From fa1a22b60b56e24b088ccc632c85f104df1bba45 Mon Sep 17 00:00:00 2001 From: 96RadhikaJadhav Date: Mon, 30 Nov 2020 12:07:04 +0530 Subject: [PATCH] tests(dppl): removed folder --- numba_dppy/tests/dppl/__init__.py | 6 - numba_dppy/tests/dppl/test_arg_accessor.py | 70 ----- numba_dppy/tests/dppl/test_arg_types.py | 96 ------ numba_dppy/tests/dppl/test_atomic_op.py | 284 ----------------- numba_dppy/tests/dppl/test_barrier.py | 81 ----- numba_dppy/tests/dppl/test_black_scholes.py | 121 ------- numba_dppy/tests/dppl/test_caching.py | 37 --- .../tests/dppl/test_device_array_args.py | 49 --- numba_dppy/tests/dppl/test_dpctl_api.py | 25 -- numba_dppy/tests/dppl/test_dpnp_functions.py | 218 ------------- numba_dppy/tests/dppl/test_dppl_fallback.py | 56 ---- numba_dppy/tests/dppl/test_dppl_func.py | 63 ---- numba_dppy/tests/dppl/test_math_functions.py | 142 --------- .../test_numpy_bit_twiddling_functions.py | 111 ------- .../dppl/test_numpy_comparison_functions.py | 193 ------------ .../dppl/test_numpy_floating_functions.py | 95 ------ .../tests/dppl/test_numpy_math_functions.py | 296 ------------------ .../test_numpy_trigonomteric_functions.py | 222 ------------- .../tests/dppl/test_parfor_lower_message.py | 36 --- numba_dppy/tests/dppl/test_prange.py | 155 --------- numba_dppy/tests/dppl/test_print.py | 39 --- numba_dppy/tests/dppl/test_sum_reduction.py | 48 --- numba_dppy/tests/dppl/test_vectorize.py | 43 --- numba_dppy/tests/dppl/test_with_context.py | 120 ------- 24 files changed, 2606 deletions(-) delete mode 100644 numba_dppy/tests/dppl/__init__.py delete mode 100644 numba_dppy/tests/dppl/test_arg_accessor.py delete mode 100644 numba_dppy/tests/dppl/test_arg_types.py delete mode 100644 numba_dppy/tests/dppl/test_atomic_op.py delete mode 100644 numba_dppy/tests/dppl/test_barrier.py delete mode 100644 numba_dppy/tests/dppl/test_black_scholes.py delete mode 100644 numba_dppy/tests/dppl/test_caching.py delete mode 100644 numba_dppy/tests/dppl/test_device_array_args.py delete mode 100644 numba_dppy/tests/dppl/test_dpctl_api.py delete mode 100644 numba_dppy/tests/dppl/test_dpnp_functions.py delete mode 100644 numba_dppy/tests/dppl/test_dppl_fallback.py delete mode 100644 numba_dppy/tests/dppl/test_dppl_func.py delete mode 100644 numba_dppy/tests/dppl/test_math_functions.py delete mode 100644 numba_dppy/tests/dppl/test_numpy_bit_twiddling_functions.py delete mode 100644 numba_dppy/tests/dppl/test_numpy_comparison_functions.py delete mode 100644 numba_dppy/tests/dppl/test_numpy_floating_functions.py delete mode 100644 numba_dppy/tests/dppl/test_numpy_math_functions.py delete mode 100644 numba_dppy/tests/dppl/test_numpy_trigonomteric_functions.py delete mode 100644 numba_dppy/tests/dppl/test_parfor_lower_message.py delete mode 100644 numba_dppy/tests/dppl/test_prange.py delete mode 100644 numba_dppy/tests/dppl/test_print.py delete mode 100644 numba_dppy/tests/dppl/test_sum_reduction.py delete mode 100644 numba_dppy/tests/dppl/test_vectorize.py delete mode 100644 numba_dppy/tests/dppl/test_with_context.py diff --git a/numba_dppy/tests/dppl/__init__.py b/numba_dppy/tests/dppl/__init__.py deleted file mode 100644 index cff5a36cc2..0000000000 --- a/numba_dppy/tests/dppl/__init__.py +++ /dev/null @@ -1,6 +0,0 @@ -from numba.testing import SerialSuite -from numba.testing import load_testsuite -import os - -def load_tests(loader, tests, pattern): - return SerialSuite(load_testsuite(loader, os.path.dirname(__file__))) diff --git a/numba_dppy/tests/dppl/test_arg_accessor.py b/numba_dppy/tests/dppl/test_arg_accessor.py deleted file mode 100644 index ecc5d839bb..0000000000 --- a/numba_dppy/tests/dppl/test_arg_accessor.py +++ /dev/null @@ -1,70 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - - -@dppl.kernel(access_types={"read_only": ['a', 'b'], "write_only": ['c'], "read_write": []}) -def sum_with_accessor(a, b, c): - i = dppl.get_global_id(0) - c[i] = a[i] + b[i] - -@dppl.kernel -def sum_without_accessor(a, b, c): - i = dppl.get_global_id(0) - c[i] = a[i] + b[i] - -def call_kernel(global_size, local_size, - A, B, C, func): - func[global_size, dppl.DEFAULT_LOCAL_SIZE](A, B, C) - - -global_size = 10 -local_size = 1 -N = global_size * local_size - -A = np.array(np.random.random(N), dtype=np.float32) -B = np.array(np.random.random(N), dtype=np.float32) -D = A + B - - -@unittest.skipUnless(dpctl.has_cpu_queues(), 'test only on CPU system') -class TestDPPLArgAccessorCPU(DPPLTestCase): - def test_arg_with_accessor(self): - C = np.ones_like(A) - with dpctl.device_context("opencl:cpu") as cpu_queue: - call_kernel(global_size, local_size, - A, B, C, sum_with_accessor) - self.assertTrue(np.all(D == C)) - - def test_arg_without_accessor(self): - C = np.ones_like(A) - with dpctl.device_context("opencl:cpu") as cpu_queue: - call_kernel(global_size, local_size, - A, B, C, sum_without_accessor) - self.assertTrue(np.all(D == C)) - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLArgAccessorGPU(DPPLTestCase): - def test_arg_with_accessor(self): - C = np.ones_like(A) - with dpctl.device_context("opencl:gpu") as gpu_queue: - call_kernel(global_size, local_size, - A, B, C, sum_with_accessor) - self.assertTrue(np.all(D == C)) - - def test_arg_without_accessor(self): - C = np.ones_like(A) - with dpctl.device_context("opencl:gpu") as gpu_queue: - call_kernel(global_size, local_size, - A, B, C, sum_without_accessor) - self.assertTrue(np.all(D == C)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_arg_types.py b/numba_dppy/tests/dppl/test_arg_types.py deleted file mode 100644 index fc2eae105d..0000000000 --- a/numba_dppy/tests/dppl/test_arg_types.py +++ /dev/null @@ -1,96 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - - -@dppl.kernel -def mul_kernel(A, B, test): - i = dppl.get_global_id(0) - B[i] = A[i] * test - -def call_mul_device_kernel(global_size, A, B, test): - mul_kernel[global_size, dppl.DEFAULT_LOCAL_SIZE](A, B, test) - - -global_size = 10 -N = global_size -A = np.array(np.random.random(N), dtype=np.float32) -B = np.array(np.random.random(N), dtype=np.float32) - - -@unittest.skipUnless(dpctl.has_cpu_queues(), 'test only on CPU system') -class TestDPPLArrayArgCPU(DPPLTestCase): - def test_integer_arg(self): - x = np.int32(2) - with dpctl.device_context("opencl:cpu") as cpu_queue: - call_mul_device_kernel(global_size, A, B, x) - self.assertTrue(np.all((A * x) == B)) - - def test_float_arg(self): - x = np.float32(2.0) - with dpctl.device_context("opencl:cpu") as cpu_queue: - call_mul_device_kernel(global_size, A, B, x) - self.assertTrue(np.all(A * x == B)) - - x = np.float64(3.0) - call_mul_device_kernel(global_size, A, B, x) - self.assertTrue(np.all(A * x == B)) - - def test_bool_arg(self): - @dppl.kernel - def check_bool_kernel(A, test): - if test: - A[0] = 111 - else: - A[0] = 222 - - A = np.array([0], dtype='float64') - - with dpctl.device_context("opencl:cpu") as cpu_queue: - check_bool_kernel[global_size, dppl.DEFAULT_LOCAL_SIZE](A, True) - self.assertTrue(A[0] == 111) - check_bool_kernel[global_size, dppl.DEFAULT_LOCAL_SIZE](A, False) - self.assertTrue(A[0] == 222) - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLArrayArgGPU(DPPLTestCase): - def test_integer_arg(self): - x = np.int32(2) - with dpctl.device_context("opencl:gpu") as gpu_queue: - call_mul_device_kernel(global_size, A, B, x) - self.assertTrue(np.all((A * x) == B)) - - def test_float_arg(self): - x = np.float32(2.0) - with dpctl.device_context("opencl:gpu") as gpu_queue: - call_mul_device_kernel(global_size, A, B, x) - self.assertTrue(np.all(A * x == B)) - - x = np.float64(3.0) - call_mul_device_kernel(global_size, A, B, x) - self.assertTrue(np.all(A * x == B)) - - def test_bool_arg(self): - @dppl.kernel - def check_bool_kernel(A, test): - if test: - A[0] = 111 - else: - A[0] = 222 - - A = np.array([0], dtype='float64') - - with dpctl.device_context("opencl:gpu") as gpu_queue: - check_bool_kernel[global_size, dppl.DEFAULT_LOCAL_SIZE](A, True) - self.assertTrue(A[0] == 111) - check_bool_kernel[global_size, dppl.DEFAULT_LOCAL_SIZE](A, False) - self.assertTrue(A[0] == 222) - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_atomic_op.py b/numba_dppy/tests/dppl/test_atomic_op.py deleted file mode 100644 index 9825c707d1..0000000000 --- a/numba_dppy/tests/dppl/test_atomic_op.py +++ /dev/null @@ -1,284 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -import numba -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - -def atomic_add_int32(ary): - tid = dppl.get_local_id(0) - lm = dppl.local.static_alloc(32, numba.uint32) - lm[tid] = 0 - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - bin = ary[tid] % 32 - dppl.atomic.add(lm, bin, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[tid] = lm[tid] - - -def atomic_sub_int32(ary): - tid = dppl.get_local_id(0) - lm = dppl.local.static_alloc(32, numba.uint32) - lm[tid] = 0 - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - bin = ary[tid] % 32 - dppl.atomic.sub(lm, bin, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[tid] = lm[tid] - - -def atomic_add_float32(ary): - lm = dppl.local.static_alloc(1, numba.float32) - lm[0] = ary[0] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.add(lm, 0, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[0] = lm[0] - - -def atomic_sub_float32(ary): - lm = dppl.local.static_alloc(1, numba.float32) - lm[0] = ary[0] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.sub(lm, 0, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[0] = lm[0] - - -def atomic_add_int64(ary): - lm = dppl.local.static_alloc(1, numba.int64) - lm[0] = ary[0] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.add(lm, 0, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[0] = lm[0] - - -def atomic_sub_int64(ary): - lm = dppl.local.static_alloc(1, numba.int64) - lm[0] = ary[0] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.sub(lm, 0, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[0] = lm[0] - - -def atomic_add_float64(ary): - lm = dppl.local.static_alloc(1, numba.float64) - lm[0] = ary[0] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.add(lm, 0, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[0] = lm[0] - - -def atomic_sub_float64(ary): - lm = dppl.local.static_alloc(1, numba.float64) - lm[0] = ary[0] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.sub(lm, 0, 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[0] = lm[0] - - -def atomic_add2(ary): - tx = dppl.get_local_id(0) - ty = dppl.get_local_id(1) - lm = dppl.local.static_alloc((4, 8), numba.uint32) - lm[tx, ty] = ary[tx, ty] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.add(lm, (tx, ty), 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[tx, ty] = lm[tx, ty] - - -def atomic_add3(ary): - tx = dppl.get_local_id(0) - ty = dppl.get_local_id(1) - lm = dppl.local.static_alloc((4, 8), numba.uint32) - lm[tx, ty] = ary[tx, ty] - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - dppl.atomic.add(lm, (tx, numba.uint64(ty)), 1) - dppl.barrier(dppl.CLK_GLOBAL_MEM_FENCE) - ary[tx, ty] = lm[tx, ty] - - -def call_fn_for_datatypes(fn, result, input, global_size): - dtypes = [np.int32, np.int64, np.float32, np.double] - - for dtype in dtypes: - a = np.array(input, dtype=dtype) - - with dpctl.device_context("opencl:gpu") as gpu_queue: - # TODO: dpctl needs to expose this functions - #if dtype == np.double and not device_env.device_support_float64_atomics(): - # continue - #if dtype == np.int64 and not device_env.device_support_int64_atomics(): - # continue - fn[global_size, dppl.DEFAULT_LOCAL_SIZE](a) - - assert(a[0] == result) - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -@unittest.skipUnless(numba_dppy.ocl.atomic_support_present(), 'test only when atomic support is present') -class TestAtomicOp(DPPLTestCase): - def test_atomic_add_global(self): - @dppl.kernel - def atomic_add(B): - dppl.atomic.add(B, 0, 1) - - N = 100 - B = np.array([0]) - - call_fn_for_datatypes(atomic_add, N, B, N) - - - def test_atomic_sub_global(self): - @dppl.kernel - def atomic_sub(B): - dppl.atomic.sub(B, 0, 1) - - N = 100 - B = np.array([100]) - - call_fn_for_datatypes(atomic_sub, 0, B, N) - - - def test_atomic_add_local_int32(self): - ary = np.random.randint(0, 32, size=32).astype(np.uint32) - orig = ary.copy() - - #dppl_atomic_add = dppl.kernel('void(uint32[:])')(atomic_add_int32) - dppl_atomic_add = dppl.kernel(atomic_add_int32) - with dpctl.device_context("opencl:gpu") as gpu_queue: - dppl_atomic_add[32, dppl.DEFAULT_LOCAL_SIZE](ary) - - gold = np.zeros(32, dtype=np.uint32) - for i in range(orig.size): - gold[orig[i]] += 1 - - self.assertTrue(np.all(ary == gold)) - - - def test_atomic_sub_local_int32(self): - ary = np.random.randint(0, 32, size=32).astype(np.uint32) - orig = ary.copy() - - #dppl_atomic_sub = dppl.kernel('void(uint32[:])')(atomic_sub_int32) - dppl_atomic_sub = dppl.kernel(atomic_sub_int32) - with dpctl.device_context("opencl:gpu") as gpu_queue: - dppl_atomic_sub[32, dppl.DEFAULT_LOCAL_SIZE](ary) - - gold = np.zeros(32, dtype=np.uint32) - for i in range(orig.size): - gold[orig[i]] -= 1 - - self.assertTrue(np.all(ary == gold)) - - - def test_atomic_add_local_float32(self): - ary = np.array([0], dtype=np.float32) - - #dppl_atomic_add = dppl.kernel('void(float32[:])')(atomic_add_float32) - dppl_atomic_add = dppl.kernel(atomic_add_float32) - with dpctl.device_context("opencl:gpu") as gpu_queue: - dppl_atomic_add[32, dppl.DEFAULT_LOCAL_SIZE](ary) - - self.assertTrue(ary[0] == 32) - - - def test_atomic_sub_local_float32(self): - ary = np.array([32], dtype=np.float32) - - #dppl_atomic_sub = dppl.kernel('void(float32[:])')(atomic_sub_float32) - dppl_atomic_sub = dppl.kernel(atomic_sub_float32) - with dpctl.device_context("opencl:gpu") as gpu_queue: - - dppl_atomic_sub[32, dppl.DEFAULT_LOCAL_SIZE](ary) - - self.assertTrue(ary[0] == 0) - - - def test_atomic_add_local_int64(self): - ary = np.array([0], dtype=np.int64) - - #dppl_atomic_add = dppl.kernel('void(int64[:])')(atomic_add_int64) - dppl_atomic_add = dppl.kernel(atomic_add_int64) - with dpctl.device_context("opencl:gpu") as gpu_queue: - # TODO: dpctl needs to expose this functions - #if device_env.device_support_int64_atomics(): - dppl_atomic_add[32, dppl.DEFAULT_LOCAL_SIZE](ary) - self.assertTrue(ary[0] == 32) - #else: - # return - - - def test_atomic_sub_local_int64(self): - ary = np.array([32], dtype=np.int64) - - #fn = dppl.kernel('void(int64[:])')(atomic_sub_int64) - fn = dppl.kernel(atomic_sub_int64) - with dpctl.device_context("opencl:gpu") as gpu_queue: - # TODO: dpctl needs to expose this functions - #if device_env.device_support_int64_atomics(): - fn[32, dppl.DEFAULT_LOCAL_SIZE](ary) - self.assertTrue(ary[0] == 0) - #else: - # return - - - def test_atomic_add_local_float64(self): - ary = np.array([0], dtype=np.double) - - #fn = dppl.kernel('void(float64[:])')(atomic_add_float64) - fn = dppl.kernel(atomic_add_float64) - with dpctl.device_context("opencl:gpu") as gpu_queue: - # TODO: dpctl needs to expose this functions - #if device_env.device_support_float64_atomics(): - fn[32, dppl.DEFAULT_LOCAL_SIZE](ary) - self.assertTrue(ary[0] == 32) - #else: - # return - - - def test_atomic_sub_local_float64(self): - ary = np.array([32], dtype=np.double) - - #fn = dppl.kernel('void(float64[:])')(atomic_sub_int64) - fn = dppl.kernel(atomic_sub_int64) - with dpctl.device_context("opencl:gpu") as gpu_queue: - # TODO: dpctl needs to expose this functions - #if device_env.device_support_float64_atomics(): - fn[32, dppl.DEFAULT_LOCAL_SIZE](ary) - self.assertTrue(ary[0] == 0) - #else: - # return - - - def test_atomic_add2(self): - ary = np.random.randint(0, 32, size=32).astype(np.uint32).reshape(4, 8) - orig = ary.copy() - #dppl_atomic_add2 = dppl.kernel('void(uint32[:,:])')(atomic_add2) - dppl_atomic_add2 = dppl.kernel(atomic_add2) - with dpctl.device_context("opencl:gpu") as gpu_queue: - dppl_atomic_add2[(4, 8), dppl.DEFAULT_LOCAL_SIZE](ary) - self.assertTrue(np.all(ary == orig + 1)) - - - def test_atomic_add3(self): - ary = np.random.randint(0, 32, size=32).astype(np.uint32).reshape(4, 8) - orig = ary.copy() - #dppl_atomic_add3 = dppl.kernel('void(uint32[:,:])')(atomic_add3) - dppl_atomic_add3 = dppl.kernel(atomic_add3) - with dpctl.device_context("opencl:gpu") as gpu_queue: - dppl_atomic_add3[(4, 8), dppl.DEFAULT_LOCAL_SIZE](ary) - - self.assertTrue(np.all(ary == orig + 1)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_barrier.py b/numba_dppy/tests/dppl/test_barrier.py deleted file mode 100644 index aeff16dd40..0000000000 --- a/numba_dppy/tests/dppl/test_barrier.py +++ /dev/null @@ -1,81 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -from numba import float32 -import numba_dppy, numba_dppy as dppl -import dpctl - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestBarrier(unittest.TestCase): - def test_proper_lowering(self): - #@dppl.kernel("void(float32[::1])") - @dppl.kernel - def twice(A): - i = dppl.get_global_id(0) - d = A[i] - dppl.barrier(dppl.CLK_LOCAL_MEM_FENCE) # local mem fence - A[i] = d * 2 - - N = 256 - arr = np.random.random(N).astype(np.float32) - orig = arr.copy() - - with dpctl.device_context("opencl:gpu") as gpu_queue: - twice[N, N//2](arr) - - # The computation is correct? - np.testing.assert_allclose(orig * 2, arr) - - def test_no_arg_barrier_support(self): - #@dppl.kernel("void(float32[::1])") - @dppl.kernel - def twice(A): - i = dppl.get_global_id(0) - d = A[i] - # no argument defaults to global mem fence - dppl.barrier() - A[i] = d * 2 - - N = 256 - arr = np.random.random(N).astype(np.float32) - orig = arr.copy() - - with dpctl.device_context("opencl:gpu") as gpu_queue: - twice[N, dppl.DEFAULT_LOCAL_SIZE](arr) - - # The computation is correct? - np.testing.assert_allclose(orig * 2, arr) - - - def test_local_memory(self): - blocksize = 10 - - #@dppl.kernel("void(float32[::1])") - @dppl.kernel - def reverse_array(A): - lm = dppl.local.static_alloc(shape=10, dtype=float32) - i = dppl.get_global_id(0) - - # preload - lm[i] = A[i] - # barrier local or global will both work as we only have one work group - dppl.barrier(dppl.CLK_LOCAL_MEM_FENCE) # local mem fence - # write - A[i] += lm[blocksize - 1 - i] - - arr = np.arange(blocksize).astype(np.float32) - orig = arr.copy() - - with dpctl.device_context("opencl:gpu") as gpu_queue: - reverse_array[blocksize, dppl.DEFAULT_LOCAL_SIZE](arr) - - expected = orig[::-1] + orig - np.testing.assert_allclose(expected, arr) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_black_scholes.py b/numba_dppy/tests/dppl/test_black_scholes.py deleted file mode 100644 index 3d9581bb54..0000000000 --- a/numba_dppy/tests/dppl/test_black_scholes.py +++ /dev/null @@ -1,121 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np -import math -import time - -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - - -RISKFREE = 0.02 -VOLATILITY = 0.30 - -A1 = 0.31938153 -A2 = -0.356563782 -A3 = 1.781477937 -A4 = -1.821255978 -A5 = 1.330274429 -RSQRT2PI = 0.39894228040143267793994605993438 - - -def cnd(d): - K = 1.0 / (1.0 + 0.2316419 * np.abs(d)) - ret_val = (RSQRT2PI * np.exp(-0.5 * d * d) * - (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5)))))) - return np.where(d > 0, 1.0 - ret_val, ret_val) - -def black_scholes(callResult, putResult, stockPrice, optionStrike, optionYears, - Riskfree, Volatility): - S = stockPrice - X = optionStrike - T = optionYears - R = Riskfree - V = Volatility - sqrtT = np.sqrt(T) - d1 = (np.log(S / X) + (R + 0.5 * V * V) * T) / (V * sqrtT) - d2 = d1 - V * sqrtT - cndd1 = cnd(d1) - cndd2 = cnd(d2) - - expRT = np.exp(- R * T) - callResult[:] = (S * cndd1 - X * expRT * cndd2) - putResult[:] = (X * expRT * (1.0 - cndd2) - S * (1.0 - cndd1)) - -def randfloat(rand_var, low, high): - return (1.0 - rand_var) * low + rand_var * high - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLBlackScholes(DPPLTestCase): - def test_black_scholes(self): - OPT_N = 400 - iterations = 2 - - stockPrice = randfloat(np.random.random(OPT_N), 5.0, 30.0) - optionStrike = randfloat(np.random.random(OPT_N), 1.0, 100.0) - optionYears = randfloat(np.random.random(OPT_N), 0.25, 10.0) - - callResultNumpy = np.zeros(OPT_N) - putResultNumpy = -np.ones(OPT_N) - - callResultNumbapro = np.zeros(OPT_N) - putResultNumbapro = -np.ones(OPT_N) - - # numpy - for i in range(iterations): - black_scholes(callResultNumpy, putResultNumpy, stockPrice, - optionStrike, optionYears, RISKFREE, VOLATILITY) - - - @dppl.kernel - def black_scholes_dppl(callResult, putResult, S, X, T, R, V): - i = dppl.get_global_id(0) - if i >= S.shape[0]: - return - sqrtT = math.sqrt(T[i]) - d1 = (math.log(S[i] / X[i]) + (R + 0.5 * V * V) * T[i]) / (V * sqrtT) - d2 = d1 - V * sqrtT - - K = 1.0 / (1.0 + 0.2316419 * math.fabs(d1)) - cndd1 = (RSQRT2PI * math.exp(-0.5 * d1 * d1) * - (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5)))))) - if d1 > 0: - cndd1 = 1.0 - cndd1 - - K = 1.0 / (1.0 + 0.2316419 * math.fabs(d2)) - cndd2 = (RSQRT2PI * math.exp(-0.5 * d2 * d2) * - (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5)))))) - if d2 > 0: - cndd2 = 1.0 - cndd2 - - expRT = math.exp((-1. * R) * T[i]) - callResult[i] = (S[i] * cndd1 - X[i] * expRT * cndd2) - putResult[i] = (X[i] * expRT * (1.0 - cndd2) - S[i] * (1.0 - cndd1)) - - # numbapro - time0 = time.time() - blockdim = 512, 1 - griddim = int(math.ceil(float(OPT_N) / blockdim[0])), 1 - - with dpctl.device_context("opencl:gpu") as gpu_queue: - time1 = time.time() - for i in range(iterations): - black_scholes_dppl[blockdim, griddim]( - callResultNumbapro, putResultNumbapro, stockPrice, optionStrike, - optionYears, RISKFREE, VOLATILITY) - - - dt = (time1 - time0) - - delta = np.abs(callResultNumpy - callResultNumbapro) - L1norm = delta.sum() / np.abs(callResultNumpy).sum() - - max_abs_err = delta.max() - self.assertTrue(L1norm < 1e-13) - self.assertTrue(max_abs_err < 1e-13) - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_caching.py b/numba_dppy/tests/dppl/test_caching.py deleted file mode 100644 index 6a6a7967a5..0000000000 --- a/numba_dppy/tests/dppl/test_caching.py +++ /dev/null @@ -1,37 +0,0 @@ -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -import numba_dppy, numba_dppy as dppl -import dpctl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -def data_parallel_sum(a, b, c): - i = dppl.get_global_id(0) - c[i] = a[i] + b[i] - - -class TestCaching(DPPLTestCase): - def test_caching_kernel(self): - global_size = 10 - N = global_size - - a = np.array(np.random.random(N), dtype=np.float32) - b = np.array(np.random.random(N), dtype=np.float32) - c = np.ones_like(a) - - - with dpctl.device_context("opencl:gpu") as gpu_queue: - func = dppl.kernel(data_parallel_sum) - caching_kernel = func[global_size, dppl.DEFAULT_LOCAL_SIZE].specialize(a, b, c) - - for i in range(10): - cached_kernel = func[global_size, dppl.DEFAULT_LOCAL_SIZE].specialize(a, b, c) - self.assertIs(caching_kernel, cached_kernel) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_device_array_args.py b/numba_dppy/tests/dppl/test_device_array_args.py deleted file mode 100644 index 024e3723a9..0000000000 --- a/numba_dppy/tests/dppl/test_device_array_args.py +++ /dev/null @@ -1,49 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -import numba_dppy, numba_dppy as dppl -import dpctl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - -@dppl.kernel -def data_parallel_sum(a, b, c): - i = dppl.get_global_id(0) - c[i] = a[i] + b[i] - - -global_size = 64 -N = global_size - -a = np.array(np.random.random(N), dtype=np.float32) -b = np.array(np.random.random(N), dtype=np.float32) -d = a + b - - -@unittest.skipUnless(dpctl.has_cpu_queues(), 'test only on CPU system') -class TestDPPLDeviceArrayArgsGPU(DPPLTestCase): - def test_device_array_args_cpu(self): - c = np.ones_like(a) - - with dpctl.device_context("opencl:cpu") as cpu_queue: - data_parallel_sum[global_size, dppl.DEFAULT_LOCAL_SIZE](a, b, c) - - self.assertTrue(np.all(c == d)) - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLDeviceArrayArgsCPU(DPPLTestCase): - def test_device_array_args_gpu(self): - c = np.ones_like(a) - - with dpctl.device_context("opencl:gpu") as gpu_queue: - data_parallel_sum[global_size, dppl.DEFAULT_LOCAL_SIZE](a, b, c) - - self.assertTrue(np.all(c == d)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_dpctl_api.py b/numba_dppy/tests/dppl/test_dpctl_api.py deleted file mode 100644 index bb72a35cf2..0000000000 --- a/numba_dppy/tests/dppl/test_dpctl_api.py +++ /dev/null @@ -1,25 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPCTLAPI(DPPLTestCase): - def test_dpctl_api(self): - with dpctl.device_context("opencl:gpu") as gpu_queue: - dpctl.dump() - dpctl.get_current_queue() - dpctl.get_num_platforms() - dpctl.get_num_activated_queues() - dpctl.has_cpu_queues() - dpctl.has_gpu_queues() - dpctl.has_sycl_platforms() - dpctl.is_in_device_context() - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_dpnp_functions.py b/numba_dppy/tests/dppl/test_dpnp_functions.py deleted file mode 100644 index bbffb30c3f..0000000000 --- a/numba_dppy/tests/dppl/test_dpnp_functions.py +++ /dev/null @@ -1,218 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -def test_for_different_datatypes(fn, test_fn, dims, arg_count, tys, np_all=False, matrix=None): - if arg_count == 1: - for ty in tys: - if matrix and matrix[0]: - a = np.array(np.random.random(dims[0] * dims[1]), dtype=ty).reshape(dims[0], dims[1]) - else: - a = np.array(np.random.random(dims[0]), dtype=ty) - c = fn(a) - d = test_fn(a) - if np_all: - max_abs_err = np.all(c - d) - else: - max_abs_err = c - d - if not (max_abs_err < 1e-4): - return False - - elif arg_count == 2: - for ty in tys: - if matrix and matrix[0]: - a = np.array(np.random.random(dims[0] * dims[1]), dtype=ty).reshape(dims[0], dims[1]) - else: - a = np.array(np.random.random(dims[0] * dims[1]), dtype=ty) - if matrix and matrix[1]: - b = np.array(np.random.random(dims[2] * dims[3]), dtype=ty).reshape(dims[2], dims[3]) - else: - b = np.array(np.random.random(dims[2] * dims[3]), dtype=ty) - - c = fn(a, b) - d = test_fn(a, b) - if np_all: - max_abs_err = np.sum(c - d) - else: - max_abs_err = c - d - if not (max_abs_err < 1e-4): - return False - - return True - -def test_for_dimensions(fn, test_fn, dims, tys, np_all=False): - total_size = 1 - for d in dims: - total_size *= d - - for ty in tys: - a = np.array(np.random.random(total_size), dtype=ty).reshape(dims) - c = fn(a) - d = test_fn(a) - if np_all: - max_abs_err = np.all(c - d) - else: - max_abs_err = c - d - if not (max_abs_err < 1e-4): - return False - - return True - -def ensure_dpnp(): - try: - # import dpnp - from numba_dppy.dpnp_glue import dpnp_fptr_interface as dpnp_glue - return True - except: - return False - - -@unittest.skipUnless(ensure_dpnp(), 'test only when dpNP is available') -class Testdpnp_functions(DPPLTestCase): - N = 10 - - a = np.array(np.random.random(N), dtype=np.float32) - b = np.array(np.random.random(N), dtype=np.float32) - tys = [np.int32, np.uint32, np.int64, np.uint64, np.float, np.double] - - def test_sum(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.sum(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.sum, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.sum, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.sum, [10, 2, 3], self.tys)) - - def test_prod(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.prod(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.prod, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.prod, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.prod, [10, 2, 3], self.tys)) - - def test_argmax(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.argmax(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.argmax, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.argmax, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.argmax, [10, 2, 3], self.tys)) - - def test_max(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.max(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.max, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.max, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.max, [10, 2, 3], self.tys)) - - def test_argmin(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.argmin(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.argmin, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.argmin, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.argmin, [10, 2, 3], self.tys)) - - def test_min(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.min(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.min, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.min, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.min, [10, 2, 3], self.tys)) - - def test_argsort(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.argsort(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.argmin, [10], 1, self.tys, np_all=True)) - - def test_median(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.median(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.median, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.median, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.median, [10, 2, 3], self.tys)) - - def test_mean(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.mean(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.mean, [10], 1, self.tys)) - self.assertTrue(test_for_dimensions(f, np.mean, [10, 2], self.tys)) - self.assertTrue(test_for_dimensions(f, np.mean, [10, 2, 3], self.tys)) - - def test_matmul(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.matmul(a, b) - return c - - self.assertTrue(test_for_different_datatypes(f, np.matmul, [10, 5, 5, 10], 2, [np.float, np.double], np_all=True, matrix=[True, True])) - - def test_dot(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.dot(a, b) - return c - - self.assertTrue(test_for_different_datatypes(f, np.dot, [10, 1, 10, 1], 2, [np.float, np.double])) - self.assertTrue(test_for_different_datatypes(f, np.dot, [10, 1, 10, 2], 2, [np.float, np.double], matrix=[False, True], np_all=True)) - self.assertTrue(test_for_different_datatypes(f, np.dot, [2, 10, 10, 1], 2, [np.float, np.double], matrix=[True, False], np_all=True)) - self.assertTrue(test_for_different_datatypes(f, np.dot, [10, 2, 2, 10], 2, [np.float, np.double], matrix=[True, True], np_all=True)) - - - def test_cov(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.cov(a) - return c - - self.assertTrue(test_for_different_datatypes(f, np.cov, [10, 7], 1, self.tys, matrix=[True], np_all=True)) - - def test_dpnp_interacting_with_parfor(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.sum(a) - e = np.add(b, a) - #d = a + 1 - return 0 - - result = f(self.a, self.b) - #np_result = np.add((self.a + np.sum(self.a)), self.b) - - #max_abs_err = result.sum() - np_result.sum() - #self.assertTrue(max_abs_err < 1e-4) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_dppl_fallback.py b/numba_dppy/tests/dppl/test_dppl_fallback.py deleted file mode 100644 index adb7ae868b..0000000000 --- a/numba_dppy/tests/dppl/test_dppl_fallback.py +++ /dev/null @@ -1,56 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -import numba -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -from numba.tests.support import captured_stderr -import dpctl -import sys -import io - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLFallback(DPPLTestCase): - def test_dppl_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 - - with captured_stderr() as msg: - dppl = numba.njit(parallel={'offload':True})(inner_call_fallback) - dppl_result = dppl() - - ref_result = inner_call_fallback() - - np.testing.assert_array_equal(dppl_result, ref_result) - self.assertTrue('Failed to lower parfor on DPPL-device' in msg.getvalue()) - - def test_dppl_fallback_reductions(self): - def reduction(a): - return np.amax(a) - - a = np.ones(10) - with captured_stderr() as msg: - dppl = numba.njit(parallel={'offload':True})(reduction) - dppl_result = dppl(a) - - ref_result = reduction(a) - - np.testing.assert_array_equal(dppl_result, ref_result) - self.assertTrue('Failed to lower parfor on DPPL-device' in msg.getvalue()) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_dppl_func.py b/numba_dppy/tests/dppl/test_dppl_func.py deleted file mode 100644 index 0f64046082..0000000000 --- a/numba_dppy/tests/dppl/test_dppl_func.py +++ /dev/null @@ -1,63 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np - -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLFunc(DPPLTestCase): - N = 257 - - def test_dppl_func_device_array(self): - @dppl.func - def g(a): - return a + 1 - - @dppl.kernel - def f(a, b): - i = dppl.get_global_id(0) - b[i] = g(a[i]) - - a = np.ones(self.N) - b = np.ones(self.N) - - with dpctl.device_context("opencl:gpu") as gpu_queue: - f[self.N, dppl.DEFAULT_LOCAL_SIZE](a, b) - - - self.assertTrue(np.all(b == 2)) - - def test_dppl_func_ndarray(self): - @dppl.func - def g(a): - return a + 1 - - @dppl.kernel - def f(a, b): - i = dppl.get_global_id(0) - b[i] = g(a[i]) - - @dppl.kernel - def h(a, b): - i = dppl.get_global_id(0) - b[i] = g(a[i]) + 1 - - a = np.ones(self.N) - b = np.ones(self.N) - - with dpctl.device_context("opencl:gpu") as gpu_queue: - f[self.N, dppl.DEFAULT_LOCAL_SIZE](a, b) - - self.assertTrue(np.all(b == 2)) - - h[self.N, dppl.DEFAULT_LOCAL_SIZE](a, b) - - self.assertTrue(np.all(b == 3)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_math_functions.py b/numba_dppy/tests/dppl/test_math_functions.py deleted file mode 100644 index 977fe85fef..0000000000 --- a/numba_dppy/tests/dppl/test_math_functions.py +++ /dev/null @@ -1,142 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -import numba_dppy, numba_dppy as dppl -import dpctl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import math - -@dppl.kernel -def dppl_fabs(a,b): - i = dppl.get_global_id(0) - b[i] = math.fabs(a[i]) - -@dppl.kernel -def dppl_exp(a,b): - i = dppl.get_global_id(0) - b[i] = math.exp(a[i]) - -@dppl.kernel -def dppl_log(a,b): - i = dppl.get_global_id(0) - b[i] = math.log(a[i]) - -@dppl.kernel -def dppl_sqrt(a,b): - i = dppl.get_global_id(0) - b[i] = math.sqrt(a[i]) - -@dppl.kernel -def dppl_sin(a,b): - i = dppl.get_global_id(0) - b[i] = math.sin(a[i]) - -@dppl.kernel -def dppl_cos(a,b): - i = dppl.get_global_id(0) - b[i] = math.cos(a[i]) - -@dppl.kernel -def dppl_tan(a,b): - i = dppl.get_global_id(0) - b[i] = math.tan(a[i]) - -global_size = 10 -N = global_size - -a = np.array(np.random.random(N), dtype=np.float32) - -def driver(a, jitfunc): - b = np.ones_like(a) - # Device buffers - jitfunc[global_size, dppl.DEFAULT_LOCAL_SIZE](a, b) - return b - - -def test_driver(input_arr, device_ty, jitfunc): - out_actual = None - if device_ty == "GPU": - with dpctl.device_context("opencl:gpu") as gpu_queue: - out_actual = driver(input_arr, jitfunc) - elif device_ty == "CPU": - with dpctl.device_context("opencl:cpu") as cpu_queue: - out_actual = driver(input_arr, jitfunc) - else: - print("Unknown device type") - raise SystemExit() - - return out_actual - - -@unittest.skipUnless(dpctl.has_cpu_queues(), 'test only on CPU system') -class TestDPPLMathFunctionsCPU(DPPLTestCase): - def test_fabs_cpu(self): - b_actual = test_driver(a, "CPU", dppl_fabs) - b_expected = np.fabs(a) - self.assertTrue(np.all(b_actual == b_expected)) - - def test_sin_cpu(self): - b_actual = test_driver(a, "CPU", dppl_sin) - b_expected = np.sin(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_cos_cpu(self): - b_actual = test_driver(a, "CPU", dppl_cos) - b_expected = np.cos(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_exp_cpu(self): - b_actual = test_driver(a, "CPU", dppl_exp) - b_expected = np.exp(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_sqrt_cpu(self): - b_actual = test_driver(a, "CPU", dppl_sqrt) - b_expected = np.sqrt(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_log_cpu(self): - b_actual = test_driver(a, "CPU", dppl_log) - b_expected = np.log(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLMathFunctionsGPU(DPPLTestCase): - def test_fabs_gpu(self): - b_actual = test_driver(a, "GPU", dppl_fabs) - b_expected = np.fabs(a) - self.assertTrue(np.all(b_actual == b_expected)) - - def test_sin_gpu(self): - b_actual = test_driver(a, "GPU", dppl_sin) - b_expected = np.sin(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_cos_gpu(self): - b_actual = test_driver(a, "GPU", dppl_cos) - b_expected = np.cos(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_exp_gpu(self): - b_actual = test_driver(a, "GPU", dppl_exp) - b_expected = np.exp(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_sqrt_gpu(self): - b_actual = test_driver(a, "GPU", dppl_sqrt) - b_expected = np.sqrt(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - def test_log_gpu(self): - b_actual = test_driver(a, "GPU", dppl_log) - b_expected = np.log(a) - self.assertTrue(np.allclose(b_actual,b_expected)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_numpy_bit_twiddling_functions.py b/numba_dppy/tests/dppl/test_numpy_bit_twiddling_functions.py deleted file mode 100644 index 5e3cd9ba24..0000000000 --- a/numba_dppy/tests/dppl/test_numpy_bit_twiddling_functions.py +++ /dev/null @@ -1,111 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -class TestNumpy_bit_twiddling_functions(DPPLTestCase): - def test_bitwise_and(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.bitwise_and(a, b) - return c - - a = np.array([2,5,255]) - b = np.array([3,14,16]) - - c = f(a, b) - d = np.bitwise_and(a, b) - self.assertTrue(np.all(c == d)) - - - def test_bitwise_or(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.bitwise_or(a, b) - return c - - a = np.array([2,5,255]) - b = np.array([4,4,4]) - - c = f(a, b) - d = np.bitwise_or(a, b) - self.assertTrue(np.all(c == d)) - - - def test_bitwise_xor(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.bitwise_xor(a, b) - return c - - a = np.array([2,5,255]) - b = np.array([4,4,4]) - - c = f(a, b) - d = np.bitwise_xor(a, b) - self.assertTrue(np.all(c == d)) - - - def test_bitwise_not(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.bitwise_not(a) - return c - - a = np.array([2,5,255]) - - c = f(a) - d = np.bitwise_not(a) - self.assertTrue(np.all(c == d)) - - - def test_invert(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.invert(a) - return c - - a = np.array([2,5,255]) - - c = f(a) - d = np.invert(a) - self.assertTrue(np.all(c == d)) - - - def test_left_shift(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.left_shift(a, b) - return c - - a = np.array([2,3,4]) - b = np.array([1,2,3]) - - c = f(a, b) - d = np.left_shift(a, b) - self.assertTrue(np.all(c == d)) - - - def test_right_shift(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.right_shift(a, b) - return c - - a = np.array([2,3,4]) - b = np.array([1,2,3]) - - c = f(a, b) - d = np.right_shift(a, b) - self.assertTrue(np.all(c == d)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_numpy_comparison_functions.py b/numba_dppy/tests/dppl/test_numpy_comparison_functions.py deleted file mode 100644 index 0bd7dcbb69..0000000000 --- a/numba_dppy/tests/dppl/test_numpy_comparison_functions.py +++ /dev/null @@ -1,193 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - -class TestNumpy_comparison_functions(DPPLTestCase): - a = np.array([4,5,6]) - b = np.array([2,6,6]) - def test_greater(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.greater(a, b) - return c - - c = f(self.a, self.b) - d = np.greater(self.a, self.b) - self.assertTrue(np.all(c == d)) - - - def test_greater_equal(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.greater_equal(a, b) - return c - - c = f(self.a, self.b) - d = np.greater_equal(self.a, self.b) - self.assertTrue(np.all(c == d)) - - - def test_less(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.less(a, b) - return c - - c = f(self.a, self.b) - d = np.less(self.a, self.b) - self.assertTrue(np.all(c == d)) - - - def test_less_equal(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.less_equal(a, b) - return c - - c = f(self.a, self.b) - d = np.less_equal(self.a, self.b) - self.assertTrue(np.all(c == d)) - - - def test_not_equal(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.not_equal(a, b) - return c - - c = f(self.a, self.b) - d = np.not_equal(self.a, self.b) - self.assertTrue(np.all(c == d)) - - - def test_equal(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.equal(a, b) - return c - - c = f(self.a, self.b) - d = np.equal(self.a, self.b) - self.assertTrue(np.all(c == d)) - - - def test_logical_and(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.logical_and(a, b) - return c - - a = np.array([True, True, False]) - b = np.array([True, False, False]) - - c = f(a, b) - d = np.logical_and(a, b) - self.assertTrue(np.all(c == d)) - - - def test_logical_or(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.logical_or(a, b) - return c - - a = np.array([True, True, False]) - b = np.array([True, False, False]) - - c = f(a, b) - d = np.logical_or(a, b) - self.assertTrue(np.all(c == d)) - - - def test_logical_xor(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.logical_xor(a, b) - return c - - a = np.array([True, True, False]) - b = np.array([True, False, False]) - - c = f(a, b) - d = np.logical_xor(a, b) - self.assertTrue(np.all(c == d)) - - - def test_logical_not(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.logical_not(a) - return c - - a = np.array([True, True, False]) - - c = f(a) - d = np.logical_not(a) - self.assertTrue(np.all(c == d)) - - - def test_maximum(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.maximum(a, b) - return c - - a = np.array([5,6,7,np.nan], dtype=np.float32) - b = np.array([5,7,6,100], dtype=np.float32) - - c = f(a, b) - d = np.maximum(a, b) - np.testing.assert_equal(c, d) - - - def test_minimum(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.minimum(a, b) - return c - - a = np.array([5,6,7,np.nan], dtype=np.float32) - b = np.array([5,7,6,100], dtype=np.float32) - - c = f(a, b) - d = np.minimum(a, b) - np.testing.assert_equal(c, d) - - - def test_fmax(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.fmax(a, b) - return c - - a = np.array([5,6,7,np.nan], dtype=np.float32) - b = np.array([5,7,6,100], dtype=np.float32) - - c = f(a, b) - d = np.fmax(a, b) - np.testing.assert_equal(c, d) - - - def test_fmin(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.fmin(a, b) - return c - - a = np.array([5,6,7,np.nan], dtype=np.float32) - b = np.array([5,7,6,100], dtype=np.float32) - - c = f(a, b) - d = np.fmin(a, b) - np.testing.assert_equal(c, d) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_numpy_floating_functions.py b/numba_dppy/tests/dppl/test_numpy_floating_functions.py deleted file mode 100644 index 62b76b1ade..0000000000 --- a/numba_dppy/tests/dppl/test_numpy_floating_functions.py +++ /dev/null @@ -1,95 +0,0 @@ -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -class TestNumpy_floating_functions(DPPLTestCase): - def test_isfinite(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.isfinite(a) - return c - - test_arr = [np.log(-1.),1.,np.log(0)] - input_arr = np.asarray(test_arr, dtype=np.float32) - - c = f(input_arr) - d = np.isfinite(input_arr) - self.assertTrue(np.all(c == d)) - - - def test_isinf(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.isinf(a) - return c - - test_arr = [np.log(-1.),1.,np.log(0)] - input_arr = np.asarray(test_arr, dtype=np.float32) - - c = f(input_arr) - d = np.isinf(input_arr) - self.assertTrue(np.all(c == d)) - - def test_isnan(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.isnan(a) - return c - - test_arr = [np.log(-1.),1.,np.log(0)] - input_arr = np.asarray(test_arr, dtype=np.float32) - - c = f(input_arr) - d = np.isnan(input_arr) - self.assertTrue(np.all(c == d)) - - - def test_floor(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.floor(a) - return c - - input_arr = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) - - c = f(input_arr) - d = np.floor(input_arr) - self.assertTrue(np.all(c == d)) - - - def test_ceil(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.ceil(a) - return c - - input_arr = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) - - c = f(input_arr) - d = np.ceil(input_arr) - self.assertTrue(np.all(c == d)) - - - def test_trunc(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.trunc(a) - return c - - input_arr = np.array([-1.7, -1.5, -0.2, 0.2, 1.5, 1.7, 2.0]) - - c = f(input_arr) - d = np.trunc(input_arr) - self.assertTrue(np.all(c == d)) - - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_numpy_math_functions.py b/numba_dppy/tests/dppl/test_numpy_math_functions.py deleted file mode 100644 index ddbb568ede..0000000000 --- a/numba_dppy/tests/dppl/test_numpy_math_functions.py +++ /dev/null @@ -1,296 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -class TestNumpy_math_functions(DPPLTestCase): - N = 10 - a = np.array(np.random.random(N), dtype=np.float32) - b = np.array(np.random.random(N), dtype=np.float32) - - def test_add(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.add(a, b) - return c - - c = f(self.a, self.b) - d = self.a + self.b - self.assertTrue(np.all(c == d)) - - def test_subtract(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.subtract(a, b) - return c - - c = f(self.a, self.b) - d = self.a - self.b - self.assertTrue(np.all(c == d)) - - def test_multiply(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.multiply(a, b) - return c - - c = f(self.a, self.b) - d = self.a * self.b - self.assertTrue(np.all(c == d)) - - def test_divide(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.divide(a, b) - return c - - c = f(self.a, self.b) - d = self.a / self.b - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-2) - - def test_true_divide(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.true_divide(a, b) - return c - - c = f(self.a, self.b) - d = np.true_divide(self.a, self.b) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-2) - - def test_negative(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.negative(a) - return c - - c = f(self.a) - self.assertTrue(np.all(c == -self.a)) - - def test_power(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.power(a, b) - return c - - input_arr = np.random.randint(self.N, size=(self.N)) - exp = np.full((self.N), 2, dtype=np.int) - - c = f(input_arr, exp) - self.assertTrue(np.all(c == input_arr * input_arr)) - - def test_remainder(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.remainder(a, b) - return c - - input_arr = np.full((self.N), 3, dtype=np.int) - divisor = np.full((self.N), 2, dtype=np.int) - - c = f(input_arr, divisor) - self.assertTrue(np.all(c == 1)) - - def test_mod(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.mod(a, b) - return c - - input_arr = np.full((self.N), 3, dtype=np.int) - divisor = np.full((self.N), 2, dtype=np.int) - - c = f(input_arr, divisor) - self.assertTrue(np.all(c == 1)) - - def test_fmod(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.fmod(a, b) - return c - - input_arr = np.full((self.N), 3, dtype=np.float32) - divisor = np.full((self.N), 2, dtype=np.int) - - c = f(input_arr, divisor) - self.assertTrue(np.all(c == 1.)) - - def test_abs(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.abs(a) - return c - - input_arr = 5 * np.random.random_sample(self.N) - 5 - - c = f(input_arr) - self.assertTrue(np.all(c == -input_arr)) - - def test_absolute(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.absolute(a) - return c - - input_arr = 5 * np.random.random_sample(self.N) - 5 - - c = f(input_arr) - self.assertTrue(np.all(c == -input_arr)) - - - def test_fabs(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.fabs(a) - return c - - input_arr = 5 * np.random.random_sample(self.N) - 5 - - c = f(input_arr) - self.assertTrue(np.all(c == -input_arr)) - - - def test_sign(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.sign(a) - return c - - input_arr = 5 * np.random.random_sample(self.N) - 5 - - c = f(input_arr) - self.assertTrue(np.all(c == -1.)) - - def test_conj(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.conj(a) - return c - - input_arr = np.eye(self.N) + 1j * np.eye(self.N) - - c = f(input_arr) - d = np.conj(input_arr) - self.assertTrue(np.all(c == d)) - - def test_exp(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.exp(a) - return c - - input_arr = np.random.randint(self.N, size=(self.N)) - c = f(input_arr) - d = np.exp(input_arr) - self.assertTrue(np.all(c == d)) - - - def test_log(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.log(a) - return c - - input_arr = np.random.randint(1, self.N, size=(self.N)) - c = f(input_arr) - d = np.log(input_arr) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_log10(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.log10(a) - return c - - input_arr = np.random.randint(1, self.N, size=(self.N)) - c = f(input_arr) - d = np.log10(input_arr) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_expm1(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.expm1(a) - return c - - input_arr = np.random.randint(1, self.N, size=(self.N)) - c = f(input_arr) - d = np.expm1(input_arr) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_log1p(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.log1p(a) - return c - - input_arr = np.random.randint(1, self.N, size=(self.N)) - c = f(input_arr) - d = np.log1p(input_arr) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - def test_sqrt(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.sqrt(a) - return c - - c = f(self.a) - d = np.sqrt(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_square(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.square(a) - return c - - input_arr = np.random.randint(self.N, size=(self.N)) - - c = f(input_arr) - self.assertTrue(np.all(c == input_arr * input_arr)) - - def test_reciprocal(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.reciprocal(a) - return c - - input_arr = 5 * np.random.random_sample(self.N) + 5 - - c = f(input_arr) - self.assertTrue(np.all(c == 1/input_arr)) - - def test_conjugate(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.conjugate(a) - return c - - input_arr = np.eye(self.N) + 1j * np.eye(self.N) - - c = f(input_arr) - d = np.conj(input_arr) - self.assertTrue(np.all(c == d)) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_numpy_trigonomteric_functions.py b/numba_dppy/tests/dppl/test_numpy_trigonomteric_functions.py deleted file mode 100644 index 8f61f941c9..0000000000 --- a/numba_dppy/tests/dppl/test_numpy_trigonomteric_functions.py +++ /dev/null @@ -1,222 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -class TestNumpy_math_functions(DPPLTestCase): - N = 10 - a = np.array(np.random.random(N), dtype=np.float32) - b = np.array(np.random.random(N), dtype=np.float32) - - def test_sin(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.sin(a) - return c - - c = f(self.a) - d = np.sin(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_cos(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.cos(a) - return c - - c = f(self.a) - d = np.cos(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_tan(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.tan(a) - return c - - c = f(self.a) - d = np.tan(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arcsin(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.arcsin(a) - return c - - c = f(self.a) - d = np.arcsin(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arccos(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.arccos(a) - return c - - c = f(self.a) - d = np.arccos(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arctan(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.arctan(a) - return c - - c = f(self.a) - d = np.arctan(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arctan2(self): - @njit(parallel={'offload':True}) - def f(a, b): - c = np.arctan2(a, b) - return c - - c = f(self.a, self.b) - d = np.arctan2(self.a, self.b) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_sinh(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.sinh(a) - return c - - c = f(self.a) - d = np.sinh(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_cosh(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.cosh(a) - return c - - c = f(self.a) - d = np.cosh(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_tanh(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.tanh(a) - return c - - c = f(self.a) - d = np.tanh(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arcsinh(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.arcsinh(a) - return c - - c = f(self.a) - d = np.arcsinh(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arccosh(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.arccosh(a) - return c - - input_arr = np.random.randint(1, self.N, size=(self.N)) - c = f(input_arr) - d = np.arccosh(input_arr) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_arctanh(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.arctanh(a) - return c - - c = f(self.a) - d = np.arctanh(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_deg2rad(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.deg2rad(a) - return c - - c = f(self.a) - d = np.deg2rad(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - - def test_rad2deg(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.rad2deg(a) - return c - - c = f(self.a) - d = np.rad2deg(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-2) - - def test_degrees(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.degrees(a) - return c - - c = f(self.a) - d = np.degrees(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-2) - - def test_radians(self): - @njit(parallel={'offload':True}) - def f(a): - c = np.radians(a) - return c - - c = f(self.a) - d = np.radians(self.a) - max_abs_err = c.sum() - d.sum() - self.assertTrue(max_abs_err < 1e-5) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_parfor_lower_message.py b/numba_dppy/tests/dppl/test_parfor_lower_message.py deleted file mode 100644 index fe8c85d356..0000000000 --- a/numba_dppy/tests/dppl/test_parfor_lower_message.py +++ /dev/null @@ -1,36 +0,0 @@ -import numpy as np -import numba -from numba import njit, prange -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest, DPPLTestCase -from numba.tests.support import captured_stdout -import dpctl - - -def prange_example(): - n = 10 - a = np.ones((n), dtype=np.float64) - b = np.ones((n), dtype=np.float64) - c = np.ones((n), dtype=np.float64) - for i in prange(n//2): - a[i] = b[i] + c[i] - - return a - - -@unittest.skipUnless(dpctl.has_gpu_queues(), "test only on GPU system") -class TestParforMessage(DPPLTestCase): - def test_parfor_message(self): - with dpctl.device_context("opencl:gpu") as gpu_queue: - numba_dppy.compiler.DEBUG = 1 - jitted = njit(parallel={"offload": True})(prange_example) - - with captured_stdout() as got: - jitted() - - numba_dppy.compiler.DEBUG = 0 - self.assertTrue("Parfor lowered on DPPL-device" in got.getvalue()) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_prange.py b/numba_dppy/tests/dppl/test_prange.py deleted file mode 100644 index 317c2cbb2f..0000000000 --- a/numba_dppy/tests/dppl/test_prange.py +++ /dev/null @@ -1,155 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -import numba -from numba import njit, prange -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest, expectedFailureIf -from numba_dppy.testing import DPPLTestCase -from numba.tests.support import captured_stdout - - -class TestPrange(DPPLTestCase): - def test_one_prange(self): - @njit(parallel={'offload':True}) - def f(a, b): - for i in prange(4): - b[i, 0] = a[i, 0] * 10 - - m = 8 - n = 8 - a = np.ones((m, n)) - b = np.ones((m, n)) - - f(a, b) - - for i in range(4): - self.assertTrue(b[i, 0] == a[i, 0] * 10) - - - def test_nested_prange(self): - @njit(parallel={'offload':True}) - def f(a, b): - # dimensions must be provided as scalar - m, n = a.shape - for i in prange(m): - for j in prange(n): - b[i, j] = a[i, j] * 10 - - m = 8 - n = 8 - a = np.ones((m, n)) - b = np.ones((m, n)) - - f(a, b) - self.assertTrue(np.all(b == 10)) - - - def test_multiple_prange(self): - @njit(parallel={'offload':True}) - def f(a, b): - # dimensions must be provided as scalar - m, n = a.shape - for i in prange(m): - val = 10 - for j in prange(n): - b[i, j] = a[i, j] * val - - - for i in prange(m): - for j in prange(n): - a[i, j] = a[i, j] * 10 - - m = 8 - n = 8 - a = np.ones((m, n)) - b = np.ones((m, n)) - - f(a, b) - self.assertTrue(np.all(b == 10)) - self.assertTrue(np.all(a == 10)) - - - def test_three_prange(self): - @njit(parallel={'offload':True}) - def f(a, b): - # dimensions must be provided as scalar - m, n, o = a.shape - for i in prange(m): - val = 10 - for j in prange(n): - constant = 2 - for k in prange(o): - b[i, j, k] = a[i, j, k] * (val + constant) - - m = 8 - n = 8 - o = 8 - a = np.ones((m, n, o)) - b = np.ones((m, n, o)) - - f(a, b) - self.assertTrue(np.all(b == 12)) - - - @expectedFailureIf(sys.platform.startswith('win')) - def test_two_consequent_prange(self): - def prange_example(): - n = 10 - a = np.ones((n), dtype=np.float64) - b = np.ones((n), dtype=np.float64) - c = np.ones((n), dtype=np.float64) - for i in prange(n//2): - a[i] = b[i] + c[i] - - return a - - old_debug = numba_dppy.compiler.DEBUG - numba_dppy.compiler.DEBUG = 1 - - jitted = njit(parallel={'offload':True})(prange_example) - with captured_stdout() as stdout: - jitted_res = jitted() - - res = prange_example() - - numba_dppy.compiler.DEBUG = old_debug - - self.assertEqual(stdout.getvalue().count('Parfor lowered on DPPL-device'), 2, stdout.getvalue()) - self.assertEqual(stdout.getvalue().count('Failed to lower parfor on DPPL-device'), 0, stdout.getvalue()) - np.testing.assert_equal(res, jitted_res) - - - @unittest.skip('NRT required but not enabled') - def test_2d_arrays(self): - def prange_example(): - n = 10 - a = np.ones((n, n), dtype=np.float64) - b = np.ones((n, n), dtype=np.float64) - c = np.ones((n, n), dtype=np.float64) - for i in prange(n//2): - a[i] = b[i] + c[i] - - return a - - old_debug = numba_dppy.compiler.DEBUG - numba_dppy.compiler.DEBUG = 1 - - jitted = njit(parallel={'offload':True})(prange_example) - with captured_stdout() as stdout: - jitted_res = jitted() - - res = prange_example() - - numba_dppy.compiler.DEBUG = old_debug - - self.assertEqual(stdout.getvalue().count('Parfor lowered on DPPL-device'), 2, stdout.getvalue()) - self.assertEqual(stdout.getvalue().count('Failed to lower parfor on DPPL-device'), 0, stdout.getvalue()) - np.testing.assert_equal(res, jitted_res) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_print.py b/numba_dppy/tests/dppl/test_print.py deleted file mode 100644 index ca1e47978a..0000000000 --- a/numba_dppy/tests/dppl/test_print.py +++ /dev/null @@ -1,39 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit, prange -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - -import dpctl - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestPrint(DPPLTestCase): - def test_print_dppl_kernel(self): - @dppl.func - def g(a): - print("value of a:", a) - return a + 1 - - @dppl.kernel - def f(a, b): - i = dppl.get_global_id(0) - b[i] = g(a[i]) - print("value of b at:", i, "is", b[i]) - - N = 10 - - a = np.ones(N) - b = np.ones(N) - - with dpctl.device_context("opencl:gpu") as gpu_queue: - f[N, dppl.DEFAULT_LOCAL_SIZE](a, b) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_sum_reduction.py b/numba_dppy/tests/dppl/test_sum_reduction.py deleted file mode 100644 index 3095497a66..0000000000 --- a/numba_dppy/tests/dppl/test_sum_reduction.py +++ /dev/null @@ -1,48 +0,0 @@ -from __future__ import print_function, division, absolute_import - -import numpy as np -import math -import time - -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase -import dpctl - -@dppl.kernel -def reduction_kernel(A, R, stride): - i = dppl.get_global_id(0) - # sum two element - R[i] = A[i] + A[i+stride] - # store the sum to be used in nex iteration - A[i] = R[i] - - -@unittest.skipUnless(dpctl.has_gpu_queues(), 'test only on GPU system') -class TestDPPLSumReduction(DPPLTestCase): - def test_sum_reduction(self): - # This test will only work for even case - N = 1024 - self.assertTrue(N%2 == 0) - - A = np.array(np.random.random(N), dtype=np.float32) - A_copy = A.copy() - # 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: - total = N - - while (total > 1): - # call kernel - global_size = total // 2 - reduction_kernel[global_size, dppl.DEFAULT_LOCAL_SIZE](A, R, global_size) - total = total // 2 - - result = A_copy.sum() - max_abs_err = result - R[0] - self.assertTrue(max_abs_err < 1e-4) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_vectorize.py b/numba_dppy/tests/dppl/test_vectorize.py deleted file mode 100644 index 12dc7b5ed3..0000000000 --- a/numba_dppy/tests/dppl/test_vectorize.py +++ /dev/null @@ -1,43 +0,0 @@ -#! /usr/bin/env python -from __future__ import print_function -from timeit import default_timer as time - -import sys -import numpy as np -from numba import njit, vectorize -import numba_dppy, numba_dppy as dppl -from numba_dppy.testing import unittest -from numba_dppy.testing import DPPLTestCase - - -class TestVectorize(DPPLTestCase): - def test_vectorize(self): - - @vectorize(nopython=True) - def axy(a, x, y): - return a * x + y - - @njit(parallel={'offload':True}) - def f(a0, a1): - return np.cos(axy(a0, np.sin(a1) - 1., 1.) ) - - def f_np(a0, a1): - sin_res = np.sin(a1) - res = [] - for i in range(len(a0)): - res.append(axy(a0[i], sin_res[i] - 1., 1.)) - return np.cos(np.array(res)) - - - A = np.random.random(10) - B = np.random.random(10) - - expected = f(A, B) - actual = f_np(A, B) - - max_abs_err = expected.sum() - actual.sum() - self.assertTrue(max_abs_err < 1e-5) - - -if __name__ == '__main__': - unittest.main() diff --git a/numba_dppy/tests/dppl/test_with_context.py b/numba_dppy/tests/dppl/test_with_context.py deleted file mode 100644 index 0749ff3e89..0000000000 --- a/numba_dppy/tests/dppl/test_with_context.py +++ /dev/null @@ -1,120 +0,0 @@ -import sys -import numba -import numpy as np -from numba import njit -import numba_dppy, numba_dppy as dppl -from numba.core import errors -from numba.tests.support import captured_stdout -from numba_dppy.testing import DPPLTestCase, unittest, expectedFailureIf -import dpctl - - -class TestWithDPPLContext(DPPLTestCase): - - @unittest.skipIf(not dpctl.has_gpu_queues(), "No GPU platforms available") - @expectedFailureIf(sys.platform.startswith('win')) - def test_with_dppl_context_gpu(self): - - @njit - def nested_func(a, b): - np.sin(a, b) - - @njit - def func(b): - a = np.ones((64), dtype=np.float64) - nested_func(a, b) - - numba_dppy.compiler.DEBUG = 1 - expected = np.ones((64), dtype=np.float64) - got_gpu = np.ones((64), dtype=np.float64) - - with captured_stdout() as got_gpu_message: - with dpctl.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 DPPL-device' in got_gpu_message.getvalue()) - - @unittest.skipIf(not dpctl.has_cpu_queues(), "No CPU platforms available") - @unittest.expectedFailure - def test_with_dppl_context_cpu(self): - - @njit - def nested_func(a, b): - np.sin(a, b) - - @njit - def func(b): - a = np.ones((64), dtype=np.float64) - nested_func(a, b) - - numba_dppy.compiler.DEBUG = 1 - expected = np.ones((64), dtype=np.float64) - got_cpu = np.ones((64), dtype=np.float64) - - with captured_stdout() as got_cpu_message: - with dpctl.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 DPPL-device' in got_cpu_message.getvalue()) - - - @unittest.skipIf(not dpctl.has_gpu_queues(), "No GPU platforms available") - def test_with_dppl_context_target(self): - - @njit(target='cpu') - def nested_func_target(a, b): - np.sin(a, b) - - @njit(target='gpu') - def func_target(b): - a = np.ones((64), dtype=np.float64) - nested_func_target(a, b) - - @njit - def func_no_target(b): - a = np.ones((64), dtype=np.float64) - nested_func_target(a, b) - - @njit(parallel=False) - def func_no_parallel(b): - a = np.ones((64), dtype=np.float64) - return a - - - a = np.ones((64), dtype=np.float64) - b = np.ones((64), dtype=np.float64) - - with self.assertRaises(errors.UnsupportedError) as raises_1: - with dpctl.device_context("opencl:gpu"): - nested_func_target(a, b) - - with self.assertRaises(errors.UnsupportedError) as raises_2: - with dpctl.device_context("opencl:gpu"): - func_target(a) - - with self.assertRaises(errors.UnsupportedError) as raises_3: - with dpctl.device_context("opencl:gpu"): - func_no_target(a) - - with self.assertRaises(errors.UnsupportedError) as raises_4: - with dpctl.device_context("opencl:gpu"): - func_no_parallel(a) - - msg_1 = "Can't use 'with' context with explicitly specified target" - msg_2 = "Can't use 'with' context with parallel option" - self.assertTrue(msg_1 in str(raises_1.exception)) - self.assertTrue(msg_1 in str(raises_2.exception)) - self.assertTrue(msg_1 in str(raises_3.exception)) - self.assertTrue(msg_2 in str(raises_4.exception)) - - -if __name__ == '__main__': - unittest.main()