Merge pull request Apress#6 from breyerml/kernel_offset_fix

removed implicit offset in favor of explicit one since offsets were deprecated in SYCL 2020

samples/Ch14_common_parallel_patterns/fig_14_14_stencil.cpp

@@ -32,10 +32,9 @@ int main() {
       accessor output{ output_buf, h };
 
       // Compute the average of each cell and its immediate neighbors
-      id<2> offset(1, 1);
-      h.parallel_for(stencil_range, offset, [=](id<2> idx) {
-        int i = idx[0];
-        int j = idx[1];
+      h.parallel_for(stencil_range, [=](id<2> idx) {
+        int i = idx[0] + 1;
+        int j = idx[1] + 1;
 
         float self = input[i][j];
         float north = input[i - 1][j];

samples/Ch14_common_parallel_patterns/fig_14_15_local_stencil.cpp

@@ -42,9 +42,8 @@ int main() {
       auto tile = local_accessor<float, 2>(tile_size, h);
 
       // Compute the average of each cell and its immediate neighbors
-      id<2> offset(1, 1);
       h.parallel_for(
-          nd_range<2>(stencil_range, local_range, offset), [=](nd_item<2> it) {
+          nd_range<2>(stencil_range, local_range), [=](nd_item<2> it) {
             // Load this tile into work-group local memory
             id<2> lid = it.get_local_id();
             range<2> lrange = it.get_local_range();
@@ -58,8 +57,8 @@ int main() {
             it.barrier(access::fence_space::local_space);
 
             // Compute the stencil using values from local memory
-            int gi = it.get_global_id(0);
-            int gj = it.get_global_id(1);
+            int gi = it.get_global_id(0) + 1;
+            int gj = it.get_global_id(1) + 1;
 
             int ti = it.get_local_id(0) + 1;
             int tj = it.get_local_id(1) + 1;