From 85748f673734d8325eed87d19859013f0776b46e Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Tue, 18 Apr 2023 10:49:53 -0700
Subject: [PATCH] [SYCL] Rename detail::dim_loop -> detail::loop and make it
 constexpr-friendly

Can be used like this now:

```
  loop<2>([](auto i) {
    if constexpr (i == 0)
      std::cout << "constexpr 0\n";
    std::cout << i << std::endl;
  });
```
---
 sycl/include/sycl/accessor.hpp        | 10 +++++-----
 sycl/include/sycl/detail/helpers.hpp  |  8 ++++----
 sycl/include/sycl/group_algorithm.hpp |  2 +-
 3 files changed, 10 insertions(+), 10 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 3fadaca031899..3f6cbaf429f78 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -1084,7 +1084,7 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   template <int Dims = Dimensions> size_t getLinearIndex(id<Dims> Id) const {
 
     size_t Result = 0;
-    detail::dim_loop<Dims>([&, this](size_t I) {
+    detail::loop<Dims>([&, this](size_t I) {
       Result = Result * getMemoryRange()[I] + Id[I];
       // We've already adjusted for the accessor's offset in the __init, so
       // don't include it here in case of device.
@@ -1147,7 +1147,7 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
               range<AdjustedDim> MemRange, id<AdjustedDim> Offset) {
     MData = Ptr;
-    detail::dim_loop<AdjustedDim>([&, this](size_t I) {
+    detail::loop<AdjustedDim>([&, this](size_t I) {
       if constexpr (!(PropertyListT::template has_property<
                         sycl::ext::oneapi::property::no_offset>())) {
         getOffset()[I] = Offset[I];
@@ -2252,7 +2252,7 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 #ifdef __SYCL_DEVICE_ONLY__
   size_t getTotalOffset() const noexcept {
     size_t TotalOffset = 0;
-    detail::dim_loop<Dimensions>([&, this](size_t I) {
+    detail::loop<Dimensions>([&, this](size_t I) {
       TotalOffset = TotalOffset * impl.MemRange[I];
       if constexpr (!(PropertyListT::template has_property<
                         sycl::ext::oneapi::property::no_offset>())) {
@@ -2515,7 +2515,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
               range<AdjustedDim>, id<AdjustedDim>) {
     MData = Ptr;
-    detail::dim_loop<AdjustedDim>(
+    detail::loop<AdjustedDim>(
         [&, this](size_t I) { getSize()[I] = AccessRange[I]; });
   }
 
@@ -2570,7 +2570,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   // Method which calculates linear offset for the ID using Range and Offset.
   template <int Dims = AdjustedDim> size_t getLinearIndex(id<Dims> Id) const {
     size_t Result = 0;
-    detail::dim_loop<Dims>(
+    detail::loop<Dims>(
         [&, this](size_t I) { Result = Result * getSize()[I] + Id[I]; });
     return Result;
   }
diff --git a/sycl/include/sycl/detail/helpers.hpp b/sycl/include/sycl/detail/helpers.hpp
index 9cd481a27f69c..ad173f45fa6e6 100644
--- a/sycl/include/sycl/detail/helpers.hpp
+++ b/sycl/include/sycl/detail/helpers.hpp
@@ -241,12 +241,12 @@ getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace,
 
 // To ensure loop unrolling is done when processing dimensions.
 template <size_t... Inds, class F>
-void dim_loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
-  (f(Inds), ...);
+void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
+  (f(std::integral_constant<size_t, Inds>{}), ...);
 }
 
-template <size_t count, class F> void dim_loop(F &&f) {
-  dim_loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
+template <size_t count, class F> void loop(F &&f) {
+  loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
 }
 
 } // namespace detail
diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp
index 0bedd60122548..ce3613238704f 100644
--- a/sycl/include/sycl/group_algorithm.hpp
+++ b/sycl/include/sycl/group_algorithm.hpp
@@ -253,7 +253,7 @@ reduce_over_group(Group g, sycl::vec<T, N> x, BinaryOperation binary_op) {
       "Result type of binary_op must match reduction accumulation type.");
   sycl::vec<T, N> result;
 
-  detail::dim_loop<N>(
+  detail::loop<N>(
       [&](size_t s) { result[s] = reduce_over_group(g, x[s], binary_op); });
   return result;
 }