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.

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

NdRange does not produce correct results depending on global size #1308

Closed
ZzEeKkAa opened this issue Feb 2, 2024 · 3 comments · Fixed by #1357
Closed

NdRange does not produce correct results depending on global size #1308

ZzEeKkAa opened this issue Feb 2, 2024 · 3 comments · Fixed by #1357

Comments

@ZzEeKkAa
Copy link
Contributor

ZzEeKkAa commented Feb 2, 2024

This test fails on windows.

import dpnp
import numpy as np

import numba_dpex as dpex
import numba_dpex.experimental as dpex_exp
from numba_dpex.kernel_api import NdItem

_SIZE = 10

@dpex_exp.kernel
def set_ones_nd_item(nd_item: NdItem, a):
    i = nd_item.get_global_id(0)
    a[i] = 1

a = dpnp.zeros(_SIZE, dtype=dpnp.float32)
dpex_exp.call_kernel(
    set_ones_nd_item, dpex.NdRange((a.size,), (a.size,)), a
)

assert np.array_equal(a.asnumpy(), np.ones(a.size, dtype=np.float32))
@ZzEeKkAa
Copy link
Contributor Author

ZzEeKkAa commented Feb 3, 2024

UPD: it fails on linux as well, if _SIZE > 8192

@diptorupd diptorupd changed the title NdRange does not work on Windows NdRange does not produce correct results depending on global size Feb 3, 2024
@diptorupd
Copy link
Contributor

The problem is in the test code. The NdRange created by you dpex.NdRange((a.size,), (a.size,)) means that you want a.size work items in total (global size) and also want a.size work items per work group. That means your kernel is launched on exactly one work group.

The linux error is easy to explain, if your a.size is greater than 8192 it exceeds the max_work_item_sizes1d for that work group after which the behavior is undefined.

>>> import dpctl
>>> cd = dpctl.SyclDevice("cpu")
>>> cd.max_work_item_sizes1d
(8192,)

My hunch will be that the Windows error also comes up for the same reason. I will test and confirm.

@ZzEeKkAa
Copy link
Contributor Author

ZzEeKkAa commented Feb 9, 2024

After further investigation - it seems to be an issue with github CI. Local tests shows that it behaves the same way on linux and windows.

UPD: I mean issue is on our side, but revealed on github CI hardware.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging a pull request may close this issue.

2 participants