-
Notifications
You must be signed in to change notification settings - Fork 788
[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
Conversation
@cmc-rep, could you please add a short description with an explanation what was done and how? |
e00874d
to
5bbc10c
Compare
More description added in the commit comment |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
There was a problem hiding this 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?
5bbc10c
to
8393201
Compare
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]>
8393201
to
281f0e0
Compare
@@ -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, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
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)) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
void kernel2() __attribute__((sycl_device)) { | |
// TODO: add executable tests for slm load/stores with verification | |
void kernel2() __attribute__((sycl_device)) { |
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
Signed-off-by: Gang Y Chen [email protected]