diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp
index 843db64dea661..c66f9b94dbdc1 100644
--- a/sycl/include/sycl/detail/cg_types.hpp
+++ b/sycl/include/sycl/detail/cg_types.hpp
@@ -69,6 +69,101 @@ enum class CGType : unsigned int {
   AsyncFree = 29,
 };
 
+// The structure represents NDRange - global, local sizes, global offset and
+// number of dimensions.
+
+// TODO: A lot of tests rely on particular values to be set for dimensions that
+// are not used. To clarify, for example, if a 2D kernel is invoked, in
+// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
+// depending on which constructor is used for no clear reason.
+// Instead, only sensible defaults should be used and tests should be updated
+// to reflect this.
+class NDRDescT {
+
+public:
+  NDRDescT() = default;
+  NDRDescT(const NDRDescT &Desc) = default;
+  NDRDescT(NDRDescT &&Desc) = default;
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
+    if (SetNumWorkGroups) {
+      for (size_t I = 0; I < Dims_; ++I) {
+        NumWorkGroups[I] = N[I];
+      }
+    } else {
+      for (size_t I = 0; I < Dims_; ++I) {
+        GlobalSize[I] = N[I];
+      }
+
+      for (int I = Dims_; I < 3; ++I) {
+        GlobalSize[I] = 1;
+      }
+    }
+  }
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
+           sycl::id<Dims_> Offset)
+      : Dims{size_t(Dims_)} {
+    for (size_t I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = NumWorkItems[I];
+      LocalSize[I] = LocalSizes[I];
+      GlobalOffset[I] = Offset[I];
+    }
+
+    for (int I = Dims_; I < 3; ++I) {
+      LocalSize[I] = LocalSizes[0] ? 1 : 0;
+    }
+
+    for (int I = Dims_; I < 3; ++I) {
+      GlobalSize[I] = 1;
+    }
+  }
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
+      : Dims{size_t(Dims_)} {
+    for (size_t I = 0; I < Dims_; ++I) {
+      GlobalSize[I] = NumWorkItems[I];
+      GlobalOffset[I] = Offset[I];
+    }
+  }
+
+  template <int Dims_>
+  NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
+      : NDRDescT(ExecutionRange.get_global_range(),
+                 ExecutionRange.get_local_range(),
+                 ExecutionRange.get_offset()) {}
+
+  template <int Dims_>
+  NDRDescT(sycl::range<Dims_> Range)
+      : NDRDescT(Range, /*SetNumWorkGroups=*/false) {}
+
+  template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
+    if (this->Dims != size_t(Dims_)) {
+      throw std::runtime_error(
+          "Dimensionality of cluster, global and local ranges must be same");
+    }
+
+    for (int I = 0; I < Dims_; ++I)
+      ClusterDimensions[I] = N[I];
+  }
+
+  NDRDescT &operator=(const NDRDescT &Desc) = default;
+  NDRDescT &operator=(NDRDescT &&Desc) = default;
+
+  std::array<size_t, 3> GlobalSize{0, 0, 0};
+  std::array<size_t, 3> LocalSize{0, 0, 0};
+  std::array<size_t, 3> GlobalOffset{0, 0, 0};
+  /// 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
+  std::array<size_t, 3> NumWorkGroups{0, 0, 0};
+  std::array<size_t, 3> ClusterDimensions{1, 1, 1};
+  size_t Dims = 0;
+};
+
 template <typename, typename T> struct check_fn_signature {
   static_assert(std::integral_constant<T, false>::value,
                 "Second template parameter is required to be of function type");
diff --git a/sycl/source/detail/cg.hpp b/sycl/source/detail/cg.hpp
index 2eb44926f7381..c7fcf235ed0c8 100644
--- a/sycl/source/detail/cg.hpp
+++ b/sycl/source/detail/cg.hpp
@@ -60,101 +60,6 @@ class ArgDesc {
   int MIndex;
 };
 
-// The structure represents NDRange - global, local sizes, global offset and
-// number of dimensions.
-
-// TODO: A lot of tests rely on particular values to be set for dimensions that
-// are not used. To clarify, for example, if a 2D kernel is invoked, in
-// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
-// depending on which constructor is used for no clear reason.
-// Instead, only sensible defaults should be used and tests should be updated
-// to reflect this.
-class NDRDescT {
-
-public:
-  NDRDescT() = default;
-  NDRDescT(const NDRDescT &Desc) = default;
-  NDRDescT(NDRDescT &&Desc) = default;
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
-    if (SetNumWorkGroups) {
-      for (size_t I = 0; I < Dims_; ++I) {
-        NumWorkGroups[I] = N[I];
-      }
-    } else {
-      for (size_t I = 0; I < Dims_; ++I) {
-        GlobalSize[I] = N[I];
-      }
-
-      for (int I = Dims_; I < 3; ++I) {
-        GlobalSize[I] = 1;
-      }
-    }
-  }
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
-           sycl::id<Dims_> Offset)
-      : Dims{size_t(Dims_)} {
-    for (size_t I = 0; I < Dims_; ++I) {
-      GlobalSize[I] = NumWorkItems[I];
-      LocalSize[I] = LocalSizes[I];
-      GlobalOffset[I] = Offset[I];
-    }
-
-    for (int I = Dims_; I < 3; ++I) {
-      LocalSize[I] = LocalSizes[0] ? 1 : 0;
-    }
-
-    for (int I = Dims_; I < 3; ++I) {
-      GlobalSize[I] = 1;
-    }
-  }
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
-      : Dims{size_t(Dims_)} {
-    for (size_t I = 0; I < Dims_; ++I) {
-      GlobalSize[I] = NumWorkItems[I];
-      GlobalOffset[I] = Offset[I];
-    }
-  }
-
-  template <int Dims_>
-  NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
-      : NDRDescT(ExecutionRange.get_global_range(),
-                 ExecutionRange.get_local_range(),
-                 ExecutionRange.get_offset()) {}
-
-  template <int Dims_>
-  NDRDescT(sycl::range<Dims_> Range)
-      : NDRDescT(Range, /*SetNumWorkGroups=*/false) {}
-
-  template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
-    if (this->Dims != size_t(Dims_)) {
-      throw std::runtime_error(
-          "Dimensionality of cluster, global and local ranges must be same");
-    }
-
-    for (int I = 0; I < Dims_; ++I)
-      ClusterDimensions[I] = N[I];
-  }
-
-  NDRDescT &operator=(const NDRDescT &Desc) = default;
-  NDRDescT &operator=(NDRDescT &&Desc) = default;
-
-  std::array<size_t, 3> GlobalSize{0, 0, 0};
-  std::array<size_t, 3> LocalSize{0, 0, 0};
-  std::array<size_t, 3> GlobalOffset{0, 0, 0};
-  /// 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
-  std::array<size_t, 3> NumWorkGroups{0, 0, 0};
-  std::array<size_t, 3> ClusterDimensions{1, 1, 1};
-  size_t Dims = 0;
-};
-
 /// Base class for all types of command groups.
 class CG {
 public: