From 6add8fab02cc719c9e44e8169fb20d6bd86b03fe Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Wed, 9 Jul 2025 13:00:10 +0000 Subject: [PATCH] [SYCL] Move the NDRDescT class to a public header --- sycl/include/sycl/queue.hpp | 95 ++++++++++++++++++ sycl/source/detail/cg.hpp | 99 +------------------ .../detail/error_handling/error_handling.cpp | 10 +- .../detail/error_handling/error_handling.hpp | 2 +- sycl/source/detail/graph/graph_impl.cpp | 4 +- sycl/source/detail/graph/graph_impl.hpp | 3 +- sycl/source/detail/graph/node_impl.hpp | 4 +- sycl/source/detail/handler_impl.hpp | 2 +- .../program_manager/program_manager.cpp | 6 +- .../program_manager/program_manager.hpp | 6 +- sycl/source/detail/scheduler/commands.cpp | 14 +-- sycl/source/detail/scheduler/commands.hpp | 6 +- sycl/source/handler.cpp | 36 +++---- .../scheduler/SchedulerTestUtils.hpp | 2 +- 14 files changed, 144 insertions(+), 145 deletions(-) diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index a0dbdf5c540e8..5e01e725e1215 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -149,6 +149,101 @@ class __SYCL_EXPORT SubmissionInfo { ext::oneapi::experimental::event_mode_enum::none; }; +// 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 + NDRDescT(sycl::range 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 + NDRDescT(sycl::range NumWorkItems, sycl::range LocalSizes, + sycl::id 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 + NDRDescT(sycl::range NumWorkItems, sycl::id Offset) + : Dims{size_t(Dims_)} { + for (size_t I = 0; I < Dims_; ++I) { + GlobalSize[I] = NumWorkItems[I]; + GlobalOffset[I] = Offset[I]; + } + } + + template + NDRDescT(sycl::nd_range ExecutionRange) + : NDRDescT(ExecutionRange.get_global_range(), + ExecutionRange.get_local_range(), + ExecutionRange.get_offset()) {} + + template + NDRDescT(sycl::range Range) + : NDRDescT(Range, /*SetNumWorkGroups=*/false) {} + + template void setClusterDimensions(sycl::range 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 GlobalSize{0, 0, 0}; + std::array LocalSize{0, 0, 0}; + std::array 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 NumWorkGroups{0, 0, 0}; + std::array ClusterDimensions{1, 1, 1}; + size_t Dims = 0; +}; + } // namespace v1 } // namespace detail diff --git a/sycl/source/detail/cg.hpp b/sycl/source/detail/cg.hpp index f48f6ace13ddd..307e4df55aaf3 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 - NDRDescT(sycl::range 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 - NDRDescT(sycl::range NumWorkItems, sycl::range LocalSizes, - sycl::id 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 - NDRDescT(sycl::range NumWorkItems, sycl::id Offset) - : Dims{size_t(Dims_)} { - for (size_t I = 0; I < Dims_; ++I) { - GlobalSize[I] = NumWorkItems[I]; - GlobalOffset[I] = Offset[I]; - } - } - - template - NDRDescT(sycl::nd_range ExecutionRange) - : NDRDescT(ExecutionRange.get_global_range(), - ExecutionRange.get_local_range(), - ExecutionRange.get_offset()) {} - - template - NDRDescT(sycl::range Range) - : NDRDescT(Range, /*SetNumWorkGroups=*/false) {} - - template void setClusterDimensions(sycl::range 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 GlobalSize{0, 0, 0}; - std::array LocalSize{0, 0, 0}; - std::array 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 NumWorkGroups{0, 0, 0}; - std::array ClusterDimensions{1, 1, 1}; - size_t Dims = 0; -}; - /// Base class for all types of command groups. class CG { public: @@ -248,7 +153,7 @@ class CG { class CGExecKernel : public CG { public: /// Stores ND-range description. - NDRDescT MNDRDesc; + detail::v1::NDRDescT MNDRDesc; std::shared_ptr MHostKernel; std::shared_ptr MSyclKernel; std::shared_ptr MKernelBundle; @@ -265,7 +170,7 @@ class CGExecKernel : public CG { bool MKernelUsesClusterLaunch = false; size_t MKernelWorkGroupMemorySize = 0; - CGExecKernel(NDRDescT NDRDesc, std::shared_ptr HKernel, + CGExecKernel(detail::v1::NDRDescT NDRDesc, std::shared_ptr HKernel, std::shared_ptr SyclKernel, std::shared_ptr KernelBundle, CG::StorageInitHelper CGData, std::vector Args, diff --git a/sycl/source/detail/error_handling/error_handling.cpp b/sycl/source/detail/error_handling/error_handling.cpp index c33852cd08cd8..7dcef0d8590a2 100644 --- a/sycl/source/detail/error_handling/error_handling.cpp +++ b/sycl/source/detail/error_handling/error_handling.cpp @@ -22,7 +22,7 @@ inline namespace _V1 { namespace detail::enqueue_kernel_launch { void handleOutOfResources(const device_impl &DeviceImpl, - ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) { + ur_kernel_handle_t Kernel, const detail::v1::NDRDescT &NDRDesc) { sycl::platform Platform = DeviceImpl.get_platform(); sycl::backend Backend = Platform.get_backend(); if (Backend == sycl::backend::ext_oneapi_cuda) { @@ -75,7 +75,7 @@ void handleOutOfResources(const device_impl &DeviceImpl, void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, ur_kernel_handle_t Kernel, - const NDRDescT &NDRDesc) { + const detail::v1::NDRDescT &NDRDesc) { sycl::platform Platform = DeviceImpl.get_platform(); // Some of the error handling below is special for particular OpenCL @@ -349,7 +349,7 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, } void handleInvalidWorkItemSize(const device_impl &DeviceImpl, - const NDRDescT &NDRDesc) { + const detail::v1::NDRDescT &NDRDesc) { const AdapterPtr &Adapter = DeviceImpl.getAdapter(); ur_device_handle_t Device = DeviceImpl.getHandleRef(); @@ -370,7 +370,7 @@ void handleInvalidWorkItemSize(const device_impl &DeviceImpl, } void handleInvalidValue(const device_impl &DeviceImpl, - const NDRDescT &NDRDesc) { + const detail::v1::NDRDescT &NDRDesc) { const AdapterPtr &Adapter = DeviceImpl.getAdapter(); ur_device_handle_t Device = DeviceImpl.getHandleRef(); @@ -393,7 +393,7 @@ void handleInvalidValue(const device_impl &DeviceImpl, } void handleErrorOrWarning(ur_result_t Error, const device_impl &DeviceImpl, - ur_kernel_handle_t Kernel, const NDRDescT &NDRDesc) { + ur_kernel_handle_t Kernel, const detail::v1::NDRDescT &NDRDesc) { assert(Error != UR_RESULT_SUCCESS && "Success is expected to be handled on caller side"); switch (Error) { diff --git a/sycl/source/detail/error_handling/error_handling.hpp b/sycl/source/detail/error_handling/error_handling.hpp index 2f0dbb8d783ca..9d00525832e9a 100644 --- a/sycl/source/detail/error_handling/error_handling.hpp +++ b/sycl/source/detail/error_handling/error_handling.hpp @@ -26,7 +26,7 @@ namespace enqueue_kernel_launch { /// This function actually never returns and always throws an exception with /// error description. void handleErrorOrWarning(ur_result_t, const device_impl &, ur_kernel_handle_t, - const NDRDescT &); + const detail::v1::NDRDescT &); } // namespace enqueue_kernel_launch namespace kernel_get_group_info { diff --git a/sycl/source/detail/graph/graph_impl.cpp b/sycl/source/detail/graph/graph_impl.cpp index d5555ef688767..61bdbc1ab3a20 100644 --- a/sycl/source/detail/graph/graph_impl.cpp +++ b/sycl/source/detail/graph/graph_impl.cpp @@ -1573,7 +1573,7 @@ void exec_graph_impl::populateURKernelUpdateStructs( std::vector &MemobjProps, std::vector &PtrDescs, std::vector &ValueDescs, - sycl::detail::NDRDescT &NDRDesc, + sycl::detail::v1::NDRDescT &NDRDesc, ur_exp_command_buffer_update_kernel_launch_desc_t &UpdateDesc) const { sycl::detail::context_impl &ContextImpl = *sycl::detail::getSyclObjImpl(MContext); @@ -1800,7 +1800,7 @@ void exec_graph_impl::updateURImpl( PtrDescsList(NumUpdatableNodes); std::vector> ValueDescsList(NumUpdatableNodes); - std::vector NDRDescList(NumUpdatableNodes); + std::vector NDRDescList(NumUpdatableNodes); std::vector UpdateDescList( NumUpdatableNodes); std::vector KernelBundleObjList(NumUpdatableNodes); diff --git a/sycl/source/detail/graph/graph_impl.hpp b/sycl/source/detail/graph/graph_impl.hpp index 0e257a77e5ef1..d397ce2984cde 100644 --- a/sycl/source/detail/graph/graph_impl.hpp +++ b/sycl/source/detail/graph/graph_impl.hpp @@ -31,7 +31,6 @@ class handler; namespace detail { class SYCLMemObjT; class queue_impl; -class NDRDescT; class ArgDesc; class CG; } // namespace detail @@ -869,7 +868,7 @@ class exec_graph_impl { std::vector &MemobjProps, std::vector &PtrDescs, std::vector &ValueDescs, - sycl::detail::NDRDescT &NDRDesc, + sycl::detail::v1::NDRDescT &NDRDesc, ur_exp_command_buffer_update_kernel_launch_desc_t &UpdateDesc) const; /// Execution schedule of nodes in the graph. diff --git a/sycl/source/detail/graph/node_impl.hpp b/sycl/source/detail/graph/node_impl.hpp index 11166e1eba897..b2733fea14b62 100644 --- a/sycl/source/detail/graph/node_impl.hpp +++ b/sycl/source/detail/graph/node_impl.hpp @@ -437,7 +437,7 @@ class node_impl : public std::enable_shared_from_this { "the node was originally created with."); } - NDRDesc = sycl::detail::NDRDescT{ExecutionRange}; + NDRDesc = sycl::detail::v1::NDRDescT{ExecutionRange}; } template void updateRange(range ExecutionRange) { @@ -458,7 +458,7 @@ class node_impl : public std::enable_shared_from_this { "the node was originally created with."); } - NDRDesc = sycl::detail::NDRDescT{ExecutionRange}; + NDRDesc = sycl::detail::v1::NDRDescT{ExecutionRange}; } /// Update this node with the command-group from another node. /// @param Other The other node to update, must be of the same node type. diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index 0fda3dd4f2769..b743a66a59aaf 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -155,7 +155,7 @@ class handler_impl { std::vector MAssociatedAccesors; /// Struct that encodes global size, local size, ... - detail::NDRDescT MNDRDesc; + detail::v1::NDRDescT MNDRDesc; /// Type of the command group, e.g. kernel, fill. Can also encode version. /// Use getType and setType methods to access this variable unless diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index cfd23e820b37c..a5e30f54cdd57 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -862,7 +862,7 @@ static void setSpecializationConstants(device_image_impl &InputImpl, // its ref count incremented. ur_program_handle_t ProgramManager::getBuiltURProgram( context_impl &ContextImpl, device_impl &DeviceImpl, - KernelNameStrRefT KernelName, const NDRDescT &NDRDesc) { + KernelNameStrRefT KernelName, const detail::v1::NDRDescT &NDRDesc) { device_impl *RootDevImpl; ur_bool_t MustBuildOnSubdevice = true; @@ -1124,7 +1124,7 @@ ur_program_handle_t ProgramManager::getBuiltURProgram( FastKernelCacheValPtr ProgramManager::getOrCreateKernel( context_impl &ContextImpl, device_impl &DeviceImpl, KernelNameStrRefT KernelName, - KernelNameBasedCacheT *KernelNameBasedCachePtr, const NDRDescT &NDRDesc) { + KernelNameBasedCacheT *KernelNameBasedCachePtr, const detail::v1::NDRDescT &NDRDesc) { if constexpr (DbgProgMgr > 0) { std::cerr << ">>> ProgramManager::getOrCreateKernel(" << &ContextImpl << ", " << &DeviceImpl << ", " << KernelName << ")\n"; @@ -3593,7 +3593,7 @@ std::optional checkDevSupportJointMatrixMad( std::optional checkDevSupportDeviceRequirements(const device_impl &Dev, const RTDeviceBinaryImage &Img, - const NDRDescT &NDRDesc) { + const detail::v1::NDRDescT &NDRDesc) { auto getPropIt = [&Img](const std::string &PropName) { auto &PropRange = Img.getDeviceRequirements(); RTDeviceBinaryImage::PropertyRange::ConstIterator PropIt = std::find_if( diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index bd60661949e87..9fba3f8faebda 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -62,7 +62,7 @@ bool doesDevSupportDeviceRequirements(const device_impl &Dev, std::optional checkDevSupportDeviceRequirements(const device_impl &Dev, const RTDeviceBinaryImage &BinImages, - const NDRDescT &NDRDesc = {}); + const detail::v1::NDRDescT &NDRDesc = {}); bool doesImageTargetMatchDevice(const RTDeviceBinaryImage &Img, const device_impl &DevImpl); @@ -179,7 +179,7 @@ class ProgramManager { ur_program_handle_t getBuiltURProgram(context_impl &ContextImpl, device_impl &DeviceImpl, KernelNameStrRefT KernelName, - const NDRDescT &NDRDesc = {}); + const detail::v1::NDRDescT &NDRDesc = {}); /// Builds a program from a given set of images or retrieves that program from /// cache. @@ -201,7 +201,7 @@ class ProgramManager { getOrCreateKernel(context_impl &ContextImpl, device_impl &DeviceImpl, KernelNameStrRefT KernelName, KernelNameBasedCacheT *KernelNameBasedCachePtr, - const NDRDescT &NDRDesc = {}); + const detail::v1::NDRDescT &NDRDesc = {}); ur_kernel_handle_t getCachedMaterializedKernel( KernelNameStrRefT KernelName, diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 3b90c13e9e3a1..388b0ef5006c6 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1989,7 +1989,7 @@ std::string instrumentationGetKernelName( } void instrumentationAddExtraKernelMetadata( - xpti_td *&CmdTraceEvent, const NDRDescT &NDRDesc, + xpti_td *&CmdTraceEvent, const detail::v1::NDRDescT &NDRDesc, detail::kernel_bundle_impl *KernelBundleImplPtr, KernelNameStrRefT KernelName, KernelNameBasedCacheT *KernelNameBasedCachePtr, @@ -2107,7 +2107,7 @@ std::pair emitKernelInstrumentationData( const detail::code_location &CodeLoc, bool IsTopCodeLoc, const std::string_view SyclKernelName, KernelNameBasedCacheT *KernelNameBasedCachePtr, queue_impl *Queue, - const NDRDescT &NDRDesc, detail::kernel_bundle_impl *KernelBundleImplPtr, + const detail::v1::NDRDescT &NDRDesc, detail::kernel_bundle_impl *KernelBundleImplPtr, std::vector &CGArgs) { auto XptiObjects = std::make_pair(nullptr, -1); @@ -2261,7 +2261,7 @@ std::string_view ExecCGCommand::getTypeString() const { // 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, ur_kernel_handle_t Kernel, +static void adjustNDRangePerKernel(detail::v1::NDRDescT &NDR, ur_kernel_handle_t Kernel, const device_impl &DeviceImpl) { if (NDR.GlobalSize[0] != 0) return; // GlobalSize is set - no need to adjust @@ -2292,7 +2292,7 @@ static void adjustNDRangePerKernel(NDRDescT &NDR, ur_kernel_handle_t Kernel, // Initially we keep the order of NDRDescT as it provided by the user, this // simplifies overall handling and do the reverse only when // the kernel is enqueued. -void ReverseRangeDimensionsForKernel(NDRDescT &NDR) { +void ReverseRangeDimensionsForKernel(detail::v1::NDRDescT &NDR) { if (NDR.Dims > 1) { std::swap(NDR.GlobalSize[0], NDR.GlobalSize[NDR.Dims - 1]); std::swap(NDR.LocalSize[0], NDR.LocalSize[NDR.Dims - 1]); @@ -2398,7 +2398,7 @@ static void SetArgBasedOnType( static ur_result_t SetKernelParamsAndLaunch( queue_impl &Queue, std::vector &Args, const std::shared_ptr &DeviceImageImpl, - ur_kernel_handle_t Kernel, NDRDescT &NDRDesc, + ur_kernel_handle_t Kernel, detail::v1::NDRDescT &NDRDesc, std::vector &RawEvents, detail::event_impl *OutEventImpl, const KernelArgMask *EliminatedArgMask, const std::function &getMemAllocationFunc, @@ -2672,7 +2672,7 @@ ur_result_t enqueueImpCommandBufferKernel( } void enqueueImpKernel( - queue_impl &Queue, NDRDescT &NDRDesc, std::vector &Args, + queue_impl &Queue, detail::v1::NDRDescT &NDRDesc, std::vector &Args, detail::kernel_bundle_impl *KernelBundleImplPtr, const detail::kernel_impl *MSyclKernel, KernelNameStrRefT KernelName, KernelNameBasedCacheT *KernelNameBasedCachePtr, @@ -3231,7 +3231,7 @@ ur_result_t ExecCGCommand::enqueueImpQueue() { assert(MQueue && "Kernel submissions should have an associated queue"); CGExecKernel *ExecKernel = (CGExecKernel *)MCommandGroup.get(); - NDRDescT &NDRDesc = ExecKernel->MNDRDesc; + detail::v1::NDRDescT &NDRDesc = ExecKernel->MNDRDesc; std::vector &Args = ExecKernel->MArgs; auto getMemAllocationFunc = [this](Requirement *Req) { diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 9a67d6c8fec31..9e51e28b332c5 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/sycl/source/detail/scheduler/commands.hpp @@ -620,7 +620,7 @@ class MemCpyCommandHost : public Command { }; void enqueueImpKernel( - queue_impl &Queue, NDRDescT &NDRDesc, std::vector &Args, + queue_impl &Queue, detail::v1::NDRDescT &NDRDesc, std::vector &Args, detail::kernel_bundle_impl *KernelBundleImplPtr, const detail::kernel_impl *MSyclKernel, KernelNameStrRefT KernelName, KernelNameBasedCacheT *KernelNameBasedCachePtr, @@ -690,7 +690,7 @@ std::pair emitKernelInstrumentationData( const detail::code_location &CodeLoc, bool IsTopCodeLoc, std::string_view SyclKernelName, KernelNameBasedCacheT *KernelNameBasedCachePtr, queue_impl *Queue, - const NDRDescT &NDRDesc, detail::kernel_bundle_impl *KernelBundleImplPtr, + const detail::v1::NDRDescT &NDRDesc, detail::kernel_bundle_impl *KernelBundleImplPtr, std::vector &CGArgs); #endif @@ -796,7 +796,7 @@ void applyFuncOnFilteredArgs( } } -void ReverseRangeDimensionsForKernel(NDRDescT &NDR); +void ReverseRangeDimensionsForKernel(detail::v1::NDRDescT &NDR); } // namespace detail } // namespace _V1 diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index f575885b6a24d..49101598f6402 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -2443,12 +2443,12 @@ void handler::setNDRangeDescriptorPadded(sycl::range<3> N, bool SetNumWorkGroups, int Dims) { if (Dims == 1) { sycl::range<1> Range = {N[0]}; - impl->MNDRDesc = NDRDescT{Range, SetNumWorkGroups}; + impl->MNDRDesc = detail::v1::NDRDescT{Range, SetNumWorkGroups}; } else if (Dims == 2) { sycl::range<2> Range = {N[0], N[1]}; - impl->MNDRDesc = NDRDescT{Range, SetNumWorkGroups}; + impl->MNDRDesc = detail::v1::NDRDescT{Range, SetNumWorkGroups}; } else if (Dims == 3) { - impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups}; + impl->MNDRDesc = detail::v1::NDRDescT{N, SetNumWorkGroups}; } } @@ -2457,13 +2457,13 @@ void handler::setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems, if (Dims == 1) { sycl::range<1> NumWorkItemsTrimmed = {NumWorkItems[0]}; sycl::id<1> OffsetTrimmed = {Offset[0]}; - impl->MNDRDesc = NDRDescT{NumWorkItemsTrimmed, OffsetTrimmed}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItemsTrimmed, OffsetTrimmed}; } else if (Dims == 2) { sycl::range<2> NumWorkItemsTrimmed = {NumWorkItems[0], NumWorkItems[1]}; sycl::id<2> OffsetTrimmed = {Offset[0], Offset[1]}; - impl->MNDRDesc = NDRDescT{NumWorkItemsTrimmed, OffsetTrimmed}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItemsTrimmed, OffsetTrimmed}; } else if (Dims == 3) { - impl->MNDRDesc = NDRDescT{NumWorkItems, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, Offset}; } } @@ -2475,56 +2475,56 @@ void handler::setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems, sycl::range<1> LocalSizeTrimmed = {LocalSize[0]}; sycl::id<1> OffsetTrimmed = {Offset[0]}; impl->MNDRDesc = - NDRDescT{NumWorkItemsTrimmed, LocalSizeTrimmed, OffsetTrimmed}; + detail::v1::NDRDescT{NumWorkItemsTrimmed, LocalSizeTrimmed, OffsetTrimmed}; } else if (Dims == 2) { sycl::range<2> NumWorkItemsTrimmed = {NumWorkItems[0], NumWorkItems[1]}; sycl::range<2> LocalSizeTrimmed = {LocalSize[0], LocalSize[1]}; sycl::id<2> OffsetTrimmed = {Offset[0], Offset[1]}; impl->MNDRDesc = - NDRDescT{NumWorkItemsTrimmed, LocalSizeTrimmed, OffsetTrimmed}; + detail::v1::NDRDescT{NumWorkItemsTrimmed, LocalSizeTrimmed, OffsetTrimmed}; } else if (Dims == 3) { - impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, LocalSize, Offset}; } } #endif void handler::setNDRangeDescriptor(sycl::range<3> N, bool SetNumWorkGroups) { - impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups}; + impl->MNDRDesc = detail::v1::NDRDescT{N, SetNumWorkGroups}; } void handler::setNDRangeDescriptor(sycl::range<3> NumWorkItems, sycl::id<3> Offset) { - impl->MNDRDesc = NDRDescT{NumWorkItems, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, Offset}; } void handler::setNDRangeDescriptor(sycl::range<3> NumWorkItems, sycl::range<3> LocalSize, sycl::id<3> Offset) { - impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, LocalSize, Offset}; } void handler::setNDRangeDescriptor(sycl::range<2> N, bool SetNumWorkGroups) { - impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups}; + impl->MNDRDesc = detail::v1::NDRDescT{N, SetNumWorkGroups}; } void handler::setNDRangeDescriptor(sycl::range<2> NumWorkItems, sycl::id<2> Offset) { - impl->MNDRDesc = NDRDescT{NumWorkItems, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, Offset}; } void handler::setNDRangeDescriptor(sycl::range<2> NumWorkItems, sycl::range<2> LocalSize, sycl::id<2> Offset) { - impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, LocalSize, Offset}; } void handler::setNDRangeDescriptor(sycl::range<1> N, bool SetNumWorkGroups) { - impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups}; + impl->MNDRDesc = detail::v1::NDRDescT{N, SetNumWorkGroups}; } void handler::setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::id<1> Offset) { - impl->MNDRDesc = NDRDescT{NumWorkItems, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, Offset}; } void handler::setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::range<1> LocalSize, sycl::id<1> Offset) { - impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset}; + impl->MNDRDesc = detail::v1::NDRDescT{NumWorkItems, LocalSize, Offset}; } void handler::setKernelNameBasedCachePtr( diff --git a/sycl/unittests/scheduler/SchedulerTestUtils.hpp b/sycl/unittests/scheduler/SchedulerTestUtils.hpp index 61ed99dd87cff..f8fa67dbeaea0 100644 --- a/sycl/unittests/scheduler/SchedulerTestUtils.hpp +++ b/sycl/unittests/scheduler/SchedulerTestUtils.hpp @@ -228,7 +228,7 @@ class MockHandler : public sycl::handler { using sycl::handler::impl; using sycl::handler::setNDRangeDescriptor; - sycl::detail::NDRDescT &getNDRDesc() { return impl->MNDRDesc; } + sycl::detail::v1::NDRDescT &getNDRDesc() { return impl->MNDRDesc; } sycl::detail::code_location &getCodeLoc() { return MCodeLoc; } std::vector> &getStreamStorage() { return MStreamStorage;