Skip to content

[WIP] Defer arg extraction until handler::finalize #18413

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 1 commit into from
Closed

[WIP] Defer arg extraction until handler::finalize #18413

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
18 changes: 11 additions & 7 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -749,8 +749,8 @@ class __SYCL_EXPORT handler {

/// Stores lambda to the template-free object
///
/// Also initializes kernel name, list of arguments and requirements using
/// information from the integration header/built-ins.
/// Also initializes the kernel name and prepares for arguments to
/// be extracted from the lambda in handler::finalize().
///
/// \param KernelFunc is a SYCL kernel function
/// \param ParamDescs is the vector of kernel parameter descriptors.
Expand Down Expand Up @@ -796,11 +796,11 @@ class __SYCL_EXPORT handler {
if constexpr (KernelHasName) {
// TODO support ESIMD in no-integration-header case too.

clearArgs();
extractArgsAndReqsFromLambda(MHostKernel->getPtr(),
&(detail::getKernelParamDesc<KernelName>),
detail::getKernelNumParams<KernelName>(),
detail::isKernelESIMD<KernelName>());
setKernelInfo((void *)MHostKernel->getPtr(),
detail::getKernelNumParams<KernelName>(),
&(detail::getKernelParamDesc<KernelName>),
detail::isKernelESIMD<KernelName>());

MKernelName = detail::getKernelName<KernelName>();
} else {
// In case w/o the integration header it is necessary to process
Expand Down Expand Up @@ -3761,6 +3761,10 @@ class __SYCL_EXPORT handler {
sycl::range<3> LocalSize, sycl::id<3> Offset,
int Dims);

void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
bool KernelIsESIMD);

friend class detail::HandlerAccess;

#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
Expand Down
6 changes: 6 additions & 0 deletions sycl/source/detail/handler_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -198,6 +198,12 @@ class handler_impl {

// Allocation ptr to be freed asynchronously.
void *MFreePtr = nullptr;

// Store information about the kernel arguments.
void *MKernelFuncPtr = nullptr;
int MKernelNumArgs = 0;
detail::kernel_param_desc_t (*MKernelParamDescGetter)(int) = nullptr;
bool MKernelIsESIMD = false;
Comment on lines +202 to +206
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I wonder if we can just keep that inside HostKernel class.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Probably? My original hope here was that we could get rid of HostKernel, but I didn't appreciate that we need it in order to extend the lifetime of the kernel function until handler::finalize is called.

I'm going to focus on getting a version of things that works and demonstrates the optimizations, but once I reach that point I'm open to talking about where the variables should go.

};

} // namespace detail
Expand Down
21 changes: 20 additions & 1 deletion sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -411,6 +411,16 @@ event handler::finalize() {
return MLastEvent;
MIsFinalized = true;

// Extract arguments from the kernel lambda.
// TODO: Skip this in simple cases.
const auto &type = getType();
if (type == detail::CGType::Kernel && impl->MKernelFuncPtr) {
clearArgs();
extractArgsAndReqsFromLambda((char *)impl->MKernelFuncPtr,
impl->MKernelParamDescGetter,
impl->MKernelNumArgs, impl->MKernelIsESIMD);
}

// According to 4.7.6.9 of SYCL2020 spec, if a placeholder accessor is passed
// to a command without being bound to a command group, an exception should
// be thrown.
Expand Down Expand Up @@ -448,7 +458,6 @@ event handler::finalize() {
}
}

const auto &type = getType();
if (type == detail::CGType::Kernel) {
// If there were uses of set_specialization_constant build the kernel_bundle
std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImpPtr =
Expand Down Expand Up @@ -2254,6 +2263,16 @@ void handler::setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems,
impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset, Dims};
}

void handler::setKernelInfo(
void *KernelFuncPtr, int KernelNumArgs,
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
bool KernelIsESIMD) {
impl->MKernelFuncPtr = KernelFuncPtr;
impl->MKernelNumArgs = KernelNumArgs;
impl->MKernelParamDescGetter = KernelParamDescGetter;
impl->MKernelIsESIMD = KernelIsESIMD;
}

void handler::saveCodeLoc(detail::code_location CodeLoc, bool IsTopCodeLoc) {
MCodeLoc = CodeLoc;
impl->MIsTopCodeLoc = IsTopCodeLoc;
Expand Down
1 change: 1 addition & 0 deletions sycl/test/abi/sycl_symbols_linux.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3523,6 +3523,7 @@ _ZN4sycl3_V17handler12addReductionERKSt10shared_ptrIKvE
_ZN4sycl3_V17handler12setArgHelperEiRNS0_3ext6oneapi12experimental6detail30dynamic_work_group_memory_baseE
_ZN4sycl3_V17handler12setArgHelperEiRNS0_6detail22work_group_memory_implE
_ZN4sycl3_V17handler13getKernelNameEv
_ZN4sycl3_V17handler13setKernelInfoEPviPFNS0_6detail19kernel_param_desc_tEiEb
_ZN4sycl3_V17handler14addAccessorReqESt10shared_ptrINS0_6detail16AccessorImplHostEE
_ZN4sycl3_V17handler14setNDRangeUsedEb
_ZN4sycl3_V17handler15ext_oneapi_copyENS0_3ext6oneapi12experimental16image_mem_handleENS0_5rangeILi3EEERKNS4_16image_descriptorEPvS7_S7_S7_
Expand Down
1 change: 1 addition & 0 deletions sycl/test/abi/sycl_symbols_windows.dump
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4377,6 +4377,7 @@
?setHandlerKernelBundle@handler@_V1@sycl@@AEAAXVkernel@23@@Z
?setKernelCacheConfig@handler@_V1@sycl@@AEAAXW4StableKernelCacheConfig@123@@Z
?setKernelClusterLaunch@handler@_V1@sycl@@AEAAXV?$range@$02@23@H@Z
?setKernelInfo@handler@_V1@sycl@@AEAAXPEAXHP6A?AUkernel_param_desc_t@detail@23@H@Z_N@Z
?setKernelIsCooperative@handler@_V1@sycl@@AEAAX_N@Z
?setKernelWorkGroupMem@handler@_V1@sycl@@AEAAX_K@Z
?setLocalAccessorArgHelper@handler@_V1@sycl@@AEAAXHAEAVLocalAccessorBaseHost@detail@23@@Z
Expand Down
Loading