Skip to content

Commit a7f5e93

Browse files
authored
[SYCL] Optimize handler::setNDRangeDescriptor (#17969)
1 parent 08d11bc commit a7f5e93

File tree

4 files changed

+47
-8
lines changed

4 files changed

+47
-8
lines changed

sycl/include/sycl/detail/array.hpp

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -58,6 +58,10 @@ template <int dimensions = 1> class array {
5858
return result;
5959
}
6060

61+
void reset(size_t value) {
62+
std::fill(common_array, common_array + dimensions, value);
63+
}
64+
6165
size_t get(int dimension) const {
6266
check_dimension(dimension);
6367
return common_array[dimension];

sycl/include/sycl/handler.hpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3718,7 +3718,8 @@ class __SYCL_EXPORT handler {
37183718
bool HasAssociatedAccessor(detail::AccessorImplHost *Req,
37193719
access::target AccessTarget) const;
37203720

3721-
template <int Dims> static sycl::range<3> padRange(sycl::range<Dims> Range) {
3721+
template <int Dims>
3722+
static sycl::range<3> padRange(const sycl::range<Dims> &Range) {
37223723
if constexpr (Dims == 3) {
37233724
return Range;
37243725
} else {
@@ -3729,7 +3730,7 @@ class __SYCL_EXPORT handler {
37293730
}
37303731
}
37313732

3732-
template <int Dims> static sycl::id<3> padId(sycl::id<Dims> Id) {
3733+
template <int Dims> static sycl::id<3> padId(const sycl::id<Dims> &Id) {
37333734
if constexpr (Dims == 3) {
37343735
return Id;
37353736
} else {

sycl/source/detail/cg.hpp

Lines changed: 37 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -108,9 +108,6 @@ class NDRDescT {
108108
setNDRangeLeftover();
109109
}
110110

111-
NDRDescT(sycl::range<3> NumWorkItems, sycl::id<3> Offset, int DimsArg)
112-
: GlobalSize{NumWorkItems}, GlobalOffset{Offset}, Dims{size_t(DimsArg)} {}
113-
114111
NDRDescT(sycl::range<3> NumWorkItems, sycl::range<3> LocalSize,
115112
sycl::id<3> Offset, int DimsArg)
116113
: GlobalSize{NumWorkItems}, LocalSize{LocalSize}, GlobalOffset{Offset},
@@ -144,6 +141,43 @@ class NDRDescT {
144141
ClusterDimensions[I] = (I < Dims) ? N[I] : 1;
145142
}
146143

144+
void reset(const sycl::range<3> &N, bool SetNumWorkGroups, int DimsArg) {
145+
if (SetNumWorkGroups) {
146+
GlobalSize.reset(0);
147+
NumWorkGroups = N;
148+
} else {
149+
GlobalSize = N;
150+
NumWorkGroups.reset(0);
151+
}
152+
LocalSize.reset(0);
153+
GlobalOffset.reset(0);
154+
ClusterDimensions.reset(1);
155+
Dims = size_t(DimsArg);
156+
setNDRangeLeftover();
157+
}
158+
159+
void reset(const sycl::range<3> &NumWorkItems, const sycl::id<3> &Offset,
160+
int DimsArg) {
161+
GlobalSize = NumWorkItems;
162+
GlobalOffset = Offset;
163+
NumWorkGroups.reset(0);
164+
LocalSize.reset(0);
165+
ClusterDimensions.reset(1);
166+
Dims = size_t(DimsArg);
167+
}
168+
169+
void reset(const sycl::range<3> &NumWorkItems,
170+
const sycl::range<3> &LocalSizeArg, const sycl::id<3> &Offset,
171+
int DimsArg) {
172+
GlobalSize = NumWorkItems;
173+
LocalSize = LocalSizeArg;
174+
GlobalOffset = Offset;
175+
NumWorkGroups.reset(0);
176+
ClusterDimensions.reset(1);
177+
Dims = size_t(DimsArg);
178+
setNDRangeLeftover();
179+
}
180+
147181
NDRDescT &operator=(const NDRDescT &Desc) = default;
148182
NDRDescT &operator=(NDRDescT &&Desc) = default;
149183

sycl/source/handler.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2218,16 +2218,16 @@ sycl::detail::CGType handler::getType() const { return impl->MCGType; }
22182218

22192219
void handler::setNDRangeDescriptorPadded(sycl::range<3> N,
22202220
bool SetNumWorkGroups, int Dims) {
2221-
impl->MNDRDesc = NDRDescT{N, SetNumWorkGroups, Dims};
2221+
impl->MNDRDesc.reset(N, SetNumWorkGroups, Dims);
22222222
}
22232223
void handler::setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems,
22242224
sycl::id<3> Offset, int Dims) {
2225-
impl->MNDRDesc = NDRDescT{NumWorkItems, Offset, Dims};
2225+
impl->MNDRDesc.reset(NumWorkItems, Offset, Dims);
22262226
}
22272227
void handler::setNDRangeDescriptorPadded(sycl::range<3> NumWorkItems,
22282228
sycl::range<3> LocalSize,
22292229
sycl::id<3> Offset, int Dims) {
2230-
impl->MNDRDesc = NDRDescT{NumWorkItems, LocalSize, Offset, Dims};
2230+
impl->MNDRDesc.reset(NumWorkItems, LocalSize, Offset, Dims);
22312231
}
22322232

22332233
void handler::saveCodeLoc(detail::code_location CodeLoc, bool IsTopCodeLoc) {

0 commit comments

Comments
 (0)