Skip to content

Commit 20ea292

Browse files
vinser52vinser52
committed
Move Kernel specific data from handler_impl to a separate data structure to use it in handler-based and handler-less submission paths
1 parent 3717903 commit 20ea292

File tree

10 files changed

+133
-58
lines changed

10 files changed

+133
-58
lines changed

‎sycl/include/sycl/handler.hpp‎

Lines changed: 5 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -856,14 +856,11 @@ class __SYCL_EXPORT handler {
856856
// Force hasSpecialCaptures to be evaluated at compile-time.
857857
constexpr bool HasSpecialCapt = detail::hasSpecialCaptures<KernelName>();
858858
setKernelInfo((void *)MHostKernel->getPtr(),
859-
detail::getKernelNumParams<KernelName>(),
860-
&(detail::getKernelParamDesc<KernelName>),
861-
detail::isKernelESIMD<KernelName>(), HasSpecialCapt);
859+
&detail::getDeviceKernelInfo<KernelName>());
862860

863861
constexpr std::string_view KernelNameStr =
864862
detail::getKernelName<KernelName>();
865863
MKernelName = KernelNameStr;
866-
setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<KernelName>());
867864
} else {
868865
// In case w/o the integration header it is necessary to process
869866
// accessors from the list(which are associated with this handler) as
@@ -3658,10 +3655,13 @@ class __SYCL_EXPORT handler {
36583655
void setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::id<1> Offset);
36593656
void setNDRangeDescriptor(sycl::range<1> NumWorkItems,
36603657
sycl::range<1> LocalSize, sycl::id<1> Offset);
3661-
3658+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
36623659
void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
36633660
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
36643661
bool KernelIsESIMD, bool KernelHasSpecialCaptures);
3662+
#endif
3663+
void setKernelInfo(void *KernelFuncPtr,
3664+
detail::DeviceKernelInfo *DeviceKernelInfoPtr);
36653665

36663666
void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr);
36673667

@@ -3689,7 +3689,6 @@ class __SYCL_EXPORT handler {
36893689
void setKernelNameBasedCachePtr(
36903690
detail::KernelNameBasedCacheT *KernelNameBasedCachePtr);
36913691
#endif
3692-
void setDeviceKernelInfoPtr(detail::DeviceKernelInfo *DeviceKernelInfoPtr);
36933692

36943693
queue getQueue();
36953694

‎sycl/source/detail/handler_impl.hpp‎

Lines changed: 2 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,7 @@
1111
#include "sycl/handler.hpp"
1212
#include <detail/cg.hpp>
1313
#include <detail/kernel_bundle_impl.hpp>
14+
#include <detail/kernel_data.hpp>
1415
#include <memory>
1516
#include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
1617

@@ -236,16 +237,7 @@ class handler_impl {
236237
// Allocation ptr to be freed asynchronously.
237238
void *MFreePtr = nullptr;
238239

239-
// Store information about the kernel arguments.
240-
void *MKernelFuncPtr = nullptr;
241-
int MKernelNumArgs = 0;
242-
detail::kernel_param_desc_t (*MKernelParamDescGetter)(int) = nullptr;
243-
bool MKernelIsESIMD = false;
244-
bool MKernelHasSpecialCaptures = true;
245-
246-
// A pointer to device kernel information. Cached on the application side in
247-
// headers or retrieved from program manager.
248-
DeviceKernelInfo *MDeviceKernelInfoPtr = nullptr;
240+
KernelData MKernelData;
249241
};
250242

251243
} // namespace detail

‎sycl/source/detail/kernel_data.hpp‎

Lines changed: 81 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,81 @@
1+
//==---------------- kernel_data.hpp - SYCL handler -----------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <sycl/detail/kernel_desc.hpp>
12+
#include <sycl/detail/kernel_name_based_cache.hpp>
13+
14+
namespace sycl {
15+
inline namespace _V1 {
16+
namespace detail {
17+
18+
class KernelData {
19+
public:
20+
using KernelParamDescGetterT = detail::kernel_param_desc_t (*)(int);
21+
22+
KernelData() = default;
23+
~KernelData() = default;
24+
KernelData(const KernelData &) = default;
25+
KernelData(KernelData &&) = default;
26+
KernelData &operator=(const KernelData &) = default;
27+
KernelData &operator=(KernelData &&) = default;
28+
29+
void *getKernelFuncPtr() const { return MKernelFuncPtr; }
30+
31+
size_t getKernelNumArgs() const { return MDeviceKernelInfoPtr->NumParams; }
32+
33+
KernelParamDescGetterT getKernelParamDescGetter() const {
34+
return MDeviceKernelInfoPtr->ParamDescGetter;
35+
}
36+
37+
bool isESIMD() const { return MDeviceKernelInfoPtr->IsESIMD; }
38+
39+
bool hasSpecialCaptures() const {
40+
return MDeviceKernelInfoPtr->HasSpecialCaptures;
41+
}
42+
43+
DeviceKernelInfo *getDeviceKernelInfoPtr() const {
44+
return MDeviceKernelInfoPtr;
45+
}
46+
47+
void setDeviceKernelInfoPtr(DeviceKernelInfo *Ptr) {
48+
MDeviceKernelInfoPtr = Ptr;
49+
}
50+
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
51+
void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
52+
KernelParamDescGetterT KernelParamDescGetter,
53+
bool KernelIsESIMD, bool KernelHasSpecialCaptures) {
54+
MKernelFuncPtr = KernelFuncPtr;
55+
MDeviceKernelInfoPtr->NumParams = KernelNumArgs;
56+
MDeviceKernelInfoPtr->ParamDescGetter = KernelParamDescGetter;
57+
MDeviceKernelInfoPtr->IsESIMD = KernelIsESIMD;
58+
MDeviceKernelInfoPtr->HasSpecialCaptures = KernelHasSpecialCaptures;
59+
}
60+
#endif
61+
62+
void setKernelInfo(void *KernelFuncPtr,
63+
detail::DeviceKernelInfo *DeviceKernelInfoPtr) {
64+
MKernelFuncPtr = KernelFuncPtr;
65+
MDeviceKernelInfoPtr = DeviceKernelInfoPtr;
66+
}
67+
68+
bool usesAssert() const { return MDeviceKernelInfoPtr->usesAssert(); }
69+
70+
private:
71+
// Store information about the kernel arguments.
72+
void *MKernelFuncPtr = nullptr;
73+
74+
// A pointer to device kernel information. Cached on the application side in
75+
// headers or retrieved from program manager.
76+
DeviceKernelInfo *MDeviceKernelInfoPtr = nullptr;
77+
};
78+
79+
} // namespace detail
80+
} // namespace _V1
81+
} // namespace sycl

‎sycl/source/detail/queue_impl.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -872,7 +872,7 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
872872
// Kernel only uses assert if it's non interop one
873873
KernelUsesAssert =
874874
(!Handler.MKernel || Handler.MKernel->hasSYCLMetadata()) &&
875-
Handler.impl->MDeviceKernelInfoPtr->usesAssert();
875+
Handler.impl->MKernelData.usesAssert();
876876

877877
auto &PostProcess = *PostProcessorFunc;
878878
PostProcess(IsKernel, KernelUsesAssert, Event);

‎sycl/source/handler.cpp‎

Lines changed: 35 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -496,12 +496,13 @@ event handler::finalize() {
496496

497497
// Extract arguments from the kernel lambda, if required.
498498
// Skipping this is currently limited to simple kernels on the fast path.
499-
if (type == detail::CGType::Kernel && impl->MKernelFuncPtr &&
500-
(!KernelFastPath || impl->MKernelHasSpecialCaptures)) {
499+
if (type == detail::CGType::Kernel && impl->MKernelData.getKernelFuncPtr() &&
500+
(!KernelFastPath || impl->MKernelData.hasSpecialCaptures())) {
501501
clearArgs();
502-
extractArgsAndReqsFromLambda((char *)impl->MKernelFuncPtr,
503-
impl->MKernelParamDescGetter,
504-
impl->MKernelNumArgs, impl->MKernelIsESIMD);
502+
extractArgsAndReqsFromLambda((char *)impl->MKernelData.getKernelFuncPtr(),
503+
impl->MKernelData.getKernelParamDescGetter(),
504+
impl->MKernelData.getKernelNumArgs(),
505+
impl->MKernelData.isESIMD());
505506
}
506507

507508
// According to 4.7.6.9 of SYCL2020 spec, if a placeholder accessor is passed
@@ -542,16 +543,17 @@ event handler::finalize() {
542543
}
543544

544545
if (type == detail::CGType::Kernel) {
545-
if (impl->MDeviceKernelInfoPtr) {
546+
if (impl->MKernelData.getDeviceKernelInfoPtr()) {
546547
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
547-
impl->MDeviceKernelInfoPtr->initIfNeeded(toKernelNameStrT(MKernelName));
548+
impl->MKernelData.getDeviceKernelInfoPtr()->initIfNeeded(
549+
toKernelNameStrT(MKernelName));
548550
#endif
549551
} else {
550552
// Fetch the device kernel info pointer if it hasn't been set (e.g.
551553
// in kernel bundle or free function cases).
552-
impl->MDeviceKernelInfoPtr =
554+
impl->MKernelData.setDeviceKernelInfoPtr(
553555
&detail::ProgramManager::getInstance().getOrCreateDeviceKernelInfo(
554-
toKernelNameStrT(MKernelName));
556+
toKernelNameStrT(MKernelName)));
555557
}
556558
// If there were uses of set_specialization_constant build the kernel_bundle
557559
detail::kernel_bundle_impl *KernelBundleImpPtr =
@@ -627,7 +629,7 @@ event handler::finalize() {
627629
if (DiscardEvent) {
628630
// Kernel only uses assert if it's non interop one
629631
bool KernelUsesAssert = !(MKernel && MKernel->isInterop()) &&
630-
impl->MDeviceKernelInfoPtr->usesAssert();
632+
impl->MKernelData.usesAssert();
631633
DiscardEvent = !KernelUsesAssert;
632634
}
633635

@@ -647,7 +649,7 @@ event handler::finalize() {
647649
if (xptiEnabled) {
648650
std::tie(CmdTraceEvent, InstanceID) = emitKernelInstrumentationData(
649651
detail::GSYCLStreamID, MKernel, MCodeLoc, impl->MIsTopCodeLoc,
650-
MKernelName.data(), *impl->MDeviceKernelInfoPtr,
652+
MKernelName.data(), *impl->MKernelData.getDeviceKernelInfoPtr(),
651653
impl->get_queue_or_null(), impl->MNDRDesc, KernelBundleImpPtr,
652654
impl->MArgs);
653655
detail::emitInstrumentationGeneral(detail::GSYCLStreamID, InstanceID,
@@ -664,11 +666,14 @@ event handler::finalize() {
664666
enqueueImpKernel(
665667
impl->get_queue(), impl->MNDRDesc, impl->MArgs, KernelBundleImpPtr,
666668
MKernel.get(), toKernelNameStrT(MKernelName),
667-
*impl->MDeviceKernelInfoPtr, RawEvents, ResultEvent.get(), nullptr,
668-
impl->MKernelCacheConfig, impl->MKernelIsCooperative,
669-
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
670-
BinImage, impl->MKernelFuncPtr, impl->MKernelNumArgs,
671-
impl->MKernelParamDescGetter, impl->MKernelHasSpecialCaptures);
669+
*impl->MKernelData.getDeviceKernelInfoPtr(), RawEvents,
670+
ResultEvent.get(), nullptr, impl->MKernelCacheConfig,
671+
impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
672+
impl->MKernelWorkGroupMemorySize, BinImage,
673+
impl->MKernelData.getKernelFuncPtr(),
674+
impl->MKernelData.getKernelNumArgs(),
675+
impl->MKernelData.getKernelParamDescGetter(),
676+
impl->MKernelData.hasSpecialCaptures());
672677
#ifdef XPTI_ENABLE_INSTRUMENTATION
673678
if (xptiEnabled) {
674679
// Emit signal only when event is created
@@ -726,7 +731,7 @@ event handler::finalize() {
726731
impl->MNDRDesc, std::move(MHostKernel), std::move(MKernel),
727732
std::move(impl->MKernelBundle), std::move(impl->CGData),
728733
std::move(impl->MArgs), toKernelNameStrT(MKernelName),
729-
*impl->MDeviceKernelInfoPtr, std::move(MStreamStorage),
734+
*impl->MKernelData.getDeviceKernelInfoPtr(), std::move(MStreamStorage),
730735
std::move(impl->MAuxiliaryResources), getType(),
731736
impl->MKernelCacheConfig, impl->MKernelIsCooperative,
732737
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
@@ -2610,26 +2615,25 @@ void handler::setNDRangeDescriptor(sycl::range<1> NumWorkItems,
26102615
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
26112616
void handler::setKernelNameBasedCachePtr(
26122617
sycl::detail::KernelNameBasedCacheT *KernelNameBasedCachePtr) {
2613-
setDeviceKernelInfoPtr(reinterpret_cast<sycl::detail::DeviceKernelInfo *>(
2614-
KernelNameBasedCachePtr));
2615-
}
2616-
#endif
2617-
2618-
void handler::setDeviceKernelInfoPtr(
2619-
sycl::detail::DeviceKernelInfo *DeviceKernelInfoPtr) {
2620-
assert(!impl->MDeviceKernelInfoPtr && "Already set!");
2621-
impl->MDeviceKernelInfoPtr = DeviceKernelInfoPtr;
2618+
assert(!impl->MKernelData.getDeviceKernelInfoPtr() && "Already set!");
2619+
impl->MKernelData.setDeviceKernelInfoPtr(
2620+
reinterpret_cast<sycl::detail::DeviceKernelInfo *>(
2621+
KernelNameBasedCachePtr));
26222622
}
26232623

26242624
void handler::setKernelInfo(
26252625
void *KernelFuncPtr, int KernelNumArgs,
26262626
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
26272627
bool KernelIsESIMD, bool KernelHasSpecialCaptures) {
2628-
impl->MKernelFuncPtr = KernelFuncPtr;
2629-
impl->MKernelNumArgs = KernelNumArgs;
2630-
impl->MKernelParamDescGetter = KernelParamDescGetter;
2631-
impl->MKernelIsESIMD = KernelIsESIMD;
2632-
impl->MKernelHasSpecialCaptures = KernelHasSpecialCaptures;
2628+
impl->MKernelData.setKernelInfo(KernelFuncPtr, KernelNumArgs,
2629+
KernelParamDescGetter, KernelIsESIMD,
2630+
KernelHasSpecialCaptures);
2631+
}
2632+
#endif
2633+
2634+
void handler::setKernelInfo(void *KernelFuncPtr,
2635+
detail::DeviceKernelInfo *DeviceKernelInfoPtr) {
2636+
impl->MKernelData.setKernelInfo(KernelFuncPtr, DeviceKernelInfoPtr);
26332637
}
26342638

26352639
void handler::instantiateKernelOnHost(void *InstantiateKernelOnHostPtr) {

‎sycl/test/abi/sycl_symbols_linux.dump‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3592,7 +3592,6 @@ _ZN4sycl3_V17handler21setKernelWorkGroupMemEm
35923592
_ZN4sycl3_V17handler21setUserFacingNodeTypeENS0_3ext6oneapi12experimental9node_typeE
35933593
_ZN4sycl3_V17handler22ext_oneapi_fill2d_implEPvmPKvmmm
35943594
_ZN4sycl3_V17handler22memcpyFromDeviceGlobalEPvPKvbmm
3595-
_ZN4sycl3_V17handler22setDeviceKernelInfoPtrEPNS0_6detail16DeviceKernelInfoE
35963595
_ZN4sycl3_V17handler22setHandlerKernelBundleENS0_6kernelE
35973596
_ZN4sycl3_V17handler22setHandlerKernelBundleERKSt10shared_ptrINS0_6detail18kernel_bundle_implEE
35983597
_ZN4sycl3_V17handler22setKernelClusterLaunchENS0_5rangeILi1EEE

‎sycl/test/abi/sycl_symbols_windows.dump‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4409,7 +4409,6 @@
44094409
?setArgsHelper@handler@_V1@sycl@@AEAAXH@Z
44104410
?setArgsToAssociatedAccessors@handler@_V1@sycl@@AEAAXXZ
44114411
?setDevice@HostProfilingInfo@detail@_V1@sycl@@QEAAXPEAVdevice_impl@234@@Z
4412-
?setDeviceKernelInfoPtr@handler@_V1@sycl@@AEAAXPEAVDeviceKernelInfo@detail@23@@Z
44134412
?setHandlerKernelBundle@handler@_V1@sycl@@AEAAXAEBV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@@Z
44144413
?setHandlerKernelBundle@handler@_V1@sycl@@AEAAXVkernel@23@@Z
44154414
?setKernelCacheConfig@handler@_V1@sycl@@AEAAXW4StableKernelCacheConfig@123@@Z

‎sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,7 @@ class MockHandler : public sycl::handler {
148148
std::move(impl->MNDRDesc), std::move(CGH->MHostKernel),
149149
std::move(CGH->MKernel), std::move(impl->MKernelBundle),
150150
std::move(impl->CGData), std::move(impl->MArgs),
151-
CGH->MKernelName.data(), *impl->MDeviceKernelInfoPtr,
151+
CGH->MKernelName.data(), *impl->MKernelData.getDeviceKernelInfo(),
152152
std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources),
153153
impl->MCGType, {}, impl->MKernelIsCooperative,
154154
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,

‎sycl/unittests/scheduler/SchedulerTestUtils.hpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -303,10 +303,10 @@ class MockHandlerCustomFinalize : public MockHandler {
303303
CommandGroup.reset(new sycl::detail::CGExecKernel(
304304
getNDRDesc(), std::move(getHostKernel()), getKernel(),
305305
std::move(impl->MKernelBundle), std::move(CGData), getArgs(),
306-
getKernelName(), *impl->MDeviceKernelInfoPtr, getStreamStorage(),
307-
impl->MAuxiliaryResources, getType(), {}, impl->MKernelIsCooperative,
308-
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
309-
getCodeLoc()));
306+
getKernelName(), *impl->MKernelData.getDeviceKernelInfoPtr(),
307+
getStreamStorage(), impl->MAuxiliaryResources, getType(), {},
308+
impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
309+
impl->MKernelWorkGroupMemorySize, getCodeLoc()));
310310
break;
311311
}
312312
case sycl::detail::CGType::CodeplayHostTask: {

‎sycl/unittests/scheduler/StreamInitDependencyOnHost.cpp‎

Lines changed: 4 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -33,9 +33,10 @@ class MockHandlerStreamInit : public MockHandler {
3333
detail::CG::StorageInitHelper(getArgsStorage(), getAccStorage(),
3434
getSharedPtrStorage(),
3535
getRequirements(), getEvents()),
36-
getArgs(), getKernelName(), *impl->MDeviceKernelInfoPtr,
37-
getStreamStorage(), std::move(impl->MAuxiliaryResources), getType(),
38-
{}, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
36+
getArgs(), getKernelName(),
37+
*impl->MKernelData.getDeviceKernelInfoPtr(), getStreamStorage(),
38+
std::move(impl->MAuxiliaryResources), getType(), {},
39+
impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
3940
impl->MKernelWorkGroupMemorySize, getCodeLoc()));
4041
break;
4142
}

0 commit comments

Comments
 (0)