Skip to content

Commit 68b2d3e

Browse files
crystarmcrystarm
committed
[SYCL] Add noexcept and hidden friends to nd_item
1 parent 0b52e2e commit 68b2d3e

3 files changed

Lines changed: 183 additions & 37 deletions

File tree

‎sycl/include/sycl/nd_item.hpp‎

Lines changed: 42 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -47,21 +47,21 @@ template <int Dimensions = 1> class nd_item {
4747
public:
4848
static constexpr int dimensions = Dimensions;
4949

50-
id<Dimensions> get_global_id() const {
50+
id<Dimensions> get_global_id() const noexcept {
5151
#ifdef __SYCL_DEVICE_ONLY__
5252
return __spirv::initBuiltInGlobalInvocationId<Dimensions, id<Dimensions>>();
5353
#else
5454
return {};
5555
#endif
5656
}
5757

58-
size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const {
58+
size_t __SYCL_ALWAYS_INLINE get_global_id(int Dimension) const noexcept {
5959
size_t Id = get_global_id()[Dimension];
6060
__SYCL_ASSUME_ID_RANGE(Id);
6161
return Id;
6262
}
6363

64-
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const {
64+
size_t __SYCL_ALWAYS_INLINE get_global_linear_id() const noexcept {
6565
size_t LinId = 0;
6666
id<Dimensions> Index = get_global_id();
6767
range<Dimensions> Extent = get_global_range();
@@ -78,21 +78,21 @@ template <int Dimensions = 1> class nd_item {
7878
return LinId;
7979
}
8080

81-
id<Dimensions> get_local_id() const {
81+
id<Dimensions> get_local_id() const noexcept {
8282
#ifdef __SYCL_DEVICE_ONLY__
8383
return __spirv::initBuiltInLocalInvocationId<Dimensions, id<Dimensions>>();
8484
#else
8585
return {};
8686
#endif
8787
}
8888

89-
size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const {
89+
size_t __SYCL_ALWAYS_INLINE get_local_id(int Dimension) const noexcept {
9090
size_t Id = get_local_id()[Dimension];
9191
__SYCL_ASSUME_ID_RANGE(Id);
9292
return Id;
9393
}
9494

95-
size_t get_local_linear_id() const {
95+
size_t get_local_linear_id() const noexcept {
9696
size_t LinId = 0;
9797
id<Dimensions> Index = get_local_id();
9898
range<Dimensions> Extent = get_local_range();
@@ -108,23 +108,23 @@ template <int Dimensions = 1> class nd_item {
108108
return LinId;
109109
}
110110

111-
group<Dimensions> get_group() const {
111+
group<Dimensions> get_group() const noexcept {
112112
// TODO: ideally Group object should be stateless and have a contructor with
113113
// no arguments.
114114
return detail::Builder::createGroup(get_global_range(), get_local_range(),
115115
get_group_range(), get_group_id());
116116
}
117117

118118
// Out-of-class definition in sub_group.hpp
119-
sub_group get_sub_group() const;
119+
sub_group get_sub_group() const noexcept;
120120

121-
size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
121+
size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const noexcept {
122122
size_t Id = get_group_id()[Dimension];
123123
__SYCL_ASSUME_ID_RANGE(Id);
124124
return Id;
125125
}
126126

127-
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const {
127+
size_t __SYCL_ALWAYS_INLINE get_group_linear_id() const noexcept {
128128
size_t LinId = 0;
129129
id<Dimensions> Index = get_group_id();
130130
range<Dimensions> Extent = get_group_range();
@@ -140,66 +140,66 @@ template <int Dimensions = 1> class nd_item {
140140
return LinId;
141141
}
142142

143-
range<Dimensions> get_group_range() const {
143+
range<Dimensions> get_group_range() const noexcept {
144144
#ifdef __SYCL_DEVICE_ONLY__
145145
return __spirv::initBuiltInNumWorkgroups<Dimensions, range<Dimensions>>();
146146
#else
147147
return {};
148148
#endif
149149
}
150150

151-
size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const {
151+
size_t __SYCL_ALWAYS_INLINE get_group_range(int Dimension) const noexcept {
152152
size_t Range = get_group_range()[Dimension];
153153
__SYCL_ASSUME_ID_RANGE(Range);
154154
return Range;
155155
}
156156

157-
range<Dimensions> get_global_range() const {
157+
range<Dimensions> get_global_range() const noexcept {
158158
#ifdef __SYCL_DEVICE_ONLY__
159159
return __spirv::initBuiltInGlobalSize<Dimensions, range<Dimensions>>();
160160
#else
161161
return {};
162162
#endif
163163
}
164164

165-
size_t get_global_range(int Dimension) const {
165+
size_t get_global_range(int Dimension) const noexcept {
166166
size_t Val = get_global_range()[Dimension];
167167
__SYCL_ASSUME_ID_RANGE(Val);
168168
return Val;
169169
}
170170

171-
range<Dimensions> get_local_range() const {
171+
range<Dimensions> get_local_range() const noexcept {
172172
#ifdef __SYCL_DEVICE_ONLY__
173173
return __spirv::initBuiltInWorkgroupSize<Dimensions, range<Dimensions>>();
174174
#else
175175
return {};
176176
#endif
177177
}
178178

179-
size_t get_local_range(int Dimension) const {
179+
size_t get_local_range(int Dimension) const noexcept {
180180
size_t Id = get_local_range()[Dimension];
181181
__SYCL_ASSUME_ID_RANGE(Id);
182182
return Id;
183183
}
184184

185185
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
186-
id<Dimensions> get_offset() const {
186+
id<Dimensions> get_offset() const noexcept {
187187
#ifdef __SYCL_DEVICE_ONLY__
188188
return __spirv::initBuiltInGlobalOffset<Dimensions, id<Dimensions>>();
189189
#else
190190
return {};
191191
#endif
192192
}
193193

194-
nd_range<Dimensions> get_nd_range() const {
194+
nd_range<Dimensions> get_nd_range() const noexcept {
195195
return nd_range<Dimensions>(get_global_range(), get_local_range(),
196196
get_offset());
197197
}
198198

199199
#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
200200
__SYCL2020_DEPRECATED("use sycl::group_barrier() free function instead")
201201
void barrier([[maybe_unused]] access::fence_space accessSpace =
202-
access::fence_space::global_and_local) const {
202+
access::fence_space::global_and_local) const noexcept {
203203
#ifdef __SYCL_DEVICE_ONLY__
204204
uint32_t flags = _V1::detail::getSPIRVMemorySemanticsMask(accessSpace);
205205
__spirv_ControlBarrier(__spv::Scope::Workgroup, __spv::Scope::Workgroup,
@@ -217,7 +217,7 @@ template <int Dimensions = 1> class nd_item {
217217
accessMode == access::mode::write ||
218218
accessMode == access::mode::read_write,
219219
access::fence_space>
220-
accessSpace = access::fence_space::global_and_local) const {
220+
accessSpace = access::fence_space::global_and_local) const noexcept {
221221
#if __SYCL_DEVICE_ONLY__
222222
uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
223223
// TODO: currently, there is no good way in SPIR-V to set the memory
@@ -248,7 +248,7 @@ template <int Dimensions = 1> class nd_item {
248248
src,
249249
[[maybe_unused]] size_t numElements,
250250
[[maybe_unused]] size_t srcStride)
251-
const {
251+
const noexcept {
252252
#ifdef __SYCL_DEVICE_ONLY__
253253
__ocl_event_t E = __spirv_GroupAsyncCopy(
254254
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -274,7 +274,7 @@ template <int Dimensions = 1> class nd_item {
274274
[[maybe_unused]] local_ptr<dataT> src,
275275
[[maybe_unused]] size_t numElements,
276276
[[maybe_unused]] size_t destStride)
277-
const {
277+
const noexcept {
278278
#ifdef __SYCL_DEVICE_ONLY__
279279
__ocl_event_t E = __spirv_GroupAsyncCopy(
280280
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -299,7 +299,7 @@ template <int Dimensions = 1> class nd_item {
299299
async_work_group_copy([[maybe_unused]] decorated_local_ptr<DestDataT> dest,
300300
[[maybe_unused]] decorated_global_ptr<SrcDataT> src,
301301
[[maybe_unused]] size_t numElements,
302-
[[maybe_unused]] size_t srcStride) const {
302+
[[maybe_unused]] size_t srcStride) const noexcept {
303303
#ifdef __SYCL_DEVICE_ONLY__
304304
__ocl_event_t E = __spirv_GroupAsyncCopy(
305305
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -324,7 +324,7 @@ template <int Dimensions = 1> class nd_item {
324324
async_work_group_copy([[maybe_unused]] decorated_global_ptr<DestDataT> dest,
325325
[[maybe_unused]] decorated_local_ptr<SrcDataT> src,
326326
[[maybe_unused]] size_t numElements,
327-
[[maybe_unused]] size_t destStride) const {
327+
[[maybe_unused]] size_t destStride) const noexcept {
328328
#ifdef __SYCL_DEVICE_ONLY__
329329
__ocl_event_t E = __spirv_GroupAsyncCopy(
330330
__spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
@@ -352,7 +352,7 @@ template <int Dimensions = 1> class nd_item {
352352
access::decorated::legacy>
353353
Src,
354354
size_t NumElements,
355-
size_t Stride) const {
355+
size_t Stride) const noexcept {
356356
static_assert(sizeof(bool) == sizeof(uint8_t),
357357
"Async copy to/from bool memory is not supported.");
358358
auto DestP = multi_ptr<uint8_t, DestS, access::decorated::legacy>(
@@ -378,7 +378,7 @@ template <int Dimensions = 1> class nd_item {
378378
access::decorated::legacy>
379379
Src,
380380
size_t NumElements,
381-
size_t Stride) const {
381+
size_t Stride) const noexcept {
382382
static_assert(sizeof(bool) == sizeof(uint8_t),
383383
"Async copy to/from bool memory is not supported.");
384384
using VecT = detail::change_base_type_t<T, uint8_t>;
@@ -401,7 +401,7 @@ template <int Dimensions = 1> class nd_item {
401401
device_event>
402402
async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
403403
multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
404-
size_t NumElements, size_t Stride) const {
404+
size_t NumElements, size_t Stride) const noexcept {
405405
static_assert(sizeof(bool) == sizeof(uint8_t),
406406
"Async copy to/from bool memory is not supported.");
407407
using QualSrcT =
@@ -429,7 +429,7 @@ template <int Dimensions = 1> class nd_item {
429429
device_event>
430430
async_work_group_copy(multi_ptr<DestT, DestS, access::decorated::yes> Dest,
431431
multi_ptr<SrcT, SrcS, access::decorated::yes> Src,
432-
size_t NumElements, size_t Stride) const {
432+
size_t NumElements, size_t Stride) const noexcept {
433433
static_assert(sizeof(bool) == sizeof(uint8_t),
434434
"Async copy to/from bool memory is not supported.");
435435
using VecT = detail::change_base_type_t<DestT, uint8_t>;
@@ -455,7 +455,7 @@ template <int Dimensions = 1> class nd_item {
455455
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
456456
device_event
457457
async_work_group_copy(local_ptr<dataT> dest, global_ptr<dataT> src,
458-
size_t numElements) const {
458+
size_t numElements) const noexcept {
459459
return async_work_group_copy(dest, src, numElements, 1);
460460
}
461461

@@ -468,7 +468,7 @@ template <int Dimensions = 1> class nd_item {
468468
__SYCL2020_DEPRECATED("Use decorated multi_ptr arguments instead")
469469
device_event
470470
async_work_group_copy(global_ptr<dataT> dest, local_ptr<dataT> src,
471-
size_t numElements) const {
471+
size_t numElements) const noexcept {
472472
return async_work_group_copy(dest, src, numElements, 1);
473473
}
474474

@@ -483,7 +483,7 @@ template <int Dimensions = 1> class nd_item {
483483
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
484484
async_work_group_copy(decorated_local_ptr<DestDataT> dest,
485485
decorated_global_ptr<SrcDataT> src,
486-
size_t numElements) const {
486+
size_t numElements) const noexcept {
487487
return async_work_group_copy(dest, src, numElements, 1);
488488
}
489489

@@ -498,11 +498,12 @@ template <int Dimensions = 1> class nd_item {
498498
std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>, device_event>
499499
async_work_group_copy(decorated_global_ptr<DestDataT> dest,
500500
decorated_local_ptr<SrcDataT> src,
501-
size_t numElements) const {
501+
size_t numElements) const noexcept {
502502
return async_work_group_copy(dest, src, numElements, 1);
503503
}
504504

505-
template <typename... eventTN> void wait_for(eventTN... events) const {
505+
template <typename... eventTN>
506+
void wait_for(eventTN... events) const noexcept {
506507
(events.wait(), ...);
507508
}
508509

@@ -517,16 +518,21 @@ template <int Dimensions = 1> class nd_item {
517518
nd_item &operator=(const nd_item &rhs) = default;
518519
nd_item &operator=(nd_item &&rhs) = default;
519520

520-
bool operator==(const nd_item &) const { return true; }
521-
bool operator!=(const nd_item &rhs) const { return !((*this) == rhs); }
521+
friend bool operator==(const nd_item &, const nd_item &) noexcept {
522+
return true;
523+
}
524+
525+
friend bool operator!=(const nd_item &lhs, const nd_item &rhs) noexcept {
526+
return !(lhs == rhs);
527+
}
522528

523529
protected:
524530
friend class detail::Builder;
525531
nd_item() {}
526532
nd_item(const item<Dimensions, true> &, const item<Dimensions, false> &,
527533
const group<Dimensions> &) {}
528534

529-
id<Dimensions> get_group_id() const {
535+
id<Dimensions> get_group_id() const noexcept {
530536
#ifdef __SYCL_DEVICE_ONLY__
531537
return __spirv::initBuiltInWorkgroupId<Dimensions, id<Dimensions>>();
532538
#else

‎sycl/include/sycl/sub_group.hpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -197,7 +197,8 @@ struct sub_group {
197197
sub_group() = default;
198198
};
199199

200-
template <int Dimensions> sub_group nd_item<Dimensions>::get_sub_group() const {
200+
template <int Dimensions>
201+
sub_group nd_item<Dimensions>::get_sub_group() const noexcept {
201202
return sub_group();
202203
}
203204

0 commit comments

Comments
 (0)