Skip to content

[SYCL][ESIMD] reduce restriction on some esimd API based upon user feedback #2526

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from

Conversation

cmc-rep
Copy link
Contributor

@cmc-rep cmc-rep commented Sep 23, 2020

Signed-off-by: Gang Y Chen [email protected]

@cmc-rep cmc-rep requested a review from a team as a code owner September 23, 2020 18:00
@dm-vodopyanov
Copy link
Contributor

@cmc-rep, could you please add a short description with an explanation what was done and how?

@cmc-rep
Copy link
Contributor Author

cmc-rep commented Sep 24, 2020

@cmc-rep, could you please add a short description with an explanation what was done and how?

More description added in the commit comment

Copy link
Contributor

@dm-vodopyanov dm-vodopyanov left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@cmc-rep it's better to put the description to the PR's description.
In common, code LGTM, but not approving yet. I would like @kbobrovs to review as well.

Copy link
Contributor

@kbobrovs kbobrovs left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@cmc-rep, can a test be added for

  • n = 1, n == 2, n == 4 added for gather scatter (newly enabled variants)
  • block size 8 and 16 for slm_block_load/store?

@dm-vodopyanov dm-vodopyanov changed the title reduce restriction on some esimd API based upon user feedback [SYCL][ESIMD] reduce restriction on some esimd API based upon user feedback Oct 2, 2020
@dm-vodopyanov dm-vodopyanov added the esimd Explicit SIMD feature label Oct 2, 2020
@cmc-rep
Copy link
Contributor Author

cmc-rep commented Oct 12, 2020

@cmc-rep, can a test be added for

  • n = 1, n == 2, n == 4 added for gather scatter (newly enabled variants)
  • block size 8 and 16 for slm_block_load/store?

Tests added

…edback

- allow 256-byte block load for slm
- allow vector-length of 1/2/4/8/16/32 for gather and scatter

Signed-off-by: Gang Y Chen <[email protected]>
@@ -9,7 +9,22 @@
using namespace sycl::INTEL::gpu;
using namespace cl::sycl;

void kernel(accessor<int, 1, access::mode::read_write, access::target::global_buffer> &buf) __attribute__((sycl_device)) {
void kernel0(accessor<int, 1, access::mode::read_write,
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
void kernel0(accessor<int, 1, access::mode::read_write,
// TODO: add executable tests for gather/scatter with verification
void kernel0(accessor<int, 1, access::mode::read_write,

@@ -18,3 +18,13 @@ void kernel() __attribute__((sycl_device)) {

slm_block_store<int, 32>(0, v0);
}

void kernel2() __attribute__((sycl_device)) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
void kernel2() __attribute__((sycl_device)) {
// TODO: add executable tests for slm load/stores with verification
void kernel2() __attribute__((sycl_device)) {

@cmc-rep cmc-rep closed this Jan 5, 2021
@cmc-rep cmc-rep deleted the cmc_rep-change2 branch January 5, 2021 19:55
iclsrc pushed a commit that referenced this pull request Apr 25, 2024
Don't attempt to print the object itself (which results in binary
noise) but print the relevant information from the SPIRVDecorate*
objects.

Original commit:
KhronosGroup/SPIRV-LLVM-Translator@eb0d9f6b1e755fa
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
esimd Explicit SIMD feature
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants