[SYCL][NFCI][ABI-Break] Move handler members to impl

sycl/include/sycl/detail/cg_types.hpp

@@ -37,117 +37,34 @@ class handler;
 namespace detail {
 class HostTask;
 
-// The structure represents kernel argument.
-class ArgDesc {
-public:
-  ArgDesc(sycl::detail::kernel_param_kind_t Type, void *Ptr, int Size,
-          int Index)
-      : MType(Type), MPtr(Ptr), MSize(Size), MIndex(Index) {}
-
-  sycl::detail::kernel_param_kind_t MType;
-  void *MPtr;
-  int MSize;
-  int MIndex;
-};
-
-// The structure represents NDRange - global, local sizes, global offset,
-// number of dimensions, and the cluster dimensions if applicable.
-class NDRDescT {
-  // The method initializes all sizes for dimensions greater than the passed one
-  // to the default values, so they will not affect execution.
-  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;
-    }
-  }
-
-public:
-  NDRDescT()
-      : GlobalSize{0, 0, 0}, LocalSize{0, 0, 0}, NumWorkGroups{0, 0, 0},
-        Dims{0} {}
-
-  template <int Dims_> void set(sycl::range<Dims_> NumWorkItems) {
-    for (int I = 0; I < Dims_; ++I) {
-      GlobalSize[I] = NumWorkItems[I];
-      LocalSize[I] = 0;
-      GlobalOffset[I] = 0;
-      NumWorkGroups[I] = 0;
-    }
-    setNDRangeLeftover(Dims_);
-    Dims = Dims_;
-  }
-
-  // Initializes this ND range descriptor with given range of work items and
-  // offset.
-  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_;
-  }
-
-  template <int Dims_> void set(sycl::nd_range<Dims_> 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_;
-  }
-
-  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_;
-  }
-
-  template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
-    if (Dims_ != 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];
-    }
-  }
-
-  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;
-  sycl::range<3> ClusterDimensions{1, 1, 1};
-  size_t Dims;
+/// Type of the command group.
+/// NOTE: Changing the values of any of these enumerators is an API-break.
+enum class CGType : unsigned int {
+  None = 0,
+  Kernel = 1,
+  CopyAccToPtr = 2,
+  CopyPtrToAcc = 3,
+  CopyAccToAcc = 4,
+  Barrier = 5,
+  BarrierWaitlist = 6,
+  Fill = 7,
+  UpdateHost = 8,
+  CopyUSM = 10,
+  FillUSM = 11,
+  PrefetchUSM = 12,
+  CodeplayHostTask = 14,
+  AdviseUSM = 15,
+  Copy2DUSM = 16,
+  Fill2DUSM = 17,
+  Memset2DUSM = 18,
+  CopyToDeviceGlobal = 19,
+  CopyFromDeviceGlobal = 20,
+  ReadWriteHostPipe = 21,
+  ExecCommandBuffer = 22,
+  CopyImage = 23,
+  SemaphoreWait = 24,
+  SemaphoreSignal = 25,
+  ProfilingTag = 26,
 };
 
 template <typename, typename T> struct check_fn_signature {
@@ -236,8 +153,6 @@ runKernelWithArg(KernelType KernelName, ArgType Arg) {
 // The pure virtual class aimed to store lambda/functors of any type.
 class HostKernelBase {
 public:
-  // The method executes lambda stored using NDRange passed.
-  virtual void call(const NDRDescT &NDRDesc, HostProfilingInfo *HPI) = 0;
   // Return pointer to the lambda object.
   // Used to extract captured variables.
   virtual char *getPtr() = 0;
@@ -255,177 +170,9 @@ class HostKernel : public HostKernelBase {
 
 public:
   HostKernel(KernelType Kernel) : MKernel(Kernel) {}
-  void call(const NDRDescT &NDRDesc, HostProfilingInfo *HPI) override {
-    // adjust ND range for serial host:
-    NDRDescT AdjustedRange = NDRDesc;
-
-    if (NDRDesc.GlobalSize[0] == 0 && NDRDesc.NumWorkGroups[0] != 0) {
-      // This is a special case - NDRange information is not complete, only the
-      // desired number of work groups is set by the user. Choose work group
-      // size (LocalSize), calculate the missing NDRange characteristics
-      // needed to invoke the kernel and adjust the NDRange descriptor
-      // accordingly. For some devices the work group size selection requires
-      // access to the device's properties, hence such late "adjustment".
-      range<3> WGsize{1, 1, 1}; // no better alternative for serial host?
-      AdjustedRange.set(NDRDesc.Dims,
-                        nd_range<3>(NDRDesc.NumWorkGroups * WGsize, WGsize));
-    }
-    // If local size for host is not set explicitly, let's adjust it to 1,
-    // so an exception for zero local size is not thrown.
-    if (AdjustedRange.LocalSize[0] == 0)
-      for (size_t I = 0; I < AdjustedRange.Dims; ++I)
-        AdjustedRange.LocalSize[I] = 1;
-    if (HPI)
-      HPI->start();
-    runOnHost(AdjustedRange);
-    if (HPI)
-      HPI->end();
-  }
 
   char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
 
-  template <class ArgT = KernelArgType>
-  typename std::enable_if_t<std::is_same_v<ArgT, void>>
-  runOnHost(const NDRDescT &) {
-    runKernelWithoutArg(MKernel);
-  }
-
-  template <class ArgT = KernelArgType>
-  typename std::enable_if_t<std::is_same_v<ArgT, sycl::id<Dims>>>
-  runOnHost(const NDRDescT &NDRDesc) {
-    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(
-        /*LowerBound=*/Offset, Stride, UpperBound,
-        [&](const sycl::id<Dims> &ID) {
-          sycl::item<Dims, /*Offset=*/true> Item =
-              IDBuilder::createItem<Dims, true>(Range, ID, Offset);
-
-          runKernelWithArg<const sycl::id<Dims> &>(MKernel, ID);
-        });
-  }
-
-  template <class ArgT = KernelArgType>
-  typename std::enable_if_t<std::is_same_v<ArgT, item<Dims, /*Offset=*/false>>>
-  runOnHost(const NDRDescT &NDRDesc) {
-    sycl::id<Dims> ID;
-    sycl::range<Dims> Range(InitializedVal<Dims, range>::template get<0>());
-    for (int I = 0; I < Dims; ++I)
-      Range[I] = NDRDesc.GlobalSize[I];
-
-    detail::NDLoop<Dims>::iterate(Range, [&](const sycl::id<Dims> ID) {
-      sycl::item<Dims, /*Offset=*/false> Item =
-          IDBuilder::createItem<Dims, false>(Range, ID);
-      sycl::item<Dims, /*Offset=*/true> ItemWithOffset = Item;
-
-      runKernelWithArg<sycl::item<Dims, /*Offset=*/false>>(MKernel, Item);
-    });
-  }
-
-  template <class ArgT = KernelArgType>
-  typename std::enable_if_t<std::is_same_v<ArgT, item<Dims, /*Offset=*/true>>>
-  runOnHost(const NDRDescT &NDRDesc) {
-    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(
-        /*LowerBound=*/Offset, Stride, UpperBound,
-        [&](const sycl::id<Dims> &ID) {
-          sycl::item<Dims, /*Offset=*/true> Item =
-              IDBuilder::createItem<Dims, true>(Range, ID, Offset);
-
-          runKernelWithArg<sycl::item<Dims, /*Offset=*/true>>(MKernel, Item);
-        });
-  }
-
-  template <class ArgT = KernelArgType>
-  typename std::enable_if_t<std::is_same_v<ArgT, nd_item<Dims>>>
-  runOnHost(const NDRDescT &NDRDesc) {
-    sycl::range<Dims> GroupSize(InitializedVal<Dims, range>::template get<0>());
-    for (int I = 0; I < Dims; ++I) {
-      if (NDRDesc.LocalSize[I] == 0 ||
-          NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0)
-        throw sycl::exception(make_error_code(errc::nd_range),
-                              "Invalid local size for global size");
-      GroupSize[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
-    }
-
-    sycl::range<Dims> LocalSize(InitializedVal<Dims, range>::template get<0>());
-    sycl::range<Dims> GlobalSize(
-        InitializedVal<Dims, range>::template get<0>());
-    sycl::id<Dims> GlobalOffset;
-    for (int I = 0; I < Dims; ++I) {
-      GlobalOffset[I] = NDRDesc.GlobalOffset[I];
-      LocalSize[I] = NDRDesc.LocalSize[I];
-      GlobalSize[I] = NDRDesc.GlobalSize[I];
-    }
-
-    detail::NDLoop<Dims>::iterate(GroupSize, [&](const id<Dims> &GroupID) {
-      sycl::group<Dims> Group = IDBuilder::createGroup<Dims>(
-          GlobalSize, LocalSize, GroupSize, GroupID);
-
-      detail::NDLoop<Dims>::iterate(LocalSize, [&](const id<Dims> &LocalID) {
-        id<Dims> GlobalID =
-            GroupID * id<Dims>{LocalSize} + LocalID + GlobalOffset;
-        const sycl::item<Dims, /*Offset=*/true> GlobalItem =
-            IDBuilder::createItem<Dims, true>(GlobalSize, GlobalID,
-                                              GlobalOffset);
-        const sycl::item<Dims, /*Offset=*/false> LocalItem =
-            IDBuilder::createItem<Dims, false>(LocalSize, LocalID);
-        const sycl::nd_item<Dims> NDItem =
-            IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
-
-        runKernelWithArg<const sycl::nd_item<Dims>>(MKernel, NDItem);
-      });
-    });
-  }
-
-  template <typename ArgT = KernelArgType>
-  std::enable_if_t<std::is_same_v<ArgT, sycl::group<Dims>>>
-  runOnHost(const NDRDescT &NDRDesc) {
-    sycl::range<Dims> NGroups(InitializedVal<Dims, range>::template get<0>());
-
-    for (int I = 0; I < Dims; ++I) {
-      if (NDRDesc.LocalSize[I] == 0 ||
-          NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0)
-        throw sycl::exception(make_error_code(errc::nd_range),
-                              "Invalid local size for global size");
-      NGroups[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
-    }
-
-    sycl::range<Dims> LocalSize(InitializedVal<Dims, range>::template get<0>());
-    sycl::range<Dims> GlobalSize(
-        InitializedVal<Dims, range>::template get<0>());
-    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, NGroups, GroupID);
-      runKernelWithArg<sycl::group<Dims>>(MKernel, Group);
-    });
-  }
-
   ~HostKernel() = default;
 };
 

sycl/include/sycl/ext/oneapi/memcpy2d.hpp

@@ -7,7 +7,6 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
-#include <sycl/detail/host_task_impl.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/queue.hpp>
 #include <sycl/usm/usm_enums.hpp>

sycl/include/sycl/group.hpp

@@ -673,13 +673,7 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
   friend class detail::Builder;
   group(const range<Dimensions> &G, const range<Dimensions> &L,
         const range<Dimensions> GroupRange, const id<Dimensions> &I)
-      : globalRange(G), localRange(L), groupRange(GroupRange), index(I) {
-    // Make sure local range divides global without remainder:
-    __SYCL_ASSERT(((G % L).size() == 0) &&
-                  "global range is not multiple of local");
-    __SYCL_ASSERT((((G / L) - GroupRange).size() == 0) &&
-                  "inconsistent group constructor arguments");
-  }
+      : globalRange(G), localRange(L), groupRange(GroupRange), index(I) {}
 };
 } // namespace _V1
 } // namespace sycl