Merge pull request #1421 from IntelPython/fix/examples-to-use-new-api

Updates few leftover examples to use new kernel api

Author: ZzEeKkAa

numba_dpex/examples/debug/sum_local_vars.py

@@ -9,7 +9,7 @@
 
 @ndpx.kernel(debug=True)
 def data_parallel_sum(item, a, b, c):
-    i = ndpx.get_global_id(0)
+    i = item.get_id(0)
     l1 = a[i] + 2.5
     l2 = b[i] * 0.3
     c[i] = l1 + l2

numba_dpex/examples/kernel/pipelining.py

@@ -23,8 +23,8 @@
 
 
 @dpex.kernel
-def async_kernel(x):
-    idx = dpex.get_global_id(0)
+def async_kernel(item, x):
+    idx = item.get_id(0)
 
     for i in range(1300):
         den = x.dtype.type(i + 1)

numba_dpex/examples/kernel/sum_reduction_recursive.py

@@ -18,11 +18,11 @@
 
 @ndpx.kernel
 def sum_reduction_kernel(nditem: kapi.NdItem, A, input_size, partial_sums, slm):
-    local_id = ndpx.get_local_id(0)
-    global_id = ndpx.get_global_id(0)
-    group_size = ndpx.get_local_size(0)
-    group_id = ndpx.get_group_id(0)
-
+    local_id = nditem.get_local_id(0)
+    global_id = nditem.get_global_id(0)
+    group_size = nditem.get_local_range(0)
+    gr = nditem.get_group()
+    group_id = gr.get_group_id(0)
     slm[local_id] = 0
 
     if global_id < input_size:
@@ -32,7 +32,7 @@ def sum_reduction_kernel(nditem: kapi.NdItem, A, input_size, partial_sums, slm):
     stride = group_size // 2
     while stride > 0:
         # Waiting for each 2x2 addition into given workgroup
-        ndpx.barrier(ndpx.LOCAL_MEM_FENCE)
+        kapi.group_barrier(gr)
 
         # Add elements 2 by 2 between local_id and local_id + stride
         if local_id < stride: