From 632f4b886c284ae3209522692432e2e06deb96a1 Mon Sep 17 00:00:00 2001 From: Yevhenii Havrylko Date: Wed, 6 Mar 2024 19:18:19 -0500 Subject: [PATCH] Extend indexing function tests to kernel simulator --- numba_dpex/tests/_helper.py | 7 + .../experimental/test_index_space_ids.py | 156 +++++++++--------- 2 files changed, 84 insertions(+), 79 deletions(-) diff --git a/numba_dpex/tests/_helper.py b/numba_dpex/tests/_helper.py index 3eca02ae06..df84553d39 100644 --- a/numba_dpex/tests/_helper.py +++ b/numba_dpex/tests/_helper.py @@ -40,6 +40,13 @@ def has_opencl_cpu(): return bool(dpctl.get_num_devices(backend="opencl", device_type="cpu")) +def has_cpu(): + """ + Checks if dpctl is able to select any CPU device. + """ + return bool(dpctl.get_num_devices(device_type="cpu")) + + def has_level_zero(): """ Checks if dpctl is able to select a Level Zero GPU device. diff --git a/numba_dpex/tests/experimental/test_index_space_ids.py b/numba_dpex/tests/experimental/test_index_space_ids.py index c8c4550839..eae0133406 100644 --- a/numba_dpex/tests/experimental/test_index_space_ids.py +++ b/numba_dpex/tests/experimental/test_index_space_ids.py @@ -7,48 +7,66 @@ import dpnp import numpy as np import pytest -from numba.core.errors import TypingError import numba_dpex as dpex import numba_dpex.experimental as dpex_exp from numba_dpex.kernel_api import Item, NdItem, NdRange from numba_dpex.kernel_api import call_kernel as kapi_call_kernel +from numba_dpex.tests._helper import has_cpu -_SIZE = 16 -_GROUP_SIZE = 4 +_SIZE = 100 +_GROUP_SIZE = 20 + + +_kernel_decorators = [(dpex_exp.call_kernel, dpex_exp.kernel)] +# run simulator tests only with arrays allocated on cpu to avoid performance +# issues +if has_cpu(): + _kernel_decorators.append((kapi_call_kernel, lambda a: a)) + + +@pytest.fixture(params=_kernel_decorators) +def call_kernel_decorator(request): + return request.param + + +@pytest.fixture +def call_kernel(call_kernel_decorator): + _call_kernel, _ = call_kernel_decorator + return _call_kernel + + +@pytest.fixture +def decorator(call_kernel_decorator): + _, _decorator = call_kernel_decorator + return _decorator -@dpex_exp.kernel def set_ones_no_item(a): a[0] = 1 -@dpex_exp.kernel def set_ones_item(item: Item, a): i = item.get_id(0) a[i] = 1 -@dpex_exp.kernel def set_last_one_item(item: Item, a): i = item.get_range(0) - 1 a[i] = 1 -@dpex_exp.kernel def set_last_one_linear_item(item: Item, a): i = item.get_linear_range() - 1 a[i] = 1 -@dpex_exp.kernel def set_last_one_linear_nd_item(nd_item: NdItem, a): i = nd_item.get_global_linear_range() - 1 a[0] = i a[i] = 1 -@dpex_exp.kernel def set_last_one_nd_item(item: NdItem, a): if item.get_global_id(0) == 0: i = item.get_global_range(0) - 1 @@ -56,21 +74,18 @@ def set_last_one_nd_item(item: NdItem, a): a[i] = 1 -@dpex_exp.kernel def set_last_group_one_linear_nd_item(nd_item: NdItem, a): i = nd_item.get_local_linear_range() - 1 a[0] = i a[i] = 1 -@dpex_exp.kernel def set_last_group_one_group_linear_nd_item(nd_item: NdItem, a): i = nd_item.get_group().get_local_linear_range() - 1 a[0] = i a[i] = 1 -@dpex_exp.kernel def set_last_group_one_nd_item(item: NdItem, a): if item.get_global_id(0) == 0: i = item.get_local_range(0) - 1 @@ -78,31 +93,26 @@ def set_last_group_one_nd_item(item: NdItem, a): a[i] = 1 -@dpex_exp.kernel def set_ones_nd_item(nd_item: NdItem, a): i = nd_item.get_global_id(0) a[i] = 1 -@dpex_exp.kernel def set_local_ones_nd_item(nd_item: NdItem, a): i = nd_item.get_local_id(0) a[i] = 1 -@dpex_exp.kernel def set_dimensions_item(item: Item, a): i = item.get_id(0) a[i] = item.dimensions -@dpex_exp.kernel def set_dimensions_nd_item(nd_item: NdItem, a): i = nd_item.get_global_id(0) a[i] = nd_item.dimensions -@dpex_exp.kernel def set_dimensions_group(nd_item: NdItem, a): i = nd_item.get_global_id(0) a[i] = nd_item.get_group().dimensions @@ -138,16 +148,16 @@ def _get_group_local_range_driver(nditem: NdItem, a): a[i] = g.get_local_range(0) -def test_item_get_id(): +def test_item_get_id(call_kernel, decorator): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel(set_ones_item, dpex.Range(a.size), a) + call_kernel(decorator(set_ones_item), dpex.Range(a.size), a) assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.float32)) -def test_item_get_range(): +def test_item_get_range(call_kernel, decorator): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel(set_last_one_item, dpex.Range(a.size), a) + call_kernel(decorator(set_last_one_item), dpex.Range(a.size), a) want = np.zeros(a.size, dtype=np.float32) want[-1] = 1 @@ -159,9 +169,9 @@ def test_item_get_range(): "rng", [dpex.Range(_SIZE), dpex.Range(1, _GROUP_SIZE, int(_SIZE / _GROUP_SIZE))], ) -def test_item_get_linear_range(rng): +def test_item_get_linear_range(call_kernel, decorator, rng): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel(set_last_one_linear_item, rng, a) + call_kernel(decorator(set_last_one_linear_item), rng, a) want = np.zeros(a.size, dtype=np.float32) want[-1] = 1 @@ -180,9 +190,9 @@ def test_item_get_linear_range(rng): ), ], ) -def test_nd_item_get_global_range(kernel, rng): +def test_nd_item_get_global_range(call_kernel, decorator, kernel, rng): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel(kernel, rng, a) + call_kernel(decorator(kernel), rng, a) want = np.zeros(a.size, dtype=np.float32) want[-1] = 1 @@ -213,9 +223,9 @@ def test_nd_item_get_global_range(kernel, rng): ), ], ) -def test_nd_item_get_local_range(kernel, rng): +def test_nd_item_get_local_range(call_kernel, decorator, kernel, rng): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel(kernel, rng, a) + call_kernel(decorator(kernel), rng, a) want = np.zeros(a.size, dtype=np.float32) want[_GROUP_SIZE - 1] = 1 @@ -224,20 +234,22 @@ def test_nd_item_get_local_range(kernel, rng): assert np.array_equal(a.asnumpy(), want) -def test_nd_item_get_global_id(): +def test_nd_item_get_global_id(call_kernel, decorator): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel( - set_ones_nd_item, dpex.NdRange((a.size,), (_GROUP_SIZE,)), a + call_kernel( + decorator(set_ones_nd_item), dpex.NdRange((a.size,), (_GROUP_SIZE,)), a ) assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.float32)) -def test_nd_item_get_local_id(): +def test_nd_item_get_local_id(call_kernel, decorator): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel( - set_local_ones_nd_item, dpex.NdRange((a.size,), (_GROUP_SIZE,)), a + call_kernel( + decorator(set_local_ones_nd_item), + dpex.NdRange((a.size,), (_GROUP_SIZE,)), + a, ) assert np.array_equal( @@ -250,11 +262,11 @@ def test_nd_item_get_local_id(): @pytest.mark.parametrize("dims", [1, 2, 3]) -def test_item_dimensions(dims): +def test_item_dimensions(call_kernel, decorator, dims): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) rng = [1] * dims rng[0] = a.size - dpex_exp.call_kernel(set_dimensions_item, dpex.Range(*rng), a) + call_kernel(decorator(set_dimensions_item), dpex.Range(*rng), a) assert np.array_equal(a.asnumpy(), dims * np.ones(a.size, dtype=np.float32)) @@ -263,25 +275,28 @@ def test_item_dimensions(dims): @pytest.mark.parametrize( "kernel", [set_dimensions_nd_item, set_dimensions_group] ) -def test_nd_item_dimensions(dims, kernel): +def test_nd_item_dimensions(call_kernel, decorator, dims, kernel): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) rng, grp = [1] * dims, [1] * dims rng[0], grp[0] = a.size, _GROUP_SIZE - dpex_exp.call_kernel(kernel, dpex.NdRange(rng, grp), a) + call_kernel(decorator(kernel), dpex.NdRange(rng, grp), a) assert np.array_equal(a.asnumpy(), dims * np.ones(a.size, dtype=np.float32)) -def test_error_item_get_global_id(): +def test_error_item_get_global_id(call_kernel, decorator): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - with pytest.raises(TypingError): - dpex_exp.call_kernel(set_ones_nd_item, dpex.Range(a.size), a) + with pytest.raises(Exception): + call_kernel(decorator(set_ones_nd_item), dpex.Range(a.size), a) + +def test_no_item(call_kernel, decorator): + if call_kernel == kapi_call_kernel: + pytest.skip() -def test_no_item(): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) - dpex_exp.call_kernel(set_ones_no_item, dpex.Range(a.size), a) + call_kernel(decorator(set_ones_no_item), dpex.Range(a.size), a) assert np.array_equal( a.asnumpy(), np.array([1] + [0] * (a.size - 1), dtype=np.float32) @@ -299,21 +314,18 @@ def test_no_item(): ), ], ) -def test_get_group_id(driver, rng): +def test_get_group_id(call_kernel, decorator, driver, rng): num_groups = _SIZE // _GROUP_SIZE a = dpnp.empty(_SIZE, dtype=dpnp.int32) - ka = dpnp.empty(_SIZE, dtype=dpnp.int32) expected = np.empty(_SIZE, dtype=np.int32) - dpex_exp.call_kernel(dpex_exp.kernel(driver), rng, a) - kapi_call_kernel(driver, rng, ka) + call_kernel(decorator(driver), rng, a) for gid in range(num_groups): for lid in range(_GROUP_SIZE): expected[gid * _GROUP_SIZE + lid] = gid assert np.array_equal(a.asnumpy(), expected) - assert np.array_equal(ka.asnumpy(), expected) @pytest.mark.parametrize( @@ -330,49 +342,38 @@ def test_get_group_id(driver, rng): ), ], ) -def test_get_group_range(driver, rng): +def test_get_group_range(call_kernel, decorator, driver, rng): num_groups = _SIZE // _GROUP_SIZE a = dpnp.empty(_SIZE, dtype=dpnp.int32) - ka = dpnp.empty(_SIZE, dtype=dpnp.int32) expected = np.empty(_SIZE, dtype=np.int32) - dpex_exp.call_kernel(dpex_exp.kernel(driver), rng, a) - kapi_call_kernel(driver, rng, ka) + call_kernel(decorator(driver), rng, a) for gid in range(num_groups): for lid in range(_GROUP_SIZE): expected[gid * _GROUP_SIZE + lid] = num_groups assert np.array_equal(a.asnumpy(), expected) - assert np.array_equal(ka.asnumpy(), expected) -def test_get_group_local_range(): - global_size = 100 - group_size = 20 - num_groups = global_size // group_size +def test_get_group_local_range(call_kernel, decorator): + num_groups = _SIZE // _GROUP_SIZE - a = dpnp.empty(global_size, dtype=dpnp.int32) - ka = dpnp.empty(global_size, dtype=dpnp.int32) - expected = np.empty(global_size, dtype=np.int32) - ndrange = NdRange((global_size,), (group_size,)) - dpex_exp.call_kernel( - dpex_exp.kernel(_get_group_local_range_driver), ndrange, a - ) - kapi_call_kernel(_get_group_local_range_driver, ndrange, ka) + a = dpnp.empty(_SIZE, dtype=dpnp.int32) + expected = np.empty(_SIZE, dtype=np.int32) + ndrange = NdRange((_SIZE,), (_GROUP_SIZE,)) + call_kernel(decorator(_get_group_local_range_driver), ndrange, a) for gid in range(num_groups): - for lid in range(group_size): - expected[gid * group_size + lid] = group_size + for lid in range(_GROUP_SIZE): + expected[gid * _GROUP_SIZE + lid] = _GROUP_SIZE assert np.array_equal(a.asnumpy(), expected) - assert np.array_equal(ka.asnumpy(), expected) I_SIZE, J_SIZE, K_SIZE = 2, 3, 4 -@dpex_exp.kernel def set_3d_ones_item(item: Item, a): i = item.get_id(0) j = item.get_id(1) @@ -385,7 +386,6 @@ def set_3d_ones_item(item: Item, a): a[index] = 1 -@dpex_exp.kernel def set_3d_ones_item_linear(item: Item, a): # Since we have different sizes for each dimension, wrong order will result # that some indexes will be set twice and some won't be set. @@ -394,7 +394,6 @@ def set_3d_ones_item_linear(item: Item, a): a[index] = 1 -@dpex_exp.kernel def set_3d_ones_nd_item_linear(nd_item: NdItem, a): # Since we have different sizes for each dimension, wrong order will result # that some indexes will be set twice and some won't be set. @@ -403,7 +402,6 @@ def set_3d_ones_nd_item_linear(nd_item: NdItem, a): a[index] = 1 -@dpex_exp.kernel def set_local_3d_ones_nd_item_linear(nd_item: NdItem, a): # Since we have different sizes for each dimension, wrong order will result # that some indexes will be set twice and some won't be set. @@ -413,19 +411,19 @@ def set_local_3d_ones_nd_item_linear(nd_item: NdItem, a): @pytest.mark.parametrize("kernel", [set_3d_ones_item, set_3d_ones_item_linear]) -def test_item_index_order(kernel): +def test_item_index_order(call_kernel, decorator, kernel): a = dpnp.zeros(I_SIZE * J_SIZE * K_SIZE, dtype=dpnp.int32) - dpex_exp.call_kernel(kernel, dpex.Range(I_SIZE, J_SIZE, K_SIZE), a) + call_kernel(decorator(kernel), dpex.Range(I_SIZE, J_SIZE, K_SIZE), a) assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.int32)) -def test_nd_item_index_order(): +def test_nd_item_index_order(call_kernel, decorator): a = dpnp.zeros(I_SIZE * J_SIZE * K_SIZE, dtype=dpnp.int32) - dpex_exp.call_kernel( - set_3d_ones_nd_item_linear, + call_kernel( + decorator(set_3d_ones_nd_item_linear), dpex.NdRange((I_SIZE, J_SIZE, K_SIZE), (1, 1, K_SIZE)), a, ) @@ -433,11 +431,11 @@ def test_nd_item_index_order(): assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.int32)) -def test_nd_item_local_linear_id(): +def test_nd_item_local_linear_id(call_kernel, decorator): a = dpnp.zeros(I_SIZE * J_SIZE * K_SIZE, dtype=dpnp.int32) - dpex_exp.call_kernel( - set_local_3d_ones_nd_item_linear, + call_kernel( + decorator(set_local_3d_ones_nd_item_linear), dpex.NdRange((I_SIZE, J_SIZE, K_SIZE), (1, 1, K_SIZE)), a, )