Skip to content

[SYCL][ESIMD] Disallow use of accessor::operator[] in ESIMD code #5706

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.

Sign up for GitHub

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

Merged
merged 2 commits into from
Mar 4, 2022
Merged

[SYCL][ESIMD] Disallow use of accessor::operator[] in ESIMD code #5706

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
066dab8
[SYCL][ESIMD] Disallow use of accessor::operator[] in ESIMD code
sndmitriev Mar 2, 2022
ca2c224
Merge remote-tracking branch 'intel_llvm/sycl' into sndmitriev/esimd-…
sndmitriev Mar 3, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,8 @@ using namespace llvm;
// A list of unsupported functions in ESIMD context.
static const char *IllegalFunctions[] = {
"^cl::sycl::multi_ptr<.+> cl::sycl::accessor<.+>::get_pointer<.+>\\(\\) "
"const"};
"const",
" cl::sycl::accessor<.+>::operator\\[\\]<.+>\\(.+\\) const"};

namespace {

Expand Down
9 changes: 8 additions & 1 deletion sycl/test/esimd/esimd_verify.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8,10 +8,17 @@
using namespace cl::sycl;
using namespace sycl::ext::intel::experimental::esimd;

// CHECK: error: function 'cl::sycl::multi_ptr<{{.+}}> cl::sycl::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
// CHECK-DAG: error: function 'cl::sycl::multi_ptr<{{.+}}> cl::sycl::accessor<{{.+}}>::get_pointer<{{.+}}>() const' is not supported in ESIMD context
// CHECK-DAG: error: function '{{.+}} cl::sycl::accessor<{{.+}}>::operator[]<{{.+}}>({{.+}}) const' is not supported in ESIMD context

SYCL_EXTERNAL auto
test(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
SYCL_ESIMD_FUNCTION {
return acc.get_pointer();
}

SYCL_EXTERNAL void
test1(accessor<int, 1, access::mode::read_write, access::target::device> &acc)
SYCL_ESIMD_FUNCTION {
acc[0] = 0;
}
24 changes: 14 additions & 10 deletions sycl/test/esimd/hw_compile.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,16 +22,20 @@ int main(void) {
auto accB = bufB.get_access<cl::sycl::access::mode::read>(cgh);
auto accC = bufC.get_access<cl::sycl::access::mode::write>(cgh);

cgh.parallel_for<class Test>(UnitRange * UnitRange,
[=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
// those operations below would normally be
// represented as a single vector operation
// through ESIMD vector
accC[i + 0] = accA[i + 0] + accB[i + 0];
accC[i + 1] = accA[i + 1] + accB[i + 1];
accC[i + 2] = accA[i + 2] + accB[i + 2];
accC[i + 3] = accA[i + 3] + accB[i + 3];
});
cgh.parallel_for<class Test>(
UnitRange * UnitRange, [=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
int off = i.get(0) * Size;
// those operations below would normally be
// represented as a single vector operation
// through ESIMD vector
cl::sycl::ext::intel::experimental::esimd::simd<int, Size> A(
accA, off * sizeof(int));
cl::sycl::ext::intel::experimental::esimd::simd<int, Size> B(
accB, off * sizeof(int));
cl::sycl::ext::intel::experimental::esimd::simd<int, Size> C =
A + B;
C.copy_to(accC, off * sizeof(int));
});
});
}

Expand Down
11 changes: 3 additions & 8 deletions sycl/test/esimd/spirv_intrins_trans.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,16 +21,12 @@ __attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
size_t caller() {

size_t DoNotOpt[1];
cl::sycl::buffer<size_t, 1> buf(&DoNotOpt[0], 1);
uint32_t DoNotOpt32[1];
cl::sycl::buffer<uint32_t, 1> buf32(&DoNotOpt32[0], 1);

size_t DoNotOptXYZ[3];
cl::sycl::buffer<size_t, 1> bufXYZ(&DoNotOptXYZ[0], sycl::range<1>(3));

cl::sycl::queue().submit([&](cl::sycl::handler &cgh) {
auto DoNotOptimize = buf.get_access<cl::sycl::access::mode::write>(cgh);
auto DoNotOptimize32 = buf32.get_access<cl::sycl::access::mode::write>(cgh);
auto DoNotOptimize = &DoNotOpt[0];
auto DoNotOptimize32 = &DoNotOpt32[0];

kernel<class kernel_GlobalInvocationId_x>([=]() SYCL_ESIMD_KERNEL {
DoNotOptimize[0] = __spirv_GlobalInvocationId_x();
Expand Down Expand Up @@ -193,8 +189,7 @@ size_t caller() {
// x i64> %0, i64 1 %3 = extractelement <3 x i64> %0, i64 2
// In this case we will generate 3 calls to the same GenX intrinsic,
// But -early-cse will later remove this redundancy.
auto DoNotOptimizeXYZ =
bufXYZ.get_access<cl::sycl::access::mode::write>(cgh);
auto DoNotOptimizeXYZ = &DoNotOptXYZ[0];
kernel<class kernel_LocalInvocationId_xyz>([=]() SYCL_ESIMD_KERNEL {
DoNotOptimizeXYZ[0] = __spirv_LocalInvocationId_x();
DoNotOptimizeXYZ[1] = __spirv_LocalInvocationId_y();
Expand Down