Skip to content
This repository was archived by the owner on Jan 26, 2024. It is now read-only.

Commit 08ad294

Browse files
cjdbcjdb
authored
Applied SYCL 1.2.1 renamed member functions for nd_item and nd_range: (#81)
* `nd_item::get_local` -> `nd_item::get_local_id` * `nd_item::get_global` -> `nd_item::get_global_id` * `nd_range::get_local` -> `nd_range::get_local_range` * `nd_range::get_global` -> `nd_range::get_global_range`
1 parent 6dcbb81 commit 08ad294

20 files changed

Lines changed: 52 additions & 52 deletions

‎include/sycl/algorithm/algorithm_composite_patterns.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -57,8 +57,8 @@ class ReductionStrategy {
5757
public:
5858
ReductionStrategy(int loc, int len, cl::sycl::nd_item<1> i,
5959
const local_rw_acc<T>& localmem)
60-
: localid_(i.get_local(0)),
61-
globalid_(i.get_global(0)),
60+
: localid_(i.get_local_id(0)),
61+
globalid_(i.get_global_id(0)),
6262
local_(loc),
6363
length_(len),
6464
id_(i),

‎include/sycl/algorithm/buffer_algorithms.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -286,7 +286,7 @@ B buffer_map2reduce(ExecutionPolicy &snp,
286286
cl::sycl::access::target::local>
287287
sum { cl::sycl::range<1>(d.nb_work_item), cgh };
288288
cgh.parallel_for_work_group<typename ExecutionPolicy::kernelName>(
289-
rng.get_global(), rng.get_local(), [=](cl::sycl::group<1> grp) {
289+
rng.get_global_range(), rng.get_local_range(), [=](cl::sycl::group<1> grp) {
290290
size_t group_id = grp.get_id(0);
291291
//assert(group_id < d.nb_work_group);
292292
size_t group_begin = group_id * d.size_per_work_group;

‎include/sycl/algorithm/count_if.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,7 @@ typename std::iterator_traits<InputIterator>::difference_type count_if(
6767
cl::sycl::buffer<int, 1> bufR((cl::sycl::range<1>(vectorSize)));
6868
auto length = vectorSize;
6969
auto ndRange = exec.calculateNdRange(vectorSize);
70-
const auto local = ndRange.get_local()[0];
70+
const auto local = ndRange.get_local_range()[0];
7171
int passes = 0;
7272

7373
auto f = [&passes, &length, &ndRange, local, &bufI, &bufR, unary_op, binary_op](
@@ -76,7 +76,7 @@ typename std::iterator_traits<InputIterator>::difference_type count_if(
7676
auto aR = bufR.template get_access<cl::sycl::access::mode::read_write>(h);
7777
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
7878
cl::sycl::access::target::local>
79-
scratch(ndRange.get_local(), h);
79+
scratch(ndRange.get_local_range(), h);
8080

8181
h.parallel_for<typename ExecutionPolicy::kernelName>(
8282
ndRange, [aI, aR, scratch, passes, local, length, unary_op, binary_op](
@@ -95,7 +95,7 @@ typename std::iterator_traits<InputIterator>::difference_type count_if(
9595
q.submit(f);
9696
length = length / local;
9797
ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)),
98-
ndRange.get_local()};
98+
ndRange.get_local_range()};
9999
passes++;
100100
} while (length > 1);
101101
q.wait_and_throw();

‎include/sycl/algorithm/equal.hpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -69,7 +69,7 @@ bool equal(ExecutionPolicy& exec, ForwardIt1 first1, ForwardIt1 last1,
6969

7070
auto length = size1;
7171
auto ndRange = exec.calculateNdRange(size1);
72-
const auto local = ndRange.get_local()[0];
72+
const auto local = ndRange.get_local_range()[0];
7373

7474
auto buf1 = sycl::helpers::make_const_buffer(first1, last1);
7575
auto buf2 = sycl::helpers::make_const_buffer(first2, last2);
@@ -85,7 +85,7 @@ bool equal(ExecutionPolicy& exec, ForwardIt1 first1, ForwardIt1 last1,
8585
auto aR = bufR.template get_access<cl::sycl::access::mode::read_write>(h);
8686
cl::sycl::accessor<bool, 1, cl::sycl::access::mode::read_write,
8787
cl::sycl::access::target::local>
88-
scratch(ndRange.get_local(), h);
88+
scratch(ndRange.get_local_range(), h);
8989

9090
h.parallel_for<typename ExecutionPolicy::kernelName>(
9191
ndRange, [a1, a2, aR, scratch, passes, local, length,
@@ -105,7 +105,7 @@ bool equal(ExecutionPolicy& exec, ForwardIt1 first1, ForwardIt1 last1,
105105
q.submit(f);
106106
length = length / local;
107107
ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)),
108-
ndRange.get_local()};
108+
ndRange.get_local_range()};
109109
++passes;
110110
} while (length > 1);
111111
q.wait_and_throw();

‎include/sycl/algorithm/exclusive_scan.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -82,7 +82,7 @@ OutputIterator exclusive_scan(ExecutionPolicy &sep, InputIterator b,
8282
h.parallel_for<
8383
cl::sycl::helpers::NameGen<0, typename ExecutionPolicy::kernelName> >(
8484
ndRange, [aI, aO, init, vectorSize](cl::sycl::nd_item<1> id) {
85-
size_t m_id = id.get_global(0);
85+
size_t m_id = id.get_global_id(0);
8686
if (m_id > 0) {
8787
aO[m_id] = aI[m_id - 1];
8888
} else {
@@ -106,7 +106,7 @@ OutputIterator exclusive_scan(ExecutionPolicy &sep, InputIterator b,
106106
cl::sycl::helpers::NameGen<1, typename ExecutionPolicy::kernelName> >(
107107
ndRange, [aI, aO, bop, vectorSize, i](cl::sycl::nd_item<1> id) {
108108
size_t td = 1 << (i - 1);
109-
size_t m_id = id.get_global(0);
109+
size_t m_id = id.get_global_id(0);
110110
if (m_id < vectorSize && m_id >= td) {
111111
aO[m_id] = bop(aI[m_id - td], aI[m_id]);
112112
} else {

‎include/sycl/algorithm/fill.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -57,8 +57,8 @@ void fill(ExecutionPolicy &sep, ForwardIt b, ForwardIt e, const T &value) {
5757
auto aI = bufI.template get_access<cl::sycl::access::mode::read_write>(h);
5858
h.parallel_for<typename ExecutionPolicy::kernelName>(
5959
ndRange, [aI, val, vectorSize](cl::sycl::nd_item<1> id) {
60-
if (id.get_global(0) < vectorSize) {
61-
aI[id.get_global(0)] = val;
60+
if (id.get_global_id(0) < vectorSize) {
61+
aI[id.get_global_id(0)] = val;
6262
}
6363
});
6464
};

‎include/sycl/algorithm/find.hpp‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e,
6767
}
6868

6969
auto ndRange = sep.calculateNdRange(vectorSize);
70-
const auto local = ndRange.get_local()[0];
70+
const auto local = ndRange.get_local_range()[0];
7171

7272
// map across the input testing whether they match the predicate
7373
// store the result of the predicate and the index in the array of the result
@@ -78,7 +78,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e,
7878
h.parallel_for<
7979
cl::sycl::helpers::NameGen<0, typename ExecutionPolicy::kernelName> >(
8080
ndRange, [aI, aO, vectorSize, p](cl::sycl::nd_item<1> id) {
81-
const auto m_id = id.get_global(0);
81+
const auto m_id = id.get_global_id(0);
8282
// store index or the vector length, so that we can find the
8383
// _first_ index which is true, as opposed to just "one" of them
8484
if (m_id < vectorSize) {
@@ -99,7 +99,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e,
9999
t_buf.template get_access<cl::sycl::access::mode::read_write>(h);
100100
cl::sycl::accessor<std::size_t, 1, cl::sycl::access::mode::read_write,
101101
cl::sycl::access::target::local>
102-
scratch(ndRange.get_local(), h);
102+
scratch(ndRange.get_local_range(), h);
103103

104104
h.parallel_for<
105105
cl::sycl::helpers::NameGen<1, typename ExecutionPolicy::kernelName> >(
@@ -116,7 +116,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e,
116116
q.submit(rf);
117117
length = length / local;
118118
ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)),
119-
ndRange.get_local()};
119+
ndRange.get_local_range()};
120120
} while (length > 1);
121121
q.wait_and_throw();
122122

‎include/sycl/algorithm/for_each.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -56,8 +56,8 @@ void for_each(ExecutionPolicy &sep, Iterator b, Iterator e, UnaryFunction op) {
5656
auto aI = bufI.template get_access<cl::sycl::access::mode::read_write>(h);
5757
h.parallel_for<typename ExecutionPolicy::kernelName>(
5858
ndRange, [aI, op, vectorSize](cl::sycl::nd_item<1> id) {
59-
if (id.get_global(0) < vectorSize) {
60-
op(aI[id.get_global(0)]);
59+
if (id.get_global_id(0) < vectorSize) {
60+
op(aI[id.get_global_id(0)]);
6161
}
6262
});
6363
};

‎include/sycl/algorithm/for_each_n.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -65,8 +65,8 @@ InputIterator for_each_n(ExecutionPolicy &exec, InputIterator first, Size n,
6565
auto aI = bufI.template get_access<cl::sycl::access::mode::read_write>(h);
6666
h.parallel_for<typename ExecutionPolicy::kernelName>(
6767
ndRange, [vectorSize, aI, f](cl::sycl::nd_item<1> id) {
68-
if (id.get_global(0) < vectorSize) {
69-
f(aI[id.get_global(0)]);
68+
if (id.get_global_id(0) < vectorSize) {
69+
f(aI[id.get_global_id(0)]);
7070
}
7171
});
7272
};

‎include/sycl/algorithm/generate.hpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -59,8 +59,8 @@ void generate(ExecutionPolicy &sep, ForwardIt first, ForwardIt last,
5959
const auto aI = bufI.template get_access<cl::sycl::access::mode::read_write>(h);
6060
h.parallel_for<typename ExecutionPolicy::kernelName>(
6161
ndRange, [aI, g, vectorSize](cl::sycl::nd_item<1> id) {
62-
if (id.get_global(0) < vectorSize) {
63-
aI[id.get_global(0)] = g();
62+
if (id.get_global_id(0) < vectorSize) {
63+
aI[id.get_global_id(0)] = g();
6464
}
6565
});
6666
};

0 commit comments

Comments
 (0)