Rename h_item::get_global() and h_item::get_local() to get_global_id(…

…) and get_local_id()

Better SYCL 1.2.1 compliant.

include/CL/sycl/h_item.hpp

@@ -89,42 +89,44 @@ struct h_item {
   /** Return the constituent global id representing the work-item's
       position in the global iteration space
   */
-  id<Dimensions> get_global() const { return global_index; }
+  id<Dimensions> get_global_id() const { return global_index; }
 
 
   /** Return the constituent element of the global id representing the
       work-item's position in the global iteration space in the given
       dimension
   */
-  size_t get_global(int dimension) const { return get_global()[dimension]; }
+  size_t get_global_id(int dimension) const {
+    return get_global_id()[dimension];
+  }
 
 
   /** Return the flattened id of the current work-item after subtracting
       the offset
   */
   size_t get_global_linear_id() const {
-    return detail::linear_id(get_global_range(), get_global(), get_offset());
+    return detail::linear_id(get_global_range(), get_global_id(), get_offset());
   }
 
 
   /** Return the constituent local id representing the work-item's
       position within the current work-group
   */
-  id<Dimensions> get_local() const { return local_index; }
+  id<Dimensions> get_local_id() const { return local_index; }
 
 
   /** Return the constituent element of the local id representing the
       work-item's position within the current work-group in the given
       dimension
   */
-  size_t get_local(int dimension) const { return get_local()[dimension]; }
+  size_t get_local_id(int dimension) const { return get_local_id()[dimension]; }
 
 
   /** Return the flattened id of the current work-item within the current
       work-group
    */
   size_t get_local_linear_id() const {
-    return detail::linear_id(get_local_range(), get_local());
+    return detail::linear_id(get_local_range(), get_local_id());
   }
 
 
@@ -134,7 +136,7 @@ struct h_item {
   id<Dimensions> get_group() const {
     /* Convert get_local_range() to an id<> to remove ambiguity into using
        implicit conversion either from range<> to id<> or the opposite */
-    return get_global()/id<Dimensions> { get_local_range() };
+    return get_global_id()/id<Dimensions> { get_local_range() };
   }
 
 
@@ -192,7 +194,7 @@ struct h_item {
       \todo Add to the specification
   */
   item<Dimensions> get_item() const {
-    return { get_global_range(), get_global(), get_offset() };
+    return { get_global_range(), get_global_id(), get_offset() };
   }
 
 

tests/2014-04-21-HPC-GPU_Meetup/slide_22.cpp

@@ -28,7 +28,7 @@ my_queue.submit([&](handler &cgh)
                                                              range<>(groupsize)),
                                                   [=](group<> group) {
                                                     group.parallel_for_work_item([=](h_item<1> tile) {
-                                                        out_access[tile.get_global()] = in_access[tile.get_global()] * 2;
+                                                        out_access[tile.get_global_id()] = in_access[tile.get_global_id()] * 2;
                                                       });
                                                   });
  });

tests/accessor/local_accessor_hierarchical_convolution.cpp

@@ -65,17 +65,17 @@ int test_main(int argc, char *argv[]) {
         [=](group<1> g) {
           g.parallel_for_work_item([&](h_item<1> i) {
               // Cache the convolution kernel in local memory
-              if (i.get_local(0) < a_conv_kernel.get_count())
-                a_conv_cache[i.get_local(0)] =
-                  a_conv_kernel[i.get_global()];
+              if (i.get_local_id(0) < a_conv_kernel.get_count())
+                a_conv_cache[i.get_local_id(0)] =
+                  a_conv_kernel[i.get_global_id()];
             });
           // An implicit barrier happens here
           g.parallel_for_work_item([&](h_item<1> i) {
-              a_result[i.get_global()] = 0;
+              a_result[i.get_global_id()] = 0;
               // The convolution
               for (unsigned int j = 0; j < conv_kernel.get_count(); ++j)
-                a_result[i.get_global()] +=
-                  a_a[i.get_local(0) + j]*a_conv_cache[j];
+                a_result[i.get_global_id()] +=
+                  a_a[i.get_local_id(0) + j]*a_conv_cache[j];
             });
         });
     });

tests/jacobi/include/stencil-fxd.hpp

@@ -182,21 +182,21 @@ class operation_fxd2D {
   }
 
   inline void eval_local(cl::sycl::h_item<2> it, cl::sycl::accessor<T, 2, cl::sycl::access::mode::write> out, T *local_tab, int glob_max0, int glob_max1) {
-    int i = it.get_global().get(0);
-    int j = it.get_global().get(1);
+    int i = it.get_global_id(0);
+    int j = it.get_global_id(1);
     if (i >= glob_max0 || j >= glob_max1)
       return;
     i += of0;
     j += of1;
-    int i_local = it.get_local().get(0) - st::min_ind0;
-    int j_local = it.get_local().get(1) - st::min_ind1;
+    int i_local = it.get_local_id(0) - st::min_ind0;
+    int j_local = it.get_local_id(1) - st::min_ind1;
     f(i, j, out) = stencil.template eval_local<local_dim1>(local_tab, i_local, j_local);
   }
 
   inline void store_local(T * local_tab, cl::sycl::accessor<T, 2, cl::sycl::access::mode::read> in, cl::sycl::h_item<2> it, cl::sycl::group<2> gr, int glob_max0, int glob_max1) {
     cl::sycl::range<2> l_range = it.get_local_range();
     cl::sycl::id<2> g_ind = gr.get_id(); //it.get_group_id(); error because ambiguous / operator redefinition 
-    cl::sycl::id<2> l_ind = it.get_local();
+    cl::sycl::id<2> l_ind = it.get_local_id();
 
     int l_range0 = l_range.get(0);
     int l_range1 = l_range.get(1);

tests/jacobi/include/stencil-gen-var.hpp

@@ -135,8 +135,6 @@ class operation_var2D {
 
   operation_var2D() {
     cl::sycl::range<2> rg1 = aB->get_range();
-    cl::sycl::range<2> rg2 = B->get_range();
-    assert(rg1 == rg2); //seems a bad idea, indeed ! we will not do that ...
     range = rg1-d;
     int r0 = range.get(0);
     int r1 = range.get(1);
@@ -161,22 +159,22 @@ class operation_var2D {
 
   // Not really static because of the use of global_max (which is passed by args)
   static inline void eval_local(cl::sycl::h_item<2> it, cl::sycl::accessor<T, 2, cl::sycl::access::mode::write> out, T *local_tab, cl::sycl::accessor<T, 1, cl::sycl::access::mode::read> coef, int glob_max0, int glob_max1) {
-    int i = it.get_global().get(0);
-    int j = it.get_global().get(1);
+    int i = it.get_global_id(0);
+    int j = it.get_global_id(1);
     if (i >= glob_max0 || j >= glob_max1)
       return;
     i += of0;
     j += of1;
-    int i_local = it.get_local().get(0) - st::min_ind0;
-    int j_local = it.get_local().get(1) - st::min_ind1;
+    int i_local = it.get_local_id(0) - st::min_ind0;
+    int j_local = it.get_local_id(1) - st::min_ind1;
     f(i, j, out) = st::template eval_local<T, local_dim1, b_f, c_f, r_f>(local_tab, coef, i, j, i_local, j_local);
   }
 
   // Not really static because of the use of global_max (which is passed by args)
   static inline void store_local(T * local_tab, cl::sycl::accessor<T, 2, cl::sycl::access::mode::read> in, cl::sycl::h_item<2> it, cl::sycl::group<2> gr, int glob_max0, int glob_max1) {
     cl::sycl::range<2> l_range = it.get_local_range();
     cl::sycl::id<2> g_ind = gr.get_id(); //it.get_group_id(); error because ambiguous / operator redefinition 
-    cl::sycl::id<2> l_ind = it.get_local();
+    cl::sycl::id<2> l_ind = it.get_local_id();
 
     int l_range0 = l_range.get(0);
     int l_range1 = l_range.get(1);

tests/jacobi/include/stencil-var.hpp

@@ -161,22 +161,22 @@ class operation_var2D {
 
   // Not really static because of the use of global_max (which is passed by args)
   static inline void eval_local(cl::sycl::h_item<2> it, cl::sycl::accessor<T, 2, cl::sycl::access::mode::write> out, T *local_tab, cl::sycl::accessor<T, 1, cl::sycl::access::mode::read> coef, int glob_max0, int glob_max1) {
-    int i = it.get_global().get(0);
-    int j = it.get_global().get(1);
+    int i = it.get_global_id(0);
+    int j = it.get_global_id(1);
     if (i >= glob_max0 || j >= glob_max1)
       return;
     i += of0;
     j += of1;
-    int i_local = it.get_local().get(0) - st::min_ind0;
-    int j_local = it.get_local().get(1) - st::min_ind1;
+    int i_local = it.get_local_id(0) - st::min_ind0;
+    int j_local = it.get_local_id(1) - st::min_ind1;
     f(i, j, out) = st::template eval_local<T, local_dim1, b_f>(local_tab, coef, i, j, i_local, j_local);
   }
 
   // Not really static because of the use of global_max (which is passed by args)
   static inline void store_local(T * local_tab, cl::sycl::accessor<T, 2, cl::sycl::access::mode::read> in, cl::sycl::h_item<2> it, cl::sycl::group<2> gr, int glob_max0, int glob_max1) {
     cl::sycl::range<2> l_range = it.get_local_range();
     cl::sycl::id<2> g_ind = gr.get_id(); //it.get_group_id(); error because ambiguous / operator redefinition 
-    cl::sycl::id<2> l_ind = it.get_local();
+    cl::sycl::id<2> l_ind = it.get_local_id();
 
     int l_range0 = l_range.get(0);
     int l_range1 = l_range.get(1);

tests/jacobi/jacobi2d-tile.cpp

@@ -59,8 +59,8 @@ int main(int argc, char **argv) {
               // tab_local has to be passed by reference
              group.parallel_for_work_item([=,&local](sycl::h_item<2> it){
                   sycl::range<2> l_range = it.get_local_range();
-                  sycl::id<2> g_ind = it.get_global();
-                  sycl::id<2> l_ind = it.get_local();
+                  sycl::id<2> g_ind = it.get_global_id();
+                  sycl::id<2> l_ind = it.get_local_id();
                   sycl::id<2> offset = it.get_offset();
                   sycl::id<2> id1(sycl::range<2> {0,1});
                   sycl::id<2> id2(sycl::range<2> {1,0});
@@ -83,8 +83,8 @@ int main(int argc, char **argv) {
               begin_op = counters::clock_type::now();
 
               group.parallel_for_work_item([=,&local](sycl::h_item<2> it){
-                  sycl::id<2> g_ind = it.get_global();
-                  sycl::id<2> l_ind = it.get_local();
+                  sycl::id<2> g_ind = it.get_global_id();
+                  sycl::id<2> l_ind = it.get_local_id();
                   sycl::id<2> offset = it.get_offset();
                   sycl::id<2> id1(sycl::range<2> {0,1});
                   sycl::id<2> id2(sycl::range<2> {1,0});

tests/multiple_compilation_units/parallel_for.cc

@@ -30,8 +30,8 @@ int test_main(int argc, char *argv[]) {
         nd_range<> { range<> { size }, range<> { groupsize } },
         [=](group<> group) {
           group.parallel_for_work_item([=](h_item<1> tile) {
-              out_access[tile.get_global()] =
-                1000*group.get_id(0) + tile.get_local(0);
+              out_access[tile.get_global_id()] =
+                1000*group.get_id(0) + tile.get_local_id(0);
             });
         });
     });

tests/parallel_for/hierarchical.cpp

@@ -51,9 +51,9 @@ my_queue.submit([&](handler &cgh)
 
     group.parallel_for_work_item([=](h_item<1> tile)
     {
-      std::cout << "Local id = " << tile.get_local(0)
-                << " (global id = " << tile.get_global(0) << ")" << std::endl;
-      out_access[tile.get_global()] = in_access[tile.get_global()] * 2;
+      std::cout << "Local id = " << tile.get_local_id(0)
+                << " (global id = " << tile.get_global_id(0) << ")" << std::endl;
+      out_access[tile.get_global_id()] = in_access[tile.get_global_id()] * 2;
     });
   });
 });

tests/parallel_for/hierarchical_new.cpp

@@ -58,14 +58,14 @@ my_queue.submit([&](handler &cgh) {
     std::cout << "Group id = " << group.get_id(0) << std::endl;
 
     group.parallel_for_work_item([=](h_item<1> tile) {
-      std::cout << "Local id = " << tile.get_local(0)
-                << " (global id = " << tile.get_global(0) << ")" << std::endl;
-      out_access[tile.get_global()] = in_access[tile.get_global()] * 2;
+      std::cout << "Local id = " << tile.get_local_id(0)
+                << " (global id = " << tile.get_global_id(0) << ")" << std::endl;
+      out_access[tile.get_global_id()] = in_access[tile.get_global_id()] * 2;
     });
 
     group.parallel_for_work_item([=](h_item<1> tile) {
-      std::cout << "Global id = " << tile.get_global(0) << std::endl;
-      out_access[tile.get_global()] = in_access[tile.get_global()] * 2;
+      std::cout << "Global id = " << tile.get_global_id(0) << std::endl;
+      out_access[tile.get_global_id()] = in_access[tile.get_global_id()] * 2;
     });
   });
 });

tests/pipe/1pipe_read_reserve.cpp

@@ -76,7 +76,7 @@ int test_main(int argc, char *argv[]) {
                    launch the work-items in this work-group to do the
                    reading in parallel */
                 group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                    ac[start + i.get_global(0)] = r[i.get_global(0)];
+                    ac[start + i.get_global_id(0)] = r[i.get_global_id(0)];
                   });
               }
               // Here the reservation object goes out of scope: commit

tests/pipe/1pipe_read_write_reserve.cpp

@@ -65,7 +65,7 @@ int test_main(int argc, char *argv[]) {
                    launch the work-items in this work-group to do the
                    writing in parallel */
                 group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                    r[i.get_global(0)] = aa[start + i.get_global(0)];
+                    r[i.get_global_id(0)] = aa[start + i.get_global_id(0)];
                   });
               }
               // Here the reservation object goes out of scope: commit
@@ -104,7 +104,7 @@ int test_main(int argc, char *argv[]) {
                the work-items in this work-group to do the reading in
                parallel */
             group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                ac[start + i.get_global(0)] = r[i.get_global(0)];
+                ac[start + i.get_global_id(0)] = r[i.get_global_id(0)];
             });
             /** Explicit commit requested here. Note that in this
                 simple example, since there is nothing useful after

tests/pipe/1pipe_write_reserve.cpp

@@ -61,7 +61,7 @@ int test_main(int argc, char *argv[]) {
                    launch the work-items in this work-group to do the
                    writing in parallel */
                 group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                    r[i.get_global(0)] = aa[start + i.get_global(0)];
+                    r[i.get_global_id(0)] = aa[start + i.get_global_id(0)];
                   });
               }
               // Here the reservation object goes out of scope: commit

tests/pipe/3pipes_reserve_producer_consumer.cpp

@@ -69,7 +69,7 @@ int test_main(int argc, char *argv[]) {
                    launch the work-items in this work-group to do the
                    writing in parallel */
                 group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                    r[i.get_global(0)] = aa[start + i.get_global(0)];
+                    r[i.get_global_id(0)] = aa[start + i.get_global_id(0)];
                   });
               // Here the reservation object goes out of scope: commit if ok
             }

tests/pipe/blocking_pipe_read_write_reserve.cpp

@@ -52,7 +52,7 @@ int test_main(int argc, char *argv[]) {
           for (int start = 0; start != N; start += WI) {
             auto r = apa.reserve(WI);
             group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                r[i.get_global(0)] = aa[start + i.get_global(0)];
+                r[i.get_global_id(0)] = aa[start + i.get_global_id(0)];
               });
             // Here the reservation object goes out of scope: commit
           }
@@ -76,7 +76,7 @@ int test_main(int argc, char *argv[]) {
           for (int start = 0; start != N; start += WI) {
             auto r = apa.reserve(WI);
             group.parallel_for_work_item([=] (cl::sycl::h_item<> i) {
-                ac[start + i.get_global(0)] = r[i.get_global(0)];
+                ac[start + i.get_global_id(0)] = r[i.get_global_id(0)];
               });
             // Here the reservation object goes out of scope: commit
           }