-
Notifications
You must be signed in to change notification settings - Fork 796
[SYCL][NFC] Move kernel launch property parsing out of sycl::handler
#19589
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.
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
base: sycl
Are you sure you want to change the base?
Conversation
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Pull Request Overview
This PR refactors kernel launch property parsing to move it out of the sycl::handler
class into a separate helper structure for better code organization and reusability. The motivation is to enable reuse of kernel launch property parsing code in no-handler submission paths.
Key Changes:
- Creates
KernelLaunchPropertyWrapper
struct to encapsulate property parsing logic - Converts kernel launch properties to use
std::optional
to track if properties have been set - Updates all references to kernel properties to use the new structure with dereferencing
Reviewed Changes
Copilot reviewed 10 out of 10 changed files in this pull request and generated 3 comments.
Show a summary per file
File | Description |
---|---|
sycl/include/sycl/detail/kernel_launch_helper.hpp | Implements new KernelLaunchPropertyWrapper with property parsing methods |
sycl/include/sycl/handler.hpp | Updates handler to use new property wrapper and adds helper methods |
sycl/source/detail/handler_impl.hpp | Makes handler_impl inherit from KernelLaunchPropertyWrapper |
sycl/source/handler.cpp | Updates property handling to use new structure and adds property setter method |
sycl/unittests/scheduler/*.cpp | Updates mock handlers to dereference optional kernel properties |
sycl/test/virtual-functions/properties-negative.cpp | Updates expected error location from handler.hpp to kernel_launch_helper.hpp |
sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp | Updates expected error location |
sycl/test/include_deps/sycl_detail_core.hpp.cpp | Adjusts include dependency ordering |
sycl/unittests/program_manager/arg_mask/EliminatedArgMask.cpp | Updates mock handler to use dereferenced properties |
MKernelClusterSize.fill(0); | ||
} | ||
|
||
KernelLaunchPropertiesT &operator=(const KernelLaunchPropertiesT &Other) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
The custom assignment operator should have a corresponding custom copy constructor, destructor, and move operations following the Rule of Five, or use the default implementations if possible.
Copilot uses AI. Check for mistakes.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yeah, why not = default;
?
typename KernelType> | ||
static void parseProperties([[maybe_unused]] PropertyProcessor h, | ||
[[maybe_unused]] const KernelType &KernelFunc) { | ||
// Changing values in this will break ABI/API. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Needs a test under sycl/test/abi/**
then.
// In some code paths, kernel launch properties are set multiple times | ||
// for the same kernel, i.e why using std::optional to avoid overriding | ||
// previously set properties. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
// In some code paths, kernel launch properties are set multiple times | |
// for the same kernel, i.e why using std::optional to avoid overriding | |
// previously set properties. | |
// In some code paths, kernel launch properties are set multiple times | |
// for the same kernel, that is why using std::optional to avoid overriding | |
// previously set properties. |
However, why not fix multiple parsing instead?
MKernelClusterSize.fill(0); | ||
} | ||
|
||
KernelLaunchPropertiesT &operator=(const KernelLaunchPropertiesT &Other) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yeah, why not = default;
?
KLProps.MKernelCacheConfig = UR_KERNEL_CACHE_CONFIG_DEFAULT; | ||
KLProps.MKernelIsCooperative = false; | ||
KLProps.MKernelWorkGroupMemorySize = 0; | ||
KLProps.MKernelUsesClusterLaunch = false; | ||
KLProps.MKernelClusterDims = 0; | ||
KLProps.MKernelClusterSize.fill(0); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
MKernel*
made sense when it was part of handler
, now that the class itself is specific for the kernel properties, you can name properties without "Kernel" in them, e.g. MIsCooperative
.
|
||
template <sycl::ext::oneapi::experimental::detail::UnsupportedGraphFeatures | ||
FeatureT> | ||
static void throwIfGraphAssociated(bool HasGraph) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
This looks ugly now, but I'm not yet sure what would be a better approach.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I can try inline throwIfGraphAssociated
in processLaunchProperties
as it's just used twice.
/// Process runtime kernel properties. | ||
/// | ||
/// Stores information about kernel properties into the handler. | ||
template <typename PropertiesT, typename PropertyProcessorPtrT> |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
PropertyProcessorPtrT
needs more documentation in the comment above. What you could also do is
template <typename Self>
struct PropertiesProcessor {
void foo() {
using self = static_cast<Self *>(this);
self->bar();
}
};
struct handler : public PropertiesProcessor<handler> { ... };
IIUC what you're trying to achieve here.
template < | ||
bool IsESIMDKernel, typename PropertyProcessorPtrT, | ||
typename PropertiesT = ext::oneapi::experimental::empty_properties_t> | ||
static KernelLaunchPropertiesT processProperties(PropertiesT Props, |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I was thinking about inlining launch properties processing before, but it makes even less sense to have them separate now.
@@ -1098,6 +1100,8 @@ class __SYCL_EXPORT handler { | |||
// supported gcc version is bumped. | |||
std::tuple<std::array<size_t, 3>, bool> getMaxWorkGroups_v2(); | |||
|
|||
detail::device_impl &getDeviceImpl(); |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I wanted to add that when working with postProcess
refactoring, but ultimately decided it was a wrong abstraction. Wouldn't be surprised if the same is true here.
if (auto prop = | ||
detail::KernelLaunchPropertyWrapper::parseProperties<NameT>( | ||
KernelFunc, this)) { | ||
setKernelLaunchProperties(*prop); | ||
} |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
optional
part needs a dedicated comment (near parseProperties
, I'd think). Is that a performance optimization?
@@ -3483,6 +3516,7 @@ class __SYCL_EXPORT handler { | |||
bool IsDeviceImageScoped, size_t NumBytes, | |||
size_t Offset); | |||
|
|||
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Please make sure to have a local e2e run with split build/run mode where build is done with the sycl-toolchain
built before this PR and run is using libsycl.so with this PR's changes.
Motivation is to decouple kernel launch property parsing from
sycl::handler
so that we can reuse this code for no-handler submission path.