[SYCL] Fix launching a kernel with a default constructed local access…

…or (#13382)

According to [4.7.6.11.1. Interface for local
accessors](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_interface_for_local_accessors),
passing a default constructed local accessor to a kernel function is
allowed, as long as no accesses to the data elements are made.

When default constructing a `local_accessor`, it will default construct
its `local_accessor_base`


https://github.com/intel/llvm/blob/3f841ff1d21bac2601beaf087bc3d09170af6d35/sycl/include/sycl/accessor.hpp#L2803-L2805

leading its `LocalAccessorImplHost` to be constructed with `MDims == 0`.
When setting the kernel arguments corresponding to the local accessor,
this caused `piKernelSetArg` to be called with `arg_size == 0`, which
then causes an error. This PR fixes this by checking if `MDims` would be
0 and then setting an appropriate value.

sycl/source/handler.cpp

@@ -775,7 +775,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
  // to a single kernel argument set above.
  if (!IsESIMD && !IsKernelCreatedFromSource) {
  ++IndexShift;
- const size_t SizeAccField = Dims * sizeof(Size[0]);
+ const size_t SizeAccField = (Dims == 0 ? 1 : Dims) * sizeof(Size[0]);
  MArgs.emplace_back(kernel_param_kind_t::kind_std_layout, &Size,
  SizeAccField, Index + IndexShift);
  ++IndexShift;

sycl/test-e2e/Regression/default-constructed-local-accessor.cpp

@@ -0,0 +1,20 @@
+// -O0 is necessary; on higher levels of optimization, an error
+// would not occur because of dead argument elimination of the local_accessor.
+// RUN: %{build} -o %t.out -O0
+// RUN: %{run} %t.out
+
+#include <sycl/detail/core.hpp>
+
+using namespace sycl;
+
+using acc_t = local_accessor<int, 1>;
+
+struct foo {
+ acc_t acc;
+ void operator()(nd_item<1>) const {}
+};
+
+int main() {
+ queue q;
+ q.submit([&](handler &cgh) { cgh.parallel_for(nd_range<1>(1, 1), foo{}); });
+}