diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 012a257bb6707..3f313d2438397 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -839,8 +839,17 @@ class __SYCL_SPECIAL_CLASS accessor : size_t Result = 0; // Unroll the following loop for both host and device code __SYCL_UNROLL(3) - for (int I = 0; I < Dims; ++I) - Result = Result * getMemoryRange()[I] + getOffset()[I] + Id[I]; + for (int I = 0; I < Dims; ++I) { + Result = Result * getMemoryRange()[I] + Id[I]; +#if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>())) { + Result += getOffset()[I]; + } +#else + Result += getOffset()[I]; +#endif + } return Result; } @@ -897,14 +906,28 @@ class __SYCL_SPECIAL_CLASS accessor : MData = Ptr; #pragma unroll for (int I = 0; I < AdjustedDim; ++I) { +#if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>())) { + getOffset()[I] = Offset[I]; + } +#else getOffset()[I] = Offset[I]; +#endif getAccessRange()[I] = AccessRange[I]; getMemoryRange()[I] = MemRange[I]; } // In case of 1D buffer, adjust pointer during initialization rather // then each time in operator[] or get_pointer functions. if (1 == AdjustedDim) +#if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>())) { + MData += Offset[0]; + } +#else MData += Offset[0]; +#endif } // __init variant used by the device compiler for ESIMD kernels. @@ -1530,6 +1553,12 @@ class __SYCL_SPECIAL_CLASS accessor : template 0)>> id get_offset() const { +#if __cplusplus >= 201703L + static_assert( + !(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>()), + "Accessor has no_offset property, get_offset() can not be used"); +#endif return detail::convertToArrayOfN(getOffset()); } diff --git a/sycl/test/check_device_code/no_offset.cpp b/sycl/test/check_device_code/no_offset.cpp new file mode 100644 index 0000000000000..b2329b010e11e --- /dev/null +++ b/sycl/test/check_device_code/no_offset.cpp @@ -0,0 +1,46 @@ +// RUN: %clangxx -fsycl-device-only -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -S -emit-llvm -o - %s | FileCheck %s + +#include + +inline constexpr int size = 100; + +int main() { + { + sycl::buffer a{sycl::range{size}}; + sycl::buffer b{sycl::range{size}}; + + sycl::queue q; + + q.submit([&](sycl::handler &cgh) { + sycl::ext::oneapi::accessor_property_list PL{sycl::ext::oneapi::no_offset, sycl::no_init}; + sycl::accessor acc_a(a, cgh, sycl::write_only, PL); + sycl::accessor acc_b{b, cgh, sycl::read_only}; + // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE_clES2_EUlT_E_(i32 addrspace(1)* %_arg_, i32 addrspace(1)* readonly %_arg_4, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 %_arg_8) + cgh.parallel_for(size, [=](auto i) { + acc_a[i] = acc_b[i]; + }); + }); + + q.wait(); + } + + { + sycl::buffer a{sycl::range{size}}; + sycl::buffer b{sycl::range{size}}; + + sycl::queue q; + + q.submit([&](sycl::handler &cgh) { + sycl::accessor acc_a(a, cgh, sycl::write_only); + sycl::accessor acc_b{b, cgh, sycl::read_only}; + // CHECK: define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN2cl4sycl7handlerEE0_clES2_EUlT_E_(i32 addrspace(1)* %_arg_, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 %_arg_3, i32 addrspace(1)* readonly %_arg_4, %"class.cl::sycl::id"* byval(%"class.cl::sycl::id") align 8 %_arg_8) + cgh.parallel_for(size, [=](auto i) { + acc_a[i] = acc_b[i]; + }); + }); + + q.wait(); + } + + return 0; +} diff --git a/sycl/test/check_device_code/no_offset_error.cpp b/sycl/test/check_device_code/no_offset_error.cpp new file mode 100644 index 0000000000000..9754c830b91fa --- /dev/null +++ b/sycl/test/check_device_code/no_offset_error.cpp @@ -0,0 +1,21 @@ +// RUN: %clangxx -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note -emit-llvm -o - %s + +#include + +inline constexpr int size = 100; + +int main() { + + sycl::buffer a{sycl::range{size}}; + sycl::queue q; + + q.submit([&](sycl::handler &cgh) { + sycl::ext::oneapi::accessor_property_list PL{sycl::ext::oneapi::no_offset, sycl::no_init}; + sycl::accessor acc_a(a, cgh, sycl::write_only, PL); + // expected-error@CL/sycl/accessor.hpp:* {{static_assert failed due to requirement '!(accessor_property_list, sycl::property::no_init>::has_property())' "Accessor has no_offset property, get_offset() can not be used"}} + auto b = acc_a.get_offset(); + }); + + q.wait(); + return 0; +}