From feac6d2e6060bec92b667b61f67786fbcd9905ee Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Mon, 3 Oct 2022 08:41:11 -0700 Subject: [PATCH 1/2] [NFC][SYCL] Replace #pragma unroll with dim_loop in accessor.hpp The utility was introduced in https://github.com/intel/llvm/pull/6560 because "#pragma unroll" doesn't always work and template-based solution is much more reliable. Original PR only changed the loops that resulted in immediate performance difference but other occurrences were missed. This PR updates remaining ones. Note that I've found them by looking into the LLVM IR produced by our device compiler and having the loop really unrolled improves readability of such dumps (and most likely codesize/perf, although not significantly). --- sycl/include/sycl/accessor.hpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 3510d9281684c..0e64b90084d53 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -1082,8 +1082,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : void __init(ConcreteASPtrType Ptr, range AccessRange, range MemRange, id Offset) { MData = Ptr; -#pragma unroll - for (int I = 0; I < AdjustedDim; ++I) { + detail::dim_loop([&, this](size_t I) { #if __cplusplus >= 201703L if constexpr (!(PropertyListT::template has_property< sycl::ext::oneapi::property::no_offset>())) { @@ -1094,7 +1093,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor : #endif getAccessRange()[I] = AccessRange[I]; getMemoryRange()[I] = MemRange[I]; - } + }); // Adjust for offsets as that part is invariant for all invocations of // operator[]. Will have to re-adjust in get_pointer. @@ -2362,9 +2361,9 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : void __init(ConcreteASPtrType Ptr, range AccessRange, range, id) { MData = Ptr; -#pragma unroll - for (int I = 0; I < AdjustedDim; ++I) + detail::dim_loop([&, this](size_t I) { getSize()[I] = AccessRange[I]; + }); } public: From f09bcbae8af96d1ff2c7f3b3bb51fc3c6e3a323f Mon Sep 17 00:00:00 2001 From: Andrei Elovikov Date: Mon, 3 Oct 2022 08:50:09 -0700 Subject: [PATCH 2/2] clang-format --- sycl/include/sycl/accessor.hpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp index 0e64b90084d53..c060625d3329a 100644 --- a/sycl/include/sycl/accessor.hpp +++ b/sycl/include/sycl/accessor.hpp @@ -2361,9 +2361,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base : void __init(ConcreteASPtrType Ptr, range AccessRange, range, id) { MData = Ptr; - detail::dim_loop([&, this](size_t I) { - getSize()[I] = AccessRange[I]; - }); + detail::dim_loop( + [&, this](size_t I) { getSize()[I] = AccessRange[I]; }); } public: