Skip to content

Commit 3cba04f

Browse files
aobolenskaobolensk
committed
Fix deprecation macro and add deprecation tests
1 parent 5af7ed6 commit 3cba04f

File tree

2 files changed

+89
-73
lines changed

2 files changed

+89
-73
lines changed

‎sycl/include/CL/sycl/group.hpp

Lines changed: 80 additions & 73 deletions
Original file line numberDiff line numberDiff line change
@@ -104,10 +104,10 @@ template <int Dimensions = 1> class group {
104104

105105
group() = delete;
106106

107-
__SYCL_DEPRECATED("use sycl::group::get_group_id() instead")
107+
__SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead")
108108
id<Dimensions> get_id() const { return index; }
109109

110-
__SYCL_DEPRECATED("use sycl::group::get_group_id() instead")
110+
__SYCL2020_DEPRECATED("use sycl::group::get_group_id() instead")
111111
size_t get_id(int dimension) const { return index[dimension]; }
112112

113113
id<Dimensions> get_group_id() const { return index; }
@@ -141,25 +141,8 @@ template <int Dimensions = 1> class group {
141141

142142
size_t get_local_range(int dimension) const { return localRange[dimension]; }
143143

144-
template <int dims = Dimensions>
145-
typename detail::enable_if_t<(dims == 1), size_t>
146-
get_local_linear_range() const {
147-
auto localRange = get_local_range();
148-
return localRange[0];
149-
}
150-
151-
template <int dims = Dimensions>
152-
typename detail::enable_if_t<(dims == 2), size_t>
153-
get_local_linear_range() const {
154-
auto localRange = get_local_range();
155-
return localRange[0] * localRange[1];
156-
}
157-
158-
template <int dims = Dimensions>
159-
typename detail::enable_if_t<(dims == 3), size_t>
160-
get_local_linear_range() const {
161-
auto localRange = get_local_range();
162-
return localRange[0] * localRange[1] * localRange[2];
144+
size_t get_local_linear_range() const {
145+
return get_local_linear_range_impl();
163146
}
164147

165148
range<Dimensions> get_group_range() const { return groupRange; }
@@ -168,65 +151,18 @@ template <int Dimensions = 1> class group {
168151
return get_group_range()[dimension];
169152
}
170153

171-
template <int dims = Dimensions>
172-
typename detail::enable_if_t<(dims == 1), size_t>
173-
get_group_linear_range() const {
174-
auto groupRange = get_group_range();
175-
return groupRange[0];
176-
}
177-
178-
template <int dims = Dimensions>
179-
typename detail::enable_if_t<(dims == 2), size_t>
180-
get_group_linear_range() const {
181-
auto groupRange = get_group_range();
182-
return groupRange[0] * groupRange[1];
183-
}
184-
185-
template <int dims = Dimensions>
186-
typename detail::enable_if_t<(dims == 3), size_t>
187-
get_group_linear_range() const {
188-
auto groupRange = get_group_range();
189-
return groupRange[0] * groupRange[1] * groupRange[2];
154+
size_t get_group_linear_range() const {
155+
return get_group_linear_range_impl();
190156
}
191157

192158
range<Dimensions> get_max_local_range() const { return get_local_range(); }
193159

194160
size_t operator[](int dimension) const { return index[dimension]; }
195161

196-
template <int dims = Dimensions>
197-
__SYCL_DEPRECATED("use sycl::group::get_group_linear_id() instead")
198-
size_t get_linear_id() const {
199-
return get_group_linear_id<dims>();
200-
}
201-
202-
template <int dims = Dimensions>
203-
typename detail::enable_if_t<(dims == 1), size_t>
204-
get_group_linear_id() const {
205-
return index[0];
206-
}
162+
__SYCL2020_DEPRECATED("use sycl::group::get_group_linear_id() instead")
163+
size_t get_linear_id() const { return get_group_linear_id(); }
207164

208-
template <int dims = Dimensions>
209-
typename detail::enable_if_t<(dims == 2), size_t>
210-
get_group_linear_id() const {
211-
return index[0] * groupRange[1] + index[1];
212-
}
213-
214-
// SYCL specification 1.2.1rev5, section 4.7.6.5 "Buffer accessor":
215-
// Whenever a multi-dimensional index is passed to a SYCL accessor the
216-
// linear index is calculated based on the index {id1, id2, id3} provided
217-
// and the range of the SYCL accessor {r1, r2, r3} according to row-major
218-
// ordering as follows:
219-
// id3 + (id2 · r3) + (id1 · r3 · r2) (4.3)
220-
// section 4.8.1.8 "group class":
221-
// size_t get_linear_id()const
222-
// Get a linearized version of the work-group id. Calculating a linear
223-
// work-group id from a multi-dimensional index follows the equation 4.3.
224-
template <int dims = Dimensions>
225-
typename detail::enable_if_t<(dims == 3), size_t>
226-
get_group_linear_id() const {
227-
return (index[0] * groupRange[1] * groupRange[2]) +
228-
(index[1] * groupRange[2]) + index[2];
229-
}
165+
size_t get_group_linear_id() const { return get_group_linear_id_impl(); }
230166

231167
bool leader() const { return (get_local_linear_id() == 0); }
232168

@@ -497,6 +433,77 @@ template <int Dimensions = 1> class group {
497433
(localId[1] * groupRange[2]) + localId[2];
498434
}
499435

436+
template <int dims = Dimensions>
437+
typename detail::enable_if_t<(dims == 1), size_t>
438+
get_local_linear_range_impl() const {
439+
auto localRange = get_local_range();
440+
return localRange[0];
441+
}
442+
443+
template <int dims = Dimensions>
444+
typename detail::enable_if_t<(dims == 2), size_t>
445+
get_local_linear_range_impl() const {
446+
auto localRange = get_local_range();
447+
return localRange[0] * localRange[1];
448+
}
449+
450+
template <int dims = Dimensions>
451+
typename detail::enable_if_t<(dims == 3), size_t>
452+
get_local_linear_range_impl() const {
453+
auto localRange = get_local_range();
454+
return localRange[0] * localRange[1] * localRange[2];
455+
}
456+
457+
template <int dims = Dimensions>
458+
typename detail::enable_if_t<(dims == 1), size_t>
459+
get_group_linear_range_impl() const {
460+
auto groupRange = get_group_range();
461+
return groupRange[0];
462+
}
463+
464+
template <int dims = Dimensions>
465+
typename detail::enable_if_t<(dims == 2), size_t>
466+
get_group_linear_range_impl() const {
467+
auto groupRange = get_group_range();
468+
return groupRange[0] * groupRange[1];
469+
}
470+
471+
template <int dims = Dimensions>
472+
typename detail::enable_if_t<(dims == 3), size_t>
473+
get_group_linear_range_impl() const {
474+
auto groupRange = get_group_range();
475+
return groupRange[0] * groupRange[1] * groupRange[2];
476+
}
477+
478+
template <int dims = Dimensions>
479+
typename detail::enable_if_t<(dims == 1), size_t>
480+
get_group_linear_id_impl() const {
481+
return index[0];
482+
}
483+
484+
template <int dims = Dimensions>
485+
typename detail::enable_if_t<(dims == 2), size_t>
486+
get_group_linear_id_impl() const {
487+
return index[0] * groupRange[1] + index[1];
488+
}
489+
490+
// SYCL specification 1.2.1rev5, section 4.7.6.5 "Buffer accessor":
491+
// Whenever a multi-dimensional index is passed to a SYCL accessor the
492+
// linear index is calculated based on the index {id1, id2, id3} provided
493+
// and the range of the SYCL accessor {r1, r2, r3} according to row-major
494+
// ordering as follows:
495+
// id3 + (id2 · r3) + (id1 · r3 · r2) (4.3)
496+
// section 4.8.1.8 "group class":
497+
// size_t get_linear_id()const
498+
// Get a linearized version of the work-group id. Calculating a linear
499+
// work-group id from a multi-dimensional index follows the equation 4.3.
500+
template <int dims = Dimensions>
501+
typename detail::enable_if_t<(dims == 3), size_t>
502+
get_group_linear_id_impl() const {
503+
return (index[0] * groupRange[1] * groupRange[2]) +
504+
(index[1] * groupRange[2]) + index[2];
505+
}
506+
500507
void waitForHelper() const {}
501508

502509
void waitForHelper(device_event Event) const {

‎sycl/test/warnings/sycl_2020_deprecations.cpp

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -179,5 +179,14 @@ int main() {
179179
// expected-warning@+1 {{'atomic<int, sycl::access::address_space::global_space>' is deprecated: sycl::atomic is deprecated since SYCL 2020}}
180180
cl::sycl::atomic<int> b(a);
181181

182+
cl::sycl::group<1> group =
183+
cl::sycl::detail::Builder::createGroup<1>({8}, {4}, {1});
184+
// expected-warning@+1{{'get_id' is deprecated: use sycl::group::get_group_id() instead}}
185+
group.get_id();
186+
// expected-warning@+1{{'get_id' is deprecated: use sycl::group::get_group_id() instead}}
187+
group.get_id(1);
188+
// expected-warning@+1{{'get_linear_id' is deprecated: use sycl::group::get_group_linear_id() instead}}
189+
group.get_linear_id();
190+
182191
return 0;
183192
}

0 commit comments

Comments
 (0)