Add kernel launcher allocation alignment

Author: ZzEeKkAa

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,

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)