[SYCL] implement no_offset property for accessor_property_list class (#4920)

Author: HabKaffee

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 <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
   id<Dimensions> 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<Dimensions, 0>(getOffset());
   }
 

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 <CL/sycl.hpp>
+
+inline constexpr int size = 100;
+
+int main() {
+    {
+        sycl::buffer<int> a{sycl::range{size}};
+        sycl::buffer<int> 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<int> a{sycl::range{size}};
+        sycl::buffer<int> 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;
+}

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 <CL/sycl.hpp>
+
+inline constexpr int size = 100;
+
+int main() {
+
+    sycl::buffer<int> 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::ext::oneapi::property::no_offset::instance<true>, 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;
+}