Skip to content

Commit 560a214

Browse files
dsawferdsawfer
authored
[SYCL] Align id and range class implementation with SYCL specification (#4538)
The DPC++ "range" implementation has a conversion to "id", which is not in the spec, and "id" implementation has a conversion to "range", which is also not in the spec. As for operators: The 16 binary operators "+, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=" are defined as either member functions or hidden friend functions depending on the arguments. However, the spec defines these operators as hidden friend functions for all argument variations. The 16 binary operators are defined as templates that take any integral type, but the spec defines these operators only for "size_t". The 10 assignment operators "+=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=" are defined as member functions, but the spec defines them as hidden friend functions. In addition, DPC++ is missing the following hidden friend functions which are defined in the spec: Unary "+" and "-". Prefix "++" and "--" Postfix "++" and "--"
1 parent b79ae69 commit 560a214

File tree

5 files changed

+222
-34
lines changed

5 files changed

+222
-34
lines changed

‎sycl/include/CL/sycl/id.hpp

Lines changed: 62 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -92,6 +92,7 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
9292
id(ParamTy<N, 3, const item<dimensions, with_offset>> &item)
9393
: base(item.get_id(0), item.get_id(1), item.get_id(2)) {}
9494

95+
__SYCL_DEPRECATED("range() conversion is deprecated")
9596
explicit operator range<dimensions>() const {
9697
range<dimensions> result(
9798
detail::InitializedVal<dimensions, range>::template get<0>());
@@ -150,10 +151,11 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
150151

151152
// OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
152153
#define __SYCL_GEN_OPT_BASE(op) \
153-
id<dimensions> operator op(const id<dimensions> &rhs) const { \
154+
friend id<dimensions> operator op(const id<dimensions> &lhs, \
155+
const id<dimensions> &rhs) { \
154156
id<dimensions> result; \
155157
for (int i = 0; i < dimensions; ++i) { \
156-
result.common_array[i] = this->common_array[i] op rhs.common_array[i]; \
158+
result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
157159
} \
158160
return result; \
159161
}
@@ -163,10 +165,11 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
163165
#define __SYCL_GEN_OPT(op) \
164166
__SYCL_GEN_OPT_BASE(op) \
165167
template <typename T> \
166-
EnableIfIntegral<T, id<dimensions>> operator op(const T &rhs) const { \
168+
friend EnableIfIntegral<T, id<dimensions>> operator op( \
169+
const id<dimensions> &lhs, const T &rhs) { \
167170
id<dimensions> result; \
168171
for (int i = 0; i < dimensions; ++i) { \
169-
result.common_array[i] = this->common_array[i] op rhs; \
172+
result.common_array[i] = lhs.common_array[i] op rhs; \
170173
} \
171174
return result; \
172175
} \
@@ -182,10 +185,11 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
182185
#else
183186
#define __SYCL_GEN_OPT(op) \
184187
__SYCL_GEN_OPT_BASE(op) \
185-
id<dimensions> operator op(const size_t &rhs) const { \
188+
friend id<dimensions> operator op(const id<dimensions> &lhs, \
189+
const size_t &rhs) { \
186190
id<dimensions> result; \
187191
for (int i = 0; i < dimensions; ++i) { \
188-
result.common_array[i] = this->common_array[i] op rhs; \
192+
result.common_array[i] = lhs.common_array[i] op rhs; \
189193
} \
190194
return result; \
191195
} \
@@ -221,17 +225,18 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
221225

222226
// OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
223227
#define __SYCL_GEN_OPT(op) \
224-
id<dimensions> &operator op(const id<dimensions> &rhs) { \
228+
friend id<dimensions> &operator op(id<dimensions> &lhs, \
229+
const id<dimensions> &rhs) { \
225230
for (int i = 0; i < dimensions; ++i) { \
226-
this->common_array[i] op rhs.common_array[i]; \
231+
lhs.common_array[i] op rhs.common_array[i]; \
227232
} \
228-
return *this; \
233+
return lhs; \
229234
} \
230-
id<dimensions> &operator op(const size_t &rhs) { \
235+
friend id<dimensions> &operator op(id<dimensions> &lhs, const size_t &rhs) { \
231236
for (int i = 0; i < dimensions; ++i) { \
232-
this->common_array[i] op rhs; \
237+
lhs.common_array[i] op rhs; \
233238
} \
234-
return *this; \
239+
return lhs; \
235240
}
236241

237242
__SYCL_GEN_OPT(+=)
@@ -247,6 +252,51 @@ template <int dimensions = 1> class id : public detail::array<dimensions> {
247252

248253
#undef __SYCL_GEN_OPT
249254

255+
// OP is unary +, -
256+
#define __SYCL_GEN_OPT(op) \
257+
friend id<dimensions> operator op(const id<dimensions> &rhs) { \
258+
id<dimensions> result; \
259+
for (int i = 0; i < dimensions; ++i) { \
260+
result.common_array[i] = (op rhs.common_array[i]); \
261+
} \
262+
return result; \
263+
}
264+
265+
__SYCL_GEN_OPT(+)
266+
__SYCL_GEN_OPT(-)
267+
268+
#undef __SYCL_GEN_OPT
269+
270+
// OP is prefix ++, --
271+
#define __SYCL_GEN_OPT(op) \
272+
friend id<dimensions> &operator op(id<dimensions> &rhs) { \
273+
for (int i = 0; i < dimensions; ++i) { \
274+
op rhs.common_array[i]; \
275+
} \
276+
return rhs; \
277+
}
278+
279+
__SYCL_GEN_OPT(++)
280+
__SYCL_GEN_OPT(--)
281+
282+
#undef __SYCL_GEN_OPT
283+
284+
// OP is postfix ++, --
285+
#define __SYCL_GEN_OPT(op) \
286+
friend id<dimensions> operator op(id<dimensions> &lhs, int) { \
287+
id<dimensions> old_lhs; \
288+
for (int i = 0; i < dimensions; ++i) { \
289+
old_lhs.common_array[i] = lhs.common_array[i]; \
290+
op lhs.common_array[i]; \
291+
} \
292+
return old_lhs; \
293+
}
294+
295+
__SYCL_GEN_OPT(++)
296+
__SYCL_GEN_OPT(--)
297+
298+
#undef __SYCL_GEN_OPT
299+
250300
private:
251301
// Friend to get access to private method set_allowed_range().
252302
template <typename, int, typename> friend class detail::RoundedRangeKernel;

‎sycl/include/CL/sycl/range.hpp

Lines changed: 88 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -47,14 +47,6 @@ template <int dimensions = 1> class range : public detail::array<dimensions> {
4747
size_t dim2)
4848
: base(dim0, dim1, dim2) {}
4949

50-
explicit operator id<dimensions>() const {
51-
id<dimensions> result;
52-
for (int i = 0; i < dimensions; ++i) {
53-
result[i] = this->get(i);
54-
}
55-
return result;
56-
}
57-
5850
size_t size() const {
5951
size_t size = 1;
6052
for (int i = 0; i < dimensions; ++i) {
@@ -70,19 +62,26 @@ template <int dimensions = 1> class range : public detail::array<dimensions> {
7062
range() = delete;
7163

7264
// OP is: +, -, *, /, %, <<, >>, &, |, ^, &&, ||, <, >, <=, >=
73-
#define __SYCL_GEN_OPT(op) \
74-
range<dimensions> operator op(const range<dimensions> &rhs) const { \
75-
range<dimensions> result(*this); \
65+
#define __SYCL_GEN_OPT_BASE(op) \
66+
friend range<dimensions> operator op(const range<dimensions> &lhs, \
67+
const range<dimensions> &rhs) { \
68+
range<dimensions> result(lhs); \
7669
for (int i = 0; i < dimensions; ++i) { \
77-
result.common_array[i] = this->common_array[i] op rhs.common_array[i]; \
70+
result.common_array[i] = lhs.common_array[i] op rhs.common_array[i]; \
7871
} \
7972
return result; \
80-
} \
73+
}
74+
75+
#ifndef __SYCL_DISABLE_ID_TO_INT_CONV__
76+
// Enable operators with integral types only
77+
#define __SYCL_GEN_OPT(op) \
78+
__SYCL_GEN_OPT_BASE(op) \
8179
template <typename T> \
82-
IntegralType<T, range<dimensions>> operator op(const T &rhs) const { \
83-
range<dimensions> result(*this); \
80+
friend IntegralType<T, range<dimensions>> operator op( \
81+
const range<dimensions> &lhs, const T &rhs) { \
82+
range<dimensions> result(lhs); \
8483
for (int i = 0; i < dimensions; ++i) { \
85-
result.common_array[i] = this->common_array[i] op rhs; \
84+
result.common_array[i] = lhs.common_array[i] op rhs; \
8685
} \
8786
return result; \
8887
} \
@@ -95,6 +94,26 @@ template <int dimensions = 1> class range : public detail::array<dimensions> {
9594
} \
9695
return result; \
9796
}
97+
#else
98+
#define __SYCL_GEN_OPT(op) \
99+
__SYCL_GEN_OPT_BASE(op) \
100+
friend range<dimensions> operator op(const range<dimensions> &lhs, \
101+
const size_t &rhs) { \
102+
range<dimensions> result(lhs); \
103+
for (int i = 0; i < dimensions; ++i) { \
104+
result.common_array[i] = lhs.common_array[i] op rhs; \
105+
} \
106+
return result; \
107+
} \
108+
friend range<dimensions> operator op(const size_t &lhs, \
109+
const range<dimensions> &rhs) { \
110+
range<dimensions> result(rhs); \
111+
for (int i = 0; i < dimensions; ++i) { \
112+
result.common_array[i] = lhs op rhs.common_array[i]; \
113+
} \
114+
return result; \
115+
}
116+
#endif // __SYCL_DISABLE_ID_TO_INT_CONV__
98117

99118
__SYCL_GEN_OPT(+)
100119
__SYCL_GEN_OPT(-)
@@ -114,20 +133,23 @@ template <int dimensions = 1> class range : public detail::array<dimensions> {
114133
__SYCL_GEN_OPT(>=)
115134

116135
#undef __SYCL_GEN_OPT
136+
#undef __SYCL_GEN_OPT_BASE
117137

118138
// OP is: +=, -=, *=, /=, %=, <<=, >>=, &=, |=, ^=
119139
#define __SYCL_GEN_OPT(op) \
120-
range<dimensions> &operator op(const range<dimensions> &rhs) { \
140+
friend range<dimensions> &operator op(range<dimensions> &lhs, \
141+
const range<dimensions> &rhs) { \
121142
for (int i = 0; i < dimensions; ++i) { \
122-
this->common_array[i] op rhs[i]; \
143+
lhs.common_array[i] op rhs[i]; \
123144
} \
124-
return *this; \
145+
return lhs; \
125146
} \
126-
range<dimensions> &operator op(const size_t &rhs) { \
147+
friend range<dimensions> &operator op(range<dimensions> &lhs, \
148+
const size_t &rhs) { \
127149
for (int i = 0; i < dimensions; ++i) { \
128-
this->common_array[i] op rhs; \
150+
lhs.common_array[i] op rhs; \
129151
} \
130-
return *this; \
152+
return lhs; \
131153
}
132154

133155
__SYCL_GEN_OPT(+=)
@@ -143,6 +165,50 @@ template <int dimensions = 1> class range : public detail::array<dimensions> {
143165

144166
#undef __SYCL_GEN_OPT
145167

168+
// OP is unary +, -
169+
#define __SYCL_GEN_OPT(op) \
170+
friend range<dimensions> operator op(const range<dimensions> &rhs) { \
171+
range<dimensions> result(rhs); \
172+
for (int i = 0; i < dimensions; ++i) { \
173+
result.common_array[i] = (op rhs.common_array[i]); \
174+
} \
175+
return result; \
176+
}
177+
178+
__SYCL_GEN_OPT(+)
179+
__SYCL_GEN_OPT(-)
180+
181+
#undef __SYCL_GEN_OPT
182+
183+
// OP is prefix ++, --
184+
#define __SYCL_GEN_OPT(op) \
185+
friend range<dimensions> &operator op(range<dimensions> &rhs) { \
186+
for (int i = 0; i < dimensions; ++i) { \
187+
op rhs.common_array[i]; \
188+
} \
189+
return rhs; \
190+
}
191+
192+
__SYCL_GEN_OPT(++)
193+
__SYCL_GEN_OPT(--)
194+
195+
#undef __SYCL_GEN_OPT
196+
197+
// OP is postfix ++, --
198+
#define __SYCL_GEN_OPT(op) \
199+
friend range<dimensions> operator op(range<dimensions> &lhs, int) { \
200+
range<dimensions> old_lhs(lhs); \
201+
for (int i = 0; i < dimensions; ++i) { \
202+
op lhs.common_array[i]; \
203+
} \
204+
return old_lhs; \
205+
}
206+
207+
__SYCL_GEN_OPT(++)
208+
__SYCL_GEN_OPT(--)
209+
210+
#undef __SYCL_GEN_OPT
211+
146212
private:
147213
friend class handler;
148214
friend class detail::Builder;

‎sycl/test/basic_tests/id.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -439,4 +439,37 @@ int main() {
439439

440440
#undef numValue
441441
#endif // __SYCL_DISABLE_ID_TO_INT_CONV__
442+
443+
{
444+
cl::sycl::id<1> one_dim_id(64);
445+
cl::sycl::id<1> one_dim_id_neg(-64);
446+
cl::sycl::id<1> one_dim_id_copy(64);
447+
cl::sycl::id<2> two_dim_id(64, 1);
448+
cl::sycl::id<2> two_dim_id_neg(-64, -1);
449+
cl::sycl::id<2> two_dim_id_copy(64, 1);
450+
cl::sycl::id<3> three_dim_id(64, 1, 2);
451+
cl::sycl::id<3> three_dim_id_neg(-64, -1, -2);
452+
cl::sycl::id<3> three_dim_id_copy(64, 1, 2);
453+
454+
assert((+one_dim_id) == one_dim_id);
455+
assert(-one_dim_id == one_dim_id_neg);
456+
assert((+two_dim_id) == two_dim_id);
457+
assert(-two_dim_id == two_dim_id_neg);
458+
assert((+three_dim_id) == three_dim_id);
459+
assert(-three_dim_id == three_dim_id_neg);
460+
461+
assert((++one_dim_id) == (one_dim_id_copy + 1));
462+
assert((--one_dim_id) == (one_dim_id_copy));
463+
assert((++two_dim_id) == (two_dim_id_copy + 1));
464+
assert((--two_dim_id) == (two_dim_id_copy));
465+
assert((++three_dim_id) == (three_dim_id_copy + 1));
466+
assert((--three_dim_id) == (three_dim_id_copy));
467+
468+
assert((one_dim_id++) == (one_dim_id_copy));
469+
assert((one_dim_id--) == (one_dim_id_copy + 1));
470+
assert((two_dim_id++) == (two_dim_id_copy));
471+
assert((two_dim_id--) == (two_dim_id_copy + 1));
472+
assert((three_dim_id++) == (three_dim_id_copy));
473+
assert((three_dim_id--) == (three_dim_id_copy + 1));
474+
}
442475
}

‎sycl/test/basic_tests/range.cpp

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -34,4 +34,32 @@ int main() {
3434
assert(three_dim_range.get(2) ==2);
3535
assert(three_dim_range[2] ==2);
3636
cout << "three_dim_range passed " << endl;
37+
38+
cl::sycl::range<1> one_dim_range_neg(-64);
39+
cl::sycl::range<1> one_dim_range_copy(64);
40+
cl::sycl::range<2> two_dim_range_neg(-64, -1);
41+
cl::sycl::range<2> two_dim_range_copy(64, 1);
42+
cl::sycl::range<3> three_dim_range_copy(64, 1, 2);
43+
cl::sycl::range<3> three_dim_range_neg(-64, -1, -2);
44+
45+
assert((+one_dim_range) == one_dim_range);
46+
assert(-one_dim_range == one_dim_range_neg);
47+
assert((+two_dim_range) == two_dim_range);
48+
assert(-two_dim_range == two_dim_range_neg);
49+
assert((+three_dim_range) == three_dim_range);
50+
assert(-three_dim_range == three_dim_range_neg);
51+
52+
assert((++one_dim_range) == (one_dim_range_copy + 1));
53+
assert((--one_dim_range) == (one_dim_range_copy));
54+
assert((++two_dim_range) == (two_dim_range_copy + 1));
55+
assert((--two_dim_range) == (two_dim_range_copy));
56+
assert((++three_dim_range) == (three_dim_range_copy + 1));
57+
assert((--three_dim_range) == (three_dim_range_copy));
58+
59+
assert((one_dim_range++) == (one_dim_range_copy));
60+
assert((one_dim_range--) == (one_dim_range_copy + 1));
61+
assert((two_dim_range++) == (two_dim_range_copy));
62+
assert((two_dim_range--) == (two_dim_range_copy + 1));
63+
assert((three_dim_range++) == (three_dim_range_copy));
64+
assert((three_dim_range--) == (three_dim_range_copy + 1));
3765
}

‎sycl/test/warnings/conversion_deprecation.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
// RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
2+
3+
#include <CL/sycl.hpp>
4+
5+
int main() {
6+
cl::sycl::id<1> id_obj(64);
7+
// expected-warning@+1 {{'operator range' is deprecated: range() conversion is deprecated}}
8+
(void)(cl::sycl::range<1>) id_obj;
9+
10+
return 0;
11+
}

0 commit comments

Comments
 (0)