Tests/port existing kernel tests to experimental

numba_dpex/tests/kernel_tests/test_complex_array.py → numba_dpex/tests/kernel_tests/test_complex_array_kernel_arg.py

@@ -6,21 +6,21 @@
 import numpy
 import pytest
 
-import numba_dpex as dpex
+import numba_dpex.experimental as dpex
 from numba_dpex.tests._helper import get_all_dtypes
 
 N = 1024
 
 
 @dpex.kernel
-def kernel_scalar(a, b, c):
-    i = dpex.get_global_id(0)
+def kernel_scalar(item, a, b, c):
+    i = item.get_id(0)
     b[i] = a[i] * c
 
 
 @dpex.kernel
-def kernel_array(a, b, c):
-    i = dpex.get_global_id(0)
+def kernel_array(item, a, b, c):
+    i = item.get_id(0)
     b[i] = a[i] * c[i]
 
 

numba_dpex/tests/kernel_tests/test_dpnp_ndarray_args.py

@@ -3,10 +3,10 @@
 # SPDX-License-Identifier: Apache-2.0
 
 import dpnp
-from numpy.testing import assert_almost_equal
 
-import numba_dpex as ndpx
+import numba_dpex.experimental as ndpx
 from numba_dpex import float32
+from numba_dpex import kernel_api as kapi
 
 COEFFICIENTS = dpnp.asarray(
     [
@@ -64,9 +64,11 @@
 
 
 @ndpx.kernel()
-def _kernel(coefficients):
-    c = ndpx.private.array(4, dtype=float32)
-    gr_id = ndpx.get_group_id(0)
+def _kernel(nditem: kapi.NdItem, coefficients):
+    c = kapi.PrivateArray(4, dtype=float32)
+    gr: kapi.Group = nditem.get_group()
+    gr_id = gr.get_group_id(0)
+
     c[0] = coefficients[gr_id][0]
     c[1] = coefficients[gr_id][1]
     c[2] = coefficients[gr_id][2]

numba_dpex/tests/kernel_tests/test_func.py

@@ -5,17 +5,17 @@
 import dpnp
 import numpy
 
-import numba_dpex as dpex
+import numba_dpex.experimental as dpex
 
 
-@dpex.func
+@dpex.device_func
 def g(a):
     return a + dpnp.float32(1)
 
 
 @dpex.kernel
-def f(a, b):
-    i = dpex.get_global_id(0)
+def f(item, a, b):
+    i = item.get_id(0)
     b[i] = g(a[i])
 
 

numba_dpex/tests/kernel_tests/test_func_qualname_disambiguation.py

@@ -2,26 +2,24 @@
 #
 # SPDX-License-Identifier: Apache-2.0
 
-import dpctl.tensor as dpt
 import dpnp
 import numpy as np
 
-import numba_dpex as ndpx
+import numba_dpex.experimental as ndpx
 
 
 def make_write_values_kernel(n_rows):
-    """Uppermost kernel to set 1s in a certain way.
-    The uppermost kernel function invokes two levels
-    of inner functions to set 1s in an empty matrix
+    """Uppermost kernel to set 1s in a certain way. The uppermost kernel
+    function invokes two levels of inner functions to set 1s in an empty matrix
     in a certain way.
 
     Args:
         n_rows (int): Number of rows to iterate.
 
     Returns:
         numba_dpex.core.kernel_interface.dispatcher.JitKernel:
-            A JitKernel object that encapsulates a @kernel
-            decorated numba_dpex compiled kernel object.
+            A KernelDispatcher object that encapsulates a @kernel decorated
+            numba_dpex compiled kernel object.
     """
     write_values = make_write_values_kernel_func()
 
@@ -31,7 +29,7 @@ def write_values_kernel(array_in):
             is_even = (row_idx % 2) == 0
             write_values(array_in, row_idx, is_even)
 
-    return write_values_kernel[ndpx.NdRange(ndpx.Range(1), ndpx.Range(1))]
+    return write_values_kernel
 
 
 def make_write_values_kernel_func():
@@ -47,7 +45,7 @@ def make_write_values_kernel_func():
     write_when_odd = make_write_values_kernel_func_inner(1)
     write_when_even = make_write_values_kernel_func_inner(3)
 
-    @ndpx.func
+    @ndpx.device_func
     def write_values(array_in, row_idx, is_even):
         if is_even:
             write_when_even(array_in, row_idx)
@@ -70,7 +68,7 @@ def make_write_values_kernel_func_inner(n_cols):
             numba_dpex compiled function object.
     """
 
-    @ndpx.func
+    @ndpx.device_func
     def write_values_inner(array_in, row_idx):
         for idx in range(n_cols):
             array_in[row_idx, idx] = 1
@@ -89,9 +87,9 @@ def test_qualname_basic():
         else:
             ans[i, 0] = 1
 
-    a = dpnp.zeros((10, 10), dtype=dpt.int64)
+    a = dpnp.zeros((10, 10), dtype=dpnp.int64)
 
     kernel = make_write_values_kernel(10)
-    kernel(a)
+    ndpx.call_kernel(kernel, ndpx.NdRange((1,), (1,)), a)
 
     assert np.array_equal(dpnp.asnumpy(a), ans)

numba_dpex/tests/kernel_tests/test_invalid_kernel_args.py

@@ -5,15 +5,15 @@
 import numpy
 import pytest
 
-import numba_dpex as dpex
-from numba_dpex.core.exceptions import UnsupportedKernelArgumentError
+import numba_dpex.experimental as dpex
+from numba_dpex import kernel_api as kapi
 
 N = 1024
 
 
 @dpex.kernel
-def vecadd_kernel(a, b, c):
-    i = dpex.get_global_id(0)
+def vecadd_kernel(item: kapi.Item, a, b, c):
+    i = item.get_id(0)
     c[i] = a[i] + b[i]
 
 
@@ -25,5 +25,5 @@ def test_passing_numpy_arrays_as_kernel_args():
     b = numpy.ones(N)
     c = numpy.zeros(N)
 
-    with pytest.raises(UnsupportedKernelArgumentError):
+    with pytest.raises(Exception):
         dpex.call_kernel(vecadd_kernel, dpex.Range(N), a, b, c)

numba_dpex/tests/kernel_tests/test_math_functions.py

@@ -8,7 +8,7 @@
 import numpy
 import pytest
 
-import numba_dpex as dpex
+import numba_dpex.experimental as dpex
 from numba_dpex.tests._helper import get_all_dtypes
 
 list_of_unary_ops = ["fabs", "exp", "log", "sqrt", "sin", "cos", "tan"]
@@ -39,8 +39,8 @@ def test_binary_ops(unary_op, input_arrays):
     dpnp_uop = getattr(dpnp, unary_op)
 
     @dpex.kernel
-    def f(a, b):
-        i = dpex.get_global_id(0)
+    def f(item, a, b):
+        i = item.get_id(0)
         b[i] = uop(a[i])
 
     dpex.call_kernel(f, dpex.Range(a.size), a, b)

numba_dpex/tests/kernel_tests/test_print.py

@@ -5,9 +5,9 @@
 import dpctl
 import dpnp
 import pytest
-from numba.core.errors import LoweringError
+from numba.core.errors import TypingError
 
-import numba_dpex as dpex
+import numba_dpex.experimental as dpex
 
 list_of_dtypes = [
     dpnp.int32,
@@ -84,9 +84,11 @@ def print_string(a):
 
     a = input_arrays
 
-    with pytest.raises(LoweringError):
+    with pytest.raises(TypingError) as ex_info:
         dpex.call_kernel(print_string, dpex.Range(1), a)
 
+    assert "LoweringError" in ex_info.value.args[0]
+
 
 @skip_on_gpu
 def test_print_array(input_arrays):
@@ -100,5 +102,7 @@ def print_string(a):
 
     a = input_arrays
 
-    with pytest.raises(LoweringError):
+    with pytest.raises(TypingError) as ex_info:
         dpex.call_kernel(print_string, dpex.Range(1), a)
+
+    assert "LoweringError" in ex_info.value.args[0]

numba_dpex/tests/kernel_tests/test_scalar_arg_types.py

@@ -6,21 +6,21 @@
 import numpy
 import pytest
 
-import numba_dpex as dpex
+import numba_dpex.experimental as dpex
 from numba_dpex.tests._helper import get_all_dtypes
 
 N = 1024
 
 
 @dpex.kernel
-def scaling_kernel(a, b, c):
-    i = dpex.get_global_id(0)
+def scaling_kernel(item, a, b, c):
+    i = item.get_id(0)
     b[i] = a[i] * c
 
 
 @dpex.kernel
-def kernel_with_bool_arg(a, b, test):
-    i = dpex.get_global_id(0)
+def kernel_with_bool_arg(item, a, b, test):
+    i = item.get_id(0)
     if test:
         b[i] = a[i] + a[i]
     else:

numba_dpex/tests/kernel_tests/test_usm_ndarray_args.py

@@ -0,0 +1,61 @@
+# SPDX-FileCopyrightText: 2020 - 2024 Intel Corporation
+#
+# SPDX-License-Identifier: Apache-2.0
+
+import dpctl.tensor as dpt
+import numpy
+import pytest
+
+import numba_dpex.experimental as dpex
+from numba_dpex.tests._helper import get_all_dtypes
+
+list_of_dtype = get_all_dtypes(
+    no_bool=True, no_float16=True, no_none=True, no_complex=True
+)
+
+
+@pytest.fixture(params=list_of_dtype)
+def dtype(request):
+    return request.param
+
+
+list_of_usm_type = [
+    "shared",
+    "device",
+    "host",
+]
+
+
+@pytest.fixture(params=list_of_usm_type)
+def usm_type(request):
+    return request.param
+
+
+def test_consuming_usm_ndarray(dtype, usm_type):
+    @dpex.kernel
+    def data_parallel_sum(item, a, b, c):
+        """
+        Vector addition using the ``kernel`` decorator.
+        """
+        i = item.get_id(0)
+        j = item.get_id(1)
+        c[i, j] = a[i, j] + b[i, j]
+
+    N = 1000
+    global_size = N * N
+
+    a = dpt.arange(global_size, dtype=dtype, usm_type=usm_type)
+    a = dpt.reshape(a, shape=(N, N))
+
+    b = dpt.arange(global_size, dtype=dtype, usm_type=usm_type)
+    b = dpt.reshape(b, shape=(N, N))
+
+    c = dpt.empty_like(a)
+
+    dpex.call_kernel(data_parallel_sum, dpex.Range(N, N), a, b, c)
+
+    na = dpt.asnumpy(a)
+    nb = dpt.asnumpy(b)
+    nc = dpt.asnumpy(c)
+
+    assert numpy.array_equal(nc, na + nb)