[SYCL] Fix launching a kernel with a default constructed local accessor (#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.

Author: jzc

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