From 70b7810ef270958251d2ac026029d7a3238ef646 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 14:40:42 -0500 Subject: [PATCH 01/21] Add more test cases to test_llvm_registration.py --- .../core/runtime/test_llvm_registration.py | 38 ++++++++++++++++--- 1 file changed, 32 insertions(+), 6 deletions(-) diff --git a/numba_dpex/tests/core/runtime/test_llvm_registration.py b/numba_dpex/tests/core/runtime/test_llvm_registration.py index a5801ab401..20c9f9a5c3 100644 --- a/numba_dpex/tests/core/runtime/test_llvm_registration.py +++ b/numba_dpex/tests/core/runtime/test_llvm_registration.py @@ -8,12 +8,8 @@ def test_llvm_symbol_registered(): - """ "Register the helper function in _dpexrt_python so that we can insert calls to them via llvmlite. - - 1. DPEXRT_sycl_usm_ndarray_from_python - - 2.DPEXRT_sycl_usm_ndarray_to_python_acqref - + """Checks if the functions in the _dpexrt_python module are accessible + using llvmlite. """ assert ( llb.address_of_symbol("DPEXRT_sycl_usm_ndarray_from_python") @@ -29,3 +25,33 @@ def test_llvm_symbol_registered(): llb.address_of_symbol("NRT_ExternalAllocator_new_for_usm") == runtime._dpexrt_python.NRT_ExternalAllocator_new_for_usm ) + + assert ( + llb.address_of_symbol("DPEXRT_sycl_queue_from_python") + == runtime._dpexrt_python.DPEXRT_sycl_queue_from_python + ) + + assert ( + llb.address_of_symbol("DPEXRT_sycl_queue_to_python") + == runtime._dpexrt_python.DPEXRT_sycl_queue_to_python + ) + + assert ( + llb.address_of_symbol("DPEXRTQueue_CreateFromFilterString") + == runtime._dpexrt_python.DPEXRTQueue_CreateFromFilterString + ) + + assert ( + llb.address_of_symbol("DpexrtQueue_SubmitRange") + == runtime._dpexrt_python.DpexrtQueue_SubmitRange + ) + + assert ( + llb.address_of_symbol("DPEXRT_MemInfo_alloc") + == runtime._dpexrt_python.DPEXRT_MemInfo_alloc + ) + + assert ( + llb.address_of_symbol("DPEXRT_MemInfo_fill") + == runtime._dpexrt_python.DPEXRT_MemInfo_fill + ) From b2bbbdf080180d3c4b0a3340c1c5f7f9f78d1525 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 14:41:41 -0500 Subject: [PATCH 02/21] Remove redundant unit tests. --- .../tests/core/runtime/test_nrt_python.py | 17 ----------------- numba_dpex/tests/core/runtime/test_runtime.py | 18 ------------------ 2 files changed, 35 deletions(-) delete mode 100644 numba_dpex/tests/core/runtime/test_nrt_python.py delete mode 100644 numba_dpex/tests/core/runtime/test_runtime.py diff --git a/numba_dpex/tests/core/runtime/test_nrt_python.py b/numba_dpex/tests/core/runtime/test_nrt_python.py deleted file mode 100644 index c2c5da28ec..0000000000 --- a/numba_dpex/tests/core/runtime/test_nrt_python.py +++ /dev/null @@ -1,17 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -"""Tests for numba.core.runtime._nrt_python""" - - -def test_c_helpers(): - from numba.core.runtime._nrt_python import c_helpers - - functions = [ - "MemInfo_release", - ] - - for fn_name in functions: - assert fn_name in c_helpers - assert isinstance(c_helpers[fn_name], int) diff --git a/numba_dpex/tests/core/runtime/test_runtime.py b/numba_dpex/tests/core/runtime/test_runtime.py deleted file mode 100644 index 2521d36671..0000000000 --- a/numba_dpex/tests/core/runtime/test_runtime.py +++ /dev/null @@ -1,18 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -from numba_dpex.core.runtime import _dpexrt_python - - -def test_pointers_exposed(): - """This test is to check attributts in _dpexrt_python.""" - - def exposed(function_name): - assert hasattr(_dpexrt_python, function_name) - assert isinstance(getattr(_dpexrt_python, function_name), int) - - exposed("DPEXRT_sycl_usm_ndarray_from_python") - exposed("DPEXRT_sycl_usm_ndarray_to_python_acqref") - exposed("DPEXRT_MemInfo_alloc") - exposed("NRT_ExternalAllocator_new_for_usm") From a036e4e6c57b0621168ccfd8480b84a88b0fb755 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 14:54:28 -0500 Subject: [PATCH 03/21] Remove njit tests for NumPy ufuncs --- .../njit_tests/test_numpy_bitwise_ops.py | 87 ------------- .../tests/njit_tests/test_numpy_logic_ops.py | 94 ------------- .../test_numpy_transcedental_functions.py | 123 ------------------ .../test_numpy_trigonometric_functions.py | 111 ---------------- 4 files changed, 415 deletions(-) delete mode 100644 numba_dpex/tests/njit_tests/test_numpy_bitwise_ops.py delete mode 100644 numba_dpex/tests/njit_tests/test_numpy_logic_ops.py delete mode 100644 numba_dpex/tests/njit_tests/test_numpy_transcedental_functions.py delete mode 100644 numba_dpex/tests/njit_tests/test_numpy_trigonometric_functions.py diff --git a/numba_dpex/tests/njit_tests/test_numpy_bitwise_ops.py b/numba_dpex/tests/njit_tests/test_numpy_bitwise_ops.py deleted file mode 100644 index d9f855d53c..0000000000 --- a/numba_dpex/tests/njit_tests/test_numpy_bitwise_ops.py +++ /dev/null @@ -1,87 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np -import pytest -from numba import njit - -from numba_dpex.tests._helper import assert_auto_offloading, filter_strings - -list_of_binary_ops = [ - "bitwise_and", - "bitwise_or", - "bitwise_xor", - "left_shift", - "right_shift", -] - - -@pytest.fixture(params=list_of_binary_ops) -def binary_op(request): - return request.param - - -list_of_unary_ops = [ - "bitwise_not", - "invert", -] - - -@pytest.fixture(params=list_of_unary_ops) -def unary_op(request): - return request.param - - -list_of_dtypes = [ - np.int32, - np.int64, -] - - -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - # The size of input and out arrays to be used - N = 2048 - a = np.array(np.random.random(N), request.param) - b = np.array(np.random.random(N), request.param) - return a, b - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_binary_ops(filter_str, binary_op, input_arrays): - a, b = input_arrays - binop = getattr(np, binary_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - @njit - def f(a, b): - return binop(a, b) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a, b) - - expected = binop(a, b) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_unary_ops(filter_str, unary_op, input_arrays): - a = input_arrays[0] - uop = getattr(np, unary_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - @njit - def f(a): - return uop(a) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a) - - expected = uop(a) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) diff --git a/numba_dpex/tests/njit_tests/test_numpy_logic_ops.py b/numba_dpex/tests/njit_tests/test_numpy_logic_ops.py deleted file mode 100644 index 1e7be67e15..0000000000 --- a/numba_dpex/tests/njit_tests/test_numpy_logic_ops.py +++ /dev/null @@ -1,94 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np -import pytest -from numba import njit - -from numba_dpex.tests._helper import assert_auto_offloading, filter_strings - -list_of_binary_ops = [ - "greater", - "greater_equal", - "less", - "less_equal", - "not_equal", - "equal", - "logical_and", - "logical_or", - "logical_xor", -] - - -@pytest.fixture(params=list_of_binary_ops) -def binary_op(request): - return request.param - - -list_of_unary_ops = [ - "isinf", - "isfinite", - "isnan", -] - - -@pytest.fixture(params=list_of_unary_ops) -def unary_op(request): - return request.param - - -list_of_dtypes = [ - np.int32, - np.int64, - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - # The size of input and out arrays to be used - N = 2048 - a = np.array(np.random.random(N), request.param) - b = np.array(np.random.random(N), request.param) - return a, b - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_binary_ops(filter_str, binary_op, input_arrays): - a, b = input_arrays - binop = getattr(np, binary_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - @njit - def f(a, b): - return binop(a, b) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a, b) - - expected = binop(a, b) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_unary_ops(filter_str, unary_op, input_arrays): - a = input_arrays[0] - uop = getattr(np, unary_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - @njit - def f(a): - return uop(a) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a) - - expected = uop(a) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) diff --git a/numba_dpex/tests/njit_tests/test_numpy_transcedental_functions.py b/numba_dpex/tests/njit_tests/test_numpy_transcedental_functions.py deleted file mode 100644 index 0fb35160bf..0000000000 --- a/numba_dpex/tests/njit_tests/test_numpy_transcedental_functions.py +++ /dev/null @@ -1,123 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np -import pytest -from numba import njit - -from numba_dpex.tests._helper import ( - assert_auto_offloading, - filter_strings, - is_gen12, -) - -list_of_binary_ops = [ - "add", - "subtract", - "multiply", - "divide", - "true_divide", - "power", - "remainder", - "mod", - "fmod", - "hypot", - "maximum", - "minimum", - "fmax", - "fmin", -] - - -@pytest.fixture(params=list_of_binary_ops) -def binary_op(request): - return request.param - - -list_of_unary_ops = [ - "negative", - "abs", - "absolute", - "fabs", - "sign", - "conj", - "exp", - "exp2", - "log", - "log2", - "log10", - "expm1", - "log1p", - "sqrt", - "square", - "reciprocal", - "conjugate", - "floor", - "ceil", - "trunc", -] - - -@pytest.fixture(params=list_of_unary_ops) -def unary_op(request): - return request.param - - -list_of_dtypes = [ - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - # The size of input and out arrays to be used - N = 2048 - a = np.array(np.random.random(N), request.param) - b = np.array(np.random.random(N), request.param) - return a, b - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_binary_ops(filter_str, binary_op, input_arrays): - a, b = input_arrays - binop = getattr(np, binary_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - @njit - def f(a, b): - return binop(a, b) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a, b) - - expected = binop(a, b) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_unary_ops(filter_str, unary_op, input_arrays): - # FIXME: Why does sign fail on Gen12 discrete graphics card? - skip_ops = ["sign", "log", "log2", "log10", "expm1"] - if unary_op in skip_ops and is_gen12(filter_str): - pytest.skip() - - a = input_arrays[0] - uop = getattr(np, unary_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - @njit - def f(a): - return uop(a) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a) - - expected = uop(a) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) diff --git a/numba_dpex/tests/njit_tests/test_numpy_trigonometric_functions.py b/numba_dpex/tests/njit_tests/test_numpy_trigonometric_functions.py deleted file mode 100644 index d3702e5edf..0000000000 --- a/numba_dpex/tests/njit_tests/test_numpy_trigonometric_functions.py +++ /dev/null @@ -1,111 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np -import pytest -from numba import njit - -from numba_dpex.tests._helper import ( - assert_auto_offloading, - filter_strings, - is_gen12, -) - -list_of_filter_strs = [ - "opencl:gpu:0", - "level_zero:gpu:0", - "opencl:cpu:0", -] - - -@pytest.fixture(params=list_of_filter_strs) -def filter_str(request): - return request.param - - -list_of_trig_ops = [ - "sin", - "cos", - "tan", - "arcsin", - "arccos", - "arctan", - "arctan2", - "sinh", - "cosh", - "tanh", - "arcsinh", - "arccosh", - "arctanh", - "deg2rad", - "rad2deg", - "degrees", - "radians", -] - - -@pytest.fixture(params=list_of_trig_ops) -def trig_op(request): - return request.param - - -list_of_dtypes = [ - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_trig_ops) -def dtype(request): - return request.param - - -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - # The size of input and out arrays to be used - N = 2048 - # Note: These inputs do not work for all of the functions and - # can result in warnings. E.g. arccosh needs the range of values - # to be greater than 0, while arccos needs them to be [-1,1]. - # These warnings are relatively benign as NumPy will return "nan" - # for such cases. - a = np.array(np.random.random(N), request.param) - b = np.array(np.random.random(N), request.param) - return a, b - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_trigonometric_fn(filter_str, trig_op, input_arrays): - # FIXME: Why does archcosh fail on Gen12 discrete graphics card? - if trig_op == "arccosh" and is_gen12(filter_str): - pytest.skip() - - a, b = input_arrays - trig_fn = getattr(np, trig_op) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - if trig_op == "arctan2": - - @njit - def f(a, b): - return trig_fn(a, b) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a, b) - expected = trig_fn(a, b) - else: - - @njit - def f(a): - return trig_fn(a) - - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), assert_auto_offloading(): - actual = f(a) - expected = trig_fn(a) - - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) From fc1355e5b89e0474f594df35766bfe372817a94e Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:07:04 -0500 Subject: [PATCH 04/21] Remove test_numpy_array_creation.py, superseded by dpjit_tests/dpnp/* --- .../dpnp/test_numpy_array_creation.py | 137 ------------------ 1 file changed, 137 deletions(-) delete mode 100644 numba_dpex/tests/njit_tests/dpnp/test_numpy_array_creation.py diff --git a/numba_dpex/tests/njit_tests/dpnp/test_numpy_array_creation.py b/numba_dpex/tests/njit_tests/dpnp/test_numpy_array_creation.py deleted file mode 100644 index 81ea27cc7a..0000000000 --- a/numba_dpex/tests/njit_tests/dpnp/test_numpy_array_creation.py +++ /dev/null @@ -1,137 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np -import pytest -from numba import njit - -from numba_dpex.tests._helper import dpnp_debug, filter_strings, skip_no_dpnp - -from ._helper import args_string, wrapper_function - -pytestmark = skip_no_dpnp - -list_of_dtypes = [ - np.int32, - np.int64, - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_dtypes) -def input_array(request): - # The size of input and out arrays to be used - N = 10 - a = np.array(np.random.random(N), request.param) - return a - - -list_of_shape = [ - (10), - (5, 2), -] - - -@pytest.fixture(params=list_of_shape) -def get_shape(request): - return request.param - - -list_of_unary_op = [ - "copy", - "trace", -] - -list_of_binary_op = [ - "ones_like", - "zeros_like", -] - - -@pytest.fixture(params=list_of_unary_op) -def unary_op(request): - return request.param - - -@pytest.fixture(params=list_of_binary_op) -def binary_op(request): - return request.param - - -def get_op_fn(name, nargs): - args = args_string(nargs) - return wrapper_function(args, f"np.{name}({args})", globals()) - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_unary_ops(filter_str, unary_op, input_array, capfd): - a = input_array - if unary_op == "trace": - a = input_array.reshape((2, 5)) - fn = get_op_fn(unary_op, 1) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - f = njit(fn) - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), dpnp_debug(): - actual = f(a) - captured = capfd.readouterr() - assert "dpnp implementation" in captured.out - - expected = fn(a) - np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0) - - -@pytest.fixture(params=list_of_dtypes + [None]) -def dtype(request): - return request.param - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_binary_op(filter_str, binary_op, input_array, dtype, get_shape, capfd): - a = np.reshape(input_array, get_shape) - fn = get_op_fn(binary_op, 2) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - f = njit(fn) - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), dpnp_debug(): - actual = f(a, dtype) - captured = capfd.readouterr() - assert "dpnp implementation" in captured.out - - expected = fn(a, dtype) - np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0) - - -list_of_full = [ - "full_like", -] - - -@pytest.fixture(params=list_of_full) -def full_name(request): - return request.param - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_full(filter_str, full_name, input_array, get_shape, capfd): - a = np.reshape(input_array, get_shape) - fn = get_op_fn(full_name, 2) - actual = np.empty(shape=a.shape, dtype=a.dtype) - expected = np.empty(shape=a.shape, dtype=a.dtype) - - f = njit(fn) - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device), dpnp_debug(): - actual = f(a, np.array([2])) - captured = capfd.readouterr() - assert "dpnp implementation" in captured.out - - expected = fn(a, np.array([2])) - np.testing.assert_allclose(actual, expected, rtol=1e-3, atol=0) From c0c8d1bc1e7486c73c11343af17a64d6581a20f3 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:10:09 -0500 Subject: [PATCH 05/21] Move all tests for dpex types into core/types. --- .../tests/{ => core}/types/DpctlSyclQueue/test_box_unbox.py | 0 .../tests/{ => core}/types/DpctlSyclQueue/test_queue_ref_attr.py | 0 .../dpnp_ndarray => core/types/DpnpNdArray}/test_boxing.py | 0 .../dpnp_ndarray => core/types/DpnpNdArray}/test_bugs.py | 0 .../types/DpnpNdArray}/test_dpnp_ndarray_type.py | 0 .../dpnp_ndarray => core/types/DpnpNdArray}/test_models.py | 0 6 files changed, 0 insertions(+), 0 deletions(-) rename numba_dpex/tests/{ => core}/types/DpctlSyclQueue/test_box_unbox.py (100%) rename numba_dpex/tests/{ => core}/types/DpctlSyclQueue/test_queue_ref_attr.py (100%) rename numba_dpex/tests/{njit_tests/dpnp_ndarray => core/types/DpnpNdArray}/test_boxing.py (100%) rename numba_dpex/tests/{njit_tests/dpnp_ndarray => core/types/DpnpNdArray}/test_bugs.py (100%) rename numba_dpex/tests/{njit_tests/dpnp_ndarray => core/types/DpnpNdArray}/test_dpnp_ndarray_type.py (100%) rename numba_dpex/tests/{njit_tests/dpnp_ndarray => core/types/DpnpNdArray}/test_models.py (100%) diff --git a/numba_dpex/tests/types/DpctlSyclQueue/test_box_unbox.py b/numba_dpex/tests/core/types/DpctlSyclQueue/test_box_unbox.py similarity index 100% rename from numba_dpex/tests/types/DpctlSyclQueue/test_box_unbox.py rename to numba_dpex/tests/core/types/DpctlSyclQueue/test_box_unbox.py diff --git a/numba_dpex/tests/types/DpctlSyclQueue/test_queue_ref_attr.py b/numba_dpex/tests/core/types/DpctlSyclQueue/test_queue_ref_attr.py similarity index 100% rename from numba_dpex/tests/types/DpctlSyclQueue/test_queue_ref_attr.py rename to numba_dpex/tests/core/types/DpctlSyclQueue/test_queue_ref_attr.py diff --git a/numba_dpex/tests/njit_tests/dpnp_ndarray/test_boxing.py b/numba_dpex/tests/core/types/DpnpNdArray/test_boxing.py similarity index 100% rename from numba_dpex/tests/njit_tests/dpnp_ndarray/test_boxing.py rename to numba_dpex/tests/core/types/DpnpNdArray/test_boxing.py diff --git a/numba_dpex/tests/njit_tests/dpnp_ndarray/test_bugs.py b/numba_dpex/tests/core/types/DpnpNdArray/test_bugs.py similarity index 100% rename from numba_dpex/tests/njit_tests/dpnp_ndarray/test_bugs.py rename to numba_dpex/tests/core/types/DpnpNdArray/test_bugs.py diff --git a/numba_dpex/tests/njit_tests/dpnp_ndarray/test_dpnp_ndarray_type.py b/numba_dpex/tests/core/types/DpnpNdArray/test_dpnp_ndarray_type.py similarity index 100% rename from numba_dpex/tests/njit_tests/dpnp_ndarray/test_dpnp_ndarray_type.py rename to numba_dpex/tests/core/types/DpnpNdArray/test_dpnp_ndarray_type.py diff --git a/numba_dpex/tests/njit_tests/dpnp_ndarray/test_models.py b/numba_dpex/tests/core/types/DpnpNdArray/test_models.py similarity index 100% rename from numba_dpex/tests/njit_tests/dpnp_ndarray/test_models.py rename to numba_dpex/tests/core/types/DpnpNdArray/test_models.py From 9244dda3f477c77c714949c149eecddec555c149 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:16:48 -0500 Subject: [PATCH 06/21] Move test_dump_kernel_llvm.py into kernel_tests. --- numba_dpex/tests/{ => kernel_tests}/test_dump_kernel_llvm.py | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename numba_dpex/tests/{ => kernel_tests}/test_dump_kernel_llvm.py (100%) diff --git a/numba_dpex/tests/test_dump_kernel_llvm.py b/numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py similarity index 100% rename from numba_dpex/tests/test_dump_kernel_llvm.py rename to numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py From 4774ffa17c9380314339b45be561d12e633b1e3a Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:18:58 -0500 Subject: [PATCH 07/21] Move usm_ndarray tests into own core/types/USMNdArray. --- .../tests/{ => core/types/USMNdAArray}/test_usm_ndarray_type.py | 0 .../types/USMNdAArray}/test_usm_ndarray_type_exceptions.py | 0 2 files changed, 0 insertions(+), 0 deletions(-) rename numba_dpex/tests/{ => core/types/USMNdAArray}/test_usm_ndarray_type.py (100%) rename numba_dpex/tests/{ => core/types/USMNdAArray}/test_usm_ndarray_type_exceptions.py (100%) diff --git a/numba_dpex/tests/test_usm_ndarray_type.py b/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type.py similarity index 100% rename from numba_dpex/tests/test_usm_ndarray_type.py rename to numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type.py diff --git a/numba_dpex/tests/test_usm_ndarray_type_exceptions.py b/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type_exceptions.py similarity index 100% rename from numba_dpex/tests/test_usm_ndarray_type_exceptions.py rename to numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type_exceptions.py From 1bae9cd9f8c2298b2efed8d8d81e6d3a549f8b5b Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:24:55 -0500 Subject: [PATCH 08/21] Move more kernel API tests into kernel_tests --- numba_dpex/tests/kernel_tests/test_func.py | 29 +++++++++ .../tests/kernel_tests/test_private_memory.py | 2 - numba_dpex/tests/test_dppy_func.py | 61 ------------------- 3 files changed, 29 insertions(+), 63 deletions(-) create mode 100644 numba_dpex/tests/kernel_tests/test_func.py delete mode 100644 numba_dpex/tests/test_dppy_func.py diff --git a/numba_dpex/tests/kernel_tests/test_func.py b/numba_dpex/tests/kernel_tests/test_func.py new file mode 100644 index 0000000000..6fc0226f30 --- /dev/null +++ b/numba_dpex/tests/kernel_tests/test_func.py @@ -0,0 +1,29 @@ +# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + + +import dpnp +import numpy + +import numba_dpex as dpex + + +@dpex.func +def g(a): + return a + 1 + + +@dpex.kernel +def f(a, b): + i = dpex.get_global_id(0) + b[i] = g(a[i]) + + +def test_func_call_from_kernel(): + a = dpnp.ones(1024) + b = dpnp.ones(1024) + + f[dpex.Range(1024)](a, b) + nb = dpnp.asnumpy(b) + assert numpy.all(nb == 2) diff --git a/numba_dpex/tests/kernel_tests/test_private_memory.py b/numba_dpex/tests/kernel_tests/test_private_memory.py index 42ddad8f54..74f95afbec 100644 --- a/numba_dpex/tests/kernel_tests/test_private_memory.py +++ b/numba_dpex/tests/kernel_tests/test_private_memory.py @@ -2,9 +2,7 @@ # # SPDX-License-Identifier: Apache-2.0 -import platform -import dpctl import numpy as np import pytest diff --git a/numba_dpex/tests/test_dppy_func.py b/numba_dpex/tests/test_dppy_func.py deleted file mode 100644 index abf62d271a..0000000000 --- a/numba_dpex/tests/test_dppy_func.py +++ /dev/null @@ -1,61 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np - -import numba_dpex as dpex -from numba_dpex.tests._helper import skip_no_opencl_gpu - - -@skip_no_opencl_gpu -class TestFunc: - N = 257 - - def test_func_device_array(self): - @dpex.func - def g(a): - return a + 1 - - @dpex.kernel - def f(a, b): - i = dpex.get_global_id(0) - b[i] = g(a[i]) - - a = np.ones(self.N) - b = np.ones(self.N) - - device = dpctl.SyclDevice("opencl:gpu") - with dpctl.device_context(device): - f[self.N, dpex.DEFAULT_LOCAL_SIZE](a, b) - - assert np.all(b == 2) - - def test_func_ndarray(self): - @dpex.func - def g(a): - return a + 1 - - @dpex.kernel - def f(a, b): - i = dpex.get_global_id(0) - b[i] = g(a[i]) - - @dpex.kernel - def h(a, b): - i = dpex.get_global_id(0) - b[i] = g(a[i]) + 1 - - a = np.ones(self.N) - b = np.ones(self.N) - - device = dpctl.SyclDevice("opencl:gpu") - with dpctl.device_context(device): - f[self.N, dpex.DEFAULT_LOCAL_SIZE](a, b) - - assert np.all(b == 2) - - h[self.N, dpex.DEFAULT_LOCAL_SIZE](a, b) - - assert np.all(b == 3) From 99a7c07e4c844e202ae4cd06739a2393949fdcbf Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:27:34 -0500 Subject: [PATCH 09/21] Move passes unit tests into core. --- .../tests/{ => core}/passes/test_parfor_legalize_cfd_pass.py | 0 .../tests/{ => core}/passes/test_rename_numpy_function_pass.py | 0 2 files changed, 0 insertions(+), 0 deletions(-) rename numba_dpex/tests/{ => core}/passes/test_parfor_legalize_cfd_pass.py (100%) rename numba_dpex/tests/{ => core}/passes/test_rename_numpy_function_pass.py (100%) diff --git a/numba_dpex/tests/passes/test_parfor_legalize_cfd_pass.py b/numba_dpex/tests/core/passes/test_parfor_legalize_cfd_pass.py similarity index 100% rename from numba_dpex/tests/passes/test_parfor_legalize_cfd_pass.py rename to numba_dpex/tests/core/passes/test_parfor_legalize_cfd_pass.py diff --git a/numba_dpex/tests/passes/test_rename_numpy_function_pass.py b/numba_dpex/tests/core/passes/test_rename_numpy_function_pass.py similarity index 100% rename from numba_dpex/tests/passes/test_rename_numpy_function_pass.py rename to numba_dpex/tests/core/passes/test_rename_numpy_function_pass.py From 7c6895c508c4363e024dae3e34d1b7c404aae599 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:50:20 -0500 Subject: [PATCH 10/21] Move relevant kernel interop tests into kernel_tests. --- numba_dpex/tests/integration/DuckUSMArray.py | 49 -------- .../tests/integration/test_dpnp_interop.py | 41 ------ .../test_sycl_usm_array_iface_interop.py | 80 ------------ .../integration/test_usm_ndarray_interop.py | 88 ------------- .../test_sycl_usm_array_iface_interop.py | 117 ++++++++++++++++++ .../kernel_tests/test_usm_ndarray_interop.py | 63 ++++++++++ 6 files changed, 180 insertions(+), 258 deletions(-) delete mode 100644 numba_dpex/tests/integration/DuckUSMArray.py delete mode 100644 numba_dpex/tests/integration/test_dpnp_interop.py delete mode 100644 numba_dpex/tests/integration/test_sycl_usm_array_iface_interop.py delete mode 100644 numba_dpex/tests/integration/test_usm_ndarray_interop.py create mode 100644 numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py create mode 100644 numba_dpex/tests/kernel_tests/test_usm_ndarray_interop.py diff --git a/numba_dpex/tests/integration/DuckUSMArray.py b/numba_dpex/tests/integration/DuckUSMArray.py deleted file mode 100644 index 1f7831f4e5..0000000000 --- a/numba_dpex/tests/integration/DuckUSMArray.py +++ /dev/null @@ -1,49 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl.memory as dpmem -import dpctl.tensor as dpt -import numpy as np - - -class DuckUSMArray: - """A Python class that defines a __sycl_usm_array_interface__ attribute.""" - - def __init__(self, shape, dtype="d", host_buffer=None): - _tensor = dpt.empty(shape, dtype=dtype, usm_type="shared") - shmem = _tensor.usm_data - if isinstance(host_buffer, np.ndarray): - shmem.copy_from_host(host_buffer.view(dtype="|u1")) - self.arr = np.ndarray(shape, dtype=dtype, buffer=shmem) - - def __getitem__(self, indx): - return self.arr[indx] - - def __setitem__(self, indx, val): - self.arr.__setitem__(indx, val) - - @property - def __sycl_usm_array_interface__(self): - iface = self.arr.__array_interface__ - b = self.arr.base - iface["syclobj"] = b.__sycl_usm_array_interface__["syclobj"] - iface["version"] = 1 - return iface - - -class PseudoDuckUSMArray: - """A Python class that defines an attributed called - __sycl_usm_array_interface__, but is not actually backed by USM memory. - - """ - - def __init__(self): - pass - - @property - def __sycl_usm_array_interface__(self): - iface = {} - iface["syclobj"] = None - iface["version"] = 0 - return iface diff --git a/numba_dpex/tests/integration/test_dpnp_interop.py b/numba_dpex/tests/integration/test_dpnp_interop.py deleted file mode 100644 index f406a22e20..0000000000 --- a/numba_dpex/tests/integration/test_dpnp_interop.py +++ /dev/null @@ -1,41 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import numpy as np -import pytest - -import numba_dpex as dpex - -dpnp = pytest.importorskip("dpnp", reason="DPNP is not installed") - - -list_of_dtype = [ - np.int32, - np.int64, - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_dtype) -def dtype(request): - return request.param - - -def test_consuming_array_from_dpnp(dtype): - @dpex.kernel - def data_parallel_sum(a, b, c): - """ - Vector addition using the ``kernel`` decorator. - """ - i = dpex.get_global_id(0) - c[i] = a[i] + b[i] - - global_size = 1021 - - a = dpnp.arange(global_size, dtype=dtype) - b = dpnp.arange(global_size, dtype=dtype) - c = dpnp.ones_like(a) - - data_parallel_sum[global_size](a, b, c) diff --git a/numba_dpex/tests/integration/test_sycl_usm_array_iface_interop.py b/numba_dpex/tests/integration/test_sycl_usm_array_iface_interop.py deleted file mode 100644 index a0428125ce..0000000000 --- a/numba_dpex/tests/integration/test_sycl_usm_array_iface_interop.py +++ /dev/null @@ -1,80 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import DuckUSMArray as duckArrs -import numpy as np -import pytest - -import numba_dpex as dpex - - -@dpex.kernel -def vecadd(a, b, c): - i = dpex.get_global_id(0) - c[i] = a[i] + b[i] - - -dtypes = [ - "i4", - "i8", - "f4", - "f8", -] - - -@pytest.fixture(params=dtypes) -def dtype(request): - return request.param - - -def test_kernel_valid_usm_obj(dtype): - """Test if a ``numba_dpex.kernel`` function accepts a DuckUSMArray argument. - - The ``DuckUSMArray`` uses ``dpctl.memory`` to allocate a Python object that - defines a ``__sycl_usm_array_interface__`` attribute. We test if - ``numba_dpex`` recognizes the ``DuckUSMArray`` as a valid USM-backed Python - object and accepts it as a kernel argument. - - """ - N = 1024 - - buffA = np.arange(0, N, dtype=dtype) - A = duckArrs.DuckUSMArray(shape=buffA.shape, dtype=dtype, host_buffer=buffA) - - buffB = np.arange(0, N, dtype=dtype) - B = duckArrs.DuckUSMArray(shape=buffB.shape, dtype=dtype, host_buffer=buffB) - - buffC = np.zeros(N, dtype=dtype) - C = duckArrs.DuckUSMArray(shape=buffC.shape, dtype=dtype, host_buffer=buffC) - - try: - vecadd[N, dpex.DEFAULT_LOCAL_SIZE](A, B, C) - except Exception: - pytest.fail( - "Could not pass Python object with sycl_usm_array_interface" - + " to a kernel." - ) - - -def test_kernel_invalid_usm_obj(dtype): - """Test if a ``numba_dpex.kernel`` function rejects a PseudoDuckUSMArray - argument. - - The ``PseudoDuckUSMArray`` defines a fake attribute called - __sycl_usm_array__interface__. We test if - ``numba_dpex`` correctly recognizes and rejects the ``PseudoDuckUSMArray``. - - """ - N = 1024 - - A = duckArrs.PseudoDuckUSMArray() - - B = duckArrs.PseudoDuckUSMArray() - - C = duckArrs.PseudoDuckUSMArray() - - with pytest.raises(Exception): - with dpctl.device_context(dpctl.select_default_device()): - vecadd[N, dpex.DEFAULT_LOCAL_SIZE](A, B, C) diff --git a/numba_dpex/tests/integration/test_usm_ndarray_interop.py b/numba_dpex/tests/integration/test_usm_ndarray_interop.py deleted file mode 100644 index 1f79799008..0000000000 --- a/numba_dpex/tests/integration/test_usm_ndarray_interop.py +++ /dev/null @@ -1,88 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import dpctl.tensor as dpt -import numpy as np -import pytest -from numba import njit - -import numba_dpex as dpex -from numba_dpex.tests._helper import filter_strings - -list_of_dtype = [ - np.int32, - np.int64, - np.float32, - np.float64, -] - - -@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 - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_consuming_usm_ndarray(filter_str, dtype, usm_type): - @dpex.kernel - def data_parallel_sum(a, b, c): - """ - Vector addition using the ``kernel`` decorator. - """ - i = dpex.get_global_id(0) - j = dpex.get_global_id(1) - c[i, j] = a[i, j] + b[i, j] - - N = 1021 - global_size = N * N - - a = np.array(np.random.random(global_size), dtype=dtype).reshape(N, N) - b = np.array(np.random.random(global_size), dtype=dtype).reshape(N, N) - - got = np.ones_like(a) - - with dpctl.device_context(filter_str) as gpu_queue: - da = dpt.usm_ndarray( - a.shape, - dtype=a.dtype, - buffer=usm_type, - buffer_ctor_kwargs={"queue": gpu_queue}, - ) - da.usm_data.copy_from_host(a.reshape((-1)).view("|u1")) - - db = dpt.usm_ndarray( - b.shape, - dtype=b.dtype, - buffer=usm_type, - buffer_ctor_kwargs={"queue": gpu_queue}, - ) - db.usm_data.copy_from_host(b.reshape((-1)).view("|u1")) - - dc = dpt.usm_ndarray( - got.shape, - dtype=got.dtype, - buffer=usm_type, - buffer_ctor_kwargs={"queue": gpu_queue}, - ) - - data_parallel_sum[(N, N), dpex.DEFAULT_LOCAL_SIZE](da, db, dc) - - dc.usm_data.copy_to_host(got.reshape((-1)).view("|u1")) - - expected = a + b - - assert np.array_equal(got, expected) 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 new file mode 100644 index 0000000000..0ee27fba9f --- /dev/null +++ b/numba_dpex/tests/kernel_tests/test_sycl_usm_array_iface_interop.py @@ -0,0 +1,117 @@ +# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + +import numpy as np +import pytest +from dpctl import tensor as dpt + +import numba_dpex as dpex + + +class DuckUSMArray: + """A Python class that defines a __sycl_usm_array_interface__ attribute.""" + + def __init__(self, shape, dtype="d", host_buffer=None): + _tensor = dpt.empty(shape, dtype=dtype, usm_type="shared") + shmem = _tensor.usm_data + if isinstance(host_buffer, np.ndarray): + shmem.copy_from_host(host_buffer.view(dtype="|u1")) + self.arr = np.ndarray(shape, dtype=dtype, buffer=shmem) + + def __getitem__(self, indx): + return self.arr[indx] + + def __setitem__(self, indx, val): + self.arr.__setitem__(indx, val) + + @property + def __sycl_usm_array_interface__(self): + iface = self.arr.__array_interface__ + b = self.arr.base + iface["syclobj"] = b.__sycl_usm_array_interface__["syclobj"] + iface["version"] = 1 + return iface + + +class PseudoDuckUSMArray: + """A Python class that defines an attributed called + __sycl_usm_array_interface__, but is not actually backed by USM memory. + + """ + + def __init__(self): + pass + + @property + def __sycl_usm_array_interface__(self): + iface = {} + iface["syclobj"] = None + iface["version"] = 0 + return iface + + +@dpex.kernel +def vecadd(a, b, c): + i = dpex.get_global_id(0) + c[i] = a[i] + b[i] + + +dtypes = [ + "i4", + "i8", + "f4", + "f8", +] + + +@pytest.fixture(params=dtypes) +def dtype(request): + return request.param + + +def test_kernel_valid_usm_obj(dtype): + """Test if a ``numba_dpex.kernel`` function accepts a DuckUSMArray argument. + + The ``DuckUSMArray`` uses ``dpctl.memory`` to allocate a Python object that + defines a ``__sycl_usm_array_interface__`` attribute. We test if + ``numba_dpex`` recognizes the ``DuckUSMArray`` as a valid USM-backed Python + object and accepts it as a kernel argument. + + """ + N = 1024 + + buffA = np.arange(0, N, dtype=dtype) + buffB = np.arange(0, N, dtype=dtype) + buffC = np.zeros(N, dtype=dtype) + + A = DuckUSMArray(shape=buffA.shape, dtype=dtype, host_buffer=buffA) + B = DuckUSMArray(shape=buffB.shape, dtype=dtype, host_buffer=buffB) + C = DuckUSMArray(shape=buffC.shape, dtype=dtype, host_buffer=buffC) + + try: + vecadd[dpex.Range(N)](A, B, C) + except Exception: + pytest.fail( + "Could not pass Python object with sycl_usm_array_interface" + + " to a kernel." + ) + + +def test_kernel_invalid_usm_obj(dtype): + """Test if a ``numba_dpex.kernel`` function rejects a PseudoDuckUSMArray + argument. + + The ``PseudoDuckUSMArray`` defines a fake attribute called + __sycl_usm_array__interface__. We test if + ``numba_dpex`` correctly recognizes and rejects the ``PseudoDuckUSMArray``. + + """ + N = 1024 + + A = PseudoDuckUSMArray() + B = PseudoDuckUSMArray() + C = PseudoDuckUSMArray() + + with pytest.raises(Exception): + 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 new file mode 100644 index 0000000000..6289ed1679 --- /dev/null +++ b/numba_dpex/tests/kernel_tests/test_usm_ndarray_interop.py @@ -0,0 +1,63 @@ +# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + +import dpctl.tensor as dpt +import numpy +import pytest + +import numba_dpex as dpex + +list_of_dtype = [ + numpy.int32, + numpy.int64, + numpy.float32, + numpy.float64, +] + + +@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(a, b, c): + """ + Vector addition using the ``kernel`` decorator. + """ + i = dpex.get_global_id(0) + j = dpex.get_global_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) + + 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) From 847885e9199c591924693e645e8a423d4d61716d Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:51:33 -0500 Subject: [PATCH 11/21] Parfor fallback is about to be removed. Delete unit tests. --- numba_dpex/tests/test_dppy_fallback.py | 62 -------------------------- 1 file changed, 62 deletions(-) delete mode 100644 numba_dpex/tests/test_dppy_fallback.py diff --git a/numba_dpex/tests/test_dppy_fallback.py b/numba_dpex/tests/test_dppy_fallback.py deleted file mode 100644 index 8acb025435..0000000000 --- a/numba_dpex/tests/test_dppy_fallback.py +++ /dev/null @@ -1,62 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import warnings - -import dpctl -import numba -import numpy as np -import pytest - -from numba_dpex.tests._helper import skip_no_opencl_gpu - - -@skip_no_opencl_gpu -class TestFallback: - def test_fallback_inner_call(self): - @numba.jit - def fill_value(i): - return i - - def inner_call_fallback(): - x = 10 - a = np.empty(shape=x, dtype=np.float32) - - for i in numba.prange(x): - a[i] = fill_value(i) - - return a - - device = dpctl.SyclDevice("opencl:gpu") - with warnings.catch_warnings(record=True) as w, dpctl.device_context( - device - ): - fn = numba.njit(inner_call_fallback) - result = fn() - - ref_result = inner_call_fallback() - - np.testing.assert_array_equal(result, ref_result) - assert "Failed to offload parfor " in str(w[-1].message) - - @pytest.mark.skip - def test_fallback_reductions(self): - def reduction(a): - b = 1 - for i in numba.prange(len(a)): - b += a[i] - return b - - a = np.ones(10) - device = dpctl.SyclDevice("opencl:gpu") - with warnings.catch_warnings(record=True) as w, dpctl.device_context( - device - ): - fn = numba.njit(reduction) - result = fn(a) - - ref_result = reduction(a) - - np.testing.assert_array_equal(result, ref_result) - assert "Failed to offload parfor " in str(w[-1].message) From ddf82e9be014954447e6819a2555a47e45f18910 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:54:05 -0500 Subject: [PATCH 12/21] Move test mangler tests into core. --- numba_dpex/tests/{ => core}/test_itanium_mangler_extension.py | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename numba_dpex/tests/{ => core}/test_itanium_mangler_extension.py (100%) diff --git a/numba_dpex/tests/test_itanium_mangler_extension.py b/numba_dpex/tests/core/test_itanium_mangler_extension.py similarity index 100% rename from numba_dpex/tests/test_itanium_mangler_extension.py rename to numba_dpex/tests/core/test_itanium_mangler_extension.py From 3abfe121d2f9e91d8940ea0f835fa8225135379b Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 20:57:04 -0500 Subject: [PATCH 13/21] Move version parsing tests into misc. --- numba_dpex/tests/misc/test_dpctl_version.py | 15 +++++++++++ .../{ => misc}/test_parse_sem_version.py | 0 numba_dpex/tests/test_dpctl_api.py | 27 ------------------- 3 files changed, 15 insertions(+), 27 deletions(-) create mode 100644 numba_dpex/tests/misc/test_dpctl_version.py rename numba_dpex/tests/{ => misc}/test_parse_sem_version.py (100%) delete mode 100644 numba_dpex/tests/test_dpctl_api.py diff --git a/numba_dpex/tests/misc/test_dpctl_version.py b/numba_dpex/tests/misc/test_dpctl_version.py new file mode 100644 index 0000000000..d723264de4 --- /dev/null +++ b/numba_dpex/tests/misc/test_dpctl_version.py @@ -0,0 +1,15 @@ +# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + +import dpctl + +from numba_dpex import dpctl_version + + +def test_dpctl_version(): + dpctl_v = dpctl.__version__ + computed_v = ".".join(str(n) for n in dpctl_version) + n = len(computed_v) + assert n <= len(dpctl_v) + assert computed_v == dpctl_v[:n] diff --git a/numba_dpex/tests/test_parse_sem_version.py b/numba_dpex/tests/misc/test_parse_sem_version.py similarity index 100% rename from numba_dpex/tests/test_parse_sem_version.py rename to numba_dpex/tests/misc/test_parse_sem_version.py diff --git a/numba_dpex/tests/test_dpctl_api.py b/numba_dpex/tests/test_dpctl_api.py deleted file mode 100644 index 2a9cc8eeb0..0000000000 --- a/numba_dpex/tests/test_dpctl_api.py +++ /dev/null @@ -1,27 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import pytest - -from numba_dpex import dpctl_version -from numba_dpex.tests._helper import filter_strings - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_dpctl_api(filter_str): - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device): - dpctl.lsplatform() - dpctl.get_current_queue() - dpctl.get_num_activated_queues() - dpctl.is_in_device_context() - - -def test_dpctl_version(): - dpctl_v = dpctl.__version__ - computed_v = ".".join(str(n) for n in dpctl_version) - n = len(computed_v) - assert n <= len(dpctl_v) - assert computed_v == dpctl_v[:n] From 97bfdb4e93659fda5df9ef5b196777c8700ce15f Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 22:07:58 -0500 Subject: [PATCH 14/21] Remove test for deprecated feature. --- .../tests/kernel_tests/test_arg_accessor.py | 60 ------------------- 1 file changed, 60 deletions(-) delete mode 100644 numba_dpex/tests/kernel_tests/test_arg_accessor.py diff --git a/numba_dpex/tests/kernel_tests/test_arg_accessor.py b/numba_dpex/tests/kernel_tests/test_arg_accessor.py deleted file mode 100644 index 31ab39f4e1..0000000000 --- a/numba_dpex/tests/kernel_tests/test_arg_accessor.py +++ /dev/null @@ -1,60 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import dpctl -import numpy as np -import pytest - -import numba_dpex as dpex -from numba_dpex.tests._helper import filter_strings - - -def call_kernel(global_size, local_size, A, B, C, func): - func[global_size, local_size](A, B, C) - - -global_size = 10 -local_size = 1 -N = global_size * local_size - - -def sum_kernel(a, b, c): - i = dpex.get_global_id(0) - c[i] = a[i] + b[i] - - -list_of_dtypes = [ - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - # The size of input and out arrays to be used - a = np.array(np.random.random(N), request.param) - b = np.array(np.random.random(N), request.param) - c = np.zeros_like(a) - return a, b, c - - -list_of_kernel_opt = [ - {"read_only": ["a", "b"], "write_only": ["c"], "read_write": []}, - {}, -] - - -@pytest.fixture(params=list_of_kernel_opt) -def kernel(request): - return dpex.kernel(access_types=request.param)(sum_kernel) - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_kernel_arg_accessor(filter_str, input_arrays, kernel): - a, b, actual = input_arrays - expected = a + b - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device): - call_kernel(global_size, local_size, a, b, actual, kernel) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) From 2f458d64800768783d5a44ac2ba7dbb2b0b35bbc Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 22:27:56 -0500 Subject: [PATCH 15/21] Fix test_arg_types.py --- .../tests/kernel_tests/test_arg_types.py | 106 ------------------ .../kernel_tests/test_scalar_arg_types.py | 85 ++++++++++++++ 2 files changed, 85 insertions(+), 106 deletions(-) delete mode 100644 numba_dpex/tests/kernel_tests/test_arg_types.py create mode 100644 numba_dpex/tests/kernel_tests/test_scalar_arg_types.py diff --git a/numba_dpex/tests/kernel_tests/test_arg_types.py b/numba_dpex/tests/kernel_tests/test_arg_types.py deleted file mode 100644 index 14f8689b47..0000000000 --- a/numba_dpex/tests/kernel_tests/test_arg_types.py +++ /dev/null @@ -1,106 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -import sys - -import dpctl -import dpctl.tensor as dpt -import numpy as np -import pytest - -import numba_dpex as dpex - -global_size = 1054 -local_size = 1 -N = global_size * local_size - - -def mul_kernel(a, b, c): - i = dpex.get_global_id(0) - b[i] = a[i] * c - - -list_of_dtypes = [ - np.int32, - np.int64, - np.float32, - np.float64, -] - - -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - a = np.array(np.random.random(N), request.param) - b = np.empty_like(a, request.param) - c = np.array([2], request.param) - return a, b, c[0] - - -def test_kernel_arg_types(input_arrays): - usm_type = "device" - - a, b, c = input_arrays - expected = a * c - - queue = dpctl.SyclQueue(dpctl.select_default_device()) - - da = dpt.usm_ndarray( - a.shape, - dtype=a.dtype, - buffer=usm_type, - buffer_ctor_kwargs={"queue": queue}, - ) - da.usm_data.copy_from_host(a.reshape((-1)).view("|u1")) - - db = dpt.usm_ndarray( - b.shape, - dtype=b.dtype, - buffer=usm_type, - buffer_ctor_kwargs={"queue": queue}, - ) - db.usm_data.copy_from_host(b.reshape((-1)).view("|u1")) - - kernel = dpex.kernel(mul_kernel) - kernel[dpex.NdRange(dpex.Range(global_size), dpex.Range(local_size))]( - da, db, c - ) - - result = np.zeros_like(b) - db.usm_data.copy_to_host(result.reshape((-1)).view("|u1")) - - np.testing.assert_allclose(result, expected, rtol=1e-5, atol=0) - - -def check_bool_kernel(A, test): - if test: - A[0] = 111 - else: - A[0] = 222 - - -def test_bool_type(): - usm_type = "device" - a = np.array([2], np.int64) - - queue = dpctl.SyclQueue(dpctl.select_default_device()) - - da = dpt.usm_ndarray( - a.shape, - dtype=a.dtype, - buffer=usm_type, - buffer_ctor_kwargs={"queue": queue}, - ) - da.usm_data.copy_from_host(a.reshape((-1)).view("|u1")) - - kernel = dpex.kernel(check_bool_kernel) - - kernel[dpex.Range(a.size)](da, True) - result = np.zeros_like(a) - da.usm_data.copy_to_host(result.reshape((-1)).view("|u1")) - assert result[0] == 111 - - kernel[dpex.Range(a.size)](da, False) - result = np.zeros_like(a) - da.usm_data.copy_to_host(result.reshape((-1)).view("|u1")) - assert result[0] == 222 diff --git a/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py b/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py new file mode 100644 index 0000000000..7c08667e16 --- /dev/null +++ b/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py @@ -0,0 +1,85 @@ +# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + +import dpctl +import dpctl.tensor as dpt +import dpnp +import numpy +import pytest + +import numba_dpex as dpex + +N = 1024 + + +@dpex.kernel +def scaling_kernel(a, b, c): + i = dpex.get_global_id(0) + b[i] = a[i] * c + + +@dpex.kernel +def kernel_with_bool_arg(a, b, test): + i = dpex.get_global_id(0) + if test: + b[i] = a[i] + a[i] + else: + b[i] = a[i] - a[i] + + +list_of_dtypes = [ + dpnp.int32, + dpnp.int64, + dpnp.float32, + dpnp.float64, +] + +list_of_usm_types = ["shared", "device", "host"] + + +@pytest.fixture(params=list_of_dtypes) +def input_arrays(request): + a = dpnp.ones(N, dtype=request.param) + b = dpnp.empty_like(a) + return a, b + + +def test_numeric_kernel_arg_types(input_arrays): + """Tests passing float and int type scalar arguments to a kernel function. + + Args: + input_arrays (dpnp.ndarray): Array arguments to be passed to a kernel. + """ + s = 2 + a, b = input_arrays + + scaling_kernel[dpex.Range(N)](a, b, s) + + nb = dpnp.asnumpy(b) + nexpected = numpy.full_like(nb, fill_value=2) + + assert numpy.allclose(nb, nexpected) + + +def test_bool_kernel_arg_type(input_arrays): + """Tests passing boolean arguments to a kernel function. + + Args: + input_arrays (dpnp.ndarray): Array arguments to be passed to a kernel. + """ + a, b = input_arrays + + 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) + + nb = dpnp.asnumpy(b) + nexpected_false = numpy.zeros_like(nb) + + assert numpy.allclose(nb, nexpected_false) From e3c533b9477da2d72904597b26bef4907eeb8fc8 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 22:59:08 -0500 Subject: [PATCH 16/21] Fixes the private memory allocation test case. - Test case now uses dpnp arrays and compute follows data - The previous test case was not actually testing private memory allocation as the code was being optimized out. The modified test case makes sure private memory is actually allocated. --- .../tests/kernel_tests/test_private_memory.py | 31 ----------------- .../test_private_memory_allocation.py | 34 +++++++++++++++++++ 2 files changed, 34 insertions(+), 31 deletions(-) delete mode 100644 numba_dpex/tests/kernel_tests/test_private_memory.py create mode 100644 numba_dpex/tests/kernel_tests/test_private_memory_allocation.py diff --git a/numba_dpex/tests/kernel_tests/test_private_memory.py b/numba_dpex/tests/kernel_tests/test_private_memory.py deleted file mode 100644 index 74f95afbec..0000000000 --- a/numba_dpex/tests/kernel_tests/test_private_memory.py +++ /dev/null @@ -1,31 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - - -import numpy as np -import pytest - -import numba_dpex -from numba_dpex.tests._helper import filter_strings - - -@pytest.mark.parametrize("filter_str", filter_strings) -def test_private_memory(filter_str): - @numba_dpex.kernel - def private_memory_kernel(A): - i = numba_dpex.get_global_id(0) - prvt_mem = numba_dpex.private.array(shape=1, dtype=np.float32) - prvt_mem[0] = i - numba_dpex.barrier(numba_dpex.LOCAL_MEM_FENCE) # local mem fence - A[i] = prvt_mem[0] * 2 - - N = 64 - arr = np.zeros(N).astype(np.float32) - orig = np.arange(N).astype(np.float32) - - with numba_dpex.offload_to_sycl_device(filter_str): - private_memory_kernel[N, N](arr) - - # The computation is correct? - np.testing.assert_allclose(orig * 2, arr) diff --git a/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py b/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py new file mode 100644 index 0000000000..9d12602048 --- /dev/null +++ b/numba_dpex/tests/kernel_tests/test_private_memory_allocation.py @@ -0,0 +1,34 @@ +# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation +# +# SPDX-License-Identifier: Apache-2.0 + + +import dpnp +import numpy + +import numba_dpex as dpex + + +@dpex.kernel +def kernel_with_private_memory_allocation(A): + i = dpex.get_global_id(0) + prvt_mem = dpex.private.array(shape=10, dtype=dpnp.float32) + for idx in range(10): + prvt_mem[idx] = idx * idx + dpex.barrier(dpex.LOCAL_MEM_FENCE) # local mem fence + for idx in range(10): + A[i] += prvt_mem[idx] + + +def test_private_memory_allocation(): + N = 64 + arr = dpnp.zeros(N, dtype=dpnp.float32) + kernel_with_private_memory_allocation[dpex.Range(N)](arr) + + nparr = dpnp.asnumpy(arr) + + expected = 0 + for i in range(10): + expected += i * i + + assert numpy.all(numpy.isclose(nparr, expected)) From 341951be8c8d9f08d5f438174429c4c3a20e60b2 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Mon, 24 Apr 2023 23:46:11 -0500 Subject: [PATCH 17/21] Minor typo fixes in comments. --- numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py | 5 +++-- numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py | 4 ++-- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py b/numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py index fa3dfee222..96e962e7e1 100644 --- a/numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py +++ b/numba_dpex/tests/kernel_tests/test_dump_kernel_llvm.py @@ -30,8 +30,9 @@ def _get_kernel_llvm(fn, sig, debug=False): def test_dump_file_on_dump_kernel_llvm_flag_on(): """ - Test functionality of DUMP_KERNEL_LLVM config variable. - Check llvm source is dumped in .ll file in current directory + Tests functionality of DUMP_KERNEL_LLVM config variable. + + Check llvm source is dumped into a .ll file in current directory and compare with llvm source stored in SprivKernel. """ diff --git a/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py b/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py index 8be8d5d3b8..6c1320b061 100644 --- a/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py +++ b/numba_dpex/tests/kernel_tests/test_ndrange_exceptions.py @@ -24,8 +24,8 @@ def kernel_vector_sum(a, b, c): ], ) def test_ndrange_config_error(error, ranges): - """Test if a exception is raised when calling a - ndrange kernel with unspported arguments. + """Test if a exception is raised when calling a ndrange kernel with + unsupported arguments. """ a = dpt.ones(1024, dtype=dpt.int32) From 806ef505e9a593217b3f05db6236b3182f53d528 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Tue, 25 Apr 2023 11:30:41 -0500 Subject: [PATCH 18/21] Fix and improve test_print.py --- numba_dpex/tests/kernel_tests/test_print.py | 98 +++++++++++++-------- 1 file changed, 59 insertions(+), 39 deletions(-) diff --git a/numba_dpex/tests/kernel_tests/test_print.py b/numba_dpex/tests/kernel_tests/test_print.py index cb78cfa850..d53e1059cf 100644 --- a/numba_dpex/tests/kernel_tests/test_print.py +++ b/numba_dpex/tests/kernel_tests/test_print.py @@ -2,27 +2,63 @@ # # SPDX-License-Identifier: Apache-2.0 -import dpctl -import numpy as np +import dpnp import pytest +from numba.core.errors import LoweringError import numba_dpex as dpex -from numba_dpex.tests._helper import filter_strings_opencl_gpu + +list_of_dtypes = [ + dpnp.int32, + dpnp.int64, + dpnp.float32, + dpnp.float64, +] + + +@pytest.fixture(params=list_of_dtypes) +def input_arrays(request): + a = dpnp.array([0], dtype=request.param) + a[0] = 10 + return a + + +def test_print_scalar_with_string(input_arrays, capfd): + """Tests if we can print a scalar value with a string.""" + + @dpex.kernel + def print_scalar_val(s): + print("printing ...", s[0]) + + a = input_arrays + + print_scalar_val[dpex.Range(1)](a) + captured = capfd.readouterr() + assert "printing ... 10" in captured.out -@pytest.mark.parametrize("filter_str", filter_strings_opencl_gpu) -@pytest.mark.xfail -def test_print_only_str(filter_str): - try: - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device): - pass - except Exception: - pytest.skip() +def test_print_scalar(input_arrays, capfd): + """Tests if we can print a scalar value.""" @dpex.kernel - def f(): - print("test", "test2") + def print_scalar_val(s): + print(s[0]) + + a = input_arrays + + print_scalar_val[dpex.Range(1)](a) + captured = capfd.readouterr() + assert "10" in captured.out + + +def test_print_only_str(input_arrays): + """Negative test to capture LoweringError as printing strings is + unsupported. + """ + + @dpex.kernel + def print_string(a): + print("cannot print only a string inside a kernel") # This test will fail, we currently can not print only string. # The LLVM generated for printf() function with only string gets @@ -30,36 +66,20 @@ def f(): # puts function signature right now, and would fail in general due # to lack of support for puts() in OpenCL. - with dpctl.device_context(filter_str): - f[3, dpex.DEFAULT_LOCAL_SIZE]() - - -list_of_dtypes = [ - np.int32, - np.int64, - np.float32, - np.float64, -] + a = input_arrays + with pytest.raises(LoweringError): + print_string[dpex.Range(1)](a) -@pytest.fixture(params=list_of_dtypes) -def input_arrays(request): - a = np.array([0], request.param) - a[0] = 6 - return a +def test_print_array(input_arrays): + """Negative test to capture LoweringError as printing arrays is unsupported.""" -@pytest.mark.parametrize("filter_str", filter_strings_opencl_gpu) -def test_print(filter_str, input_arrays, capfd): @dpex.kernel - def f(a): - print("test", a[0]) + def print_string(a): + print(a) a = input_arrays - global_size = 3 - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device): - f[global_size, dpex.DEFAULT_LOCAL_SIZE](a) - captured = capfd.readouterr() - assert "test" in captured.out + with pytest.raises(LoweringError): + print_string[dpex.Range(1)](a) From 496b8299b07f91d4dda94418f20b3c4a54a02c9e Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Tue, 25 Apr 2023 15:58:32 -0500 Subject: [PATCH 19/21] Updates to kernel tests. --- .../test_func_qualname_disambiguation.py | 36 ++------------- .../kernel_tests/test_func_specialization.py | 46 +++++++++---------- .../test_kernel_has_return_value_error.py | 9 ++-- .../tests/kernel_tests/test_math_functions.py | 31 ++++++------- .../kernel_tests/test_scalar_arg_types.py | 2 - 5 files changed, 47 insertions(+), 77 deletions(-) diff --git a/numba_dpex/tests/kernel_tests/test_func_qualname_disambiguation.py b/numba_dpex/tests/kernel_tests/test_func_qualname_disambiguation.py index f6f7571b8c..507fef3740 100644 --- a/numba_dpex/tests/kernel_tests/test_func_qualname_disambiguation.py +++ b/numba_dpex/tests/kernel_tests/test_func_qualname_disambiguation.py @@ -2,13 +2,11 @@ # # SPDX-License-Identifier: Apache-2.0 -import dpctl import dpctl.tensor as dpt +import dpnp import numpy as np -import pytest import numba_dpex as ndpx -from numba_dpex.tests._helper import filter_strings def make_write_values_kernel(n_rows): @@ -80,8 +78,7 @@ def write_values_inner(array_in, row_idx): return write_values_inner -@pytest.mark.parametrize("offload_device", filter_strings) -def test_qualname_basic(offload_device): +def test_qualname_basic(): """A basic test function to test qualified name disambiguation. """ @@ -92,32 +89,9 @@ def test_qualname_basic(offload_device): else: ans[i, 0] = 1 - a = np.zeros((10, 10), dtype=dpt.int64) - - device = dpctl.SyclDevice(offload_device) - queue = dpctl.SyclQueue(device) - - da = dpt.usm_ndarray( - a.shape, - dtype=a.dtype, - buffer="device", - buffer_ctor_kwargs={"queue": queue}, - ) - da.usm_data.copy_from_host(a.reshape((-1)).view("|u1")) + a = dpnp.zeros((10, 10), dtype=dpt.int64) kernel = make_write_values_kernel(10) - kernel(da) - - result = np.zeros_like(a) - da.usm_data.copy_to_host(result.reshape((-1)).view("|u1")) - - print(ans) - print(result) - - assert np.array_equal(result, ans) - + kernel(a) -if __name__ == "__main__": - test_qualname_basic("level_zero:gpu:0") - test_qualname_basic("opencl:gpu:0") - test_qualname_basic("opencl:cpu:0") + assert np.array_equal(dpnp.asnumpy(a), ans) diff --git a/numba_dpex/tests/kernel_tests/test_func_specialization.py b/numba_dpex/tests/kernel_tests/test_func_specialization.py index 53ebe7e0d9..386535ba18 100644 --- a/numba_dpex/tests/kernel_tests/test_func_specialization.py +++ b/numba_dpex/tests/kernel_tests/test_func_specialization.py @@ -2,7 +2,7 @@ # # SPDX-License-Identifier: Apache-2.0 -import dpctl.tensor as dpt +import dpnp import numpy as np import pytest @@ -32,12 +32,12 @@ def kernel_function(a, b): k = dpex.kernel(kernel_function) - a = dpt.ones(N) - b = dpt.ones(N) + a = dpnp.ones(N) + b = dpnp.ones(N) - k[N](a, b) + k[dpex.Range(N)](a, b) - assert np.array_equal(dpt.asnumpy(b), dpt.asnumpy(a) + 1) + assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) def test_single_signature(): @@ -53,19 +53,19 @@ def kernel_function(a, b): k = dpex.kernel(kernel_function) # Test with int32, should work - a = dpt.ones(N, dtype=dpt.int32) - b = dpt.ones(N, dtype=dpt.int32) + a = dpnp.ones(N, dtype=dpnp.int32) + b = dpnp.ones(N, dtype=dpnp.int32) - k[N](a, b) + k[dpex.Range(N)](a, b) - assert np.array_equal(dpt.asnumpy(b), dpt.asnumpy(a) + 1) + assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) # Test with int64, should fail - a = dpt.ones(N, dtype=dpt.int64) - b = dpt.ones(N, dtype=dpt.int64) + a = dpnp.ones(N, dtype=dpnp.int64) + b = dpnp.ones(N, dtype=dpnp.int64) with pytest.raises(Exception) as e: - k[N](a, b) + k[dpex.Range(N)](a, b) assert " >>> (int64)" in e.value.args[0] @@ -83,26 +83,26 @@ def kernel_function(a, b): k = dpex.kernel(kernel_function) # Test with int32, should work - a = dpt.ones(N, dtype=dpt.int32) - b = dpt.ones(N, dtype=dpt.int32) + a = dpnp.ones(N, dtype=dpnp.int32) + b = dpnp.ones(N, dtype=dpnp.int32) - k[N](a, b) + k[dpex.Range(N)](a, b) - assert np.array_equal(dpt.asnumpy(b), dpt.asnumpy(a) + 1) + assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) # Test with float32, should work - a = dpt.ones(N, dtype=dpt.float32) - b = dpt.ones(N, dtype=dpt.float32) + a = dpnp.ones(N, dtype=dpnp.float32) + b = dpnp.ones(N, dtype=dpnp.float32) - k[N](a, b) + k[dpex.Range(N)](a, b) - assert np.array_equal(dpt.asnumpy(b), dpt.asnumpy(a) + 1) + assert np.array_equal(dpnp.asnumpy(b), dpnp.asnumpy(a) + 1) # Test with int64, should fail - a = dpt.ones(N, dtype=dpt.int64) - b = dpt.ones(N, dtype=dpt.int64) + a = dpnp.ones(N, dtype=dpnp.int64) + b = dpnp.ones(N, dtype=dpnp.int64) with pytest.raises(Exception) as e: - k[N](a, b) + k[dpex.Range(N)](a, b) assert " >>> (int64)" in e.value.args[0] 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 3864e55859..60a403ea2f 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 @@ -2,8 +2,7 @@ # # SPDX-License-Identifier: Apache-2.0 -import dpctl.tensor as dpt -import numpy as np +import dpnp import pytest import numba_dpex as dpex @@ -28,8 +27,8 @@ def sig(request): def test_return(sig): - a = dpt.arange(1024, dtype=dpt.int32, device="0") + a = dpnp.arange(1024, dtype=dpnp.int32) with pytest.raises(dpex.core.exceptions.KernelHasReturnValueError): - kernel = dpex.kernel(sig)(f) - kernel[a.size, dpex.DEFAULT_LOCAL_SIZE](a) + kernel_fn = dpex.kernel(sig)(f) + kernel_fn[dpex.Range(a.size)](a) diff --git a/numba_dpex/tests/kernel_tests/test_math_functions.py b/numba_dpex/tests/kernel_tests/test_math_functions.py index 52316b8139..7c9eb5c754 100644 --- a/numba_dpex/tests/kernel_tests/test_math_functions.py +++ b/numba_dpex/tests/kernel_tests/test_math_functions.py @@ -4,12 +4,11 @@ import math -import dpctl -import numpy as np +import dpnp +import numpy import pytest import numba_dpex as dpex -from numba_dpex.tests._helper import filter_strings list_of_unary_ops = ["fabs", "exp", "log", "sqrt", "sin", "cos", "tan"] @@ -20,8 +19,8 @@ def unary_op(request): list_of_dtypes = [ - np.float32, - np.float64, + dpnp.float32, + dpnp.float64, ] @@ -29,26 +28,26 @@ def unary_op(request): def input_arrays(request): # The size of input and out arrays to be used N = 2048 - a = np.array(np.random.random(N), request.param) - b = np.array(np.random.random(N), request.param) + a = dpnp.arange(N, dtype=request.param) + b = dpnp.arange(N, dtype=request.param) return a, b -@pytest.mark.parametrize("filter_str", filter_strings) -def test_binary_ops(filter_str, unary_op, input_arrays): - a, actual = input_arrays +def test_binary_ops(unary_op, input_arrays): + a, b = input_arrays uop = getattr(math, unary_op) - np_uop = getattr(np, unary_op) + dpnp_uop = getattr(dpnp, unary_op) @dpex.kernel def f(a, b): i = dpex.get_global_id(0) b[i] = uop(a[i]) - device = dpctl.SyclDevice(filter_str) - with dpctl.device_context(device): - f[a.size, dpex.DEFAULT_LOCAL_SIZE](a, actual) + f[dpex.Range(a.size)](a, b) - expected = np_uop(a) + expected = dpnp_uop(a) - np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) + np_expected = dpnp.asnumpy(expected) + np_actual = dpnp.asnumpy(b) + + assert numpy.allclose(np_expected, np_actual) 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 7c08667e16..6fbc80a062 100644 --- a/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py +++ b/numba_dpex/tests/kernel_tests/test_scalar_arg_types.py @@ -2,8 +2,6 @@ # # SPDX-License-Identifier: Apache-2.0 -import dpctl -import dpctl.tensor as dpt import dpnp import numpy import pytest From e21f04a9a553835c39c4937823943b4e85bc16d2 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Tue, 25 Apr 2023 20:02:31 -0500 Subject: [PATCH 20/21] XFail test_print test cases on OCL CPU. --- numba_dpex/tests/kernel_tests/test_print.py | 12 +++++++++++- 1 file changed, 11 insertions(+), 1 deletion(-) diff --git a/numba_dpex/tests/kernel_tests/test_print.py b/numba_dpex/tests/kernel_tests/test_print.py index d53e1059cf..e337c40915 100644 --- a/numba_dpex/tests/kernel_tests/test_print.py +++ b/numba_dpex/tests/kernel_tests/test_print.py @@ -2,6 +2,7 @@ # # SPDX-License-Identifier: Apache-2.0 +import dpctl import dpnp import pytest from numba.core.errors import LoweringError @@ -26,6 +27,9 @@ def input_arrays(request): def test_print_scalar_with_string(input_arrays, capfd): """Tests if we can print a scalar value with a string.""" + if dpctl.SyclDevice().device_type == dpctl.device_type.cpu: + pytest.xfail("Printing scalars on OpenCL CPU devices is unsupported.") + @dpex.kernel def print_scalar_val(s): print("printing ...", s[0]) @@ -40,6 +44,9 @@ def print_scalar_val(s): def test_print_scalar(input_arrays, capfd): """Tests if we can print a scalar value.""" + if dpctl.SyclDevice().device_type == dpctl.device_type.cpu: + pytest.xfail("Printing scalars on OpenCL CPU devices is unsupported.") + @dpex.kernel def print_scalar_val(s): print(s[0]) @@ -48,6 +55,7 @@ def print_scalar_val(s): print_scalar_val[dpex.Range(1)](a) captured = capfd.readouterr() + assert "10" in captured.out @@ -73,7 +81,9 @@ def print_string(a): def test_print_array(input_arrays): - """Negative test to capture LoweringError as printing arrays is unsupported.""" + """Negative test to capture LoweringError as printing arrays + is unsupported. + """ @dpex.kernel def print_string(a): From 5655ab7b01901b9fa3464c86d331fcf2c2f330f2 Mon Sep 17 00:00:00 2001 From: Diptorup Deb Date: Tue, 25 Apr 2023 20:47:56 -0500 Subject: [PATCH 21/21] Unit tests for USMNdArray creation. Remove test_usm_ndarray_type_exceptions --- .../USMNdAArray/test_usm_ndarray_creation.py | 79 +++++++++++++++++++ .../test_usm_ndarray_type_exceptions.py | 31 -------- 2 files changed, 79 insertions(+), 31 deletions(-) create mode 100644 numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_creation.py delete mode 100644 numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type_exceptions.py diff --git a/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_creation.py b/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_creation.py new file mode 100644 index 0000000000..007b287190 --- /dev/null +++ b/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_creation.py @@ -0,0 +1,79 @@ +import dpctl +import pytest + +from numba_dpex.core.types import USMNdArray + +"""Negative tests for expected exceptions raised during USMNdArray creation. + +""" + + +def test_default_type_construction(): + """Tests call USMNdArray constructor with no device or queue args.""" + usma = USMNdArray(1, queue=None) + + assert usma.ndim == 1 + assert usma.layout == "C" + assert usma.addrspace == 1 + assert usma.usm_type == "device" + + default_device = dpctl.SyclDevice() + cached_queue = dpctl.get_device_cached_queue(default_device) + + assert usma.device == default_device.filter_string + assert usma.queue == cached_queue + + +def test_type_creation_with_device(): + """Tests creating a USMNdArray with a device arg and no queue""" + + default_device_str = dpctl.SyclDevice().filter_string + + usma = USMNdArray(1, device=default_device_str, queue=None) + + assert usma.ndim == 1 + assert usma.layout == "C" + assert usma.addrspace == 1 + assert usma.usm_type == "device" + + assert usma.device == default_device_str + + cached_queue = dpctl.get_device_cached_queue(default_device_str) + + assert usma.queue == cached_queue + + +def test_type_creation_with_queue(): + """Tests creating a USMNdArray with a queue arg and no device""" + queue = dpctl.SyclQueue() + usma = USMNdArray(1, queue=queue) + + assert usma.ndim == 1 + assert usma.layout == "C" + assert usma.addrspace == 1 + assert usma.usm_type == "device" + + assert usma.device == queue.sycl_device.filter_string + assert usma.queue == queue + + +def test_exception_when_both_device_and_queue_arg_specified(): + """Tests if TypeError is raised when both queue and device specified""" + + queue = dpctl.SyclQueue() + with pytest.raises(TypeError): + USMNdArray(1, device="cpu", queue=queue) + + +def test_improper_queue_type(): + """Tests if TypeError is raised if queue argument is of invalid type""" + + with pytest.raises(TypeError): + USMNdArray(1, queue="cpu") + + +def test_improper_device_type(): + """Tests if TypeError is raised if device argument is of invalid type""" + + with pytest.raises(TypeError): + USMNdArray(1, device=0) diff --git a/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type_exceptions.py b/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type_exceptions.py deleted file mode 100644 index 62aa17e663..0000000000 --- a/numba_dpex/tests/core/types/USMNdAArray/test_usm_ndarray_type_exceptions.py +++ /dev/null @@ -1,31 +0,0 @@ -# SPDX-FileCopyrightText: 2020 - 2023 Intel Corporation -# -# SPDX-License-Identifier: Apache-2.0 - -"""This is to test USMNdArray raising correct exceptions.""" - -import dpnp -import pytest -from numba import njit -from numba.core.errors import TypingError - -from numba_dpex.dpctl_iface import get_current_queue - -arguments = [ - ("shape=10", 'device="cpu"', "queue=a.sycl_queue"), - ("shape=10", "device=10", "queue=a.sycl_queue"), - ("shape=10", 'device="cpu"', "queue=test"), - ("shape=10", 'device="dpu"'), -] - - -@pytest.mark.parametrize("argument", arguments) -def test_usm_ndarray_type_exceptions(argument): - a = dpnp.ndarray(10) - - @njit - def func(a): - dpnp.empty(argument) - - with pytest.raises(TypingError): - func(a)