From 7c9426fe6cfbcb74b2726f66f08c9a1a4211f81f Mon Sep 17 00:00:00 2001 From: HabKaffee Date: Tue, 9 Nov 2021 13:43:01 +0300 Subject: [PATCH 1/4] [SYCL] implement no_offset property for accessor_property_list class --- sycl/include/CL/sycl/accessor.hpp | 30 ++++++++++-- sycl/test/check_device_code/no_offset.cpp | 46 +++++++++++++++++++ .../check_device_code/no_offset_error.cpp | 21 +++++++++ 3 files changed, 94 insertions(+), 3 deletions(-) create mode 100644 sycl/test/check_device_code/no_offset.cpp create mode 100644 sycl/test/check_device_code/no_offset_error.cpp diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 012a257bb6707..1ea5f95e67594 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -839,8 +839,16 @@ 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())) { + Result += getOffset()[I]; + } + #else + Result += getOffset()[I]; + #endif + } return Result; } @@ -897,14 +905,26 @@ class __SYCL_SPECIAL_CLASS accessor : MData = Ptr; #pragma unroll for (int I = 0; I < AdjustedDim; ++I) { - getOffset()[I] = Offset[I]; + #if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property())) { + 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())) { + MData += Offset[0]; + } + #else MData += Offset[0]; + #endif } // __init variant used by the device compiler for ESIMD kernels. @@ -1530,6 +1550,10 @@ class __SYCL_SPECIAL_CLASS accessor : template 0)>> id get_offset() const { + #if __cplusplus >= 201703L + static_assert (!(PropertyListT::template has_property()), + "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..95e10974aea87 --- /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; +} \ No newline at end of file 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..6c02286da2609 --- /dev/null +++ b/sycl/test/check_device_code/no_offset_error.cpp @@ -0,0 +1,21 @@ +// RUN: %clangxx -fsycl-device-only -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -Xclang -verify -Xclang -verify-ignore-unexpected=note -S -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; +} \ No newline at end of file From ff8e0cc61a4836379800030c2ccb7d0a7bfe25dc Mon Sep 17 00:00:00 2001 From: HabKaffee Date: Tue, 9 Nov 2021 13:53:56 +0300 Subject: [PATCH 2/4] Clang format --- sycl/include/CL/sycl/accessor.hpp | 45 +++++++++++++++++-------------- 1 file changed, 25 insertions(+), 20 deletions(-) diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 1ea5f95e67594..3f313d2438397 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -841,13 +841,14 @@ class __SYCL_SPECIAL_CLASS accessor : __SYCL_UNROLL(3) for (int I = 0; I < Dims; ++I) { Result = Result * getMemoryRange()[I] + Id[I]; - #if __cplusplus >= 201703L - if constexpr (!(PropertyListT::template has_property())) { - Result += getOffset()[I]; - } - #else +#if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>())) { Result += getOffset()[I]; - #endif + } +#else + Result += getOffset()[I]; +#endif } return Result; } @@ -905,26 +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())) { - getOffset()[I] = Offset[I]; - } - #else +#if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>())) { getOffset()[I] = Offset[I]; - #endif + } +#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())) { +#if __cplusplus >= 201703L + if constexpr (!(PropertyListT::template has_property< + sycl::ext::oneapi::property::no_offset>())) { MData += Offset[0]; } - #else +#else MData += Offset[0]; - #endif +#endif } // __init variant used by the device compiler for ESIMD kernels. @@ -1550,10 +1553,12 @@ class __SYCL_SPECIAL_CLASS accessor : template 0)>> id get_offset() const { - #if __cplusplus >= 201703L - static_assert (!(PropertyListT::template has_property()), - "Accessor has no_offset property, get_offset() can not be used"); - #endif +#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()); } From cf14d29d8759794a8025ea29aec12c49c0618ce8 Mon Sep 17 00:00:00 2001 From: Naumov Nikita Date: Thu, 11 Nov 2021 15:54:01 +0300 Subject: [PATCH 3/4] Update sycl/test/check_device_code/no_offset_error.cpp Co-authored-by: vladimirlaz --- sycl/test/check_device_code/no_offset_error.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/no_offset_error.cpp b/sycl/test/check_device_code/no_offset_error.cpp index 6c02286da2609..35ee12f5c0ba2 100644 --- a/sycl/test/check_device_code/no_offset_error.cpp +++ b/sycl/test/check_device_code/no_offset_error.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl-device-only -fsycl-early-optimizations -fsycl-dead-args-optimization -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -Xclang -verify -Xclang -verify-ignore-unexpected=note -S -emit-llvm -o - %s +// RUN: %clangxx -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note -emit-llvm -o - %s #include From a1cef531b9c81bcd226f5f9943f6018b9f0f3f56 Mon Sep 17 00:00:00 2001 From: Naumov Nikita Date: Thu, 11 Nov 2021 21:04:51 +0300 Subject: [PATCH 4/4] Apply suggestions from code review Co-authored-by: Steffen Larsen --- sycl/test/check_device_code/no_offset.cpp | 2 +- sycl/test/check_device_code/no_offset_error.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/check_device_code/no_offset.cpp b/sycl/test/check_device_code/no_offset.cpp index 95e10974aea87..b2329b010e11e 100644 --- a/sycl/test/check_device_code/no_offset.cpp +++ b/sycl/test/check_device_code/no_offset.cpp @@ -43,4 +43,4 @@ int main() { } return 0; -} \ No newline at end of file +} diff --git a/sycl/test/check_device_code/no_offset_error.cpp b/sycl/test/check_device_code/no_offset_error.cpp index 35ee12f5c0ba2..9754c830b91fa 100644 --- a/sycl/test/check_device_code/no_offset_error.cpp +++ b/sycl/test/check_device_code/no_offset_error.cpp @@ -18,4 +18,4 @@ int main() { q.wait(); return 0; -} \ No newline at end of file +}