Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add kernel launcher array allocation alignment #1357

Merged
merged 1 commit into from
Feb 23, 2024
Merged

Add kernel launcher array allocation alignment #1357

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 5 additions & 1 deletion numba_dpex/core/utils/kernel_launcher.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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,
Expand Down
17 changes: 0 additions & 17 deletions numba_dpex/tests/experimental/test_index_space_ids.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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(
Expand All @@ -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(
Expand All @@ -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(
Expand All @@ -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)

Expand Down Expand Up @@ -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
Expand All @@ -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
Expand All @@ -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
Expand Down Expand Up @@ -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)

Expand Down
Loading