diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 48ee02d661c05..6e22fd8323379 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -196,9 +196,6 @@ class accessor_common {
   template <int SubDims> class AccessorSubscript {
     static constexpr int Dims = Dimensions;
 
-    template <bool B, class T = void>
-    using enable_if_t = typename std::enable_if<B, T>::type;
-
     mutable id<Dims> MIDs;
     AccType MAccessor;
 
@@ -212,21 +209,22 @@ class accessor_common {
       MIDs[0] = Index;
     }
 
-    template <int CurDims = SubDims, typename = enable_if_t<(CurDims > 1)>>
+    template <int CurDims = SubDims,
+              typename = detail::enable_if_t<(CurDims > 1)>>
     AccessorSubscript<CurDims - 1> operator[](size_t Index) {
       MIDs[Dims - CurDims] = Index;
       return AccessorSubscript<CurDims - 1>(MAccessor, MIDs);
     }
 
     template <int CurDims = SubDims,
-              typename = enable_if_t<IsAccessAnyWrite && CurDims == 1>>
+              typename = detail::enable_if_t<IsAccessAnyWrite && CurDims == 1>>
     RefType operator[](size_t Index) const {
       MIDs[Dims - CurDims] = Index;
       return MAccessor[MIDs];
     }
 
     template <int CurDims = SubDims,
-              typename = enable_if_t<IsAccessReadOnly && CurDims == 1>>
+              typename = detail::enable_if_t<IsAccessReadOnly && CurDims == 1>>
     DataT operator[](size_t Index) const {
       MIDs[Dims - SubDims] = Index;
       return MAccessor[MIDs];
@@ -250,9 +248,6 @@ class accessor :
                  AccessTarget == access::target::host_buffer),
                 "Expected buffer type");
 
-  template <bool B, class T = void>
-  using enable_if_t = typename std::enable_if<B, T>::type;
-
   using AccessorCommonT = detail::accessor_common<DataT, Dimensions, AccessMode,
                                                   AccessTarget, IsPlaceholder>;
 
@@ -336,7 +331,7 @@ class accessor :
   using const_reference = const DataT &;
 
   template <int Dims = Dimensions>
-  accessor(enable_if_t<((!IsPlaceH && IsHostBuf) ||
+  accessor(detail::enable_if_t<((!IsPlaceH && IsHostBuf) ||
                         (IsPlaceH && (IsGlobalBuf || IsConstantBuf))) &&
                        Dims == 0, buffer<DataT, 1> > &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
@@ -358,7 +353,7 @@ class accessor :
   template <int Dims = Dimensions>
   accessor(
       buffer<DataT, 1> &BufferRef,
-      enable_if_t<(!IsPlaceH && (IsGlobalBuf || IsConstantBuf)) && Dims == 0,
+      detail::enable_if_t<(!IsPlaceH && (IsGlobalBuf || IsConstantBuf)) && Dims == 0,
                    handler> &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<AdjustedDim>(), BufferRef.get_range(), BufferRef.MemRange) {
@@ -374,11 +369,11 @@ class accessor :
   }
 #endif
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<((!IsPlaceH && IsHostBuf) ||
-                              (IsPlaceH && (IsGlobalBuf || IsConstantBuf))) &&
-                             (Dims > 0)>>
+  template <int Dims = Dimensions,
+            typename = detail::enable_if_t<((!IsPlaceH && IsHostBuf) ||
+                                            (IsPlaceH &&
+                                             (IsGlobalBuf || IsConstantBuf))) &&
+                                           (Dims > 0)>>
   accessor(buffer<DataT, Dimensions> &BufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(id<Dimensions>(), BufferRef.get_range(), BufferRef.MemRange) {
@@ -398,7 +393,7 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<
+            typename = detail::enable_if_t<
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf)) && (Dims > 0)>>
   accessor(buffer<DataT, Dimensions> &BufferRef, handler &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
@@ -415,11 +410,11 @@ class accessor :
   }
 #endif
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<((!IsPlaceH && IsHostBuf) ||
-                              (IsPlaceH && (IsGlobalBuf || IsConstantBuf))) &&
-                             (Dims > 0)>>
+  template <int Dims = Dimensions,
+            typename = detail::enable_if_t<((!IsPlaceH && IsHostBuf) ||
+                                            (IsPlaceH &&
+                                             (IsGlobalBuf || IsConstantBuf))) &&
+                                           (Dims > 0)>>
   accessor(buffer<DataT, Dimensions> &BufferRef, range<Dimensions> AccessRange,
            id<Dimensions> AccessOffset = {})
 #ifdef __SYCL_DEVICE_ONLY__
@@ -427,10 +422,10 @@ class accessor :
   }
 #else
       : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
-                           detail::convertToArrayOfN<3, 1>(AccessRange),
-                           detail::convertToArrayOfN<3, 1>(BufferRef.MemRange),
-                           AccessMode, detail::getSyclObjImpl(BufferRef).get(),
-                           Dimensions, sizeof(DataT)) {
+                         detail::convertToArrayOfN<3, 1>(AccessRange),
+                         detail::convertToArrayOfN<3, 1>(BufferRef.MemRange),
+                         AccessMode, detail::getSyclObjImpl(BufferRef).get(),
+                         Dimensions, sizeof(DataT)) {
     detail::EventImplPtr Event =
         detail::Scheduler::getInstance().addHostAccessor(
             AccessorBaseHost::impl.get());
@@ -439,7 +434,7 @@ class accessor :
 #endif
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<
+            typename = detail::enable_if_t<
                 (!IsPlaceH && (IsGlobalBuf || IsConstantBuf)) && (Dims > 0)>>
   accessor(buffer<DataT, Dimensions> &BufferRef, handler &CommandGroupHandler,
            range<Dimensions> AccessRange, id<Dimensions> AccessOffset = {})
@@ -448,10 +443,10 @@ class accessor :
   }
 #else
       : AccessorBaseHost(detail::convertToArrayOfN<3, 0>(AccessOffset),
-                           detail::convertToArrayOfN<3, 1>(AccessRange),
-                           detail::convertToArrayOfN<3, 1>(BufferRef.MemRange),
-                           AccessMode, detail::getSyclObjImpl(BufferRef).get(),
-                           Dimensions, sizeof(DataT)) {
+                         detail::convertToArrayOfN<3, 1>(AccessRange),
+                         detail::convertToArrayOfN<3, 1>(BufferRef.MemRange),
+                         AccessMode, detail::getSyclObjImpl(BufferRef).get(),
+                         Dimensions, sizeof(DataT)) {
     CommandGroupHandler.associateWithHandler(*this);
   }
 #endif
@@ -462,53 +457,53 @@ class accessor :
 
   size_t get_count() const { return getMemoryRange().size(); }
 
-  template <int Dims = Dimensions, typename = enable_if_t<(Dims > 0)>>
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
   range<Dimensions> get_range() const {
     return detail::convertToArrayOfN<Dimensions, 1>(getAccessRange());
   }
 
-  template <int Dims = Dimensions, typename = enable_if_t<(Dims > 0)>>
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
   id<Dimensions> get_offset() const {
     return detail::convertToArrayOfN<Dimensions, 0>(getOffset());
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessAnyWrite && Dims == 0>>
+            typename = detail::enable_if_t<IsAccessAnyWrite && Dims == 0>>
   operator RefType() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return *(getQualifiedPtr() + LinearIndex);
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessAnyWrite && (Dims > 0)>>
+            typename = detail::enable_if_t<IsAccessAnyWrite && (Dims > 0)>>
   RefType operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return getQualifiedPtr()[LinearIndex];
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessAnyWrite && Dims == 1>>
+            typename = detail::enable_if_t<IsAccessAnyWrite && Dims == 1>>
   RefType operator[](size_t Index) const {
     const size_t LinearIndex = getLinearIndex(id<Dimensions>(Index));
     return getQualifiedPtr()[LinearIndex];
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessReadOnly && Dims == 0>>
+            typename = detail::enable_if_t<IsAccessReadOnly && Dims == 0>>
   operator DataT() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return *(getQualifiedPtr() + LinearIndex);
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessReadOnly && (Dims > 0)>>
+            typename = detail::enable_if_t<IsAccessReadOnly && (Dims > 0)>>
   DataT operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return getQualifiedPtr()[LinearIndex];
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessReadOnly && Dims == 1>>
+            typename = detail::enable_if_t<IsAccessReadOnly && Dims == 1>>
   DataT operator[](size_t Index) const {
     const size_t LinearIndex = getLinearIndex(id<Dimensions>(Index));
     return getQualifiedPtr()[LinearIndex];
@@ -533,39 +528,39 @@ class accessor :
   }
 
   template <int Dims = Dimensions>
-  typename enable_if_t<AccessMode == access::mode::atomic && Dims == 1,
-                       atomic<DataT, AS>>::type
+  typename detail::enable_if_t<AccessMode == access::mode::atomic && Dims == 1,
+                               atomic<DataT, AS>>::type
   operator[](size_t Index) const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
   }
 
-  template <int Dims = Dimensions, typename = enable_if_t<(Dims > 1)>>
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
   typename AccessorCommonT::template AccessorSubscript<Dims - 1>
   operator[](size_t Index) const {
     return AccessorSubscript<Dims - 1>(*this, Index);
   }
 
-  template <
-      access::target AccessTarget_ = AccessTarget,
-      typename = enable_if_t<AccessTarget_ == access::target::host_buffer>>
+  template <access::target AccessTarget_ = AccessTarget,
+            typename = detail::enable_if_t<AccessTarget_ ==
+                                           access::target::host_buffer>>
   DataT *get_pointer() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return getQualifiedPtr() + LinearIndex;
   }
 
-  template <
-      access::target AccessTarget_ = AccessTarget,
-      typename = enable_if_t<AccessTarget_ == access::target::global_buffer>>
+  template <access::target AccessTarget_ = AccessTarget,
+            typename = detail::enable_if_t<AccessTarget_ ==
+                                           access::target::global_buffer>>
   global_ptr<DataT> get_pointer() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return global_ptr<DataT>(getQualifiedPtr() + LinearIndex);
   }
 
-  template <
-      access::target AccessTarget_ = AccessTarget,
-      typename = enable_if_t<AccessTarget_ == access::target::constant_buffer>>
+  template <access::target AccessTarget_ = AccessTarget,
+            typename = detail::enable_if_t<AccessTarget_ ==
+                                           access::target::constant_buffer>>
   constant_ptr<DataT> get_pointer() const {
     const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
     return constant_ptr<DataT>(getQualifiedPtr() + LinearIndex);
@@ -601,9 +596,6 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
   using RefType = typename detail::PtrValueType<DataT, AS>::type &;
   using PtrType = typename detail::PtrValueType<DataT, AS>::type *;
 
-  template <bool B, class T = void>
-  using enable_if_t = typename std::enable_if<B, T>::type;
-
 #ifdef __SYCL_DEVICE_ONLY__
   detail::LocalAccessorBaseDevice<AdjustedDim> impl;
 
@@ -646,7 +638,7 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
   using reference = DataT &;
   using const_reference = const DataT &;
 
-  template <int Dims = Dimensions, typename = enable_if_t<Dims == 0>>
+  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
   accessor(handler &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(range<AdjustedDim>{1}) {
@@ -656,7 +648,7 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
   }
 #endif
 
-  template <int Dims = Dimensions, typename = enable_if_t<(Dims > 0)>>
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
   accessor(range<Dimensions> AllocationSize, handler &CommandGroupHandler)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AllocationSize) {
@@ -672,48 +664,48 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
   size_t get_count() const { return getSize().size(); }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessAnyWrite && Dims == 0>>
+            typename = detail::enable_if_t<IsAccessAnyWrite && Dims == 0>>
   operator RefType() const {
     return *getQualifiedPtr();
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessAnyWrite && (Dims > 0)>>
+            typename = detail::enable_if_t<IsAccessAnyWrite && (Dims > 0)>>
   RefType operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return getQualifiedPtr()[LinearIndex];
   }
 
   template <int Dims = Dimensions,
-            typename = enable_if_t<IsAccessAnyWrite && Dims == 1>>
+            typename = detail::enable_if_t<IsAccessAnyWrite && Dims == 1>>
   RefType operator[](size_t Index) const {
     return getQualifiedPtr()[Index];
   }
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<AccessMode == access::mode::atomic && Dims == 0>>
+  template <int Dims = Dimensions,
+            typename = detail::enable_if_t<AccessMode == access::mode::atomic &&
+                                           Dims == 0>>
   operator atomic<DataT, AS>() const {
     return atomic<DataT, AS>(multi_ptr<DataT, AS>(getQualifiedPtr()));
   }
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<AccessMode == access::mode::atomic && (Dims > 0)>>
+  template <int Dims = Dimensions,
+            typename = detail::enable_if_t<AccessMode == access::mode::atomic &&
+                                           (Dims > 0)>>
   atomic<DataT, AS> operator[](id<Dimensions> Index) const {
     const size_t LinearIndex = getLinearIndex(Index);
     return atomic<DataT, AS>(
         multi_ptr<DataT, AS>(getQualifiedPtr() + LinearIndex));
   }
 
-  template <
-      int Dims = Dimensions,
-      typename = enable_if_t<AccessMode == access::mode::atomic && Dims == 1>>
+  template <int Dims = Dimensions,
+            typename = detail::enable_if_t<AccessMode == access::mode::atomic &&
+                                           Dims == 1>>
   atomic<DataT, AS> operator[](size_t Index) const {
     return atomic<DataT, AS>(multi_ptr<DataT, AS>(getQualifiedPtr() + Index));
   }
 
-  template <int Dims = Dimensions, typename = enable_if_t<(Dims > 1)>>
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
   typename AccessorCommonT::template AccessorSubscript<Dims - 1>
   operator[](size_t Index) const {
     return AccessorSubscript<Dims - 1>(*this, Index);
diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp
index 1a8eb3b98e336..73c0619e4e206 100644
--- a/sycl/include/CL/sycl/detail/cg.hpp
+++ b/sycl/include/CL/sycl/detail/cg.hpp
@@ -9,8 +9,10 @@
 #pragma once
 
 #include <CL/sycl/detail/accessor_impl.hpp>
+#include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/helpers.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
+#include <CL/sycl/group.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/nd_item.hpp>
@@ -45,11 +47,12 @@ class ArgDesc {
 class NDRDescT {
   // The method initializes all sizes for dimensions greater than the passed one
   // to the default values, so they will not affect execution.
-  template <int Dims_> void setNDRangeLeftover() {
+  void setNDRangeLeftover(int Dims_) {
     for (int I = Dims_; I < 3; ++I) {
       GlobalSize[I] = 1;
       LocalSize[I] = LocalSize[0] ? 1 : 0;
       GlobalOffset[I] = 0;
+      NumWorkGroups[I] = 0;
     }
   }
 
@@ -61,9 +64,21 @@ class NDRDescT {
       GlobalSize[I] = NumWorkItems[I];
       LocalSize[I] = 0;
       GlobalOffset[I] = 0;
+      NumWorkGroups[I] = 0;
     }
+    setNDRangeLeftover(Dims_);
+    Dims = Dims_;
+  }
 
-    setNDRangeLeftover<Dims_>();
+  template <int Dims_>
+  void set(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset) {
+    for (int I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = NumWorkItems[I];
+      LocalSize[I] = 0;
+      GlobalOffset[I] = Offset[I];
+      NumWorkGroups[I] = 0;
+    }
+    setNDRangeLeftover(Dims_);
     Dims = Dims_;
   }
 
@@ -72,14 +87,42 @@ class NDRDescT {
       GlobalSize[I] = ExecutionRange.get_global_range()[I];
       LocalSize[I] = ExecutionRange.get_local_range()[I];
       GlobalOffset[I] = ExecutionRange.get_offset()[I];
+      NumWorkGroups[I] = 0;
     }
-    setNDRangeLeftover<Dims_>();
+    setNDRangeLeftover(Dims_);
+    Dims = Dims_;
+  }
+
+  void set(int Dims_, sycl::nd_range<3> ExecutionRange) {
+    for (int I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = ExecutionRange.get_global_range()[I];
+      LocalSize[I] = ExecutionRange.get_local_range()[I];
+      GlobalOffset[I] = ExecutionRange.get_offset()[I];
+      NumWorkGroups[I] = 0;
+    }
+    setNDRangeLeftover(Dims_);
+    Dims = Dims_;
+  }
+
+  template <int Dims_> void setNumWorkGroups(sycl::range<Dims_> N) {
+    for (int I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = 0;
+      // '0' is a mark to adjust before kernel launch when there is enough info:
+      LocalSize[I] = 0;
+      GlobalOffset[I] = 0;
+      NumWorkGroups[I] = N[I];
+    }
+    setNDRangeLeftover(Dims_);
     Dims = Dims_;
   }
 
   sycl::range<3> GlobalSize;
   sycl::range<3> LocalSize;
   sycl::id<3> GlobalOffset;
+  /// Number of workgroups, used to record the number of workgroups from the
+  /// simplest form of parallel_for_work_group. If set, all other fields must be
+  /// zero
+  sycl::range<3> NumWorkGroups;
   size_t Dims;
 };
 
@@ -102,7 +145,20 @@ class HostKernel : public HostKernelBase {
 
 public:
   HostKernel(KernelType Kernel) : MKernel(Kernel) {}
-  void call(const NDRDescT &NDRDesc) override { runOnHost(NDRDesc); }
+  void call(const NDRDescT &NDRDesc) override {
+    // adjust ND range for serial host:
+    NDRDescT R1;
+    bool Adjust = false;
+
+    if (NDRDesc.GlobalSize[0] == 0 && NDRDesc.NumWorkGroups[0] != 0) {
+      // LocalSize - which is the workgroup size - is not set; need to adjust
+      range<3> WGsize = {1, 1, 1}; // no better alternative for serial host?
+      R1.set(NDRDesc.Dims, nd_range<3>(NDRDesc.NumWorkGroups * WGsize, WGsize));
+      Adjust = true;
+    }
+    const NDRDescT &R = Adjust ? R1 : NDRDesc;
+    runOnHost(R);
+  }
 
   char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
 
@@ -165,6 +221,7 @@ class HostKernel : public HostKernelBase {
     sycl::id<3> GroupSize;
     for (int I = 0; I < 3; ++I) {
       GroupSize[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
+      // TODO supoport case NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0
     }
 
     sycl::range<Dims> GlobalSize;
@@ -217,6 +274,30 @@ class HostKernel : public HostKernelBase {
       }
     }
   }
+
+  template <typename ArgT = KernelArgType>
+  enable_if_t<std::is_same<ArgT, cl::sycl::group<Dims>>::value>
+  runOnHost(const NDRDescT &NDRDesc) {
+    sycl::id<Dims> NGroups;
+
+    for (int I = 0; I < Dims; ++I) {
+      NGroups[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
+      assert(NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] == 0);
+    }
+    sycl::range<Dims> GlobalSize;
+    sycl::range<Dims> LocalSize;
+
+    for (int I = 0; I < Dims; ++I) {
+      LocalSize[I] = NDRDesc.LocalSize[I];
+      GlobalSize[I] = NDRDesc.GlobalSize[I];
+    }
+    detail::NDLoop<Dims>::iterate(NGroups, [&](const id<Dims> &GroupID) {
+      sycl::group<Dims> Group =
+          IDBuilder::createGroup<Dims>(GlobalSize, LocalSize, GroupID);
+      MKernel(Group);
+    });
+  }
+
   ~HostKernel() = default;
 };
 
diff --git a/sycl/include/CL/sycl/detail/cnri.h b/sycl/include/CL/sycl/detail/cnri.h
index 5dc0d79e4d121..1a4ba5dd321bf 100644
--- a/sycl/include/CL/sycl/detail/cnri.h
+++ b/sycl/include/CL/sycl/detail/cnri.h
@@ -93,6 +93,7 @@ typedef cl_context cnri_context;
 typedef cl_event cnri_event;
 typedef cl_program cnri_program;
 typedef cl_kernel cnri_kernel;
+typedef cl_device_id cnri_device;
 
 enum { CNRI_SUCCESS = CL_SUCCESS };
 
diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp
index 4031e138f8d47..1dd2994198cee 100644
--- a/sycl/include/CL/sycl/detail/common.hpp
+++ b/sycl/include/CL/sycl/detail/common.hpp
@@ -103,6 +103,105 @@ template <class T> T createSyclObjFromImpl(decltype(T::impl) ImplObj) {
   return T(ImplObj);
 }
 
+// Produces N-dimensional object of type T whose all components are initialized
+// to given integer value.
+template <int N, template <int> class T> struct InitializedVal {
+  template <int Val> static T<N> &&get();
+};
+
+// Specialization for a one-dimensional type.
+template <template <int> class T> struct InitializedVal<1, T> {
+  template <int Val> static T<1> &&get() { return T<1>{Val}; }
+};
+
+// Specialization for a two-dimensional type.
+template <template <int> class T> struct InitializedVal<2, T> {
+  template <int Val> static T<2> &&get() { return T<2>{Val, Val}; }
+};
+
+// Specialization for a three-dimensional type.
+template <template <int> class T> struct InitializedVal<3, T> {
+  template <int Val> static T<3> &&get() { return T<3>{Val, Val, Val}; }
+};
+
+// Fills the lack of enable_if_t in C++11.
+template <bool Cond, typename T = void>
+using enable_if_t = typename std::enable_if<Cond, T>::type;
+
+/// Helper class for the \c NDLoop.
+template <int NDIMS, int DIM, template <int> typename LoopBoundTy,
+          typename FuncTy, template <int> typename LoopIndexTy>
+struct NDLoopIterateImpl {
+  NDLoopIterateImpl(const LoopIndexTy<NDIMS> &LowerBound,
+                    const LoopBoundTy<NDIMS> &Stride,
+                    const LoopBoundTy<NDIMS> &UpperBound, FuncTy f,
+                    LoopIndexTy<NDIMS> &Index) {
+
+    for (Index[DIM] = LowerBound[DIM]; Index[DIM] < UpperBound[DIM];
+         Index[DIM] += Stride[DIM]) {
+
+      NDLoopIterateImpl<NDIMS, DIM - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
+          LowerBound, Stride, UpperBound, f, Index};
+    }
+  }
+};
+
+// spcialization for DIM=0 to terminate recursion
+template <int NDIMS, template <int> typename LoopBoundTy, typename FuncTy,
+          template <int> typename LoopIndexTy>
+struct NDLoopIterateImpl<NDIMS, 0, LoopBoundTy, FuncTy, LoopIndexTy> {
+  NDLoopIterateImpl(const LoopIndexTy<NDIMS> &LowerBound,
+                    const LoopBoundTy<NDIMS> &Stride,
+                    const LoopBoundTy<NDIMS> &UpperBound, FuncTy f,
+                    LoopIndexTy<NDIMS> &Index) {
+
+    for (Index[0] = LowerBound[0]; Index[0] < UpperBound[0];
+         Index[0] += Stride[0]) {
+
+      f(Index);
+    }
+  }
+};
+
+/// Generates an NDIMS-dimensional perfect loop nest. The purpose of this class
+/// is to better support handling of situations where there must be a loop nest
+/// over a multi-dimensional space - it allows to avoid generating unnecessary
+/// outer loops like 'for (int z=0; z<1; z++)' in case of 1D and 2D iteration
+/// spaces or writing specializations of the algorithms for 1D, 2D and 3D cases.
+template <int NDIMS> struct NDLoop {
+  /// Generates ND loop nest with {0,..0} .. \c UpperBound bounds with unit
+  /// stride. Applies \c f at each iteration, passing current index of
+  /// \c LoopIndexTy<NDIMS> type as the parameter.
+  template <template <int> class LoopBoundTy, typename FuncTy,
+            template <int> class LoopIndexTy = LoopBoundTy>
+  static ALWAYS_INLINE void iterate(const LoopBoundTy<NDIMS> &UpperBound,
+                                    FuncTy f) {
+    const LoopIndexTy<NDIMS> LowerBound =
+        InitializedVal<NDIMS, LoopIndexTy>::template get<0>();
+    const LoopBoundTy<NDIMS> Stride =
+        InitializedVal<NDIMS, LoopBoundTy>::template get<1>();
+    LoopIndexTy<NDIMS> Index; // initialized down the call stack
+
+    NDLoopIterateImpl<NDIMS, NDIMS-1, LoopBoundTy, FuncTy, LoopIndexTy>{
+        LowerBound, Stride, UpperBound, f, Index};
+  }
+
+  /// Generates ND loop nest with \c LowerBound .. \c UpperBound bounds and
+  /// stride \c Stride. Applies \c f at each iteration, passing current index of
+  /// \c LoopIndexTy<NDIMS> type as the parameter.
+  template <template <int> class LoopBoundTy, typename FuncTy,
+            template <int> class LoopIndexTy = LoopBoundTy>
+  static ALWAYS_INLINE void iterate(const LoopIndexTy<NDIMS> &LowerBound,
+                                    const LoopBoundTy<NDIMS> &Stride,
+                                    const LoopBoundTy<NDIMS> &UpperBound,
+                                    FuncTy f) {
+    LoopIndexTy<NDIMS> Index; // initialized down the call stack
+    NDLoopIterateImpl<NDIMS, NDIMS-1, LoopBoundTy, FuncTy, LoopIndexTy>{
+        LowerBound, Stride, UpperBound, f, Index};
+  }
+};
+
 } // namespace detail
 } // namespace sycl
 } // namespace cl
+
diff --git a/sycl/include/CL/sycl/detail/helpers.hpp b/sycl/include/CL/sycl/detail/helpers.hpp
index 8fd73d5ef1cfa..cb349ecc6fe8c 100644
--- a/sycl/include/CL/sycl/detail/helpers.hpp
+++ b/sycl/include/CL/sycl/detail/helpers.hpp
@@ -11,7 +11,6 @@
 #include <CL/__spirv/spirv_types.hpp>
 #include <CL/sycl/access/access.hpp>
 #include <CL/sycl/detail/common.hpp>
-
 #include <memory>
 #include <stdexcept>
 #include <type_traits>
@@ -27,6 +26,8 @@ template <int dimensions> class range;
 template <int dimensions> struct id;
 template <int dimensions> class nd_item;
 enum class memory_order;
+template <int dimensions> class h_item;
+
 namespace detail {
 class context_impl;
 // The function returns list of events that can be passed to OpenCL API as
@@ -69,6 +70,21 @@ struct Builder {
                const cl::sycl::group<dimensions> &GR) {
     return cl::sycl::nd_item<dimensions>(GL, L, GR);
   }
+
+  template <int dimensions>
+  static h_item<dimensions>
+  createHItem(const cl::sycl::item<dimensions, false> &GlobalItem,
+              const cl::sycl::item<dimensions, false> &LocalItem) {
+    return cl::sycl::h_item<dimensions>(GlobalItem, LocalItem);
+  }
+
+  template <int dimensions>
+  static h_item<dimensions>
+  createHItem(const cl::sycl::item<dimensions, false> &GlobalItem,
+              const cl::sycl::item<dimensions, false> &LocalItem,
+              const cl::sycl::range<dimensions> &FlexRange) {
+    return cl::sycl::h_item<dimensions>(GlobalItem, LocalItem, FlexRange);
+  }
 };
 
 inline __spv::MemorySemanticsMask getSPIRVMemorySemanticsMask(memory_order) {
diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp
index 014f500699350..fd0c892501007 100644
--- a/sycl/include/CL/sycl/group.hpp
+++ b/sycl/include/CL/sycl/group.hpp
@@ -9,8 +9,12 @@
 #pragma once
 
 #include <CL/__spirv/spirv_ops.hpp>
+#include <CL/__spirv/spirv_types.hpp>
+#include <CL/__spirv/spirv_vars.hpp>
+#include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/helpers.hpp>
 #include <CL/sycl/device_event.hpp>
+#include <CL/sycl/h_item.hpp>
 #include <CL/sycl/id.hpp>
 #include <CL/sycl/pointers.hpp>
 #include <CL/sycl/range.hpp>
@@ -21,6 +25,19 @@ namespace cl {
 namespace sycl {
 namespace detail {
 struct Builder;
+
+// Implements a barrier accross work items within a work group.
+static inline void workGroupBarrier() {
+#ifdef __SYCL_DEVICE_ONLY__
+  uint32_t flags =
+      static_cast<uint32_t>(
+          __spv::MemorySemanticsMask::SequentiallyConsistent) |
+      static_cast<uint32_t>(__spv::MemorySemanticsMask::WorkgroupMemory);
+  __spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
+                         flags);
+#endif // __SYCL_DEVICE_ONLY__
+}
+
 } // namespace detail
 
 template <int dimensions = 1> class group {
@@ -66,12 +83,109 @@ template <int dimensions = 1> class group {
            index[0];
   }
 
-  // template<typename workItemFunctionT>
-  // void parallel_for_work_item(workItemFunctionT func) const;
+  template <typename workItemFunctionT>
+  void parallel_for_work_item(workItemFunctionT func) const {
+    // need barriers to enforce SYCL semantics for the work item loop -
+    // compilers are expected to optimize when possible
+    detail::workGroupBarrier();
+#ifdef __SYCL_DEVICE_ONLY__
+    range<dimensions> global_size;
+    range<dimensions> local_size;
+    id<dimensions> global_id;
+    id<dimensions> local_id;
+
+    __spirv::initGlobalSize<dimensions>(global_size);
+    __spirv::initWorkgroupSize<dimensions>(local_size);
+    __spirv::initGlobalInvocationId<dimensions>(global_id);
+    __spirv::initLocalInvocationId<dimensions>(local_id);
+
+    // no 'iterate' in the device code variant, because
+    // (1) this code is already invoked by each work item as a part of the
+    //     enclosing parallel_for_work_group kernel
+    // (2) the range this pfwi iterates over matches work group size exactly
+    item<dimensions, false> globalItem =
+        detail::Builder::createItem<dimensions, false>(global_size, global_id);
+    item<dimensions, false> localItem =
+        detail::Builder::createItem<dimensions, false>(local_size, local_id);
+    h_item<dimensions> hItem =
+        detail::Builder::createHItem<dimensions>(globalItem, localItem);
+
+    func(hItem);
+#else
+    id<dimensions> GroupStartID = index * localRange;
+
+    // ... host variant needs explicit 'iterate' because it is serial
+    detail::NDLoop<dimensions>::iterate(
+        localRange, [&](const id<dimensions> &LocalID) {
+          item<dimensions, false> GlobalItem =
+              detail::Builder::createItem<dimensions, false>(
+                  globalRange, GroupStartID + LocalID);
+          item<dimensions, false> LocalItem =
+              detail::Builder::createItem<dimensions, false>(localRange,
+                                                             LocalID);
+          h_item<dimensions> hItem =
+              detail::Builder::createHItem<dimensions>(GlobalItem, LocalItem);
+          func(hItem);
+        });
+#endif // __SYCL_DEVICE_ONLY__
+    detail::workGroupBarrier();
+  }
 
-  // template<typename workItemFunctionT>
-  // void parallel_for_work_item(range<dimensions> flexibleRange,
-  // workItemFunctionT func) const;
+  template <typename workItemFunctionT>
+  void parallel_for_work_item(range<dimensions> flexibleRange,
+                              workItemFunctionT func) const {
+    detail::workGroupBarrier();
+#ifdef __SYCL_DEVICE_ONLY__
+    range<dimensions> global_size;
+    range<dimensions> local_size;
+    id<dimensions> global_id;
+    id<dimensions> local_id;
+
+    __spirv::initGlobalSize<dimensions>(global_size);
+    __spirv::initWorkgroupSize<dimensions>(local_size);
+    __spirv::initGlobalInvocationId<dimensions>(global_id);
+    __spirv::initLocalInvocationId<dimensions>(local_id);
+
+    item<dimensions, false> globalItem =
+        detail::Builder::createItem<dimensions, false>(global_size, global_id);
+    item<dimensions, false> localItem =
+        detail::Builder::createItem<dimensions, false>(local_size, local_id);
+    h_item<dimensions> hItem =
+        detail::Builder::createHItem<dimensions>(globalItem, localItem);
+
+    // iterate over flexible range with work group size stride; each item
+    // performs flexibleRange/local_size iterations (if the former is divisible
+    // by the latter)
+    detail::NDLoop<dimensions>::iterate(
+        local_id, local_size, flexibleRange,
+        [&](const id<dimensions> &LogicalLocalID) {
+          hItem.setLogicalLocalID(LogicalLocalID);
+          func(hItem);
+        });
+#else
+    id<dimensions> GroupStartID = index * localRange;
+
+    detail::NDLoop<dimensions>::iterate(
+        localRange, [&](const id<dimensions> &LocalID) {
+          item<dimensions, false> GlobalItem =
+              detail::Builder::createItem<dimensions, false>(
+                  globalRange, GroupStartID + LocalID);
+          item<dimensions, false> LocalItem =
+              detail::Builder::createItem<dimensions, false>(localRange,
+                                                             LocalID);
+          h_item<dimensions> hItem =
+              detail::Builder::createHItem<dimensions>(GlobalItem, LocalItem);
+
+          detail::NDLoop<dimensions>::iterate(
+              LocalID, localRange, flexibleRange,
+              [&](const id<dimensions> &LogicalLocalID) {
+                hItem.setLogicalLocalID(LogicalLocalID);
+                func(hItem);
+              });
+        });
+#endif // __SYCL_DEVICE_ONLY__
+    detail::workGroupBarrier();
+  }
 
   /// Executes a work-group mem-fence with memory ordering on the local address
   /// space, global address space or both based on the value of \p accessSpace.
diff --git a/sycl/include/CL/sycl/h_item.hpp b/sycl/include/CL/sycl/h_item.hpp
new file mode 100644
index 0000000000000..c781c85072cad
--- /dev/null
+++ b/sycl/include/CL/sycl/h_item.hpp
@@ -0,0 +1,131 @@
+//==-------------- h_item.hpp - SYCL standard header file ------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/id.hpp>
+#include <CL/sycl/item.hpp>
+
+namespace cl {
+namespace sycl {
+
+namespace detail {
+class Builder;
+}
+
+template <int dimensions> class h_item {
+public:
+  h_item() = delete;
+
+  h_item(const h_item<dimensions> &hi) = default;
+
+  h_item<dimensions> &operator=(const h_item<dimensions> &hi) = default;
+
+  /* -- public interface members -- */
+  item<dimensions, false> get_global() const { return globalItem; }
+
+  item<dimensions, false> get_local() const { return get_logical_local(); }
+
+  item<dimensions, false> get_logical_local() const { return logicalLocalItem; }
+
+  item<dimensions, false> get_physical_local() const { return localItem; }
+
+  range<dimensions> get_global_range() const {
+    return get_global().get_range();
+  }
+
+  size_t get_global_range(int dimension) const {
+    return get_global().get_range(dimension);
+  }
+
+  id<dimensions> get_global_id() const { return get_global().get_id(); }
+
+  size_t get_global_id(int dimension) const {
+    return get_global().get_id(dimension);
+  }
+
+  range<dimensions> get_local_range() const { return get_local().get_range(); }
+
+  size_t get_local_range(int dimension) const {
+    return get_local().get_range(dimension);
+  }
+
+  id<dimensions> get_local_id() const { return get_local().get_id(); }
+
+  size_t get_local_id(int dimension) const {
+    return get_local().get_id(dimension);
+  }
+
+  range<dimensions> get_logical_local_range() const {
+    return get_logical_local().get_range();
+  }
+
+  size_t get_logical_local_range(int dimension) const {
+    return get_logical_local().get_range(dimension);
+  }
+
+  id<dimensions> get_logical_local_id() const {
+    return get_logical_local().get_id();
+  }
+
+  size_t get_logical_local_id(int dimension) const {
+    return get_logical_local().get_id(dimension);
+  }
+
+  range<dimensions> get_physical_local_range() const {
+    return get_physical_local().get_range();
+  }
+
+  size_t get_physical_local_range(int dimension) const {
+    return get_physical_local().get_range(dimension);
+  }
+
+  id<dimensions> get_physical_local_id() const {
+    return get_physical_local().get_id();
+  }
+
+  size_t get_physical_local_id(int dimension) const {
+    return get_physical_local().get_id(dimension);
+  }
+
+  bool operator==(const h_item<dimensions> &rhs) const {
+    return (rhs.localItem == localItem) && (rhs.globalItem == globalItem) &&
+           (rhs.logicalLocalItem == logicalLocalItem);
+  }
+
+  bool operator!=(const h_item<dimensions> &rhs) const {
+    return !((*this) == rhs);
+  }
+
+protected:
+  friend class detail::Builder;
+  friend class group<dimensions>;
+  h_item(const item<dimensions, false> &GL, const item<dimensions, false> &L,
+         const range<dimensions> &flexLocalRange)
+      : globalItem(GL), localItem(L),
+        logicalLocalItem(flexLocalRange, L.get_id()) {}
+
+  h_item(const item<dimensions, false> &GL, const item<dimensions, false> &L)
+      : globalItem(GL), localItem(L),
+        logicalLocalItem(localItem.get_range(), localItem.get_id()) {}
+
+  void setLogicalLocalID(const id<dimensions> &ID) {
+    logicalLocalItem.setID(ID);
+  }
+
+private:
+  /// Describles physical workgroup range and current \c h_item position in it.
+  item<dimensions, false> localItem;
+  /// Describles global range and current \c h_item position in it.
+  item<dimensions, false> globalItem;
+  /// Describles logical flexible range and current \c h_item position in it.
+  item<dimensions, false> logicalLocalItem;
+};
+
+} // namespace sycl
+} // namespace cl
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index 62ad16e71cd35..68e6c572b8033 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -663,16 +663,51 @@ class handler {
 #endif
   }
 
-  // template <typename KernelName, typename WorkgroupFunctionType, int
-  // dimensions>
-  // void parallel_for_work_group(range<dimensions> numWorkGroups,
-  //                              WorkgroupFunctionType KernelFunc);
+  template <typename KernelName = csd::auto_name, typename KernelType, int Dims>
+  void parallel_for_work_group(range<Dims> NumWorkGroups,
+                               KernelType KernelFunc) {
+    using NameT = typename csd::get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
+#else
+    MNDRDesc.setNumWorkGroups(NumWorkGroups);
+    StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
+    MCGType = detail::CG::KERNEL;
+#endif // __SYCL_DEVICE_ONLY__
+  }
+
+#ifdef __SYCL_DEVICE_ONLY__
+  template <typename KernelName, typename KernelType, int Dims>
+  __attribute__((sycl_kernel)) void
+  kernel_parallel_for_work_group(KernelType KernelFunc) {
 
-  // template <typename KernelName, typename WorkgroupFunctionType, int
-  // dimensions>
-  // void parallel_for_work_group(range<dimensions> numWorkGroups,
-  //                              range<dimensions> workGroupSize,
-  //                              WorkgroupFunctionType KernelFunc);
+    range<Dims> GlobalSize;
+    range<Dims> LocalSize;
+    id<Dims> GroupId;
+
+    __spirv::initGlobalSize<Dims>(GlobalSize);
+    __spirv::initWorkgroupSize<Dims>(LocalSize);
+    __spirv::initWorkgroupId<Dims>(GroupId);
+
+    group<Dims> G =
+        detail::Builder::createGroup<Dims>(GlobalSize, LocalSize, GroupId);
+    KernelFunc(G);
+  }
+#endif // __SYCL_DEVICE_ONLY__
+
+  template <typename KernelName = csd::auto_name, typename KernelType, int Dims>
+  void parallel_for_work_group(range<Dims> NumWorkGroups,
+                               range<Dims> WorkGroupSize,
+                               KernelType KernelFunc) {
+    using NameT = typename csd::get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
+#else
+    MNDRDesc.set(nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize));
+    StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
+    MCGType = detail::CG::KERNEL;
+#endif // __SYCL_DEVICE_ONLY__
+  }
 
   // single_task version with a kernel represented as a sycl::kernel.
   // The kernel invocation method has no functors and cannot be called on host.
@@ -804,16 +839,42 @@ class handler {
 #endif
   }
 
-  // template <typename KernelName, typename WorkgroupFunctionType, int
-  // dimensions>
-  // void parallel_for_work_group(range<dimensions> num_work_groups, kernel
-  // SyclKernel, WorkgroupFunctionType KernelFunc);
+  /// This version of \c parallel_for_work_group takes two parameters
+  /// representing the same kernel. The first one - \c syclKernel - is a
+  /// compiled form of the second one - \c kernelFunc, which is the source form
+  /// of the kernel. The same source kernel can be compiled multiple times
+  /// yielding multiple kernel class objects accessible via the \c program class
+  /// interface.
+  template <typename KernelName = csd::auto_name, typename KernelType, int Dims>
+  void parallel_for_work_group(kernel SyclKernel, range<Dims> NumWorkGroups,
+                               KernelType KernelFunc) {
+    using NameT = typename csd::get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
+#else
+    MNDRDesc.setNumWorkGroups(NumWorkGroups);
+    MSyclKernel = detail::getSyclObjImpl(std::move(SyclKernel));
+    StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
+    MCGType = detail::CG::KERNEL;
+#endif // __SYCL_DEVICE_ONLY__
+  }
 
-  // template <typename KernelName, typename WorkgroupFunctionType, int
-  // dimensions>
-  // void parallel_for_work_group(range<dimensions> num_work_groups,
-  // range<dimensions> work_group_size, kernel SyclKernel, WorkgroupFunctionType
-  // KernelFunc);
+  /// Two-kernel version of the \c parallel_for_work_group with group and local
+  /// range.
+  template <typename KernelName = csd::auto_name, typename KernelType, int Dims>
+  void parallel_for_work_group(kernel SyclKernel, range<Dims> NumWorkGroups,
+                               range<Dims> WorkGroupSize,
+                               KernelType KernelFunc) {
+    using NameT = typename csd::get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
+#else
+    MNDRDesc.set(nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize));
+    MSyclKernel = detail::getSyclObjImpl(std::move(SyclKernel));
+    StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
+    MCGType = detail::CG::KERNEL;
+#endif // __SYCL_DEVICE_ONLY__
+  }
 
   // Explicit copy operations API
 
diff --git a/sycl/include/CL/sycl/id.hpp b/sycl/include/CL/sycl/id.hpp
index 1942f064df020..299eb2c3f6585 100644
--- a/sycl/include/CL/sycl/id.hpp
+++ b/sycl/include/CL/sycl/id.hpp
@@ -9,63 +9,61 @@
 #pragma once
 
 #include <CL/sycl/detail/array.hpp>
-#include <CL/sycl/item.hpp>
+#include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/range.hpp>
 
 namespace cl {
 namespace sycl {
 template <int dimensions> class range;
+template <int dimensions, bool with_offset> class item;
 template <int dimensions = 1> struct id : public detail::array<dimensions> {
 private:
   using base = detail::array<dimensions>;
   static_assert(dimensions >= 1 && dimensions <= 3,
                 "id can only be 1, 2, or 3 dimentional.");
+  template <int N, int val, typename T>
+  using ParamTy = detail::enable_if_t<(N == val), T>;
+
 public:
   id() = default;
 
   /* The following constructor is only available in the id struct
    * specialization where: dimensions==1 */
-  template <int N = dimensions>
-  id(typename std::enable_if<(N == 1), size_t>::type dim0) : base(dim0) {}
+  template <int N = dimensions> id(ParamTy<N, 1, size_t> dim0) : base(dim0) {}
 
   template <int N = dimensions>
-  id(typename std::enable_if<(N == 1), const range<dimensions> &>::type
-         range_size)
+  id(ParamTy<N, 1, const range<dimensions>> &range_size)
       : base(range_size.get(0)) {}
 
-  template <int N = dimensions>
-  id(typename std::enable_if<(N == 1), const item<dimensions> &>::type item)
+  template <int N = dimensions, bool with_offset = true>
+  id(ParamTy<N, 1, const item<dimensions, with_offset>> &item)
       : base(item.get_id(0)) {}
 
   /* The following constructor is only available in the id struct
    * specialization where: dimensions==2 */
   template <int N = dimensions>
-  id(typename std::enable_if<(N == 2), size_t>::type dim0, size_t dim1)
-      : base(dim0, dim1) {}
+  id(ParamTy<N, 2, size_t> dim0, size_t dim1) : base(dim0, dim1) {}
 
   template <int N = dimensions>
-  id(typename std::enable_if<(N == 2), const range<dimensions> &>::type
-         range_size)
+  id(ParamTy<N, 2, const range<dimensions>> &range_size)
       : base(range_size.get(0), range_size.get(1)) {}
 
-  template <int N = dimensions>
-  id(typename std::enable_if<(N == 2), const item<dimensions> &>::type item)
+  template <int N = dimensions, bool with_offset = true>
+  id(ParamTy<N, 2, const item<dimensions, with_offset>> &item)
       : base(item.get_id(0), item.get_id(1)) {}
 
   /* The following constructor is only available in the id struct
    * specialization where: dimensions==3 */
   template <int N = dimensions>
-  id(typename std::enable_if<(N == 3), size_t>::type dim0, size_t dim1,
-     size_t dim2)
+  id(ParamTy<N, 3, size_t> dim0, size_t dim1, size_t dim2)
       : base(dim0, dim1, dim2) {}
 
   template <int N = dimensions>
-  id(typename std::enable_if<(N == 3), const range<dimensions> &>::type
-         range_size)
+  id(ParamTy<N, 3, const range<dimensions>> &range_size)
       : base(range_size.get(0), range_size.get(1), range_size.get(2)) {}
 
-  template <int N = dimensions>
-  id(typename std::enable_if<(N == 3), const item<dimensions> &>::type item)
+  template <int N = dimensions, bool with_offset = true>
+  id(ParamTy<N, 3, const item<dimensions, with_offset>> &item)
       : base(item.get_id(0), item.get_id(1), item.get_id(2)) {}
 
   explicit operator range<dimensions>() const {
diff --git a/sycl/include/CL/sycl/item.hpp b/sycl/include/CL/sycl/item.hpp
index 6f6c4e210299b..d9b60258471b8 100644
--- a/sycl/include/CL/sycl/item.hpp
+++ b/sycl/include/CL/sycl/item.hpp
@@ -20,6 +20,8 @@ struct Builder;
 }
 template <int dimensions> struct id;
 template <int dimensions> class range;
+template <int dimensions> class h_item;
+
 template <int dimensions = 1, bool with_offset = true> struct item {
 
   item() = delete;
@@ -86,6 +88,7 @@ template <int dimensions = 1, bool with_offset = true> struct item {
   // For call constructor inside conversion operator
   friend struct item<dimensions, false>;
   friend struct item<dimensions, true>;
+  friend struct h_item<dimensions>;
   friend struct detail::Builder;
 
   template <size_t W = with_offset>
@@ -98,6 +101,8 @@ template <int dimensions = 1, bool with_offset = true> struct item {
        const id<dimensions> &I)
       : extent(R), index(I), offset() {}
 
+  void setID(const id<dimensions> &ID) { index = ID; }
+
 private:
   range<dimensions> extent;
   id<dimensions> index;
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 312b1af7ca059..eadb06567b8d7 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -10,6 +10,7 @@
 #include <CL/cl.h>
 #include <CL/sycl/detail/event_impl.hpp>
 #include <CL/sycl/detail/kernel_desc.hpp>
+#include <CL/sycl/detail/kernel_info.hpp>
 #include <CL/sycl/detail/memory_manager.hpp>
 #include <CL/sycl/detail/program_manager/program_manager.hpp>
 #include <CL/sycl/detail/queue_impl.hpp>
@@ -254,6 +255,44 @@ cl_int MemCpyCommandHost::enqueueImp() {
   return CL_SUCCESS;
 }
 
+// SYCL has a parallel_for_work_group variant where the only NDRange
+// characteristics set by a user is the number of work groups. This does not map
+// to the OpenCL clEnqueueNDRangeAPI, which requires global work size to be set
+// as well. This function determines local work size based on the device
+// characteristics and the number of work groups requested by the user, then
+// calculates the global work size.
+// SYCL specification (from 4.8.5.3):
+// The member function handler::parallel_for_work_group is parameterized by the
+// number of work - groups, such that the size of each group is chosen by the
+// runtime, or by the number of work - groups and number of work - items for
+// users who need more control.
+static void adjustNDRangePerKernel(NDRDescT &NDR, cnri_kernel Kernel,
+                                   cnri_device Device) {
+  if (NDR.GlobalSize[0] != 0)
+    return; // GlobalSize is set - no need to adjust
+  // check the prerequisites:
+  assert(NDR.NumWorkGroups[0] != 0 && NDR.LocalSize[0] == 0);
+  // TODO might be good to cache this info together with the kernel info to
+  // avoid get_kernel_work_group_info_cl on every kernel run
+  range<3> WGSize = get_kernel_work_group_info_cl<
+      range<3>,
+      cl::sycl::info::kernel_work_group::compile_work_group_size>::_(Kernel,
+                                                                     Device);
+
+  if (WGSize[0] == 0) {
+    // kernel does not request specific workgroup shape - set one
+    // TODO maximum work group size as the local size might not be the best
+    //      choice for CPU or FPGA devices
+    size_t WGSize1D = get_kernel_work_group_info_cl<
+        size_t, cl::sycl::info::kernel_work_group::work_group_size>::_(Kernel,
+                                                                       Device);
+    assert(WGSize1D != 0);
+    // TODO implement better default for 2D/3D case:
+    WGSize = {WGSize1D, 1, 1};
+  }
+  NDR.set(NDR.Dims, nd_range<3>(NDR.NumWorkGroups * WGSize, WGSize));
+}
+
 cl_int ExecCGCommand::enqueueImp() {
   std::vector<cl_event> RawEvents =
       Command::prepareEvents(detail::getSyclObjImpl(MQueue->get_context()));
@@ -383,7 +422,7 @@ cl_int ExecCGCommand::enqueueImp() {
         assert(!"Unhandled");
       }
     }
-
+    adjustNDRangePerKernel(NDRDesc, Kernel, MQueue->get_device().get());
     cl_int Error = CL_SUCCESS;
     Error = clEnqueueNDRangeKernel(
         MQueue->getHandleRef(), Kernel, NDRDesc.Dims, &NDRDesc.GlobalOffset[0],
diff --git a/sycl/test/hier_par/hier_par_basic.cpp b/sycl/test/hier_par/hier_par_basic.cpp
new file mode 100644
index 0000000000000..49657fc2cddd3
--- /dev/null
+++ b/sycl/test/hier_par/hier_par_basic.cpp
@@ -0,0 +1,183 @@
+//==- hier_par_basic.cpp --- hierarchical parallelism API test -------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lOpenCL -lstdc++
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks hierarchical parallelism invocation APIs, but without any
+// data or code with side-effects between the work group and work item scopes.
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <memory>
+
+using namespace cl::sycl;
+
+template <typename GoldFnTy>
+bool verify(int testcase, int range_length, int *ptr, GoldFnTy get_gold) {
+  int err_cnt = 0;
+
+  for (int i = 0; i < range_length; i++) {
+    int gold = get_gold(i);
+
+    if (ptr[i] != gold) {
+      if (++err_cnt < 20) {
+        std::cout << testcase << " - ERROR at " << i << ": " << ptr[i]
+                  << " != " << gold << "(expected)\n";
+      }
+    }
+  }
+  if (err_cnt > 0)
+    std::cout << "-- Failure rate: " << err_cnt << "/" << range_length << "("
+              << err_cnt / (float)range_length * 100.f << "%)\n";
+  return err_cnt == 0;
+}
+
+int main() {
+  constexpr int N_WG = 7;
+  constexpr int WG_SIZE_MAIN = 3;
+  constexpr int WG_SIZE1 = 5;
+  constexpr int WG_SIZE2 = 2;
+  constexpr int N_ITER = 2;
+
+  constexpr size_t range_length = N_WG * WG_SIZE_MAIN;
+  std::unique_ptr<int> data(new int[range_length]);
+  int *ptr = data.get();
+  bool passed = true;
+
+  try {
+    queue myQueue;
+
+    std::cout << "Running on "
+              << myQueue.get_device().get_info<cl::sycl::info::device::name>()
+              << "\n";
+    {
+      // Testcase1
+      // - handler::parallel_for_work_group w/o local size specification +
+      //   group::parallel_for_work_item w/o flexible range
+      // - h_item::get_global_id
+      // The global size is not known, so kernel code needs to adapt
+      // dynamically, hence the complex loop bound and index computation.
+      std::memset(ptr, 0, range_length * sizeof(ptr[0]));
+      buffer<int, 1> buf(ptr, range<1>(range_length));
+      myQueue.submit([&](handler &cgh) {
+        auto dev_ptr = buf.get_access<access::mode::read_write>(cgh);
+        // numer of 'buf' elements per work group:
+        size_t wg_chunk = (range_length + N_WG - 1) / N_WG;
+
+        cgh.parallel_for_work_group<class hpar_simple>(
+            range<1>(N_WG), [=](group<1> g) {
+              size_t wg_offset = wg_chunk * g.get_id(0);
+              size_t wg_size = g.get_local_range(0);
+
+              for (int cnt = 0; cnt < N_ITER; cnt++) {
+                g.parallel_for_work_item([&](h_item<1> i) {
+                  // numer of buf elements per work item:
+                  size_t wi_chunk = (wg_chunk + wg_size - 1) / wg_size;
+                  auto id = i.get_physical_local_id().get(0);
+                  if (id >= wg_chunk)
+                    return;
+                  size_t wi_offset = wg_offset + id * wi_chunk;
+                  size_t ub = cl::sycl::min(wi_offset + wi_chunk, range_length);
+
+                  for (size_t ind = wi_offset; ind < ub; ind++)
+                    dev_ptr[ind]++;
+                });
+              }
+            });
+      });
+      auto ptr1 = buf.get_access<access::mode::read>().get_pointer();
+      passed &=
+          verify(1, range_length, ptr1, [&](int i) -> int { return N_ITER; });
+    }
+
+    {
+      // Testcase2
+      // - handler::parallel_for_work_group with local size specification +
+      //   group::parallel_for_work_item with flexible range
+      // - h_item::get_global_id
+      std::memset(ptr, 0, range_length * sizeof(ptr[0]));
+      buffer<int, 1> buf(ptr, range<1>(range_length));
+      myQueue.submit([&](handler &cgh) {
+        auto dev_ptr = buf.get_access<access::mode::read_write>(cgh);
+
+        cgh.parallel_for_work_group<class hpar_flex>(
+            range<1>(N_WG), range<1>(WG_SIZE_MAIN), [=](group<1> g) {
+              for (int cnt = 0; cnt < N_ITER; cnt++) {
+                g.parallel_for_work_item(range<1>(WG_SIZE1), [&](h_item<1> i) {
+                  dev_ptr[i.get_global_id(0)]++;
+                });
+                g.parallel_for_work_item(range<1>(WG_SIZE2), [&](h_item<1> i) {
+                  dev_ptr[i.get_global_id(0)]++;
+                });
+              }
+            });
+      });
+      auto ptr1 = buf.get_access<access::mode::read>().get_pointer();
+      passed &= verify(2, range_length, ptr1, [&](int i) -> int {
+        // consider increments by the first PFWI:
+        int gold = (WG_SIZE1 - 1) / WG_SIZE_MAIN;
+        if (i % WG_SIZE_MAIN < WG_SIZE1 % WG_SIZE_MAIN) {
+          gold++;
+        }
+        // consider increments by the second PFWI:
+        if (i % WG_SIZE_MAIN < WG_SIZE2) {
+          gold++;
+        }
+        gold *= N_ITER;
+        return gold;
+      });
+    }
+
+    {
+      // Testcase3
+      // - handler::parallel_for_work_group with local size specification +
+      //   group::parallel_for_work_item with flexible range
+      // - h_item::get_logical_local_id,get_physical_local_id
+      std::memset(ptr, 0, range_length * sizeof(ptr[0]));
+      buffer<int, 1> buf(ptr, range<1>(range_length));
+      myQueue.submit([&](handler &cgh) {
+        auto dev_ptr = buf.get_access<access::mode::read_write>(cgh);
+
+        cgh.parallel_for_work_group<class hpar_hitem>(
+            range<1>(N_WG), range<1>(WG_SIZE_MAIN), [=](group<1> g) {
+              for (int cnt = 0; cnt < N_ITER; cnt++) {
+                g.parallel_for_work_item(range<1>(WG_SIZE1), [&](h_item<1> i) {
+                  int n = i.get_logical_local_id() == i.get_physical_local_id()
+                              ? 0
+                              : 1;
+                  dev_ptr[i.get_global_id(0)] += n;
+                });
+              }
+            });
+      });
+      auto ptr1 = buf.get_access<access::mode::read>().get_pointer();
+      passed &= verify(3, range_length, ptr1, [&](int i) -> int {
+        int gold = 0;
+        if (i % WG_SIZE_MAIN < WG_SIZE1 % WG_SIZE_MAIN) {
+          gold++;
+        }
+        gold *= N_ITER;
+        return gold;
+      });
+    }
+  } catch (cl::sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+    return 2;
+  }
+
+  if (passed) {
+    std::cout << "Passed\n";
+    return 0;
+  }
+  std::cout << "FAILED\n";
+  return 1;
+}