Skip to content

Commit bc8f182

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 f054815 commit bc8f182

File tree

7 files changed

+111
-39
lines changed

7 files changed

+111
-39
lines changed

‎sycl/source/detail/handler_impl.hpp‎

Lines changed: 2 additions & 9 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,15 +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 a kernel name based cache retrieved on the application side.
247-
KernelNameBasedCacheT *MKernelNameBasedCachePtr = nullptr;
240+
KernelData MKernelData;
248241
};
249242

250243
} // namespace detail

‎sycl/source/detail/kernel_data.hpp‎

Lines changed: 73 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,73 @@
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 MKernelNumArgs; }
32+
33+
KernelParamDescGetterT getKernelParamDescGetter() const {
34+
return MKernelParamDescGetter;
35+
}
36+
37+
bool isESIMD() const { return MKernelIsESIMD; }
38+
39+
bool hasSpecialCaptures() const { return MKernelHasSpecialCaptures; }
40+
41+
KernelNameBasedCacheT *getKernelNameBasedCachePtr() const {
42+
return MKernelNameBasedCachePtr;
43+
}
44+
45+
void setKernelNameBasedCachePtr(KernelNameBasedCacheT *Ptr) {
46+
MKernelNameBasedCachePtr = Ptr;
47+
}
48+
49+
void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
50+
KernelParamDescGetterT KernelParamDescGetter,
51+
bool KernelIsESIMD, bool KernelHasSpecialCaptures) {
52+
MKernelFuncPtr = KernelFuncPtr;
53+
MKernelNumArgs = KernelNumArgs;
54+
MKernelParamDescGetter = KernelParamDescGetter;
55+
MKernelIsESIMD = KernelIsESIMD;
56+
MKernelHasSpecialCaptures = KernelHasSpecialCaptures;
57+
}
58+
59+
private:
60+
// Store information about the kernel arguments.
61+
void *MKernelFuncPtr = nullptr;
62+
size_t MKernelNumArgs = 0;
63+
KernelParamDescGetterT MKernelParamDescGetter = nullptr;
64+
bool MKernelIsESIMD = false;
65+
bool MKernelHasSpecialCaptures = true;
66+
67+
// A pointer to a kernel name based cache retrieved on the application side.
68+
KernelNameBasedCacheT *MKernelNameBasedCachePtr = nullptr;
69+
};
70+
71+
} // namespace detail
72+
} // namespace _V1
73+
} // namespace sycl

‎sycl/source/detail/queue_impl.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -873,7 +873,7 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
873873
(!Handler.MKernel || Handler.MKernel->hasSYCLMetadata()) &&
874874
ProgramManager::getInstance().kernelUsesAssert(
875875
Handler.MKernelName.data(),
876-
Handler.impl->MKernelNameBasedCachePtr);
876+
Handler.impl->MKernelData.getKernelNameBasedCachePtr());
877877

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

‎sycl/source/handler.cpp‎

Lines changed: 25 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -492,12 +492,13 @@ event handler::finalize() {
492492

493493
// Extract arguments from the kernel lambda, if required.
494494
// Skipping this is currently limited to simple kernels on the fast path.
495-
if (type == detail::CGType::Kernel && impl->MKernelFuncPtr &&
496-
(!KernelFastPath || impl->MKernelHasSpecialCaptures)) {
495+
if (type == detail::CGType::Kernel && impl->MKernelData.getKernelFuncPtr() &&
496+
(!KernelFastPath || impl->MKernelData.hasSpecialCaptures())) {
497497
clearArgs();
498-
extractArgsAndReqsFromLambda((char *)impl->MKernelFuncPtr,
499-
impl->MKernelParamDescGetter,
500-
impl->MKernelNumArgs, impl->MKernelIsESIMD);
498+
extractArgsAndReqsFromLambda((char *)impl->MKernelData.getKernelFuncPtr(),
499+
impl->MKernelData.getKernelParamDescGetter(),
500+
impl->MKernelData.getKernelNumArgs(),
501+
impl->MKernelData.isESIMD());
501502
}
502503

503504
// According to 4.7.6.9 of SYCL2020 spec, if a placeholder accessor is passed
@@ -614,7 +615,8 @@ event handler::finalize() {
614615
bool KernelUsesAssert =
615616
!(MKernel && MKernel->isInterop()) &&
616617
detail::ProgramManager::getInstance().kernelUsesAssert(
617-
toKernelNameStrT(MKernelName), impl->MKernelNameBasedCachePtr);
618+
toKernelNameStrT(MKernelName),
619+
impl->MKernelData.getKernelNameBasedCachePtr());
618620
DiscardEvent = !KernelUsesAssert;
619621
}
620622

@@ -635,7 +637,8 @@ event handler::finalize() {
635637
StreamID = xptiRegisterStream(detail::SYCL_STREAM_NAME);
636638
std::tie(CmdTraceEvent, InstanceID) = emitKernelInstrumentationData(
637639
StreamID, MKernel, MCodeLoc, impl->MIsTopCodeLoc,
638-
MKernelName.data(), impl->MKernelNameBasedCachePtr,
640+
MKernelName.data(),
641+
impl->MKernelData.getKernelNameBasedCachePtr(),
639642
impl->get_queue_or_null(), impl->MNDRDesc, KernelBundleImpPtr,
640643
impl->MArgs);
641644
detail::emitInstrumentationGeneral(StreamID, InstanceID,
@@ -652,11 +655,14 @@ event handler::finalize() {
652655
enqueueImpKernel(
653656
impl->get_queue(), impl->MNDRDesc, impl->MArgs, KernelBundleImpPtr,
654657
MKernel.get(), toKernelNameStrT(MKernelName),
655-
impl->MKernelNameBasedCachePtr, RawEvents, ResultEvent.get(),
656-
nullptr, impl->MKernelCacheConfig, impl->MKernelIsCooperative,
657-
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
658-
BinImage, impl->MKernelFuncPtr, impl->MKernelNumArgs,
659-
impl->MKernelParamDescGetter, impl->MKernelHasSpecialCaptures);
658+
impl->MKernelData.getKernelNameBasedCachePtr(), RawEvents,
659+
ResultEvent.get(), nullptr, impl->MKernelCacheConfig,
660+
impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
661+
impl->MKernelWorkGroupMemorySize, BinImage,
662+
impl->MKernelData.getKernelFuncPtr(),
663+
impl->MKernelData.getKernelNumArgs(),
664+
impl->MKernelData.getKernelParamDescGetter(),
665+
impl->MKernelData.hasSpecialCaptures());
660666
#ifdef XPTI_ENABLE_INSTRUMENTATION
661667
if (xptiEnabled) {
662668
// Emit signal only when event is created
@@ -713,9 +719,9 @@ event handler::finalize() {
713719
impl->MNDRDesc, std::move(MHostKernel), std::move(MKernel),
714720
std::move(impl->MKernelBundle), std::move(impl->CGData),
715721
std::move(impl->MArgs), toKernelNameStrT(MKernelName),
716-
impl->MKernelNameBasedCachePtr, std::move(MStreamStorage),
717-
std::move(impl->MAuxiliaryResources), getType(),
718-
impl->MKernelCacheConfig, impl->MKernelIsCooperative,
722+
impl->MKernelData.getKernelNameBasedCachePtr(),
723+
std::move(MStreamStorage), std::move(impl->MAuxiliaryResources),
724+
getType(), impl->MKernelCacheConfig, impl->MKernelIsCooperative,
719725
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
720726
MCodeLoc));
721727
break;
@@ -2596,18 +2602,16 @@ void handler::setNDRangeDescriptor(sycl::range<1> NumWorkItems,
25962602

25972603
void handler::setKernelNameBasedCachePtr(
25982604
sycl::detail::KernelNameBasedCacheT *KernelNameBasedCachePtr) {
2599-
impl->MKernelNameBasedCachePtr = KernelNameBasedCachePtr;
2605+
impl->MKernelData.setKernelNameBasedCachePtr(KernelNameBasedCachePtr);
26002606
}
26012607

26022608
void handler::setKernelInfo(
26032609
void *KernelFuncPtr, int KernelNumArgs,
26042610
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
26052611
bool KernelIsESIMD, bool KernelHasSpecialCaptures) {
2606-
impl->MKernelFuncPtr = KernelFuncPtr;
2607-
impl->MKernelNumArgs = KernelNumArgs;
2608-
impl->MKernelParamDescGetter = KernelParamDescGetter;
2609-
impl->MKernelIsESIMD = KernelIsESIMD;
2610-
impl->MKernelHasSpecialCaptures = KernelHasSpecialCaptures;
2612+
impl->MKernelData.setKernelInfo(KernelFuncPtr, KernelNumArgs,
2613+
KernelParamDescGetter, KernelIsESIMD,
2614+
KernelHasSpecialCaptures);
26112615
}
26122616

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

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,8 @@ 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->MKernelNameBasedCachePtr,
151+
CGH->MKernelName.data(),
152+
impl->MKernelData.getKernelNameBasedCachePtr(),
152153
std::move(CGH->MStreamStorage), std::move(impl->MAuxiliaryResources),
153154
impl->MCGType, {}, impl->MKernelIsCooperative,
154155
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->MKernelNameBasedCachePtr, getStreamStorage(),
307-
impl->MAuxiliaryResources, getType(), {}, impl->MKernelIsCooperative,
308-
impl->MKernelUsesClusterLaunch, impl->MKernelWorkGroupMemorySize,
309-
getCodeLoc()));
306+
getKernelName(), impl->MKernelData.getKernelNameBasedCachePtr(),
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->MKernelNameBasedCachePtr,
37-
getStreamStorage(), std::move(impl->MAuxiliaryResources), getType(),
38-
{}, impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
36+
getArgs(), getKernelName(),
37+
impl->MKernelData.getKernelNameBasedCachePtr(), getStreamStorage(),
38+
std::move(impl->MAuxiliaryResources), getType(), {},
39+
impl->MKernelIsCooperative, impl->MKernelUsesClusterLaunch,
3940
impl->MKernelWorkGroupMemorySize, getCodeLoc()));
4041
break;
4142
}

0 commit comments

Comments
 (0)