Skip to content

[SYCL] Move Kernel specific data from handler_impl to a separate data structure #19843

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

Merged
sergey-semenov merged 4 commits into from
Sep 16, 2025
Merged

[SYCL] Move Kernel specific data from handler_impl to a separate data structure #19843

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
75d73a2
Move Kernel specific data from handler_impl to a separate data struct…
Aug 20, 2025
c28ef2a
Fix backward compatibility test
Sep 11, 2025
55d8ba2
Rename setKernelInfo to setKernelFunc
Sep 15, 2025
9a32a2f
Returning assert to the DeviceKernelInfo::setCompileTimeInfoIfNeeded …
Sep 15, 2025
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
5 changes: 3 additions & 2 deletions sycl/include/sycl/ext/oneapi/experimental/work_group_memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,9 +18,10 @@

namespace sycl {
inline namespace _V1 {
class handler;

namespace detail {
class KernelData;

template <typename T> struct is_unbounded_array : std::false_type {};

template <typename T> struct is_unbounded_array<T[]> : std::true_type {};
Expand All @@ -38,7 +39,7 @@ class work_group_memory_impl {

private:
size_t buffer_size;
friend class sycl::handler;
friend class KernelData;
};

} // namespace detail
Expand Down
21 changes: 11 additions & 10 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -491,9 +491,7 @@ class __SYCL_EXPORT handler {
template <class Kernel> void setDeviceKernelInfo(void *KernelFuncPtr) {
constexpr auto Info = detail::CompileTimeKernelInfo<Kernel>;
MKernelName = Info.Name;
// TODO support ESIMD in no-integration-header case too.
setKernelInfo(KernelFuncPtr, Info.NumParams, Info.ParamDescGetter,
Info.IsESIMD, Info.HasSpecialCaptures);
setKernelFunc(KernelFuncPtr);
setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<Kernel>());
setType(detail::CGType::Kernel);
}
Expand All @@ -510,23 +508,21 @@ class __SYCL_EXPORT handler {
extractArgsAndReqsFromLambda(char *LambdaPtr, size_t KernelArgsNum,
const detail::kernel_param_desc_t *KernelArgs,
bool IsESIMD);
#endif
/// Extracts and prepares kernel arguments from the lambda using information
/// from the built-ins or integration header.
void extractArgsAndReqsFromLambda(
char *LambdaPtr, detail::kernel_param_desc_t (*ParamDescGetter)(int),
size_t NumKernelParams, bool IsESIMD);

#endif
/// Extracts and prepares kernel arguments set via set_arg(s).
void extractArgsAndReqs();

#if defined(__INTEL_PREVIEW_BREAKING_CHANGES)
// TODO: processArg need not to be public
__SYCL_DLL_LOCAL
#endif
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// TODO: remove in the next ABI-breaking window.
void processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
const int Size, const size_t Index, size_t &IndexShift,
bool IsKernelCreatedFromSource, bool IsESIMD);
#endif

/// \return a string containing name of SYCL kernel.
detail::ABINeutralKernelNameStrT getKernelName();
Expand Down Expand Up @@ -3600,7 +3596,10 @@ class __SYCL_EXPORT handler {

void addArg(detail::kernel_param_kind_t ArgKind, void *Req, int AccessTarget,
int ArgIndex);
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// TODO: remove in the next ABI-breaking window
void clearArgs();
#endif
void setArgsToAssociatedAccessors();

bool HasAssociatedAccessor(detail::AccessorImplHost *Req,
Expand Down Expand Up @@ -3647,10 +3646,12 @@ class __SYCL_EXPORT handler {
void setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::id<1> Offset);
void setNDRangeDescriptor(sycl::range<1> NumWorkItems,
sycl::range<1> LocalSize, sycl::id<1> Offset);

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
bool KernelIsESIMD, bool KernelHasSpecialCaptures);
#endif
void setKernelFunc(void *KernelFuncPtr);

void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr);

Expand Down
3 changes: 2 additions & 1 deletion sycl/include/sycl/stream.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,6 +42,7 @@ inline namespace _V1 {
namespace detail {

class stream_impl;
class KernelData;

using FmtFlags = unsigned int;

Expand Down Expand Up @@ -1041,7 +1042,7 @@ class __SYCL_EXPORT __SYCL_SPECIAL_CLASS __SYCL_TYPE(stream) stream
}
#endif

friend class handler;
friend class detail::KernelData;

template <typename SYCLObjT> friend class ext::oneapi::weak_object;

Expand Down
1 change: 1 addition & 0 deletions sycl/source/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -260,6 +260,7 @@ set(SYCL_COMMON_SOURCES
"detail/device_filter.cpp"
"detail/host_pipe_map.cpp"
"detail/device_global_map.cpp"
"detail/kernel_data.cpp"
"detail/kernel_global_info.cpp"
"detail/device_global_map_entry.cpp"
"detail/device_image_impl.cpp"
Expand Down
4 changes: 4 additions & 0 deletions sycl/source/detail/device_kernel_info.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,11 @@ void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
const CompileTimeKernelInfoTy &Info) {
if (!isCompileTimeInfoSet())
CompileTimeKernelInfoTy::operator=(Info);
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
// In case of 6.3 compatibility mode the KernelSize is not passed to the
// runtime. So, it will always be 0 and this assert fails.
assert(isCompileTimeInfoSet());
Copy link
Contributor Author

@vinser52 vinser52 Sep 15, 2025

Choose a reason for hiding this comment

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

@sergey-semenov I removed this assert because in case of app was compiled with the 6.3 compiler, the kernel size is not passed from the headers to the runtime.

Copy link
Contributor Author

@vinser52 vinser52 Sep 15, 2025

Choose a reason for hiding this comment

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

After discussion with @sergey-semenov, we decided to keep this assert in the preview mode.
Corresponding comment is added.

Copy link
Contributor

@aelovikov-intel aelovikov-intel Sep 16, 2025

Choose a reason for hiding this comment

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

This deserves more explanation (probably in the PR description). New KernelData is under source/detail and doesn't cross ABI boundary. What are we changing here that has ABI effects and why can't that be done in a separate PR?

Copy link
Contributor Author

@vinser52 vinser52 Sep 16, 2025

Choose a reason for hiding this comment

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

But I added comment in place. And the issue happens only in case of 6.3 or 6.2 headers. There is nothing about ABI here.

#endif
assert(Info == *this);
}

Expand Down
2 changes: 1 addition & 1 deletion sycl/source/detail/graph/dynamic_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -343,7 +343,7 @@ void dynamic_command_group_impl::finalizeCGFList(
MCommandGroups.push_back(std::shared_ptr<sycl::detail::CG>(RawCGPtr));

// Track dynamic_parameter usage in command-group
auto &DynamicParams = Handler.impl->MDynamicParameters;
auto &DynamicParams = Handler.impl->MKernelData.getDynamicParameters();

if (DynamicParams.size() > 0 &&
Handler.getType() == sycl::detail::CGType::CodeplayHostTask) {
Expand Down
2 changes: 1 addition & 1 deletion sycl/source/detail/graph/graph_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -471,7 +471,7 @@ node_impl &graph_impl::add(std::function<void(handler &)> CGF,

// Retrieve any dynamic parameters which have been registered in the CGF and
// register the actual nodes with them.
auto &DynamicParams = Handler.impl->MDynamicParameters;
auto &DynamicParams = Handler.impl->MKernelData.getDynamicParameters();

if (NodeType != node_type::kernel && DynamicParams.size() > 0) {
throw sycl::exception(sycl::make_error_code(errc::invalid),
Expand Down
32 changes: 7 additions & 25 deletions sycl/source/detail/handler_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,7 @@
#include "sycl/handler.hpp"
#include <detail/cg.hpp>
#include <detail/kernel_bundle_impl.hpp>
#include <detail/kernel_data.hpp>
#include <memory>
#include <sycl/ext/oneapi/experimental/enqueue_types.hpp>

Expand Down Expand Up @@ -61,8 +62,7 @@ class handler_impl {
}

KernelNameStrRefT getKernelName() const {
assert(MDeviceKernelInfoPtr);
return static_cast<KernelNameStrRefT>(MDeviceKernelInfoPtr->Name);
return MKernelData.getKernelName();
}

/// Registers mutually exclusive submission states.
Expand Down Expand Up @@ -108,12 +108,6 @@ class handler_impl {
// If the pipe operation is read or write, 1 for read 0 for write.
bool HostPipeRead = true;

ur_kernel_cache_config_t MKernelCacheConfig = UR_KERNEL_CACHE_CONFIG_DEFAULT;

bool MKernelIsCooperative = false;
bool MKernelUsesClusterLaunch = false;
uint32_t MKernelWorkGroupMemorySize = 0;

// Extra information for bindless image copy
ur_image_desc_t MSrcImageDesc = {};
ur_image_desc_t MDstImageDesc = {};
Expand All @@ -138,29 +132,17 @@ class handler_impl {
sycl::ext::oneapi::experimental::node_type MUserFacingNodeType =
sycl::ext::oneapi::experimental::node_type::empty;

// Storage for any SYCL Graph dynamic parameters which have been flagged for
// registration in the CG, along with the argument index for the parameter.
std::vector<std::pair<
ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>
MDynamicParameters;

/// The storage for the arguments passed.
/// We need to store a copy of values that are passed explicitly through
/// set_arg, require and so on, because we need them to be alive after
/// we exit the method they are passed in.
detail::CG::StorageInitHelper CGData;

/// The list of arguments for the kernel.
std::vector<detail::ArgDesc> MArgs;

/// The list of associated accessors with this handler.
/// These accessors were created with this handler as argument or
/// have become required for this handler via require method.
std::vector<detail::ArgDesc> MAssociatedAccesors;

/// Struct that encodes global size, local size, ...
detail::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
/// manipulations with version are required
Expand Down Expand Up @@ -241,16 +223,16 @@ class handler_impl {
// Allocation ptr to be freed asynchronously.
void *MFreePtr = nullptr;

// Store information about the kernel arguments.
void *MKernelFuncPtr = nullptr;
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
// TODO: remove in the next ABI-breaking window
// Today they are used only in the handler::setKernelNameBasedCachePtr
Comment on lines +227 to +228
Copy link
Contributor

@aelovikov-intel aelovikov-intel Sep 16, 2025

Choose a reason for hiding this comment

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

This makes little sense to me. Everything under source/detail/ is inside the libsycl.so and doesn't cross ABI boundary.

Copy link
Contributor Author

@vinser52 vinser52 Sep 16, 2025

Choose a reason for hiding this comment

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

These vars are used only when app is compiled with 6.3 compiler. In case of 6.3 compiler StoreLambda function calls setKerneInfo() that accept the compiler info from the integration header. After that StoreLambda calls setDeviceKernelInfoPtr. The problem with compatibility with 6.3 is that we need to store the data from integration header somewhere because DeviceKernelInfoPtr is not set yet.

You are right that this members does not cross ABI boundaries, but they are used only by the API that is under #ifndef __INTEL_PREVIEW_BREAKING_CHANGES macro. So I put them under preview, so that we don’t forget to remove them together with corresponding API in the next ABI-breaking window.

int MKernelNumArgs = 0;
detail::kernel_param_desc_t (*MKernelParamDescGetter)(int) = nullptr;
bool MKernelIsESIMD = false;
bool MKernelHasSpecialCaptures = true;
#endif

// A pointer to device kernel information. Cached on the application side in
// headers or retrieved from program manager.
DeviceKernelInfo *MDeviceKernelInfoPtr = nullptr;
KernelData MKernelData;
};

} // namespace detail
Expand Down
Loading