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

This operator is not supported in ESIMD context.

Signed-off-by: Sergey Dmitriev <[email protected]>

Author: sndmitriev

llvm/lib/SYCLLowerIR/ESIMD/ESIMDVerifier.cpp

@@ -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 {
 

sycl/test/esimd/esimd_verify.cpp

@@ -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;
+}

sycl/test/esimd/hw_compile.cpp

@@ -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));
+          });
     });
   }
 

sycl/test/esimd/spirv_intrins_trans.cpp

@@ -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();
@@ -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();