Skip to content

Commit 6fc78b6

Browse files
romanovvladromanovvlad
authored
[SYCL][NFC] Refactor CG type versioning (#3579)
Removed dedicated `KERNEL_V1` value, start using setType and getType which hide manipulations with the version.
1 parent 4960e71 commit 6fc78b6

File tree

4 files changed

+53
-37
lines changed

4 files changed

+53
-37
lines changed

‎sycl/include/CL/sycl/detail/cg.hpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -164,8 +164,6 @@ class CG {
164164
PREFETCH_USM = 12,
165165
CODEPLAY_INTEROP_TASK = 13,
166166
CODEPLAY_HOST_TASK = 14,
167-
KERNEL_V1 =
168-
getVersionedCGType(KERNEL, static_cast<unsigned int>(CG_VERSION::V1)),
169167
};
170168

171169
CG(CGTYPE Type, vector_class<vector_class<char>> ArgsStorage,
@@ -192,6 +190,10 @@ class CG {
192190

193191
CGTYPE getType() { return static_cast<CGTYPE>(getUnversionedCGType(MType)); }
194192

193+
CG_VERSION getVersion() {
194+
return static_cast<CG_VERSION>(getCGTypeVersion(MType));
195+
}
196+
195197
std::shared_ptr<std::vector<ExtendedMemberT>> getExtendedMembers() {
196198
if (getCGTypeVersion(MType) == static_cast<unsigned int>(CG_VERSION::V0) ||
197199
MSharedPtrStorage.empty())
@@ -259,8 +261,7 @@ class CGExecKernel : public CG {
259261
MSyclKernel(std::move(SyclKernel)), MArgs(std::move(Args)),
260262
MKernelName(std::move(KernelName)), MOSModuleHandle(OSModuleHandle),
261263
MStreams(std::move(Streams)) {
262-
assert((getType() == RUN_ON_HOST_INTEL || getType() == KERNEL ||
263-
getType() == KERNEL_V1) &&
264+
assert((getType() == RUN_ON_HOST_INTEL || getType() == KERNEL) &&
264265
"Wrong type of exec kernel CG.");
265266
}
266267

‎sycl/include/CL/sycl/handler.hpp

Lines changed: 41 additions & 29 deletions
Original file line numberDiff line numberDiff line change
@@ -325,8 +325,18 @@ class __SYCL_EXPORT handler {
325325
return Storage;
326326
}
327327

328+
void setType(detail::CG::CGTYPE Type) {
329+
constexpr detail::CG::CG_VERSION Version = detail::CG::CG_VERSION::V1;
330+
MCGType = static_cast<detail::CG::CGTYPE>(
331+
getVersionedCGType(Type, static_cast<int>(Version)));
332+
}
333+
334+
detail::CG::CGTYPE getType() {
335+
return static_cast<detail::CG::CGTYPE>(getUnversionedCGType(MCGType));
336+
}
337+
328338
void throwIfActionIsCreated() {
329-
if (detail::CG::NONE != MCGType)
339+
if (detail::CG::NONE != getType())
330340
throw sycl::runtime_error("Attempt to set multiple actions for the "
331341
"command group. Command group must consist of "
332342
"a single kernel or explicit memory operation.",
@@ -822,7 +832,7 @@ class __SYCL_EXPORT handler {
822832
MNDRDesc.set(std::move(AdjustedRange));
823833
StoreLambda<NameWT, decltype(Wrapper), Dims, TransformedArgType>(
824834
std::move(Wrapper));
825-
MCGType = detail::CG::KERNEL_V1;
835+
setType(detail::CG::KERNEL);
826836
#endif
827837
} else
828838
#endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ && \
@@ -837,7 +847,7 @@ class __SYCL_EXPORT handler {
837847
MNDRDesc.set(std::move(NumWorkItems));
838848
StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
839849
std::move(KernelFunc));
840-
MCGType = detail::CG::KERNEL_V1;
850+
setType(detail::CG::KERNEL);
841851
#endif
842852
}
843853
}
@@ -856,7 +866,7 @@ class __SYCL_EXPORT handler {
856866
MKernel = detail::getSyclObjImpl(std::move(Kernel));
857867
detail::checkValueRange<Dims>(NumWorkItems);
858868
MNDRDesc.set(std::move(NumWorkItems));
859-
MCGType = detail::CG::KERNEL_V1;
869+
setType(detail::CG::KERNEL);
860870
extractArgsAndReqs();
861871
MKernelName = getKernelName();
862872
}
@@ -1173,7 +1183,7 @@ class __SYCL_EXPORT handler {
11731183
MNDRDesc.set(range<1>{1});
11741184

11751185
StoreLambda<NameT, KernelType, /*Dims*/ 0, void>(KernelFunc);
1176-
MCGType = detail::CG::KERNEL_V1;
1186+
setType(detail::CG::KERNEL);
11771187
#endif
11781188
}
11791189

@@ -1217,7 +1227,7 @@ class __SYCL_EXPORT handler {
12171227
MArgs = std::move(MAssociatedAccesors);
12181228
MHostKernel.reset(
12191229
new detail::HostKernel<FuncT, void, 1, void>(std::move(Func)));
1220-
MCGType = detail::CG::RUN_ON_HOST_INTEL;
1230+
setType(detail::CG::RUN_ON_HOST_INTEL);
12211231
}
12221232

12231233
template <typename FuncT>
@@ -1231,7 +1241,7 @@ class __SYCL_EXPORT handler {
12311241

12321242
MHostTask.reset(new detail::HostTask(std::move(Func)));
12331243

1234-
MCGType = detail::CG::CODEPLAY_HOST_TASK;
1244+
setType(detail::CG::CODEPLAY_HOST_TASK);
12351245
}
12361246

12371247
template <typename FuncT>
@@ -1245,7 +1255,7 @@ class __SYCL_EXPORT handler {
12451255

12461256
MHostTask.reset(new detail::HostTask(std::move(Func)));
12471257

1248-
MCGType = detail::CG::CODEPLAY_HOST_TASK;
1258+
setType(detail::CG::CODEPLAY_HOST_TASK);
12491259
}
12501260

12511261
// replace _KERNELFUNCPARAM(KernelFunc) with KernelType KernelFunc
@@ -1285,7 +1295,7 @@ class __SYCL_EXPORT handler {
12851295
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
12861296
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
12871297
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
1288-
MCGType = detail::CG::KERNEL_V1;
1298+
setType(detail::CG::KERNEL);
12891299
#endif
12901300
}
12911301

@@ -1317,7 +1327,7 @@ class __SYCL_EXPORT handler {
13171327
detail::checkValueRange<Dims>(ExecutionRange);
13181328
MNDRDesc.set(std::move(ExecutionRange));
13191329
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
1320-
MCGType = detail::CG::KERNEL_V1;
1330+
setType(detail::CG::KERNEL);
13211331
#endif
13221332
}
13231333

@@ -1550,7 +1560,7 @@ class __SYCL_EXPORT handler {
15501560
detail::checkValueRange<Dims>(NumWorkGroups);
15511561
MNDRDesc.setNumWorkGroups(NumWorkGroups);
15521562
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
1553-
MCGType = detail::CG::KERNEL_V1;
1563+
setType(detail::CG::KERNEL);
15541564
#endif // __SYCL_DEVICE_ONLY__
15551565
}
15561566

@@ -1586,7 +1596,7 @@ class __SYCL_EXPORT handler {
15861596
detail::checkValueRange<Dims>(ExecRange);
15871597
MNDRDesc.set(std::move(ExecRange));
15881598
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
1589-
MCGType = detail::CG::KERNEL_V1;
1599+
setType(detail::CG::KERNEL);
15901600
#endif // __SYCL_DEVICE_ONLY__
15911601
}
15921602

@@ -1603,7 +1613,7 @@ class __SYCL_EXPORT handler {
16031613
// known constant
16041614
MNDRDesc.set(range<1>{1});
16051615
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1606-
MCGType = detail::CG::KERNEL_V1;
1616+
setType(detail::CG::KERNEL);
16071617
extractArgsAndReqs();
16081618
MKernelName = getKernelName();
16091619
}
@@ -1636,7 +1646,7 @@ class __SYCL_EXPORT handler {
16361646
MKernel = detail::getSyclObjImpl(std::move(Kernel));
16371647
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
16381648
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
1639-
MCGType = detail::CG::KERNEL_V1;
1649+
setType(detail::CG::KERNEL);
16401650
extractArgsAndReqs();
16411651
MKernelName = getKernelName();
16421652
}
@@ -1655,7 +1665,7 @@ class __SYCL_EXPORT handler {
16551665
MKernel = detail::getSyclObjImpl(std::move(Kernel));
16561666
detail::checkValueRange<Dims>(NDRange);
16571667
MNDRDesc.set(std::move(NDRange));
1658-
MCGType = detail::CG::KERNEL_V1;
1668+
setType(detail::CG::KERNEL);
16591669
extractArgsAndReqs();
16601670
MKernelName = getKernelName();
16611671
}
@@ -1679,7 +1689,7 @@ class __SYCL_EXPORT handler {
16791689
// known constant
16801690
MNDRDesc.set(range<1>{1});
16811691
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1682-
MCGType = detail::CG::KERNEL_V1;
1692+
setType(detail::CG::KERNEL);
16831693
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
16841694
extractArgsAndReqs();
16851695
MKernelName = getKernelName();
@@ -1694,7 +1704,7 @@ class __SYCL_EXPORT handler {
16941704
template <typename FuncT> void interop_task(FuncT Func) {
16951705

16961706
MInteropTask.reset(new detail::InteropTask(std::move(Func)));
1697-
MCGType = detail::CG::CODEPLAY_INTEROP_TASK;
1707+
setType(detail::CG::CODEPLAY_INTEROP_TASK);
16981708
}
16991709

17001710
/// Defines and invokes a SYCL kernel function for the specified range.
@@ -1720,7 +1730,7 @@ class __SYCL_EXPORT handler {
17201730
detail::checkValueRange<Dims>(NumWorkItems);
17211731
MNDRDesc.set(std::move(NumWorkItems));
17221732
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1723-
MCGType = detail::CG::KERNEL_V1;
1733+
setType(detail::CG::KERNEL);
17241734
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
17251735
extractArgsAndReqs();
17261736
MKernelName = getKernelName();
@@ -1756,7 +1766,7 @@ class __SYCL_EXPORT handler {
17561766
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
17571767
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
17581768
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1759-
MCGType = detail::CG::KERNEL_V1;
1769+
setType(detail::CG::KERNEL);
17601770
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
17611771
extractArgsAndReqs();
17621772
MKernelName = getKernelName();
@@ -1792,7 +1802,7 @@ class __SYCL_EXPORT handler {
17921802
detail::checkValueRange<Dims>(NDRange);
17931803
MNDRDesc.set(std::move(NDRange));
17941804
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1795-
MCGType = detail::CG::KERNEL_V1;
1805+
setType(detail::CG::KERNEL);
17961806
if (!MIsHost && !lambdaAndKernelHaveEqualName<NameT>()) {
17971807
extractArgsAndReqs();
17981808
MKernelName = getKernelName();
@@ -1833,7 +1843,7 @@ class __SYCL_EXPORT handler {
18331843
MNDRDesc.setNumWorkGroups(NumWorkGroups);
18341844
MKernel = detail::getSyclObjImpl(std::move(Kernel));
18351845
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
1836-
MCGType = detail::CG::KERNEL_V1;
1846+
setType(detail::CG::KERNEL);
18371847
#endif // __SYCL_DEVICE_ONLY__
18381848
}
18391849

@@ -1874,7 +1884,7 @@ class __SYCL_EXPORT handler {
18741884
MNDRDesc.set(std::move(ExecRange));
18751885
MKernel = detail::getSyclObjImpl(std::move(Kernel));
18761886
StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
1877-
MCGType = detail::CG::KERNEL_V1;
1887+
setType(detail::CG::KERNEL);
18781888
#endif // __SYCL_DEVICE_ONLY__
18791889
}
18801890

@@ -1957,7 +1967,7 @@ class __SYCL_EXPORT handler {
19571967
return;
19581968
}
19591969
#endif
1960-
MCGType = detail::CG::COPY_ACC_TO_PTR;
1970+
setType(detail::CG::COPY_ACC_TO_PTR);
19611971

19621972
detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Src;
19631973
detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
@@ -1996,7 +2006,7 @@ class __SYCL_EXPORT handler {
19962006
return;
19972007
}
19982008
#endif
1999-
MCGType = detail::CG::COPY_PTR_TO_ACC;
2009+
setType(detail::CG::COPY_PTR_TO_ACC);
20002010

20012011
detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Dst;
20022012
detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
@@ -2041,7 +2051,7 @@ class __SYCL_EXPORT handler {
20412051
"The destination accessor does not fit the copied memory.");
20422052
if (copyAccToAccHelper(Src, Dst))
20432053
return;
2044-
MCGType = detail::CG::COPY_ACC_TO_ACC;
2054+
setType(detail::CG::COPY_ACC_TO_ACC);
20452055

20462056
detail::AccessorBaseHost *AccBaseSrc = (detail::AccessorBaseHost *)&Src;
20472057
detail::AccessorImplPtr AccImplSrc = detail::getSyclObjImpl(*AccBaseSrc);
@@ -2071,7 +2081,7 @@ class __SYCL_EXPORT handler {
20712081
throwIfActionIsCreated();
20722082
static_assert(isValidTargetForExplicitOp(AccessTarget),
20732083
"Invalid accessor target for the update_host method.");
2074-
MCGType = detail::CG::UPDATE_HOST;
2084+
setType(detail::CG::UPDATE_HOST);
20752085

20762086
detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Acc;
20772087
detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
@@ -2103,7 +2113,7 @@ class __SYCL_EXPORT handler {
21032113
"Invalid accessor target for the fill method.");
21042114
if (!MIsHost && (((Dims == 1) && isConstOrGlobal(AccessTarget)) ||
21052115
isImageOrImageArray(AccessTarget))) {
2106-
MCGType = detail::CG::FILL;
2116+
setType(detail::CG::FILL);
21072117

21082118
detail::AccessorBaseHost *AccBase = (detail::AccessorBaseHost *)&Dst;
21092119
detail::AccessorImplPtr AccImpl = detail::getSyclObjImpl(*AccBase);
@@ -2148,7 +2158,7 @@ class __SYCL_EXPORT handler {
21482158
/// complete state.
21492159
void barrier() {
21502160
throwIfActionIsCreated();
2151-
MCGType = detail::CG::BARRIER;
2161+
setType(detail::CG::BARRIER);
21522162
}
21532163

21542164
/// Prevents any commands submitted afterward to this queue from executing
@@ -2212,7 +2222,9 @@ class __SYCL_EXPORT handler {
22122222
string_class MKernelName;
22132223
/// Storage for a sycl::kernel object.
22142224
shared_ptr_class<detail::kernel_impl> MKernel;
2215-
/// Type of the command group, e.g. kernel, fill.
2225+
/// Type of the command group, e.g. kernel, fill. Can also encode version.
2226+
/// Use getType and setType methods to access this variable unless
2227+
/// manipulations with version are required
22162228
detail::CG::CGTYPE MCGType = detail::CG::NONE;
22172229
/// Pointer to the source host memory or accessor(depending on command type).
22182230
void *MSrcPtr = nullptr;

‎sycl/source/detail/scheduler/commands.cpp

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1882,8 +1882,7 @@ cl_int ExecCGCommand::enqueueImp() {
18821882
"Enqueueing run_on_host_intel task has failed.", Error);
18831883
}
18841884
}
1885-
case CG::CGTYPE::KERNEL:
1886-
case CG::CGTYPE::KERNEL_V1: {
1885+
case CG::CGTYPE::KERNEL: {
18871886
CGExecKernel *ExecKernel = (CGExecKernel *)MCommandGroup.get();
18881887

18891888
NDRDescT &NDRDesc = ExecKernel->MNDRDesc;

‎sycl/source/handler.cpp

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -127,9 +127,8 @@ event handler::finalize() {
127127
}
128128

129129
unique_ptr_class<detail::CG> CommandGroup;
130-
switch (MCGType) {
130+
switch (getType()) {
131131
case detail::CG::KERNEL:
132-
case detail::CG::KERNEL_V1:
133132
case detail::CG::RUN_ON_HOST_INTEL: {
134133
CommandGroup.reset(new detail::CGExecKernel(
135134
std::move(MNDRDesc), std::move(MHostKernel), std::move(MKernel),
@@ -204,6 +203,11 @@ event handler::finalize() {
204203
PI_INVALID_OPERATION);
205204
}
206205

206+
if (!CommandGroup)
207+
throw sycl::runtime_error(
208+
"Internal Error. Command group cannot be constructed.",
209+
PI_INVALID_OPERATION);
210+
207211
detail::EventImplPtr Event = detail::Scheduler::getInstance().addCG(
208212
std::move(CommandGroup), std::move(MQueue));
209213

0 commit comments

Comments
 (0)