diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 6376a37924d5e..6251f988dc850 100644 --- a/sycl/source/handler.cpp +++ b/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; diff --git a/sycl/test-e2e/Regression/default-constructed-local-accessor.cpp b/sycl/test-e2e/Regression/default-constructed-local-accessor.cpp new file mode 100644 index 0000000000000..e136e2833a59a --- /dev/null +++ b/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 + +using namespace sycl; + +using acc_t = local_accessor; + +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{}); }); +}