[SYCL] Fix argument size calculation for zero-dimensional accessor

Signed-off-by: Artur Gainullin <[email protected]>

Author: againull

sycl/include/CL/sycl/handler.hpp

@@ -256,7 +256,10 @@ class handler {
         detail::Requirement *AccImpl = static_cast<detail::Requirement *>(Ptr);
         MArgs.emplace_back(Kind, AccImpl, Size, Index + IndexShift);
         if (!IsKernelCreatedFromSource) {
-          const size_t SizeAccField = sizeof(size_t) * AccImpl->MDims;
+          // Dimensionality of the buffer is 1 when dimensionality of the
+          // accessor is 0.
+          const size_t SizeAccField =
+              sizeof(size_t) * (AccImpl->MDims == 0 ? 1 : AccImpl->MDims);
           ++IndexShift;
           MArgs.emplace_back(kind_std_layout, &AccImpl->MAccessRange[0],
                              SizeAccField, Index + IndexShift);

sycl/test/basic_tests/accessor/accessor.cpp

@@ -355,4 +355,28 @@ int main() {
       return 1;
     }
   }
+
+  // Accessor with dimensionality 0.
+  {
+    try {
+      int data = -1;
+      {
+        sycl::buffer<int, 1> b(&data, sycl::range<1>(1));
+        sycl::queue queue;
+        queue.submit([&](sycl::handler &cgh) {
+          sycl::accessor<int, 0, sycl::access::mode::read_write,
+                         sycl::access::target::global_buffer>
+              B(b, cgh);
+          cgh.single_task<class acc_with_zero_dim>([=]() {
+            auto B2 = B;
+            (int &)B2 = 399;
+          });
+        });
+      }
+      assert(data == 399);
+    } catch (sycl::exception e) {
+      std::cout << "SYCL exception caught: " << e.what();
+      return 1;
+    }
+  }
 }