[SYCL] Add offset support to parallel_for for host device (#3135)

The host device is not correctly supporting parallel_for offset argument in
runOnHost routines. Adding support and updating test coverage.

Signed-off-by: Chris Perkins chris.perkins@intel.com

Author: cperkinsintel

sycl/include/CL/sycl/detail/cg_types.hpp

@@ -208,21 +208,28 @@ class HostKernel : public HostKernelBase {
 
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
+    sycl::range<Dims> Stride(
+        InitializedVal<Dims, range>::template get<1>()); // initialized to 1
+    sycl::range<Dims> UpperBound(
+        InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I) {
       Range[I] = NDRDesc.GlobalSize[I];
       Offset[I] = NDRDesc.GlobalOffset[I];
+      UpperBound[I] = Range[I] + Offset[I];
     }
 
-    detail::NDLoop<Dims>::iterate(Range, [&](const sycl::id<Dims> &ID) {
-      sycl::item<Dims, /*Offset=*/true> Item =
-          IDBuilder::createItem<Dims, true>(Range, ID, Offset);
-
-      if (StoreLocation) {
-        store_id(&ID);
-        store_item(&Item);
-      }
-      MKernel(ID);
-    });
+    detail::NDLoop<Dims>::iterate(/*LowerBound=*/Offset, Stride, UpperBound,
+                                  [&](const sycl::id<Dims> &ID) {
+                                    sycl::item<Dims, /*Offset=*/true> Item =
+                                        IDBuilder::createItem<Dims, true>(
+                                            Range, ID, Offset);
+
+                                    if (StoreLocation) {
+                                      store_id(&ID);
+                                      store_item(&Item);
+                                    }
+                                    MKernel(ID);
+                                  });
   }
 
   template <class ArgT = KernelArgType>
@@ -259,22 +266,28 @@ class HostKernel : public HostKernelBase {
 
     sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
     sycl::id<Dims> Offset;
+    sycl::range<Dims> Stride(
+        InitializedVal<Dims, range>::template get<1>()); // initialized to 1
+    sycl::range<Dims> UpperBound(
+        InitializedVal<Dims, range>::template get<0>());
     for (int I = 0; I < Dims; ++I) {
       Range[I] = NDRDesc.GlobalSize[I];
       Offset[I] = NDRDesc.GlobalOffset[I];
+      UpperBound[I] = Range[I] + Offset[I];
     }
 
-    detail::NDLoop<Dims>::iterate(Range, [&](const sycl::id<Dims> &ID) {
-      sycl::id<Dims> OffsetID = ID + Offset;
-      sycl::item<Dims, /*Offset=*/true> Item =
-          IDBuilder::createItem<Dims, true>(Range, OffsetID, Offset);
-
-      if (StoreLocation) {
-        store_id(&OffsetID);
-        store_item(&Item);
-      }
-      MKernel(Item);
-    });
+    detail::NDLoop<Dims>::iterate(/*LowerBound=*/Offset, Stride, UpperBound,
+                                  [&](const sycl::id<Dims> &ID) {
+                                    sycl::item<Dims, /*Offset=*/true> Item =
+                                        IDBuilder::createItem<Dims, true>(
+                                            Range, ID, Offset);
+
+                                    if (StoreLocation) {
+                                      store_id(&ID);
+                                      store_item(&Item);
+                                    }
+                                    MKernel(Item);
+                                  });
   }
 
   template <class ArgT = KernelArgType>

sycl/include/CL/sycl/detail/common.hpp

@@ -277,7 +277,8 @@ template <int NDIMS> struct NDLoop {
                                            const LoopBoundTy<NDIMS> &Stride,
                                            const LoopBoundTy<NDIMS> &UpperBound,
                                            FuncTy f) {
-    LoopIndexTy<NDIMS> Index; // initialized down the call stack
+    LoopIndexTy<NDIMS> Index =
+        InitializedVal<NDIMS, LoopIndexTy>::template get<0>();
     NDLoopIterateImpl<NDIMS, NDIMS - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
         LowerBound, Stride, UpperBound, f, Index};
   }

sycl/test/basic_tests/parallel_for_range_host.cpp

@@ -16,6 +16,34 @@ int main() {
 
   queue Q(AsyncHandler);
 
+  // parallel_for, offset
+  try {
+    Q.submit([&](handler &CGH) {
+      CGH.parallel_for<class offset>(range<1>(1), id<1>(16),
+                                     [=](id<1> ID) { assert(ID == 16); });
+    });
+    Q.submit([&](handler &CGH) {
+      CGH.parallel_for<class offset_2D>(range<2>(1, 1), id<2>(16, 17),
+                                        [=](id<2> ID) {
+                                          assert(ID[0] == 16);
+                                          assert(ID[1] == 17);
+                                        });
+    });
+    Q.submit([&](handler &CGH) {
+      CGH.parallel_for<class offset_3D>(range<3>(1, 1, 1), id<3>(16, 17, 18),
+                                        [=](id<3> ID) {
+                                          assert(ID[0] == 16);
+                                          assert(ID[1] == 17);
+                                          assert(ID[2] == 18);
+                                        });
+    });
+    Q.wait_and_throw();
+  } catch (nd_range_error) {
+    std::cerr << "Test case 'offset' failed: exception has been thrown"
+              << std::endl;
+    return 1;
+  }
+
   // parallel_for, 100 global, 3 local -> fail.
   try {
     Q.submit([&](handler &CGH) {