diff --git a/numba_dppy/context.py b/numba_dppy/context.py new file mode 100644 index 0000000000..58bc750ce1 --- /dev/null +++ b/numba_dppy/context.py @@ -0,0 +1,45 @@ +# Copyright 2021 Intel Corporation +# +# Licensed under the Apache License, Version 2.0 (the "License"); +# you may not use this file except in compliance with the License. +# You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +from contextlib import contextmanager + +import dpctl +from numba._dispatcher import set_use_tls_target_stack +from numba.core.dispatcher import TargetConfig + +from numba_dppy.dppy_offload_dispatcher import DppyOffloadDispatcher + + +@contextmanager +def switch_target(retarget): + # __enter__ + tc = TargetConfig() + tc.push(retarget) + set_use_tls_target_stack(True) + yield + # __exit__ + tc.pop() + set_use_tls_target_stack(False) + + +def retarget_to_gpu(cpu_disp): + dispatcher = DppyOffloadDispatcher(cpu_disp.py_func) + return lambda *args, **kwargs: dispatcher(*args, **kwargs) + + +@contextmanager +def device_context(*args, **kwargs): + with switch_target(retarget_to_gpu): + with dpctl.device_context(*args, **kwargs) as queue: + yield queue diff --git a/numba_dppy/dispatcher.py b/numba_dppy/dispatcher.py index a05f25b8aa..44c8e76f9d 100644 --- a/numba_dppy/dispatcher.py +++ b/numba_dppy/dispatcher.py @@ -25,7 +25,7 @@ # from numba.npyufunc.deviceufunc import (UFuncMechanism, GenerializedUFunc, # GUFuncCallSteps) -from .. import dispatcher, utils, typing +from numba.core import dispatcher, utils, typing from .compiler import DPPYCompiler diff --git a/numba_dppy/dppy_passes.py b/numba_dppy/dppy_passes.py index cc30e0dee1..40918b6f9d 100644 --- a/numba_dppy/dppy_passes.py +++ b/numba_dppy/dppy_passes.py @@ -226,6 +226,7 @@ def run_pass(self, state): state.typingctx, state.flags.auto_parallel, state.flags, + state.metadata, state.parfor_diagnostics, ) diff --git a/numba_dppy/tests/kernel_tests/test_arg_accessor.py b/numba_dppy/tests/kernel_tests/test_arg_accessor.py index 2553984e7b..375bd2e7bf 100644 --- a/numba_dppy/tests/kernel_tests/test_arg_accessor.py +++ b/numba_dppy/tests/kernel_tests/test_arg_accessor.py @@ -15,7 +15,7 @@ import numpy as np import numba_dppy as dppy import pytest -import dpctl +from numba_dppy.context import device_context from numba_dppy.tests.skip_tests import skip_test @@ -77,6 +77,6 @@ def test_kernel_arg_accessor(filter_str, input_arrays, kernel): a, b, actual = input_arrays expected = a + b - with dpctl.device_context(filter_str): + with device_context(filter_str): call_kernel(global_size, local_size, a, b, actual, kernel) np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) diff --git a/numba_dppy/tests/kernel_tests/test_arg_types.py b/numba_dppy/tests/kernel_tests/test_arg_types.py index 8ee24ccfd8..fddeaa05e3 100644 --- a/numba_dppy/tests/kernel_tests/test_arg_types.py +++ b/numba_dppy/tests/kernel_tests/test_arg_types.py @@ -15,7 +15,7 @@ import numpy as np import numba_dppy as dppy import pytest -import dpctl +from numba_dppy.context import device_context from numba_dppy.tests.skip_tests import skip_test global_size = 1054 @@ -63,7 +63,7 @@ def test_kernel_arg_types(filter_str, input_arrays): kernel = dppy.kernel(mul_kernel) a, actual, c = input_arrays expected = a * c - with dpctl.device_context(filter_str): + with device_context(filter_str): kernel[global_size, local_size](a, actual, c) np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=0) @@ -82,7 +82,7 @@ def test_bool_type(filter_str): kernel = dppy.kernel(check_bool_kernel) a = np.array([2], np.int64) - with dpctl.device_context(filter_str): + with device_context(filter_str): kernel[a.size, dppy.DEFAULT_LOCAL_SIZE](a, True) assert a[0] == 111 kernel[a.size, dppy.DEFAULT_LOCAL_SIZE](a, False) diff --git a/numba_dppy/tests/kernel_tests/test_atomic_op.py b/numba_dppy/tests/kernel_tests/test_atomic_op.py index 0193cb2bb1..105560868c 100644 --- a/numba_dppy/tests/kernel_tests/test_atomic_op.py +++ b/numba_dppy/tests/kernel_tests/test_atomic_op.py @@ -17,7 +17,8 @@ import numba_dppy as dppy import pytest -import dpctl + +from numba_dppy.context import device_context global_size = 100 @@ -94,7 +95,7 @@ def test_kernel_atomic_simple(filter_str, input_arrays, kernel_result_pair): a, dtype = input_arrays kernel, expected = kernel_result_pair - with dpctl.device_context(filter_str): + with device_context(filter_str): kernel[global_size, dppy.DEFAULT_LOCAL_SIZE](a) assert a[0] == expected @@ -120,7 +121,7 @@ def test_kernel_atomic_local(filter_str, input_arrays, return_list_of_op): a, dtype = input_arrays op_type, expected = return_list_of_op kernel = get_kernel_local(op_type, dtype) - with dpctl.device_context(filter_str): + with device_context(filter_str): kernel[global_size, global_size](a) assert a[0] == expected @@ -162,6 +163,6 @@ def test_kernel_atomic_multi_dim( dim = return_list_of_dim kernel = get_kernel_multi_dim(op_type, len(dim)) a = np.zeros(dim, return_dtype) - with dpctl.device_context(filter_str): + with device_context(filter_str): kernel[global_size, dppy.DEFAULT_LOCAL_SIZE](a) assert a[0] == expected diff --git a/numba_dppy/tests/kernel_tests/test_barrier.py b/numba_dppy/tests/kernel_tests/test_barrier.py index 571765a3b2..c7b6231d82 100644 --- a/numba_dppy/tests/kernel_tests/test_barrier.py +++ b/numba_dppy/tests/kernel_tests/test_barrier.py @@ -15,7 +15,7 @@ import numpy as np import numba_dppy as dppy import pytest -import dpctl +from numba_dppy.context import device_context from numba_dppy.tests.skip_tests import skip_test @@ -46,7 +46,7 @@ def twice(A): arr = np.random.random(N).astype(np.float32) orig = arr.copy() - with dpctl.device_context(filter_str) as gpu_queue: + with device_context(filter_str) as gpu_queue: twice[N, N // 2](arr) # The computation is correct? @@ -69,7 +69,7 @@ def twice(A): arr = np.random.random(N).astype(np.float32) orig = arr.copy() - with dpctl.device_context(filter_str) as gpu_queue: + with device_context(filter_str) as gpu_queue: twice[N, dppy.DEFAULT_LOCAL_SIZE](arr) # The computation is correct? @@ -97,7 +97,7 @@ def reverse_array(A): arr = np.arange(blocksize).astype(np.float32) orig = arr.copy() - with dpctl.device_context(filter_str) as gpu_queue: + with device_context(filter_str) as gpu_queue: reverse_array[blocksize, blocksize](arr) expected = orig[::-1] + orig diff --git a/numba_dppy/tests/kernel_tests/test_caching.py b/numba_dppy/tests/kernel_tests/test_caching.py index 5a6bfe436c..584b3ca133 100644 --- a/numba_dppy/tests/kernel_tests/test_caching.py +++ b/numba_dppy/tests/kernel_tests/test_caching.py @@ -16,7 +16,7 @@ import numpy as np import numba_dppy as dppy import pytest -import dpctl +from numba_dppy.context import device_context from numba_dppy.tests.skip_tests import skip_test list_of_filter_strs = [ @@ -46,7 +46,7 @@ def test_caching_kernel(filter_str): b = np.array(np.random.random(N), dtype=np.float32) c = np.ones_like(a) - with dpctl.device_context(filter_str) as gpu_queue: + with device_context(filter_str) as gpu_queue: func = dppy.kernel(data_parallel_sum) caching_kernel = func[global_size, dppy.DEFAULT_LOCAL_SIZE].specialize(a, b, c) diff --git a/numba_dppy/tests/kernel_tests/test_math_functions.py b/numba_dppy/tests/kernel_tests/test_math_functions.py index feca0bd465..0a740522d6 100644 --- a/numba_dppy/tests/kernel_tests/test_math_functions.py +++ b/numba_dppy/tests/kernel_tests/test_math_functions.py @@ -12,7 +12,7 @@ # See the License for the specific language governing permissions and # limitations under the License. -import dpctl +from numba_dppy.context import device_context import numba_dppy as dppy import numpy as np import pytest @@ -67,7 +67,7 @@ def f(a, b): i = dppy.get_global_id(0) b[i] = uop(a[i]) - with dpctl.device_context(filter_str): + with device_context(filter_str): f[a.size, dppy.DEFAULT_LOCAL_SIZE](a, actual) expected = np_uop(a) diff --git a/numba_dppy/tests/kernel_tests/test_print.py b/numba_dppy/tests/kernel_tests/test_print.py index 4d9e8ee74b..afaf0586cb 100644 --- a/numba_dppy/tests/kernel_tests/test_print.py +++ b/numba_dppy/tests/kernel_tests/test_print.py @@ -15,7 +15,7 @@ import numpy as np import numba_dppy as dppy import pytest -import dpctl +from numba_dppy.context import device_context from numba_dppy.tests.skip_tests import skip_test list_of_filter_strs = [ @@ -31,7 +31,7 @@ def filter_str(request): @pytest.mark.xfail def test_print_only_str(filter_str): try: - with dpctl.device_context(filter_str): + with device_context(filter_str): pass except Exception: pytest.skip() @@ -45,7 +45,7 @@ def f(): # replaced by a puts() which fails due to lack of addrspace in the # 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), captured_stdout() as stdout: + with device_context(filter_str), captured_stdout() as stdout: f[3, dppy.DEFAULT_LOCAL_SIZE]() @@ -75,7 +75,7 @@ def f(a): a = input_arrays global_size = 3 - with dpctl.device_context(filter_str): + with device_context(filter_str): f[global_size, dppy.DEFAULT_LOCAL_SIZE](a) captured = capfd.readouterr() assert "test" in captured.out diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_creation.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_creation.py index 730c0d3ae2..ab1185528e 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_creation.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_creation.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -106,7 +106,7 @@ def test_unary_ops(filter_str, unary_op, input_array, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -130,7 +130,7 @@ def test_binary_op(filter_str, binary_op, input_array, dtype, get_shape, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a, dtype) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -159,7 +159,7 @@ def test_full(filter_str, full_name, input_array, get_shape, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a, np.array([2])) captured = capfd.readouterr() assert "dpnp implementation" in captured.out diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_ops.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_ops.py index 470bd959a4..d772eb26ae 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_ops.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_array_ops.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -103,7 +103,7 @@ def test_unary_ops(filter_str, unary_op, input_arrays, get_shape, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -142,7 +142,7 @@ def test_take(filter_str, input_arrays, indices, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a, indices) captured = capfd.readouterr() assert "dpnp implementation" in captured.out diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_linalg.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_linalg.py index b775a122ec..cf38a40474 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_linalg.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_linalg.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -102,7 +102,7 @@ def test_eig(filter_str, eig_input, capfd): fn = get_fn("linalg.eig", 1) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual_val, actual_vec = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -175,7 +175,7 @@ def test_dot(filter_str, dot_name, dot_input, dtype, capfd): fn = get_fn(dot_name, 2) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a, b) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -196,7 +196,7 @@ def test_matmul(filter_str, dtype, capfd): fn = get_fn("matmul", 2) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a, b) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -214,7 +214,7 @@ def test_cholesky(filter_str, dtype, capfd): fn = get_fn("linalg.cholesky", 1) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -248,7 +248,7 @@ def test_det(filter_str, det_input, dtype, capfd): fn = get_fn("linalg.det", 1) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -274,7 +274,7 @@ def fn(A, B, C, D): D = np.random.random((5, 333)) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(A, B, C, D) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -314,7 +314,7 @@ def test_matrix_power(filter_str, matrix_power_input, power, dtype, capfd): fn = get_fn("linalg.matrix_power", 2) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a, power) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -339,7 +339,7 @@ def test_matrix_rank(filter_str, matrix_rank_input, capfd): fn = get_fn("linalg.matrix_rank", 1) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(matrix_rank_input) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -359,7 +359,7 @@ def test_eigvals(filter_str, eig_input, capfd): fn = get_fn("linalg.eigvals", 1) f = njit(fn) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual_val = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_rng.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_rng.py index f7f4e926cc..2572822504 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_rng.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_rng.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -83,7 +83,7 @@ def test_one_arg_fn(filter_str, one_arg_fn, unary_size, capfd): op, params = one_arg_fn name, low, high = params f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(unary_size) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -133,7 +133,7 @@ def test_two_arg_fn(filter_str, two_arg_fn, unary_size, capfd): pytest.skip("AttributeError: 'NoneType' object has no attribute 'ravel'") op = get_two_arg_fn(op_name) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(first_arg, unary_size) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -199,7 +199,7 @@ def test_three_arg_fn(filter_str, three_arg_fn, three_arg_size, capfd): op = get_three_arg_fn(op_name) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(first_arg, second_arg, three_arg_size) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -233,7 +233,7 @@ def f(): c = np.random.rand(3, 2) return c - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f() actual = actual.ravel() @@ -251,7 +251,7 @@ def f(ngood, nbad, nsamp, size): return res ngood, nbad, nsamp = 100, 2, 10 - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(ngood, nbad, nsamp, three_arg_size) if np.isscalar(actual): diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_sort_search_count.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_sort_search_count.py index bfa3e08466..2a43710d73 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_sort_search_count.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_sort_search_count.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -92,7 +92,7 @@ def test_unary_ops(filter_str, unary_op, input_arrays, get_shape, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_statistics.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_statistics.py index 662e70a082..af521c779e 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_statistics.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_statistics.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -99,7 +99,7 @@ def test_unary_ops(filter_str, unary_op, input_arrays, get_shape, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out diff --git a/numba_dppy/tests/njit_tests/dpnp/test_numpy_transcendentals.py b/numba_dppy/tests/njit_tests/dpnp/test_numpy_transcendentals.py index e78e0fb86f..3980909449 100644 --- a/numba_dppy/tests/njit_tests/dpnp/test_numpy_transcendentals.py +++ b/numba_dppy/tests/njit_tests/dpnp/test_numpy_transcendentals.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -123,7 +123,7 @@ def test_unary_ops(filter_str, unary_op, input_array, get_shape, capfd): expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out @@ -144,7 +144,7 @@ def test_unary_nan_ops(filter_str, unary_nan_op, input_nan_array, get_shape, cap expected = np.empty(shape=a.shape, dtype=a.dtype) f = njit(op) - with dpctl.device_context(filter_str), dpnp_debug(): + with device_context(filter_str), dpnp_debug(): actual = f(a) captured = capfd.readouterr() assert "dpnp implementation" in captured.out diff --git a/numba_dppy/tests/njit_tests/skip_tests.py b/numba_dppy/tests/njit_tests/skip_tests.py index 758451dc12..8ebf8c4a8d 100644 --- a/numba_dppy/tests/njit_tests/skip_tests.py +++ b/numba_dppy/tests/njit_tests/skip_tests.py @@ -17,10 +17,11 @@ ################################################################################ import dpctl +from numba_dppy.context import device_context def is_gen12(device_type): - with dpctl.device_context(device_type): + with device_context(device_type): q = dpctl.get_current_queue() device = q.get_sycl_device() name = device.get_device_name() diff --git a/numba_dppy/tests/njit_tests/test_numpy_bitwise_ops.py b/numba_dppy/tests/njit_tests/test_numpy_bitwise_ops.py index bc9b1e76a2..74b681458f 100644 --- a/numba_dppy/tests/njit_tests/test_numpy_bitwise_ops.py +++ b/numba_dppy/tests/njit_tests/test_numpy_bitwise_ops.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -87,7 +87,7 @@ def test_binary_ops(filter_str, binary_op, input_arrays): def f(a, b): return binop(a, b) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a, b) expected = binop(a, b) @@ -107,7 +107,7 @@ def test_unary_ops(filter_str, unary_op, input_arrays): def f(a): return uop(a) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a) expected = uop(a) diff --git a/numba_dppy/tests/njit_tests/test_numpy_logic_ops.py b/numba_dppy/tests/njit_tests/test_numpy_logic_ops.py index 53308a8cf3..8d0b8976d6 100644 --- a/numba_dppy/tests/njit_tests/test_numpy_logic_ops.py +++ b/numba_dppy/tests/njit_tests/test_numpy_logic_ops.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -95,7 +95,7 @@ def test_binary_ops(filter_str, binary_op, input_arrays): def f(a, b): return binop(a, b) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a, b) expected = binop(a, b) @@ -115,7 +115,7 @@ def test_unary_ops(filter_str, unary_op, input_arrays): def f(a): return uop(a) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a) expected = uop(a) diff --git a/numba_dppy/tests/njit_tests/test_numpy_transcedental_functions.py b/numba_dppy/tests/njit_tests/test_numpy_transcedental_functions.py index e137736136..f710c89e31 100644 --- a/numba_dppy/tests/njit_tests/test_numpy_transcedental_functions.py +++ b/numba_dppy/tests/njit_tests/test_numpy_transcedental_functions.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -114,7 +114,7 @@ def test_binary_ops(filter_str, binary_op, input_arrays): def f(a, b): return binop(a, b) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a, b) expected = binop(a, b) @@ -139,7 +139,7 @@ def test_unary_ops(filter_str, unary_op, input_arrays): def f(a): return uop(a) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a) expected = uop(a) diff --git a/numba_dppy/tests/njit_tests/test_numpy_trigonometric_functions.py b/numba_dppy/tests/njit_tests/test_numpy_trigonometric_functions.py index 352e1a8648..63f679242a 100644 --- a/numba_dppy/tests/njit_tests/test_numpy_trigonometric_functions.py +++ b/numba_dppy/tests/njit_tests/test_numpy_trigonometric_functions.py @@ -16,7 +16,7 @@ # limitations under the License. ################################################################################ -import dpctl +from numba_dppy.context import device_context import numpy as np from numba import njit import pytest @@ -104,7 +104,7 @@ def test_trigonometric_fn(filter_str, trig_op, input_arrays): def f(a, b): return trig_fn(a, b) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a, b) expected = trig_fn(a, b) else: @@ -113,7 +113,7 @@ def f(a, b): def f(a): return trig_fn(a) - with dpctl.device_context(filter_str): + with device_context(filter_str): actual = f(a) expected = trig_fn(a) diff --git a/numba_dppy/tests/skip_tests.py b/numba_dppy/tests/skip_tests.py index 0ab06ef21a..7cc23af828 100644 --- a/numba_dppy/tests/skip_tests.py +++ b/numba_dppy/tests/skip_tests.py @@ -13,10 +13,11 @@ # limitations under the License. import dpctl +from numba_dppy.context import device_context def is_gen12(device_type): - with dpctl.device_context(device_type): + with device_context(device_type): q = dpctl.get_current_queue() device = q.get_sycl_device() name = device.get_device_name() @@ -41,7 +42,7 @@ def platform_not_supported(device_type): def skip_test(device_type): skip = False try: - with dpctl.device_context(device_type): + with device_context(device_type): pass except Exception: skip = True diff --git a/numba_dppy/tests/test_black_scholes.py b/numba_dppy/tests/test_black_scholes.py index 916bd08aa8..77e462cdfe 100644 --- a/numba_dppy/tests/test_black_scholes.py +++ b/numba_dppy/tests/test_black_scholes.py @@ -19,6 +19,7 @@ import numba_dppy, numba_dppy as dppy import unittest import dpctl +from numba_dppy.context import device_context RISKFREE = 0.02 @@ -129,7 +130,7 @@ def black_scholes_dppy(callResult, putResult, S, X, T, R, V): blockdim = 512, 1 griddim = int(math.ceil(float(OPT_N) / blockdim[0])), 1 - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: time1 = time.time() for i in range(iterations): black_scholes_dppy[blockdim, griddim]( diff --git a/numba_dppy/tests/test_controllable_fallback.py b/numba_dppy/tests/test_controllable_fallback.py index 8bafa8f059..1e4ce9844b 100644 --- a/numba_dppy/tests/test_controllable_fallback.py +++ b/numba_dppy/tests/test_controllable_fallback.py @@ -18,6 +18,7 @@ import numba_dppy from numba_dppy.testing import unittest import dpctl +from numba_dppy.context import device_context import warnings @@ -39,7 +40,7 @@ def inner_call_fallback(): numba_dppy.compiler.DEBUG = 1 with warnings.catch_warnings(record=True) as w: - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: dppy = numba.njit(parallel=True)(inner_call_fallback) dppy_fallback_true = dppy() @@ -49,7 +50,6 @@ def inner_call_fallback(): np.testing.assert_array_equal(dppy_fallback_true, ref_result) self.assertIn("Failed to lower parfor on DPPY-device", str(w[-1].message)) - @unittest.expectedFailure def test_dppy_fallback_false(self): @numba.jit def fill_value(i): @@ -68,7 +68,7 @@ def inner_call_fallback(): numba_dppy.compiler.DEBUG = 1 numba_dppy.config.FALLBACK_ON_CPU = 0 with warnings.catch_warnings(record=True) as w: - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: dppy = numba.njit(parallel=True)(inner_call_fallback) dppy_fallback_false = dppy() diff --git a/numba_dppy/tests/test_device_array_args.py b/numba_dppy/tests/test_device_array_args.py index 1183f684a3..8e0ddae698 100644 --- a/numba_dppy/tests/test_device_array_args.py +++ b/numba_dppy/tests/test_device_array_args.py @@ -17,6 +17,7 @@ import numpy as np import numba_dppy, numba_dppy as dppy import dpctl +from numba_dppy.context import device_context import unittest @@ -39,7 +40,7 @@ class TestDPPYDeviceArrayArgsGPU(unittest.TestCase): def test_device_array_args_cpu(self): c = np.ones_like(a) - with dpctl.device_context("opencl:cpu") as cpu_queue: + with device_context("opencl:cpu") as cpu_queue: data_parallel_sum[global_size, dppy.DEFAULT_LOCAL_SIZE](a, b, c) self.assertTrue(np.all(c == d)) @@ -50,7 +51,7 @@ class TestDPPYDeviceArrayArgsCPU(unittest.TestCase): def test_device_array_args_gpu(self): c = np.ones_like(a) - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: data_parallel_sum[global_size, dppy.DEFAULT_LOCAL_SIZE](a, b, c) self.assertTrue(np.all(c == d)) diff --git a/numba_dppy/tests/test_dpctl_api.py b/numba_dppy/tests/test_dpctl_api.py index c439b7a15c..685baac297 100644 --- a/numba_dppy/tests/test_dpctl_api.py +++ b/numba_dppy/tests/test_dpctl_api.py @@ -14,6 +14,8 @@ import pytest import dpctl +from numba_dppy.context import device_context +from numba_dppy.tests.skip_tests import skip_test list_of_filter_strs = [ @@ -29,7 +31,10 @@ def filter_str(request): def test_dpctl_api(filter_str): - with dpctl.device_context(filter_str) as gpu_queue: + if skip_test(filter_str): + pytest.skip() + + with device_context(filter_str) as gpu_queue: dpctl.dump() dpctl.get_current_queue() dpctl.get_num_platforms() diff --git a/numba_dppy/tests/test_dpnp_functions.py b/numba_dppy/tests/test_dpnp_functions.py index a3c6eab103..a5522f4ec5 100644 --- a/numba_dppy/tests/test_dpnp_functions.py +++ b/numba_dppy/tests/test_dpnp_functions.py @@ -16,6 +16,7 @@ import numpy as np from numba import njit import dpctl +from numba_dppy.context import device_context import unittest from numba_dppy.testing import ensure_dpnp diff --git a/numba_dppy/tests/test_dppy_fallback.py b/numba_dppy/tests/test_dppy_fallback.py index a899d49925..7ebd75053b 100644 --- a/numba_dppy/tests/test_dppy_fallback.py +++ b/numba_dppy/tests/test_dppy_fallback.py @@ -17,6 +17,7 @@ import numba import unittest import dpctl +from numba_dppy.context import device_context import warnings @@ -36,9 +37,7 @@ def inner_call_fallback(): return a - with warnings.catch_warnings(record=True) as w, dpctl.device_context( - "opencl:gpu" - ): + with warnings.catch_warnings(record=True) as w, device_context("opencl:gpu"): dppy = numba.njit(inner_call_fallback) dppy_result = dppy() @@ -55,9 +54,7 @@ def reduction(a): return b a = np.ones(10) - with warnings.catch_warnings(record=True) as w, dpctl.device_context( - "opencl:gpu" - ): + with warnings.catch_warnings(record=True) as w, device_context("opencl:gpu"): dppy = numba.njit(reduction) dppy_result = dppy(a) diff --git a/numba_dppy/tests/test_dppy_func.py b/numba_dppy/tests/test_dppy_func.py index 30c56e918a..b6af92904c 100644 --- a/numba_dppy/tests/test_dppy_func.py +++ b/numba_dppy/tests/test_dppy_func.py @@ -17,6 +17,7 @@ import numba_dppy, numba_dppy as dppy import unittest import dpctl +from numba_dppy.context import device_context @unittest.skipUnless(dpctl.has_gpu_queues(), "test only on GPU system") @@ -36,7 +37,7 @@ def f(a, b): a = np.ones(self.N) b = np.ones(self.N) - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: f[self.N, dppy.DEFAULT_LOCAL_SIZE](a, b) self.assertTrue(np.all(b == 2)) @@ -59,7 +60,7 @@ def h(a, b): a = np.ones(self.N) b = np.ones(self.N) - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: f[self.N, dppy.DEFAULT_LOCAL_SIZE](a, b) self.assertTrue(np.all(b == 2)) diff --git a/numba_dppy/tests/test_no_copy_usm_shared.py b/numba_dppy/tests/test_no_copy_usm_shared.py index af054bd773..dc5cf15836 100644 --- a/numba_dppy/tests/test_no_copy_usm_shared.py +++ b/numba_dppy/tests/test_no_copy_usm_shared.py @@ -20,7 +20,7 @@ import dpctl.dptensor.numpy_usm_shared as usmarray import numba_dppy.numpy_usm_shared as nus -import dpctl +from numba_dppy.context import device_context from numba_dppy.compiler import DPPYCompiler from numba.core.registry import cpu_target @@ -47,7 +47,7 @@ def test_no_copy_usm_shared(capfd): targetctx = cpu_target.target_context args = typingctx.resolve_argument_type(a) - with dpctl.device_context("opencl:gpu:0"): + with device_context("opencl:gpu:0"): cres = compiler.compile_extra( typingctx=typingctx, targetctx=targetctx, diff --git a/numba_dppy/tests/test_offload_diagnostics.py b/numba_dppy/tests/test_offload_diagnostics.py index 350e65762a..3e3f8b1e4e 100644 --- a/numba_dppy/tests/test_offload_diagnostics.py +++ b/numba_dppy/tests/test_offload_diagnostics.py @@ -20,6 +20,7 @@ from numba_dppy.testing import unittest from numba.tests.support import captured_stdout import dpctl +from numba_dppy.context import device_context @unittest.skipUnless(dpctl.has_gpu_queues(), "test only on GPU system") @@ -35,7 +36,7 @@ def prange_func(): return a - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): dppy_config.OFFLOAD_DIAGNOSTICS = 1 jitted = njit(parallel=True)(prange_func) @@ -59,7 +60,7 @@ def parallel_sum(a, b, c): b = np.array(np.random.random(N), dtype=np.float32) c = np.ones_like(a) - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): dppy_config.OFFLOAD_DIAGNOSTICS = 1 with captured_stdout() as got: diff --git a/numba_dppy/tests/test_parfor_lower_message.py b/numba_dppy/tests/test_parfor_lower_message.py index bcc7578779..2216a8ec8e 100644 --- a/numba_dppy/tests/test_parfor_lower_message.py +++ b/numba_dppy/tests/test_parfor_lower_message.py @@ -20,6 +20,7 @@ import unittest from numba.tests.support import captured_stdout import dpctl +from numba_dppy.context import device_context def prange_example(): @@ -36,7 +37,7 @@ def prange_example(): @unittest.skipUnless(dpctl.has_gpu_queues(), "test only on GPU system") class TestParforMessage(unittest.TestCase): def test_parfor_message(self): - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: numba_dppy.compiler.DEBUG = 1 jitted = njit(prange_example) diff --git a/numba_dppy/tests/test_prange.py b/numba_dppy/tests/test_prange.py index e6cf886044..55d450e3e6 100644 --- a/numba_dppy/tests/test_prange.py +++ b/numba_dppy/tests/test_prange.py @@ -17,6 +17,7 @@ import numpy as np import numba import dpctl +from numba_dppy.context import device_context from numba import njit, prange import numba_dppy import unittest @@ -37,7 +38,7 @@ def f(a, b): a = np.ones((m, n)) b = np.ones((m, n)) - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): f(a, b) for i in range(4): @@ -57,7 +58,7 @@ def f(a, b): a = np.ones((m, n)) b = np.ones((m, n)) - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): f(a, b) self.assertTrue(np.all(b == 10)) @@ -81,7 +82,7 @@ def f(a, b): a = np.ones((m, n)) b = np.ones((m, n)) - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): f(a, b) self.assertTrue(np.all(b == 10)) @@ -105,7 +106,7 @@ def f(a, b): a = np.ones((m, n, o)) b = np.ones((m, n, o)) - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): f(a, b) self.assertTrue(np.all(b == 12)) @@ -127,7 +128,7 @@ def prange_example(): jitted = njit(prange_example) - with captured_stdout() as stdout, dpctl.device_context("opencl:gpu"): + with captured_stdout() as stdout, device_context("opencl:gpu"): jitted_res = jitted() res = prange_example() @@ -163,7 +164,7 @@ def prange_example(): jitted = njit(prange_example) - with captured_stdout() as stdout, dpctl.device_context("opencl:gpu"): + with captured_stdout() as stdout, device_context("opencl:gpu"): jitted_res = jitted() res = prange_example() diff --git a/numba_dppy/tests/test_sum_reduction.py b/numba_dppy/tests/test_sum_reduction.py index 0d1e547e4d..c0a30af42d 100644 --- a/numba_dppy/tests/test_sum_reduction.py +++ b/numba_dppy/tests/test_sum_reduction.py @@ -17,6 +17,7 @@ import numba_dppy, numba_dppy as dppy import unittest import dpctl +from numba_dppy.context import device_context @dppy.kernel @@ -40,7 +41,7 @@ def test_sum_reduction(self): # at max we will require half the size of A to store sum R = np.array(np.random.random(math.ceil(N / 2)), dtype=np.float32) - with dpctl.device_context("opencl:gpu") as gpu_queue: + with device_context("opencl:gpu") as gpu_queue: total = N while total > 1: diff --git a/numba_dppy/tests/test_vectorize.py b/numba_dppy/tests/test_vectorize.py index 1035cae4e5..5ba2a899e1 100644 --- a/numba_dppy/tests/test_vectorize.py +++ b/numba_dppy/tests/test_vectorize.py @@ -16,6 +16,7 @@ import numpy as np from numba import njit, vectorize import dpctl +from numba_dppy.context import device_context import unittest @@ -40,7 +41,7 @@ def f_np(a0, a1): A = np.random.random(10) B = np.random.random(10) - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): expected = f(A, B) actual = f_np(A, B) diff --git a/numba_dppy/tests/test_with_context.py b/numba_dppy/tests/test_with_context.py index 1ecf08c232..31ab899549 100644 --- a/numba_dppy/tests/test_with_context.py +++ b/numba_dppy/tests/test_with_context.py @@ -12,13 +12,15 @@ # See the License for the specific language governing permissions and # limitations under the License. +import unittest + +import dpctl +import numba_dppy +from numba_dppy.context import device_context import numpy as np from numba import njit -import numba_dppy -import unittest from numba.core import errors from numba.tests.support import captured_stdout -import dpctl class TestWithDPPYContext(unittest.TestCase): @@ -38,14 +40,14 @@ def func(b): got_gpu = np.ones((64), dtype=np.float64) with captured_stdout() as got_gpu_message: - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): func(got_gpu) numba_dppy.compiler.DEBUG = 0 func(expected) np.testing.assert_array_equal(expected, got_gpu) - self.assertTrue("Parfor lowered on DPPY-device" in got_gpu_message.getvalue()) + self.assertIn("Parfor lowered on DPPY-device", got_gpu_message.getvalue()) @unittest.skipIf(not dpctl.has_cpu_queues(), "No CPU platforms available") def test_with_dppy_context_cpu(self): @@ -63,14 +65,14 @@ def func(b): got_cpu = np.ones((64), dtype=np.float64) with captured_stdout() as got_cpu_message: - with dpctl.device_context("opencl:cpu"): + with device_context("opencl:cpu"): func(got_cpu) numba_dppy.compiler.DEBUG = 0 func(expected) np.testing.assert_array_equal(expected, got_cpu) - self.assertTrue("Parfor lowered on DPPY-device" in got_cpu_message.getvalue()) + self.assertIn("Parfor lowered on DPPY-device", got_cpu_message.getvalue()) @unittest.skipIf(not dpctl.has_gpu_queues(), "No GPU platforms available") def test_with_dppy_context_target(self): @@ -97,19 +99,19 @@ def func_no_parallel(b): b = np.ones((64), dtype=np.float64) with self.assertRaises(errors.UnsupportedError) as raises_1: - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): nested_func_target(a, b) with self.assertRaises(errors.UnsupportedError) as raises_2: - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): func_target(a) with self.assertRaises(errors.UnsupportedError) as raises_3: - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): func_no_target(a) with self.assertRaises(errors.UnsupportedError) as raises_4: - with dpctl.device_context("opencl:gpu"): + with device_context("opencl:gpu"): func_no_parallel(a) msg_1 = "Can't use 'with' context with explicitly specified target"