From 492f4adfd15277b8e92e00e33a1bcf893eaf82f5 Mon Sep 17 00:00:00 2001 From: Yevhenii Havrylko Date: Fri, 23 Feb 2024 12:31:40 -0500 Subject: [PATCH] Add kernel launcher allocation alignment --- numba_dpex/core/utils/kernel_launcher.py | 6 +++++- .../tests/experimental/test_index_space_ids.py | 17 ----------------- 2 files changed, 5 insertions(+), 18 deletions(-) diff --git a/numba_dpex/core/utils/kernel_launcher.py b/numba_dpex/core/utils/kernel_launcher.py index 98814578ee..6d12e37de9 100644 --- a/numba_dpex/core/utils/kernel_launcher.py +++ b/numba_dpex/core/utils/kernel_launcher.py @@ -28,6 +28,8 @@ MAX_SIZE_OF_SYCL_RANGE = 3 +_ARRAY_ALIGN = 16 + # TODO: probably not best place for it. Should be in kernel_dispatcher once we # get merge experimental. Right now it will cause cyclic import @@ -364,11 +366,13 @@ def _allocate_array( Returns: An LLVM IR value pointing to the array. """ - return cgutils.alloca_once( + array = cgutils.alloca_once( self.builder, self.context.get_value_type(numba_type), size=self.context.get_constant(types.uintp, size), ) + array.align = _ARRAY_ALIGN + return array def _populate_array_from_python_list( self, diff --git a/numba_dpex/tests/experimental/test_index_space_ids.py b/numba_dpex/tests/experimental/test_index_space_ids.py index 54ceca8807..2d1edb54f2 100644 --- a/numba_dpex/tests/experimental/test_index_space_ids.py +++ b/numba_dpex/tests/experimental/test_index_space_ids.py @@ -13,7 +13,6 @@ 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 skip_windows _SIZE = 16 _GROUP_SIZE = 4 @@ -99,8 +98,6 @@ def test_item_get_range(): assert np.array_equal(a.asnumpy(), want) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_nd_item_get_global_range(): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) dpex_exp.call_kernel( @@ -114,8 +111,6 @@ def test_nd_item_get_global_range(): assert np.array_equal(a.asnumpy(), want) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_nd_item_get_local_range(): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) dpex_exp.call_kernel( @@ -129,8 +124,6 @@ def test_nd_item_get_local_range(): assert np.array_equal(a.asnumpy(), want) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_nd_item_get_global_id(): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) dpex_exp.call_kernel( @@ -140,8 +133,6 @@ def test_nd_item_get_global_id(): assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.float32)) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_nd_item_get_local_id(): a = dpnp.zeros(_SIZE, dtype=dpnp.float32) @@ -174,8 +165,6 @@ def test_no_item(): ) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_get_group_id(): global_size = 100 group_size = 20 @@ -196,8 +185,6 @@ def test_get_group_id(): assert np.array_equal(ka.asnumpy(), expected) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_get_group_range(): global_size = 100 group_size = 20 @@ -218,8 +205,6 @@ def test_get_group_range(): assert np.array_equal(ka.asnumpy(), expected) -# TODO: https://github.com/IntelPython/numba-dpex/issues/1308 -@skip_windows def test_get_group_local_range(): global_size = 100 group_size = 20 @@ -258,8 +243,6 @@ def set_3d_ones_item(item: Item, a): a[index] = 1 -# TODO: CI tests failing for some reason... Works fine locally on cpu and gpu -@pytest.mark.skip def test_index_order(): a = dpnp.zeros(I_SIZE * J_SIZE * K_SIZE, dtype=dpnp.int32)