[SYCL-MLIR]: Runtime workaround for paralllel_for (#7303)

The following test case (using a SYCL parallel_for) fails to run
correctly due to a SYCL RT function (dim_loop_impl) which uses a complex
templated implementation which "confuses" cgeist. As a workaround we can
use the older SYCL RT implementation which used a regular for loop.

Test case: 
```
void host_parallel_for(std::array<int, N> &A) {
  auto q = queue{};
  device d = q.get_device();
  std::cout << "Using " << d.get_info<info::device::name>() << "\n";
  auto range = sycl::range<1>{N};

  {
    auto buf = buffer<int, 1>{A.data(), range};
    q.submit([&](handler &cgh) {
      auto A = buf.get_access<access::mode::write>(cgh);
      cgh.parallel_for<class kernel_parallel_for>(
          range, [=](sycl::id<1> id) { A[3] = 33; });
    });
  }
}

Signed-off-by: Tiotto, Ettore <[email protected]>

Author: etiotto

sycl/include/sycl/accessor.hpp

@@ -1014,7 +1014,9 @@ class __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) {
+
+#pragma unroll
+    for (int I = 0; I < Dims; ++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.
@@ -1028,7 +1030,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
       Result += getOffset()[I];
 #endif
 #endif // __SYCL_DEVICE_ONLY__
-    });
+    }
 
     return Result;
   }
@@ -1084,7 +1086,9 @@ class __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) {
+
+#pragma unroll
+    for (int I = 0; I < AdjustedDim; ++I) {
 #if __cplusplus >= 201703L
       if constexpr (!(PropertyListT::template has_property<
                         sycl::ext::oneapi::property::no_offset>())) {
@@ -1095,7 +1099,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.
@@ -2148,7 +2152,9 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 #ifdef __SYCL_DEVICE_ONLY__
   size_t getTotalOffset() const {
     size_t TotalOffset = 0;
-    detail::dim_loop<Dimensions>([&, this](size_t I) {
+
+#pragma unroll
+    for (int I = 0; I < Dimensions; ++I) {
       TotalOffset = TotalOffset * impl.MemRange[I];
 #if __cplusplus >= 201703L
       if constexpr (!(PropertyListT::template has_property<
@@ -2158,7 +2164,7 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 #else
       TotalOffset += impl.Offset[I];
 #endif
-    });
+    }
 
     return TotalOffset;
   }
@@ -2414,8 +2420,10 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   void __init(ConcreteASPtrType Ptr, range<AdjustedDim> AccessRange,
               range<AdjustedDim>, id<AdjustedDim>) {
     MData = Ptr;
-    detail::dim_loop<AdjustedDim>(
-        [&, this](size_t I) { getSize()[I] = AccessRange[I]; });
+
+#pragma unroll
+    for (int I = 0; I < AdjustedDim; ++I)
+      getSize()[I] = AccessRange[I];
   }
 
 public: