Applied SYCL 1.2.1 renamed member functions for nd_item and `nd_ran…

…ge`: (#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`

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<T>& 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),

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<typename ExecutionPolicy::kernelName>(
-        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;

include/sycl/algorithm/count_if.hpp

@@ -67,7 +67,7 @@ typename std::iterator_traits<InputIterator>::difference_type count_if(
   cl::sycl::buffer<int, 1> 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<InputIterator>::difference_type count_if(
     auto aR = bufR.template get_access<cl::sycl::access::mode::read_write>(h);
     cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                        cl::sycl::access::target::local>
-        scratch(ndRange.get_local(), h);
+        scratch(ndRange.get_local_range(), h);
 
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         ndRange, [aI, aR, scratch, passes, local, length, unary_op, binary_op](
@@ -95,7 +95,7 @@ typename std::iterator_traits<InputIterator>::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();

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<cl::sycl::access::mode::read_write>(h);
       cl::sycl::accessor<bool, 1, cl::sycl::access::mode::read_write,
                          cl::sycl::access::target::local>
-          scratch(ndRange.get_local(), h);
+          scratch(ndRange.get_local_range(), h);
 
       h.parallel_for<typename ExecutionPolicy::kernelName>(
           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();

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 {

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<cl::sycl::access::mode::read_write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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;
           }
         });
   };

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<cl::sycl::access::mode::read_write>(h);
       cl::sycl::accessor<std::size_t, 1, cl::sycl::access::mode::read_write,
                          cl::sycl::access::target::local>
-          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();
 

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<cl::sycl::access::mode::read_write>(h);
       h.parallel_for<typename ExecutionPolicy::kernelName>(
           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)]);
             }
           });
     };

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<cl::sycl::access::mode::read_write>(h);
       h.parallel_for<typename ExecutionPolicy::kernelName>(
           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)]);
             }
           });
     };

include/sycl/algorithm/generate.hpp

@@ -59,8 +59,8 @@ void generate(ExecutionPolicy &sep, ForwardIt first, ForwardIt last,
     const auto aI = bufI.template get_access<cl::sycl::access::mode::read_write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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();
           }
         });
   };

include/sycl/algorithm/inclusive_scan.hpp

@@ -86,7 +86,7 @@ OutputIterator inclusive_scan(ExecutionPolicy &sep, InputIterator b,
       h.parallel_for<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]);

include/sycl/algorithm/inner_product.hpp

@@ -130,7 +130,7 @@ T inner_product(ExecutionPolicy &exec, InputIt1 first1, InputIt1 last1,
     cl::sycl::buffer<T, 1> 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<cl::sycl::access::mode::read_write>(h);
       cl::sycl::accessor<T, 1, cl::sycl::access::mode::read_write,
                          cl::sycl::access::target::local>
-          scratch(ndRange.get_local(), h);
+          scratch(ndRange.get_local_range(), h);
 
       h.parallel_for<typename ExecutionPolicy::kernelName>(
           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<cl::sycl::access::mode::read>();

include/sycl/algorithm/mismatch.hpp

@@ -65,7 +65,7 @@ std::pair<ForwardIt1, ForwardIt2> 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<ForwardIt1, ForwardIt2> 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<ForwardIt1, ForwardIt2> mismatch(ExecutionPolicy& exec,
         bufR.template get_access<cl::sycl::access::mode::read_write>(h);
     cl::sycl::accessor<std::size_t, 1, cl::sycl::access::mode::read_write,
                        cl::sycl::access::target::local>
-        scratch(ndRange.get_local(), h);
+        scratch(ndRange.get_local_range(), h);
 
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         ndRange, [aR, scratch, passes, local,
@@ -118,7 +118,7 @@ std::pair<ForwardIt1, ForwardIt2> 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<cl::sycl::access::mode::read>(

include/sycl/algorithm/reduce.hpp

@@ -71,13 +71,13 @@ typename std::iterator_traits<Iterator>::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<cl::sycl::access::mode::read_write>(h);
     cl::sycl::accessor<type_, 1, cl::sycl::access::mode::read_write,
                        cl::sycl::access::target::local>
-        scratch(ndRange.get_local(), h);
+        scratch(ndRange.get_local_range(), h);
 
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         ndRange, [aI, scratch, local, length, bop](cl::sycl::nd_item<1> id) {
@@ -91,7 +91,7 @@ typename std::iterator_traits<Iterator>::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<cl::sycl::access::mode::read>();

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<cl::sycl::access::mode::write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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)) {

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<cl::sycl::access::mode::read_write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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;

include/sycl/algorithm/reverse.hpp

@@ -57,7 +57,7 @@ void reverse(ExecutionPolicy &sep, BidirIt first, BidirIt last) {
     const auto aI = bufI.template get_access<cl::sycl::access::mode::read_write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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]);
           }

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<cl::sycl::access::mode::write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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];
           }

include/sycl/algorithm/transform.hpp

@@ -65,8 +65,8 @@ OutputIterator transform(ExecutionPolicy &sep, Iterator b, Iterator e,
       auto aO = bufO.template get_access<cl::sycl::access::mode::write>(h);
       h.parallel_for<typename ExecutionPolicy::kernelName>(
           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<cl::sycl::access::mode::write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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<cl::sycl::access::mode::write>(h);
     h.parallel_for<typename ExecutionPolicy::kernelName>(
         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<cl::sycl::access::mode::write>(h);
     h.parallel_for<class TransformAlgorithm>(
         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)]);
           }
         });
   };

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<cl::sycl::access::mode::read_write>(h);
       cl::sycl::accessor<T, 1, cl::sycl::access::mode::read_write,
                          cl::sycl::access::target::local>
-          scratch(ndRange.get_local(), h);
+          scratch(ndRange.get_local_range(), h);
 
       h.parallel_for<typename ExecutionPolicy::kernelName>(
           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<cl::sycl::access::mode::read>();