Skip to content

[SYCL] Optimize NDRDescT by removing sycl::range, sycl::id and padding #18851

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
merged 20 commits into from
Jul 1, 2025
Merged

[SYCL] Optimize NDRDescT by removing sycl::range, sycl::id and padding #18851

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
20 commits
Select commit Hold shift + click to select a range
b959782
[SYCL] Optimize NDRDescT by removing sycl::range, sycl::id and padding
DBDuncan May 30, 2025
520e446
Format code
DBDuncan Jun 6, 2025
907717c
Improve modification of NDRDescT in adjustNDRangePerKernel
DBDuncan Jun 11, 2025
adafe3d
Fix bug when setting LocalSize by preserving old behaviour of setting…
DBDuncan Jun 11, 2025
ef58ba7
Format and remove mistakenly committed code
DBDuncan Jun 11, 2025
86d7783
Merge remote-tracking branch 'origin/sycl' into duncan/ndrange-perf-fix
DBDuncan Jun 12, 2025
7d4175f
Fix issues with .size() being called on std::array when previously wa…
DBDuncan Jun 12, 2025
4fe9507
swap int with size_t
DBDuncan Jun 12, 2025
11fdc89
Set GlobalRange default value to 1
DBDuncan Jun 12, 2025
19e8982
Preserve previous behaviour to get HierPar/hier_par_basic.cpp to pass…
DBDuncan Jun 13, 2025
73b8e4d
Preserve old behaviour of GlobalSize being set to zero when default c…
DBDuncan Jun 13, 2025
889b4d7
Remove commented out code
DBDuncan Jun 13, 2025
9964dcf
remove setting extra global size dims to 1 when using SetNumWorkGroup…
DBDuncan Jun 16, 2025
6c51413
Reintroduce setting extra global size dims to 1 only when SetNumWorkG…
DBDuncan Jun 16, 2025
0c62f94
Merge remote-tracking branch 'origin/sycl' into duncan/ndrange-perf-fix
DBDuncan Jun 18, 2025
9e879fa
Fix formatting
DBDuncan Jun 18, 2025
ace2ae2
Update TODO text
DBDuncan Jun 18, 2025
abbeed9
Fix spelling error in comment
DBDuncan Jun 18, 2025
e084cb4
Merge remote-tracking branch 'origin/sycl' into duncan/ndrange-perf-fix
DBDuncan Jun 26, 2025
b04e914
Revert swapping throw within setClusterDimensions with assert
DBDuncan Jun 26, 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
65 changes: 25 additions & 40 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -913,7 +913,7 @@ class __SYCL_EXPORT handler {
.template get_property<
syclex::cuda::cluster_size_key<ClusterDim>>()
.get_cluster_size();
setKernelClusterLaunch(padRange(ClusterSize), ClusterDim);
setKernelClusterLaunch(ClusterSize);
}
}

Expand Down Expand Up @@ -3718,7 +3718,9 @@ class __SYCL_EXPORT handler {
void setKernelIsCooperative(bool);

// Set using cuda thread block cluster launch flag and set the launch bounds.
void setKernelClusterLaunch(sycl::range<3> ClusterSize, int Dims);
void setKernelClusterLaunch(sycl::range<3> ClusterSize);
void setKernelClusterLaunch(sycl::range<2> ClusterSize);
void setKernelClusterLaunch(sycl::range<1> ClusterSize);

// Set the request work group memory size (work_group_static ext).
void setKernelWorkGroupMem(size_t Size);
Expand Down Expand Up @@ -3819,54 +3821,37 @@ class __SYCL_EXPORT handler {
bool HasAssociatedAccessor(detail::AccessorImplHost *Req,
access::target AccessTarget) const;

template <int Dims> static sycl::range<3> padRange(sycl::range<Dims> Range) {
if constexpr (Dims == 3) {
return Range;
} else {
sycl::range<3> Res{0, 0, 0};
for (int I = 0; I < Dims; ++I)
Res[I] = Range[I];
return Res;
}
}

template <int Dims> static sycl::id<3> padId(sycl::id<Dims> Id) {
if constexpr (Dims == 3) {
return Id;
} else {
sycl::id<3> Res{0, 0, 0};
for (int I = 0; I < Dims; ++I)
Res[I] = Id[I];
return Res;
}
}

template <int Dims>
void setNDRangeDescriptor(sycl::range<Dims> N,
bool SetNumWorkGroups = false) {
return setNDRangeDescriptorPadded(padRange(N), SetNumWorkGroups, Dims);
return setNDRangeDescriptor(N, SetNumWorkGroups);
}
template <int Dims>
void setNDRangeDescriptor(sycl::range<Dims> NumWorkItems,
sycl::id<Dims> Offset) {
return setNDRangeDescriptorPadded(padRange(NumWorkItems), padId(Offset),
Dims);
return setNDRangeDescriptor(NumWorkItems, Offset);
}
template <int Dims>
void setNDRangeDescriptor(sycl::nd_range<Dims> ExecutionRange) {
return setNDRangeDescriptorPadded(
padRange(ExecutionRange.get_global_range()),
padRange(ExecutionRange.get_local_range()),
padId(ExecutionRange.get_offset()), Dims);
}

void setNDRangeDescriptorPadded(sycl::range<3> N, bool SetNumWorkGroups,
int Dims);
void setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems,
sycl::id<3> Offset, int Dims);
void setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems,
sycl::range<3> LocalSize, sycl::id<3> Offset,
int Dims);
return setNDRangeDescriptor(ExecutionRange.get_global_range(),
ExecutionRange.get_local_range(),
ExecutionRange.get_offset());
}

void setNDRangeDescriptor(sycl::range<3> N, bool SetNumWorkGroups);
void setNDRangeDescriptor(sycl::range<3> NumWorkItems, sycl::id<3> Offset);
void setNDRangeDescriptor(sycl::range<3> NumWorkItems,
sycl::range<3> LocalSize, sycl::id<3> Offset);

void setNDRangeDescriptor(sycl::range<2> N, bool SetNumWorkGroups);
void setNDRangeDescriptor(sycl::range<2> NumWorkItems, sycl::id<2> Offset);
void setNDRangeDescriptor(sycl::range<2> NumWorkItems,
sycl::range<2> LocalSize, sycl::id<2> Offset);

void setNDRangeDescriptor(sycl::range<1> N, bool SetNumWorkGroups);
void setNDRangeDescriptor(sycl::range<1> NumWorkItems, sycl::id<1> Offset);
void setNDRangeDescriptor(sycl::range<1> NumWorkItems,
sycl::range<1> LocalSize, sycl::id<1> Offset);
Comment on lines +3841 to +3854
Copy link
Contributor

@aelovikov-intel aelovikov-intel Jun 30, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

Are these necessary to be exported? Both "new" and "old".

Copy link
Contributor Author

Choose a reason for hiding this comment

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

So these functions are used in handler.cpp, handler.hpp and also SchedulerTestUtils.hpp.

There does not seem to be a straight forward way to stop these functions from being exported?
After chatting with a co-worker, I don't think there is an easy way to fix this. There are attributes to set visibility, but I strongly suspect doing so will break things. Complicated by the fact that setNDRangeDescriptor is used in SchedulerTestUtils.hpp.

I would have expected that only functions that are marked as __SYCL_EXPORT would actually be exported, but that appears to not matter? sycl_symbols_linux.dump contains loads of symbols that are in the detail namespace and should not be used directly by the user. Windows and linux have different visibility defaults but it appears both have been set to export all.


void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
detail::kernel_param_desc_t (*KernelParamDescGetter)(int),
Expand Down
119 changes: 58 additions & 61 deletions sycl/source/detail/cg.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,99 +62,96 @@ class ArgDesc {

// The structure represents NDRange - global, local sizes, global offset and
// number of dimensions.
class NDRDescT {
// The method initializes all sizes for dimensions greater than the passed one
// to the default values, so they will not affect execution.
void setNDRangeLeftover() {
for (int I = Dims; I < 3; ++I) {
GlobalSize[I] = 1;
LocalSize[I] = LocalSize[0] ? 1 : 0;
GlobalOffset[I] = 0;
NumWorkGroups[I] = 0;
}
}

template <int Dims> static sycl::range<3> padRange(sycl::range<Dims> Range) {
if constexpr (Dims == 3) {
return Range;
} else {
sycl::range<3> Res{0, 0, 0};
for (int I = 0; I < Dims; ++I)
Res[I] = Range[I];
return Res;
}
}

template <int Dims> static sycl::id<3> padId(sycl::id<Dims> Id) {
if constexpr (Dims == 3) {
return Id;
} else {
sycl::id<3> Res{0, 0, 0};
for (int I = 0; I < Dims; ++I)
Res[I] = Id[I];
return Res;
}
}
// TODO: A lot of tests rely on particular values to be set for dimensions that
// are not used. To clarify, for example, if a 2D kernel is invoked, in
// NDRDescT, the value of index 2 in GlobalSize must be set to either 1 or 0
// depending on which constructor is used for no clear reason.
// Instead, only sensible defaults should be used and tests should be updated
// to reflect this.
class NDRDescT {

public:
NDRDescT() = default;
NDRDescT(const NDRDescT &Desc) = default;
NDRDescT(NDRDescT &&Desc) = default;

NDRDescT(sycl::range<3> N, bool SetNumWorkGroups, int DimsArg)
: GlobalSize{SetNumWorkGroups ? sycl::range<3>{0, 0, 0} : N},
NumWorkGroups{SetNumWorkGroups ? N : sycl::range<3>{0, 0, 0}},
Dims{size_t(DimsArg)} {
setNDRangeLeftover();
}
template <int Dims_>
NDRDescT(sycl::range<Dims_> N, bool SetNumWorkGroups) : Dims{size_t(Dims_)} {
if (SetNumWorkGroups) {
for (size_t I = 0; I < Dims_; ++I) {
NumWorkGroups[I] = N[I];
}
} else {
for (size_t I = 0; I < Dims_; ++I) {
GlobalSize[I] = N[I];
}

NDRDescT(sycl::range<3> NumWorkItems, sycl::range<3> LocalSize,
sycl::id<3> Offset, int DimsArg)
: GlobalSize{NumWorkItems}, LocalSize{LocalSize}, GlobalOffset{Offset},
Dims{size_t(DimsArg)} {
setNDRangeLeftover();
for (int I = Dims_; I < 3; ++I) {
GlobalSize[I] = 1;
}
}
}

NDRDescT(sycl::range<3> NumWorkItems, sycl::id<3> Offset, int DimsArg)
: GlobalSize{NumWorkItems}, GlobalOffset{Offset}, Dims{size_t(DimsArg)} {}
template <int Dims_>
NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::range<Dims_> LocalSizes,
sycl::id<Dims_> Offset)
: Dims{size_t(Dims_)} {
for (size_t I = 0; I < Dims_; ++I) {
GlobalSize[I] = NumWorkItems[I];
LocalSize[I] = LocalSizes[I];
GlobalOffset[I] = Offset[I];
}

for (int I = Dims_; I < 3; ++I) {
LocalSize[I] = LocalSizes[0] ? 1 : 0;
Copy link
Contributor Author

Choose a reason for hiding this comment

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

There are a number of tests that depend on extra LocalSize dimensions higher than Dims_ being set to zero or one depending on whether LocalSizes[I] is zero or not respectively. RequiredWGSize.NoRequiredSize and RequiredWGSize.HasRequiredSize always fail if extra LocalSize dimensions are always set to 1 and various tests such as work_group_size_prop.cpp and six others fail if extra LocalSize dimensions are always set to zero. This preserves the old behaviour.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

It seems strange to me that this was introduced in the first place. It really should not matter what the value of dimensions higher than Dims_ are and should just be ignored. But now a number of tests depend on this behaviour.

Copy link
Contributor

Choose a reason for hiding this comment

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

Can we add a TODO to revisit this?

This sort of complexity will have a (small) impact on runtime, but it's also going to make it harder to make changes to NDRDescT later on. Making sure NDRDescT returns values we can't explain just to satisfy existing tests is one way to proceed -- but we could also look into whether those tests are actually useful, or rewrite them (and related functionality) to do the right thing.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I have added the TODO. I can also create a github issue if you think it is a good idea to keep track of this.

}

for (int I = Dims_; I < 3; ++I) {
GlobalSize[I] = 1;
}
}

template <int Dims_>
NDRDescT(sycl::nd_range<Dims_> ExecutionRange, int DimsArg)
: NDRDescT(padRange(ExecutionRange.get_global_range()),
padRange(ExecutionRange.get_local_range()),
padId(ExecutionRange.get_offset()), size_t(DimsArg)) {
setNDRangeLeftover();
NDRDescT(sycl::range<Dims_> NumWorkItems, sycl::id<Dims_> Offset)
: Dims{size_t(Dims_)} {
for (size_t I = 0; I < Dims_; ++I) {
GlobalSize[I] = NumWorkItems[I];
GlobalOffset[I] = Offset[I];
}
}

template <int Dims_>
NDRDescT(sycl::nd_range<Dims_> ExecutionRange)
: NDRDescT(ExecutionRange, Dims_) {}
: NDRDescT(ExecutionRange.get_global_range(),
ExecutionRange.get_local_range(),
ExecutionRange.get_offset()) {}

template <int Dims_>
NDRDescT(sycl::range<Dims_> Range)
: NDRDescT(padRange(Range), /*SetNumWorkGroups=*/false, Dims_) {}
: NDRDescT(Range, /*SetNumWorkGroups=*/false) {}

void setClusterDimensions(sycl::range<3> N, int Dims) {
if (this->Dims != size_t(Dims)) {
template <int Dims_> void setClusterDimensions(sycl::range<Dims_> N) {
if (this->Dims != size_t(Dims_)) {
throw std::runtime_error(
"Dimensionality of cluster, global and local ranges must be same");
}

for (int I = 0; I < 3; ++I)
ClusterDimensions[I] = (I < Dims) ? N[I] : 1;
for (int I = 0; I < Dims_; ++I)
ClusterDimensions[I] = N[I];
}

NDRDescT &operator=(const NDRDescT &Desc) = default;
NDRDescT &operator=(NDRDescT &&Desc) = default;

sycl::range<3> GlobalSize{0, 0, 0};
sycl::range<3> LocalSize{0, 0, 0};
sycl::id<3> GlobalOffset{0, 0, 0};
std::array<size_t, 3> GlobalSize{0, 0, 0};
std::array<size_t, 3> LocalSize{0, 0, 0};
std::array<size_t, 3> GlobalOffset{0, 0, 0};
/// Number of workgroups, used to record the number of workgroups from the
/// simplest form of parallel_for_work_group. If set, all other fields must be
/// zero
sycl::range<3> NumWorkGroups{0, 0, 0};
sycl::range<3> ClusterDimensions{1, 1, 1};
std::array<size_t, 3> NumWorkGroups{0, 0, 0};
std::array<size_t, 3> ClusterDimensions{1, 1, 1};
size_t Dims = 0;
};

Expand Down
7 changes: 5 additions & 2 deletions sycl/source/detail/scheduler/commands.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2276,8 +2276,11 @@ static void adjustNDRangePerKernel(NDRDescT &NDR, ur_kernel_handle_t Kernel,
if (WGSize[0] == 0) {
WGSize = {1, 1, 1};
}
NDR = sycl::detail::NDRDescT{nd_range<3>(NDR.NumWorkGroups * WGSize, WGSize),
static_cast<int>(NDR.Dims)};

for (size_t I = 0; I < NDR.Dims; ++I) {
NDR.GlobalSize[I] = WGSize[I] * NDR.NumWorkGroups[I];
NDR.LocalSize[I] = WGSize[I];
}
}

// We have the following mapping between dimensions with SPIR-V builtins:
Expand Down
Loading