From 08ad294ae88b1b731ab6cfafbe05d0c7f104ffd6 Mon Sep 17 00:00:00 2001 From: Christopher Di Bella Date: Tue, 31 Jul 2018 15:54:37 +0100 Subject: [PATCH] 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` --- .../algorithm_composite_patterns.hpp | 4 ++-- include/sycl/algorithm/buffer_algorithms.hpp | 2 +- include/sycl/algorithm/count_if.hpp | 6 ++--- include/sycl/algorithm/equal.hpp | 6 ++--- include/sycl/algorithm/exclusive_scan.hpp | 4 ++-- include/sycl/algorithm/fill.hpp | 4 ++-- include/sycl/algorithm/find.hpp | 8 +++---- include/sycl/algorithm/for_each.hpp | 4 ++-- include/sycl/algorithm/for_each_n.hpp | 4 ++-- include/sycl/algorithm/generate.hpp | 4 ++-- include/sycl/algorithm/inclusive_scan.hpp | 2 +- include/sycl/algorithm/inner_product.hpp | 6 ++--- include/sycl/algorithm/mismatch.hpp | 8 +++---- include/sycl/algorithm/reduce.hpp | 6 ++--- include/sycl/algorithm/replace_copy_if.hpp | 2 +- include/sycl/algorithm/replace_if.hpp | 2 +- include/sycl/algorithm/reverse.hpp | 2 +- include/sycl/algorithm/reverse_copy.hpp | 2 +- include/sycl/algorithm/transform.hpp | 22 +++++++++---------- include/sycl/algorithm/transform_reduce.hpp | 6 ++--- 20 files changed, 52 insertions(+), 52 deletions(-) diff --git a/include/sycl/algorithm/algorithm_composite_patterns.hpp b/include/sycl/algorithm/algorithm_composite_patterns.hpp index 01468d4..7b609c0 100644 --- a/include/sycl/algorithm/algorithm_composite_patterns.hpp +++ b/include/sycl/algorithm/algorithm_composite_patterns.hpp @@ -57,8 +57,8 @@ class ReductionStrategy { public: ReductionStrategy(int loc, int len, cl::sycl::nd_item<1> i, const local_rw_acc& localmem) - : localid_(i.get_local(0)), - globalid_(i.get_global(0)), + : localid_(i.get_local_id(0)), + globalid_(i.get_global_id(0)), local_(loc), length_(len), id_(i), diff --git a/include/sycl/algorithm/buffer_algorithms.hpp b/include/sycl/algorithm/buffer_algorithms.hpp index 1b34773..5c4adf2 100644 --- a/include/sycl/algorithm/buffer_algorithms.hpp +++ b/include/sycl/algorithm/buffer_algorithms.hpp @@ -286,7 +286,7 @@ B buffer_map2reduce(ExecutionPolicy &snp, cl::sycl::access::target::local> sum { cl::sycl::range<1>(d.nb_work_item), cgh }; cgh.parallel_for_work_group( - rng.get_global(), rng.get_local(), [=](cl::sycl::group<1> grp) { + rng.get_global_range(), rng.get_local_range(), [=](cl::sycl::group<1> grp) { size_t group_id = grp.get_id(0); //assert(group_id < d.nb_work_group); size_t group_begin = group_id * d.size_per_work_group; diff --git a/include/sycl/algorithm/count_if.hpp b/include/sycl/algorithm/count_if.hpp index 177abfe..7b5330e 100644 --- a/include/sycl/algorithm/count_if.hpp +++ b/include/sycl/algorithm/count_if.hpp @@ -67,7 +67,7 @@ typename std::iterator_traits::difference_type count_if( cl::sycl::buffer bufR((cl::sycl::range<1>(vectorSize))); auto length = vectorSize; auto ndRange = exec.calculateNdRange(vectorSize); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; int passes = 0; auto f = [&passes, &length, &ndRange, local, &bufI, &bufR, unary_op, binary_op]( @@ -76,7 +76,7 @@ typename std::iterator_traits::difference_type count_if( auto aR = bufR.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for( ndRange, [aI, aR, scratch, passes, local, length, unary_op, binary_op]( @@ -95,7 +95,7 @@ typename std::iterator_traits::difference_type count_if( q.submit(f); length = length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; passes++; } while (length > 1); q.wait_and_throw(); diff --git a/include/sycl/algorithm/equal.hpp b/include/sycl/algorithm/equal.hpp index 0052cc7..bc57cd9 100644 --- a/include/sycl/algorithm/equal.hpp +++ b/include/sycl/algorithm/equal.hpp @@ -69,7 +69,7 @@ bool equal(ExecutionPolicy& exec, ForwardIt1 first1, ForwardIt1 last1, auto length = size1; auto ndRange = exec.calculateNdRange(size1); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; auto buf1 = sycl::helpers::make_const_buffer(first1, last1); auto buf2 = sycl::helpers::make_const_buffer(first2, last2); @@ -85,7 +85,7 @@ bool equal(ExecutionPolicy& exec, ForwardIt1 first1, ForwardIt1 last1, auto aR = bufR.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for( ndRange, [a1, a2, aR, scratch, passes, local, length, @@ -105,7 +105,7 @@ bool equal(ExecutionPolicy& exec, ForwardIt1 first1, ForwardIt1 last1, q.submit(f); length = length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; ++passes; } while (length > 1); q.wait_and_throw(); diff --git a/include/sycl/algorithm/exclusive_scan.hpp b/include/sycl/algorithm/exclusive_scan.hpp index ab20b3a..c88c72c 100644 --- a/include/sycl/algorithm/exclusive_scan.hpp +++ b/include/sycl/algorithm/exclusive_scan.hpp @@ -82,7 +82,7 @@ OutputIterator exclusive_scan(ExecutionPolicy &sep, InputIterator b, h.parallel_for< cl::sycl::helpers::NameGen<0, typename ExecutionPolicy::kernelName> >( ndRange, [aI, aO, init, vectorSize](cl::sycl::nd_item<1> id) { - size_t m_id = id.get_global(0); + size_t m_id = id.get_global_id(0); if (m_id > 0) { aO[m_id] = aI[m_id - 1]; } else { @@ -106,7 +106,7 @@ OutputIterator exclusive_scan(ExecutionPolicy &sep, InputIterator b, cl::sycl::helpers::NameGen<1, typename ExecutionPolicy::kernelName> >( ndRange, [aI, aO, bop, vectorSize, i](cl::sycl::nd_item<1> id) { size_t td = 1 << (i - 1); - size_t m_id = id.get_global(0); + size_t m_id = id.get_global_id(0); if (m_id < vectorSize && m_id >= td) { aO[m_id] = bop(aI[m_id - td], aI[m_id]); } else { diff --git a/include/sycl/algorithm/fill.hpp b/include/sycl/algorithm/fill.hpp index fde9e71..dde9b08 100644 --- a/include/sycl/algorithm/fill.hpp +++ b/include/sycl/algorithm/fill.hpp @@ -57,8 +57,8 @@ void fill(ExecutionPolicy &sep, ForwardIt b, ForwardIt e, const T &value) { auto aI = bufI.template get_access(h); h.parallel_for( ndRange, [aI, val, vectorSize](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < vectorSize) { - aI[id.get_global(0)] = val; + if (id.get_global_id(0) < vectorSize) { + aI[id.get_global_id(0)] = val; } }); }; diff --git a/include/sycl/algorithm/find.hpp b/include/sycl/algorithm/find.hpp index 30e1b0d..61911c0 100644 --- a/include/sycl/algorithm/find.hpp +++ b/include/sycl/algorithm/find.hpp @@ -67,7 +67,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e, } auto ndRange = sep.calculateNdRange(vectorSize); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; // map across the input testing whether they match the predicate // 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, h.parallel_for< cl::sycl::helpers::NameGen<0, typename ExecutionPolicy::kernelName> >( ndRange, [aI, aO, vectorSize, p](cl::sycl::nd_item<1> id) { - const auto m_id = id.get_global(0); + const auto m_id = id.get_global_id(0); // store index or the vector length, so that we can find the // _first_ index which is true, as opposed to just "one" of them if (m_id < vectorSize) { @@ -99,7 +99,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e, t_buf.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for< cl::sycl::helpers::NameGen<1, typename ExecutionPolicy::kernelName> >( @@ -116,7 +116,7 @@ InputIt find_impl(ExecutionPolicy &sep, InputIt b, InputIt e, q.submit(rf); length = length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; } while (length > 1); q.wait_and_throw(); diff --git a/include/sycl/algorithm/for_each.hpp b/include/sycl/algorithm/for_each.hpp index 3d34616..dd08ffd 100644 --- a/include/sycl/algorithm/for_each.hpp +++ b/include/sycl/algorithm/for_each.hpp @@ -56,8 +56,8 @@ void for_each(ExecutionPolicy &sep, Iterator b, Iterator e, UnaryFunction op) { auto aI = bufI.template get_access(h); h.parallel_for( ndRange, [aI, op, vectorSize](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < vectorSize) { - op(aI[id.get_global(0)]); + if (id.get_global_id(0) < vectorSize) { + op(aI[id.get_global_id(0)]); } }); }; diff --git a/include/sycl/algorithm/for_each_n.hpp b/include/sycl/algorithm/for_each_n.hpp index 45cc994..9c5d486 100644 --- a/include/sycl/algorithm/for_each_n.hpp +++ b/include/sycl/algorithm/for_each_n.hpp @@ -65,8 +65,8 @@ InputIterator for_each_n(ExecutionPolicy &exec, InputIterator first, Size n, auto aI = bufI.template get_access(h); h.parallel_for( ndRange, [vectorSize, aI, f](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < vectorSize) { - f(aI[id.get_global(0)]); + if (id.get_global_id(0) < vectorSize) { + f(aI[id.get_global_id(0)]); } }); }; diff --git a/include/sycl/algorithm/generate.hpp b/include/sycl/algorithm/generate.hpp index 0810273..8585bbd 100644 --- a/include/sycl/algorithm/generate.hpp +++ b/include/sycl/algorithm/generate.hpp @@ -59,8 +59,8 @@ void generate(ExecutionPolicy &sep, ForwardIt first, ForwardIt last, const auto aI = bufI.template get_access(h); h.parallel_for( ndRange, [aI, g, vectorSize](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < vectorSize) { - aI[id.get_global(0)] = g(); + if (id.get_global_id(0) < vectorSize) { + aI[id.get_global_id(0)] = g(); } }); }; diff --git a/include/sycl/algorithm/inclusive_scan.hpp b/include/sycl/algorithm/inclusive_scan.hpp index 613f0b1..d6f1973 100644 --- a/include/sycl/algorithm/inclusive_scan.hpp +++ b/include/sycl/algorithm/inclusive_scan.hpp @@ -86,7 +86,7 @@ OutputIterator inclusive_scan(ExecutionPolicy &sep, InputIterator b, h.parallel_for( ndRange, [aI, aO, bop, vectorSize, i](cl::sycl::nd_item<1> id) { size_t td = 1 << (i - 1); - size_t m_id = id.get_global(0); + size_t m_id = id.get_global_id(0); if (m_id < vectorSize && m_id >= td) { aO[m_id] = bop(aI[m_id - td], aI[m_id]); diff --git a/include/sycl/algorithm/inner_product.hpp b/include/sycl/algorithm/inner_product.hpp index 2c215f8..6b81736 100644 --- a/include/sycl/algorithm/inner_product.hpp +++ b/include/sycl/algorithm/inner_product.hpp @@ -130,7 +130,7 @@ T inner_product(ExecutionPolicy &exec, InputIt1 first1, InputIt1 last1, cl::sycl::buffer bufr((cl::sycl::range<1>(vectorSize))); auto length = vectorSize; auto ndRange = exec.calculateNdRange(length); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; int passes = 0; auto cg = [&passes, &length, &ndRange, local, &buf1, &buf2, &bufr, op1, op2]( cl::sycl::handler &h) mutable { @@ -140,7 +140,7 @@ T inner_product(ExecutionPolicy &exec, InputIt1 first1, InputIt1 last1, bufr.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for( ndRange, [a1, a2, aR, scratch, length, local, passes, op1, op2]( @@ -160,7 +160,7 @@ T inner_product(ExecutionPolicy &exec, InputIt1 first1, InputIt1 last1, passes++; length = length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; } while (length > 1); // end do-while q.wait_and_throw(); auto hb = bufr.template get_access(); diff --git a/include/sycl/algorithm/mismatch.hpp b/include/sycl/algorithm/mismatch.hpp index d231c84..1d65a0c 100644 --- a/include/sycl/algorithm/mismatch.hpp +++ b/include/sycl/algorithm/mismatch.hpp @@ -65,7 +65,7 @@ std::pair mismatch(ExecutionPolicy& exec, const auto length = std::min(size1, size2); auto ndRange = exec.calculateNdRange(length); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; auto buf1 = sycl::helpers::make_const_buffer(first1, first1 + length); auto buf2 = sycl::helpers::make_const_buffer(first2, first2 + length); @@ -81,7 +81,7 @@ std::pair mismatch(ExecutionPolicy& exec, h.parallel_for< cl::sycl::helpers::NameGen<0, typename ExecutionPolicy::kernelName> >( ndRange, [a1, a2, aR, length, p](cl::sycl::nd_item<1> id) { - const auto m_id = id.get_global(0); + const auto m_id = id.get_global_id(0); if (m_id < length) { aR[m_id] = p(a1[m_id], a2[m_id]) ? length : m_id; @@ -99,7 +99,7 @@ std::pair mismatch(ExecutionPolicy& exec, bufR.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for( ndRange, [aR, scratch, passes, local, @@ -118,7 +118,7 @@ std::pair mismatch(ExecutionPolicy& exec, ++passes; current_length = current_length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(current_length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; } while (current_length > 1); q.wait_and_throw(); const auto hR = bufR.get_access( diff --git a/include/sycl/algorithm/reduce.hpp b/include/sycl/algorithm/reduce.hpp index 82155e4..6c1e24d 100644 --- a/include/sycl/algorithm/reduce.hpp +++ b/include/sycl/algorithm/reduce.hpp @@ -71,13 +71,13 @@ typename std::iterator_traits::value_type reduce( auto bufI = sycl::helpers::make_const_buffer(b, e); auto length = vectorSize; auto ndRange = sep.calculateNdRange(length); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; auto f = [&length, &ndRange, local, &bufI, bop](cl::sycl::handler &h) mutable { auto aI = bufI.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for( ndRange, [aI, scratch, local, length, bop](cl::sycl::nd_item<1> id) { @@ -91,7 +91,7 @@ typename std::iterator_traits::value_type reduce( q.submit(f); length = length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; } while (length > 1); q.wait_and_throw(); auto hI = bufI.template get_access(); diff --git a/include/sycl/algorithm/replace_copy_if.hpp b/include/sycl/algorithm/replace_copy_if.hpp index fb332ec..7cd50df 100644 --- a/include/sycl/algorithm/replace_copy_if.hpp +++ b/include/sycl/algorithm/replace_copy_if.hpp @@ -68,7 +68,7 @@ ForwardIt2 replace_copy_if(ExecutionPolicy &sep, ForwardIt1 first, const auto aO = bufO.template get_access(h); h.parallel_for( ndRange, [aI, aO, vectorSize, p, new_val](cl::sycl::nd_item<1> id) { - const auto global_id = id.get_global(0); + const auto global_id = id.get_global_id(0); const auto orig_value = aI[global_id]; if (global_id < vectorSize) { if (p(orig_value)) { diff --git a/include/sycl/algorithm/replace_if.hpp b/include/sycl/algorithm/replace_if.hpp index 70d4c1d..1384276 100644 --- a/include/sycl/algorithm/replace_if.hpp +++ b/include/sycl/algorithm/replace_if.hpp @@ -61,7 +61,7 @@ void replace_if(ExecutionPolicy &sep, ForwardIt first, ForwardIt last, const auto aI = bufI.template get_access(h); h.parallel_for( ndRange, [aI, vectorSize, p, new_val](cl::sycl::nd_item<1> id) { - const auto global_id = id.get_global(0); + const auto global_id = id.get_global_id(0); if (global_id < vectorSize) { if(p(aI[global_id])) { aI[global_id] = new_val; diff --git a/include/sycl/algorithm/reverse.hpp b/include/sycl/algorithm/reverse.hpp index 4125856..ff36fba 100644 --- a/include/sycl/algorithm/reverse.hpp +++ b/include/sycl/algorithm/reverse.hpp @@ -57,7 +57,7 @@ void reverse(ExecutionPolicy &sep, BidirIt first, BidirIt last) { const auto aI = bufI.template get_access(h); h.parallel_for( ndRange, [aI, vectorSize](cl::sycl::nd_item<1> id) { - const auto global_id = id.get_global(0); + const auto global_id = id.get_global_id(0); if (global_id < vectorSize / 2) { helpers::swap(aI[global_id], aI[vectorSize - global_id - 1]); } diff --git a/include/sycl/algorithm/reverse_copy.hpp b/include/sycl/algorithm/reverse_copy.hpp index f64050c..babf5da 100644 --- a/include/sycl/algorithm/reverse_copy.hpp +++ b/include/sycl/algorithm/reverse_copy.hpp @@ -61,7 +61,7 @@ ForwardIt reverse_copy(ExecutionPolicy &sep, BidirIt first, BidirIt last, const auto aO = bufO.template get_access(h); h.parallel_for( ndRange, [aI, aO, vectorSize](cl::sycl::nd_item<1> id) { - const auto global_id = id.get_global(0); + const auto global_id = id.get_global_id(0); if (global_id < vectorSize) { aO[global_id] = aI[vectorSize - global_id - 1]; } diff --git a/include/sycl/algorithm/transform.hpp b/include/sycl/algorithm/transform.hpp index cf7a051..bf8697f 100644 --- a/include/sycl/algorithm/transform.hpp +++ b/include/sycl/algorithm/transform.hpp @@ -65,8 +65,8 @@ OutputIterator transform(ExecutionPolicy &sep, Iterator b, Iterator e, auto aO = bufO.template get_access(h); h.parallel_for( ndRange, [aI, aO, op, vectorSize](cl::sycl::nd_item<1> id) { - if ((id.get_global(0) < vectorSize)) { - aO[id.get_global(0)] = op(aI[id.get_global(0)]); + if ((id.get_global_id(0) < vectorSize)) { + aO[id.get_global_id(0)] = op(aI[id.get_global_id(0)]); } }); }; @@ -104,9 +104,9 @@ OutputIterator transform(ExecutionPolicy &sep, InputIterator first1, auto aO = res.template get_access(h); h.parallel_for( ndRange, [a1, a2, aO, op, n](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < n) { - aO[id.get_global(0)] = - op(a1[id.get_global(0)], a2[id.get_global(0)]); + if (id.get_global_id(0) < n) { + aO[id.get_global_id(0)] = + op(a1[id.get_global_id(0)], a2[id.get_global_id(0)]); } }); }; @@ -144,9 +144,9 @@ OutputIterator transform(ExecutionPolicy &sep, cl::sycl::queue &q, auto aO = res.template get_access(h); h.parallel_for( ndRange, [a1, a2, aO, op, n](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < n) { - aO[id.get_global(0)] = - op(a1[id.get_global(0)], a2[id.get_global(0)]); + if (id.get_global_id(0) < n) { + aO[id.get_global_id(0)] = + op(a1[id.get_global_id(0)], a2[id.get_global_id(0)]); } }); }; @@ -177,9 +177,9 @@ void transform(ExecutionPolicy &sep, cl::sycl::queue &q, Buffer &buf1, auto aO = res.template get_access(h); h.parallel_for( ndRange, [a1, a2, aO, op, n](cl::sycl::nd_item<1> id) { - if (id.get_global(0) < n) { - aO[id.get_global(0)] = - op(a1[id.get_global(0)], a2[id.get_global(0)]); + if (id.get_global_id(0) < n) { + aO[id.get_global_id(0)] = + op(a1[id.get_global_id(0)], a2[id.get_global_id(0)]); } }); }; diff --git a/include/sycl/algorithm/transform_reduce.hpp b/include/sycl/algorithm/transform_reduce.hpp index e1e78af..ab9cc06 100644 --- a/include/sycl/algorithm/transform_reduce.hpp +++ b/include/sycl/algorithm/transform_reduce.hpp @@ -66,7 +66,7 @@ T transform_reduce(ExecutionPolicy& exec, InputIterator first, auto bufI = sycl::helpers::make_const_buffer(first, last); size_t length = vectorSize; auto ndRange = exec.calculateNdRange(vectorSize); - const auto local = ndRange.get_local()[0]; + const auto local = ndRange.get_local_range()[0]; int passes = 0; do { @@ -76,7 +76,7 @@ T transform_reduce(ExecutionPolicy& exec, InputIterator first, auto aR = bufR.template get_access(h); cl::sycl::accessor - scratch(ndRange.get_local(), h); + scratch(ndRange.get_local_range(), h); h.parallel_for( ndRange, [aI, aR, scratch, passes, local, length, unary_op, binary_op]( @@ -95,7 +95,7 @@ T transform_reduce(ExecutionPolicy& exec, InputIterator first, passes++; length = length / local; ndRange = cl::sycl::nd_range<1>{cl::sycl::range<1>(std::max(length, local)), - ndRange.get_local()}; + ndRange.get_local_range()}; } while (length > 1); q.wait_and_throw(); auto hR = bufR.template get_access();