diff --git a/numba_dpex/__init__.py b/numba_dpex/__init__.py index 802f502ed4..cf72ed70d1 100644 --- a/numba_dpex/__init__.py +++ b/numba_dpex/__init__.py @@ -17,6 +17,7 @@ from numba import __version__ as numba_version from numba.np.ufunc.decorators import Vectorize +from numba_dpex.core.kernel_interface.launcher import call_kernel from numba_dpex.vectorizers import Vectorize as DpexVectorize from .numba_patches import patch_arrayexpr_tree_to_ir, patch_is_ufunc @@ -145,4 +146,4 @@ def parse_sem_version(version_string: str) -> Tuple[int, int, int]: __version__ = get_versions()["version"] del get_versions -__all__ = types.__all__ + ["Range", "NdRange"] +__all__ = types.__all__ + ["Range", "NdRange", "call_kernel"] diff --git a/numba_dpex/core/kernel_interface/launcher.py b/numba_dpex/core/kernel_interface/launcher.py new file mode 100644 index 0000000000..a57f8cf3d2 --- /dev/null +++ b/numba_dpex/core/kernel_interface/launcher.py @@ -0,0 +1,18 @@ +"""Launcher package to provide the same way of calling kernel as experimental +one.""" + + +def call_kernel(kernel_fn, index_space, *kernel_args) -> None: + """Syntax sugar for calling kernel the same way as experimental one. + It is a temporary glue for the experimental kernel migration. + + Args: + kernel_fn (numba_dpex.experimental.KernelDispatcher): A + numba_dpex.kernel decorated function that is compiled to a + KernelDispatcher by numba_dpex. + index_space (Range | NdRange): A numba_dpex.Range or numba_dpex.NdRange + type object that specifies the index space for the kernel. + kernel_args : List of objects that are passed to the numba_dpex.kernel + decorated function. + """ + kernel_fn[index_space](*kernel_args) diff --git a/numba_dpex/tests/kernel_tests/test_atomic_op.py b/numba_dpex/tests/kernel_tests/test_atomic_op.py index ad3e9e89e1..a73992c267 100644 --- a/numba_dpex/tests/kernel_tests/test_atomic_op.py +++ b/numba_dpex/tests/kernel_tests/test_atomic_op.py @@ -56,7 +56,7 @@ def f(a): def test_kernel_atomic_simple(input_arrays, kernel_result_pair): a, dtype = input_arrays() kernel, expected = kernel_result_pair - kernel[dpex.Range(global_size)](a) + dpex.call_kernel(kernel, dpex.Range(global_size), a) assert a[0] == expected @@ -96,7 +96,7 @@ def test_kernel_atomic_local(input_arrays, return_list_of_op): op_type, expected = return_list_of_op f = get_func_local(op_type, dtype) kernel = dpex.kernel(f) - kernel[dpex.NdRange(dpex.Range(N), dpex.Range(N))](a) + dpex.call_kernel(kernel, dpex.NdRange(dpex.Range(N), dpex.Range(N)), a) assert a[0] == expected @@ -134,7 +134,7 @@ def test_kernel_atomic_multi_dim( dim = return_list_of_dim kernel = get_kernel_multi_dim(op_type, len(dim)) a = np.zeros(dim, dtype=return_dtype) - kernel[dpex.Range(global_size)](a) + dpex.call_kernel(kernel, dpex.Range(global_size), a) assert a[0] == expected diff --git a/numba_dpex/tests/kernel_tests/test_barrier.py b/numba_dpex/tests/kernel_tests/test_barrier.py index b438724339..9cef7d2563 100644 --- a/numba_dpex/tests/kernel_tests/test_barrier.py +++ b/numba_dpex/tests/kernel_tests/test_barrier.py @@ -25,7 +25,7 @@ def twice(A): orig = dpt.asnumpy(arr) global_size = (N,) local_size = (N // 2,) - twice[NdRange(global_size, local_size)](arr) + dpex.call_kernel(twice, NdRange(global_size, local_size), arr) after = dpt.asnumpy(arr) # The computation is correct? np.testing.assert_allclose(orig * 2, after) @@ -43,7 +43,7 @@ def twice(A): N = 256 arr = dpt.arange(N, dtype=dpt.float32) orig = dpt.asnumpy(arr) - twice[Range(N)](arr) + dpex.call_kernel(twice, Range(N), arr) after = dpt.asnumpy(arr) # The computation is correct? np.testing.assert_allclose(orig * 2, after) @@ -66,7 +66,9 @@ def reverse_array(A): arr = dpt.arange(blocksize, dtype=dpt.float32) orig = dpt.asnumpy(arr) - reverse_array[NdRange(Range(blocksize), Range(blocksize))](arr) + dpex.call_kernel( + reverse_array, NdRange(Range(blocksize), Range(blocksize)), arr + ) after = dpt.asnumpy(arr) expected = orig[::-1] + orig np.testing.assert_allclose(expected, after) diff --git a/numba_dpex/tests/kernel_tests/test_complex_array.py b/numba_dpex/tests/kernel_tests/test_complex_array.py index d71e434d57..fa4fe57d7a 100644 --- a/numba_dpex/tests/kernel_tests/test_complex_array.py +++ b/numba_dpex/tests/kernel_tests/test_complex_array.py @@ -48,7 +48,7 @@ def test_numeric_kernel_arg_complex_scalar(input_arrays): a, b, _ = input_arrays s = a.dtype.type(2 + 1j) - kernel_scalar[dpex.Range(N)](a, b, s) + dpex.call_kernel(kernel_scalar, dpex.Range(N), a, b, s) nb = dpnp.asnumpy(b) nexpected = numpy.full_like(nb, fill_value=2 + 1j) @@ -65,7 +65,7 @@ def test_numeric_kernel_arg_complex_array(input_arrays): a, b, c = input_arrays - kernel_array[dpex.Range(N)](a, b, c) + dpex.call_kernel(kernel_array, dpex.Range(N), a, b, c) nb = dpnp.asnumpy(b) nexpected = numpy.full_like(nb, fill_value=0 + 0j) diff --git a/numba_dpex/tests/kernel_tests/test_dpnp_ndarray_args.py b/numba_dpex/tests/kernel_tests/test_dpnp_ndarray_args.py index da9229bd40..57e5c4df40 100644 --- a/numba_dpex/tests/kernel_tests/test_dpnp_ndarray_args.py +++ b/numba_dpex/tests/kernel_tests/test_dpnp_ndarray_args.py @@ -85,7 +85,9 @@ def test_setting_private_from_dpnp_ndarray(): global_range = ndpx.Range(N_POINTS // N_POINTS_PER_WORK_ITEM) local_range = ndpx.Range(LOCAL_SIZE) try: - _kernel[ndpx.NdRange(global_range, local_range)](COEFFICIENTS) + ndpx.call_kernel( + _kernel, ndpx.NdRange(global_range, local_range), COEFFICIENTS + ) except Exception as e: assert ( False diff --git a/numba_dpex/tests/kernel_tests/test_func.py b/numba_dpex/tests/kernel_tests/test_func.py index 6fc0226f30..7b721dc77b 100644 --- a/numba_dpex/tests/kernel_tests/test_func.py +++ b/numba_dpex/tests/kernel_tests/test_func.py @@ -24,6 +24,6 @@ def test_func_call_from_kernel(): a = dpnp.ones(1024) b = dpnp.ones(1024) - f[dpex.Range(1024)](a, b) + dpex.call_kernel(f, dpex.Range(1024), a, b) nb = dpnp.asnumpy(b) assert numpy.all(nb == 2) diff --git a/numba_dpex/tests/kernel_tests/test_func_specialization.py b/numba_dpex/tests/kernel_tests/test_func_specialization.py index 386535ba18..c6eddd8928 100644 --- a/numba_dpex/tests/kernel_tests/test_func_specialization.py +++ b/numba_dpex/tests/kernel_tests/test_func_specialization.py @@ -35,7 +35,7 @@ def kernel_function(a, b): a = dpnp.ones(N) b = dpnp.ones(N) - k[dpex.Range(N)](a, b) + dpex.call_kernel(k, dpex.Range(N), a, b) assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) @@ -56,7 +56,7 @@ def kernel_function(a, b): a = dpnp.ones(N, dtype=dpnp.int32) b = dpnp.ones(N, dtype=dpnp.int32) - k[dpex.Range(N)](a, b) + dpex.call_kernel(k, dpex.Range(N), a, b) assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) @@ -65,7 +65,7 @@ def kernel_function(a, b): b = dpnp.ones(N, dtype=dpnp.int64) with pytest.raises(Exception) as e: - k[dpex.Range(N)](a, b) + dpex.call_kernel(k, dpex.Range(N), a, b) assert " >>> (int64)" in e.value.args[0] @@ -86,7 +86,7 @@ def kernel_function(a, b): a = dpnp.ones(N, dtype=dpnp.int32) b = dpnp.ones(N, dtype=dpnp.int32) - k[dpex.Range(N)](a, b) + dpex.call_kernel(k, dpex.Range(N), a, b) assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) @@ -94,7 +94,7 @@ def kernel_function(a, b): a = dpnp.ones(N, dtype=dpnp.float32) b = dpnp.ones(N, dtype=dpnp.float32) - k[dpex.Range(N)](a, b) + dpex.call_kernel(k, dpex.Range(N), a, b) assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) @@ -103,6 +103,6 @@ def kernel_function(a, b): b = dpnp.ones(N, dtype=dpnp.int64) with pytest.raises(Exception) as e: - k[dpex.Range(N)](a, b) + dpex.call_kernel(k, dpex.Range(N), a, b) assert " >>> (int64)" in e.value.args[0] diff --git a/numba_dpex/tests/kernel_tests/test_invalid_kernel_args.py b/numba_dpex/tests/kernel_tests/test_invalid_kernel_args.py index 7cea27c5c4..679d75f41e 100644 --- a/numba_dpex/tests/kernel_tests/test_invalid_kernel_args.py +++ b/numba_dpex/tests/kernel_tests/test_invalid_kernel_args.py @@ -26,4 +26,4 @@ def test_passing_numpy_arrays_as_kernel_args(): c = numpy.zeros(N) with pytest.raises(UnsupportedKernelArgumentError): - vecadd_kernel[dpex.Range(N)](a, b, c) + dpex.call_kernel(vecadd_kernel, dpex.Range(N), a, b, c) diff --git a/numba_dpex/tests/kernel_tests/test_kernel_has_return_value_error.py b/numba_dpex/tests/kernel_tests/test_kernel_has_return_value_error.py index 60a403ea2f..be5ae89a4f 100644 --- a/numba_dpex/tests/kernel_tests/test_kernel_has_return_value_error.py +++ b/numba_dpex/tests/kernel_tests/test_kernel_has_return_value_error.py @@ -31,4 +31,4 @@ def test_return(sig): with pytest.raises(dpex.core.exceptions.KernelHasReturnValueError): kernel_fn = dpex.kernel(sig)(f) - kernel_fn[dpex.Range(a.size)](a) + dpex.call_kernel(kernel_fn, dpex.Range(a.size), a) diff --git a/numba_dpex/tests/kernel_tests/test_kernel_specialization.py b/numba_dpex/tests/kernel_tests/test_kernel_specialization.py index afc189fb86..89a7d41385 100644 --- a/numba_dpex/tests/kernel_tests/test_kernel_specialization.py +++ b/numba_dpex/tests/kernel_tests/test_kernel_specialization.py @@ -60,7 +60,9 @@ def test_missing_specialization_error(): c = dpt.zeros(1024, dtype=dpt.int32) with pytest.raises(MissingSpecializationError): - specialized_kernel1(data_parallel_sum)[Range(1024)](a, b, c) + dpex.call_kernel( + specialized_kernel1(data_parallel_sum), Range(1024), a, b, c + ) def test_execution_of_specialized_kernel(): @@ -69,7 +71,9 @@ def test_execution_of_specialized_kernel(): b = dpt.ones(1024, dtype=dpt.int64) c = dpt.zeros(1024, dtype=dpt.int64) - specialized_kernel1(data_parallel_sum)[Range(1024)](a, b, c) + dpex.call_kernel( + specialized_kernel1(data_parallel_sum), Range(1024), a, b, c + ) npc = dpt.asnumpy(c) import numpy as np diff --git a/numba_dpex/tests/kernel_tests/test_math_functions.py b/numba_dpex/tests/kernel_tests/test_math_functions.py index 24ad66bd02..d920edc67a 100644 --- a/numba_dpex/tests/kernel_tests/test_math_functions.py +++ b/numba_dpex/tests/kernel_tests/test_math_functions.py @@ -43,7 +43,7 @@ def f(a, b): i = dpex.get_global_id(0) b[i] = uop(a[i]) - f[dpex.Range(a.size)](a, b) + dpex.call_kernel(f, dpex.Range(a.size), a, b) expected = dpnp_uop(a) diff --git a/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py b/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py index 6c1320b061..611f67804f 100644 --- a/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py +++ b/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py @@ -34,4 +34,4 @@ def test_ndrange_config_error(error, ranges): with pytest.raises(error): range = NdRange(ranges[0], ranges[1]) - kernel_vector_sum[range](a, b, c) + ndpx.call_kernel(kernel_vector_sum, range, a, b, c) diff --git a/numba_dpex/tests/kernel_tests/test_print.py b/numba_dpex/tests/kernel_tests/test_print.py index 25402421fc..a9a2e18692 100644 --- a/numba_dpex/tests/kernel_tests/test_print.py +++ b/numba_dpex/tests/kernel_tests/test_print.py @@ -42,7 +42,7 @@ def print_scalar_val(s): a = input_arrays - print_scalar_val[dpex.Range(1)](a) + dpex.call_kernel(print_scalar_val, dpex.Range(1), a) captured = capfd.readouterr() assert "printing ... 10" in captured.out @@ -60,7 +60,7 @@ def print_scalar_val(s): a = input_arrays - print_scalar_val[dpex.Range(1)](a) + dpex.call_kernel(print_scalar_val, dpex.Range(1), a) captured = capfd.readouterr() assert "10" in captured.out @@ -85,7 +85,7 @@ def print_string(a): a = input_arrays with pytest.raises(LoweringError): - print_string[dpex.Range(1)](a) + dpex.call_kernel(print_string, dpex.Range(1), a) @skip_on_gpu @@ -101,4 +101,4 @@ def print_string(a): a = input_arrays with pytest.raises(LoweringError): - print_string[dpex.Range(1)](a) + dpex.call_kernel(print_string, dpex.Range(1), a) diff --git a/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py b/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py index 9d12602048..1cbfcbff89 100644 --- a/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py +++ b/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py @@ -23,7 +23,7 @@ def kernel_with_private_memory_allocation(A): def test_private_memory_allocation(): N = 64 arr = dpnp.zeros(N, dtype=dpnp.float32) - kernel_with_private_memory_allocation[dpex.Range(N)](arr) + dpex.call_kernel(kernel_with_private_memory_allocation, dpex.Range(N), arr) nparr = dpnp.asnumpy(arr) diff --git a/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py b/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py index 8b1fb576cd..8513251a66 100644 --- a/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py +++ b/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py @@ -49,7 +49,7 @@ def test_numeric_kernel_arg_types1(input_arrays): a, b = input_arrays s = a.dtype.type(2) - scaling_kernel[dpex.Range(N)](a, b, s) + dpex.call_kernel(scaling_kernel, dpex.Range(N), a, b, s) nb = dpnp.asnumpy(b) nexpected = numpy.full_like(nb, fill_value=2) @@ -65,14 +65,14 @@ def test_bool_kernel_arg_type(input_arrays): """ a, b = input_arrays - kernel_with_bool_arg[dpex.Range(a.size)](a, b, True) + dpex.call_kernel(kernel_with_bool_arg, dpex.Range(a.size), a, b, True) nb = dpnp.asnumpy(b) nexpected_true = numpy.full_like(nb, fill_value=2) assert numpy.allclose(nb, nexpected_true) - kernel_with_bool_arg[dpex.Range(a.size)](a, b, False) + dpex.call_kernel(kernel_with_bool_arg, dpex.Range(a.size), a, b, False) nb = dpnp.asnumpy(b) nexpected_false = numpy.zeros_like(nb) diff --git a/numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py b/numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py index bd89265bb7..a405ffb3e8 100644 --- a/numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py +++ b/numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py @@ -88,7 +88,7 @@ def test_kernel_valid_usm_obj(dtype): C = DuckUSMArray(shape=buffC.shape, dtype=dtype, host_buffer=buffC) try: - vecadd[dpex.Range(N)](A, B, C) + dpex.call_kernel(vecadd, dpex.Range(N), A, B, C) except Exception: pytest.fail( "Could not pass Python object with sycl_usm_array_interface" @@ -112,4 +112,4 @@ def test_kernel_invalid_usm_obj(dtype): C = PseudoDuckUSMArray() with pytest.raises(Exception): - vecadd[dpex.Range(N)](A, B, C) + dpex.call_kernel(vecadd, dpex.Range(N), A, B, C) diff --git a/numba_dpex/tests/kernel_tests/test_usm_ndarray_interop.py b/numba_dpex/tests/kernel_tests/test_usm_ndarray_interop.py index 3591d942c8..a70aa66909 100644 --- a/numba_dpex/tests/kernel_tests/test_usm_ndarray_interop.py +++ b/numba_dpex/tests/kernel_tests/test_usm_ndarray_interop.py @@ -52,7 +52,7 @@ def data_parallel_sum(a, b, c): c = dpt.empty_like(a) - data_parallel_sum[dpex.Range(N, N)](a, b, c) + dpex.call_kernel(data_parallel_sum, dpex.Range(N, N), a, b, c) na = dpt.asnumpy(a) nb = dpt.asnumpy(b) diff --git a/numba_dpex/tests/misc/test_warnings.py b/numba_dpex/tests/misc/test_warnings.py index 71108b77ea..7f63199424 100644 --- a/numba_dpex/tests/misc/test_warnings.py +++ b/numba_dpex/tests/misc/test_warnings.py @@ -17,7 +17,7 @@ def test_opt_warning(): config.DPEX_OPT = 3 with pytest.warns(UserWarning): - foo[dpex.Range(10)](dpnp.arange(10)) + dpex.call_kernel(foo, dpex.Range(10), dpnp.arange(10)) config.DPEX_OPT = bkp @@ -27,7 +27,7 @@ def test_inline_threshold_eq_3_warning(): config.INLINE_THRESHOLD = 3 with pytest.warns(UserWarning): - foo[dpex.Range(10)](dpnp.arange(10)) + dpex.call_kernel(foo, dpex.Range(10), dpnp.arange(10)) config.INLINE_THRESHOLD = bkp @@ -37,7 +37,7 @@ def test_inline_threshold_negative_val_warning_(): config.INLINE_THRESHOLD = -1 with pytest.warns(UserWarning): - foo[dpex.Range(10)](dpnp.arange(10)) + dpex.call_kernel(foo, dpex.Range(10), dpnp.arange(10)) config.INLINE_THRESHOLD = bkp @@ -47,7 +47,7 @@ def test_inline_threshold_gt_3_warning(): config.INLINE_THRESHOLD = 4 with pytest.warns(UserWarning): - foo[dpex.Range(10)](dpnp.arange(10)) + dpex.call_kernel(foo, dpex.Range(10), dpnp.arange(10)) config.INLINE_THRESHOLD = bkp @@ -55,4 +55,4 @@ def test_inline_threshold_gt_3_warning(): def test_no_warning(): with warnings.catch_warnings(): warnings.simplefilter("error") - foo[dpex.Range(10)](dpnp.arange(10)) + dpex.call_kernel(foo, dpex.Range(10), dpnp.arange(10))