Call to numba_dpex.kernel corrups data in arguments

[oleksandr-pavlyk]

Python 3.9.10 (main, Mar  2 2022, 12:02:00)
Type 'copyright', 'credits' or 'license' for more information
IPython 8.5.0 -- An enhanced Interactive Python. Type '?' for help.

In [1]: import dpctl.tensor as dpt, numba_dpex

In [2]: x = dpt.ones(10, dtype='i8')

In [3]: y = dpt.empty_like(x)

In [4]: x.sycl_device == y.sycl_device
Out[4]: True

In [5]: dpt.asnumpy(x), dpt.asnumpy(y)
Out[5]: (array([1, 1, 1, 1, 1, 1, 1, 1, 1, 1]), array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0]))

In [6]: @numba_dpex.kernel
   ...: def foo(X, Y):
   ...:     i = numba_dpex.get_global_id(0)
   ...:     Y[i] = 2*X[i]
   ...:

In [7]: foo[10, numba_dpex.DEFAULT_LOCAL_SIZE](x, y)

In [8]: dpt.asnumpy(x), dpt.asnumpy(y)
Out[8]:
(array([     94391540026928,                   0, 7954887801964155392,
             94391483338752,      94391483338736,                   0,
                          0,                   0,                   0,
                          0]),
 array([94391548829776,              0,              0,              0,
                     0,              0,              0,              0,
                     0,              0]))

My setup:

(dppy_stack) opavlyk@opavlyk-mobl:~$ conda list dp
# packages in environment at /home/opavlyk/miniconda3/envs/dppy_stack:
#
# Name                    Version                   Build  Channel
dpbench                   0.0.1                     dev_0    <develop>
dpcpp-cpp-rt              2022.2.0             intel_8734    intel
dpcpp_cpp_rt              2022.2.0             intel_8734    intel
dpctl                     0.14.0dev0      py39h8c27c75_195    dppy/label/dev
dpnp                      0.10.2           py39h2bc3f7f_7    dppy/label/dev
mkl-dpcpp                 2022.2.0             intel_8748    intel
numba-dpex                0.18.1          py39hfc4b9b4_45    dppy/label/dev
threadpoolctl             2.2.0              pyh0d69192_0    intel

[mingjie-intel]

I have tested with:

(dpex_809) mingjie2@ansatnuc04:/localdisk/work/mingjie2/sandbox/numba_dpex_809$ conda list dp
# packages in environment at /localdisk/work/mingjie2/miniconda3/envs/dpex_809:
#
# Name                    Version                   Build  Channel
dpcpp-cpp-rt              2022.2.0             intel_8734    intel
dpcpp_cpp_rt              2022.2.0             intel_8734    intel
dpctl                     0.14.0dev0      py39h8c27c75_207    dppy/label/dev
dpnp                      0.10.2           py39h2bc3f7f_8    dppy/label/dev
mkl-dpcpp                 2022.2.0             intel_8748    intel
numba-dpex                0.18.1          py39hfc4b9b4_45    dppy/label/dev

the output seems correct:

(dpex_809) mingjie2@ansatnuc04:/localdisk/work/mingjie2/sandbox/numba_dpex_809$ python test_Sasha.py
in:
[1 1 1 1 1 1 1 1 1 1]
[0 0 0 0 0 0 0 0 0 0]
out:
[1 1 1 1 1 1 1 1 1 1]
[2 2 2 2 2 2 2 2 2 2]

The only difference I can see here is dpctl: py39h8c27c75_207 vs py39h8c27c75_195. Is the possible something fixed in the newer build?

[oleksandr-pavlyk]

The thing that I failed to state is my GPU environment:

(dppy_stack2) opavlyk@opavlyk-mobl:~/tmp$ sycl-ls
[opencl:acc:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device 1.2 [2022.14.7.0.30_160000]
[opencl:cpu:1] Intel(R) OpenCL, 11th Gen Intel(R) Core(TM) i7-1185G7 @ 3.00GHz 3.0 [2022.14.7.0.30_160000]
[opencl:gpu:2] Intel(R) OpenCL HD Graphics, Intel(R) Graphics [0x9a49] 3.0 [22.39.24347]
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Graphics [0x9a49] 1.3 [1.3.24347]
[host:host:0] SYCL host platform, SYCL host device 1.2 [1.2]

I am on Ubuntu 20.04 running kernel "5.10.102.1-microsoft-standard-WSL2" in WSL.

[diptorupd]

The thing that I failed to state is my GPU environment:

(dppy_stack2) opavlyk@opavlyk-mobl:~/tmp$ sycl-ls
[opencl:acc:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device 1.2 [2022.14.7.0.30_160000]
[opencl:cpu:1] Intel(R) OpenCL, 11th Gen Intel(R) Core(TM) i7-1185G7 @ 3.00GHz 3.0 [2022.14.7.0.30_160000]
[opencl:gpu:2] Intel(R) OpenCL HD Graphics, Intel(R) Graphics [0x9a49] 3.0 [22.39.24347]
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Graphics [0x9a49] 1.3 [1.3.24347]
[host:host:0] SYCL host platform, SYCL host device 1.2 [1.2]

I am on Ubuntu 20.04 running kernel "5.10.102.1-microsoft-standard-WSL2" in WSL.

I suspect an overflow. Does Iris Xe support fp64?

[oleksandr-pavlyk]

No, Iris Xe-LP does not support fp64:

In [1]: import dpctl

In [2]: dev = dpctl.SyclDevice()

In [3]: dev, dev.has_aspect_fp64
Out[3]:
(<dpctl.SyclDevice [backend_type.level_zero, device_type.gpu,  Intel(R) Graphics [0x9a49]] at 0x7eff719a21b0>,
 False)

[diptorupd]

But, then your dtype is i8. So, the issue may not be related to fp64. Let me investigate. There are two possibilities: the unpack routine where we cast PyObjects to ctypes is messing up the precision for i8 or worse we are somehow ignoring the dtype and passing in fp64 anyway. Let me check and get back.

[ndgrigorian]

I was able to replicate this issue locally.

Using Sasha's code:
In

(array([1, 1, 1, 1, 1, 1, 1, 1, 1, 1]),
array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0]))

Out

(array([     94921035162528,                   0, 7954887801964155392,
            94921012152992,      94921012152976,                   0,
                         0,                   0,                   0,
                         0]),
array([94921064734992,              0, 94921064738480, 94921064739312,
       94921035167728, 94921035162256, 94921035162272, 94921035162272,
       94921035168864, 94921035168872]))

My GPU environment is similar:

(numba-dpex-env) ngrigori@ngrigori-mobl:~/numba-dpex$ sycl-ls
[opencl:acc:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device 1.2 [2022.14.7.0.30_160000]
[opencl:cpu:1] Intel(R) OpenCL, 11th Gen Intel(R) Core(TM) i7-1185G7 @ 3.00GHz 3.0 [2022.14.7.0.30_160000]
[opencl:gpu:2] Intel(R) OpenCL HD Graphics, Intel(R) Graphics [0x9a49] 3.0 [22.28.23726.1]
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Graphics [0x9a49] 1.3 [1.3.23726]
[host:host:0] SYCL host platform, SYCL host device 1.2 [1.2]

(numba-dpex-env) ngrigori@ngrigori-mobl:~/numba-dpex$ uname -r | tr '[:upper:]' '[:lower:]'
5.10.16.3-microsoft-standard-wsl2

with Ubuntu 20.04

[oleksandr-pavlyk]

So this must be a bug specific to WSL. Just tried this on Linux box:

Python 3.9.12 (main, Aug 29 2022, 00:54:58)
Type 'copyright', 'credits' or 'license' for more information
IPython 8.5.0 -- An enhanced Interactive Python. Type '?' for help.

In [1]: import dpctl, numba_dpex, dpctl.tensor as dpt

In [2]: x = dpt.ones(10, dtype='i8')
   ...:

In [3]: y = dpt.empty_like(x)

In [4]: x.sycl_device == y.sycl_device
Out[4]: True

In [5]: dpt.asnumpy(x), dpt.asnumpy(y)
Out[5]: (array([1, 1, 1, 1, 1, 1, 1, 1, 1, 1]), array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0]))

In [6]: @numba_dpex.kernel
   ...: def foo(X, Y):
   ...:     i = numba_dpex.get_global_id(0)
   ...:     Y[i] = 2 * X[i]
   ...:

In [7]: foo[10, numba_dpex.DEFAULT_LOCAL_SIZE](x, y)

In [8]: dpt.asnumpy(x), dpt.asnumpy(y)
Out[8]: (array([1, 1, 1, 1, 1, 1, 1, 1, 1, 1]), array([2, 2, 2, 2, 2, 2, 2, 2, 2, 2]))

In [9]: x.sycl_device
Out[9]: <dpctl.SyclDevice [backend_type.level_zero, device_type.gpu,  Intel(R) Iris(R) Xe Graphics [0x9a49]] at 0x7fc5c4cfd5f0>

In [10]: x.sycl_device.driver_version
Out[10]: '1.3.24347'

In [11]: !uname -r
5.15.0-47-generic

In [12]: !lsb_release -a
No LSB modules are available.
Distributor ID: Ubuntu
Description:    Ubuntu 22.04.1 LTS
Release:        22.04
Codename:       jammy

[oleksandr-pavlyk]

Also, this appears to work correctly in WSL with opencl:gpu backend:

Python 3.9.12 (main, Aug 29 2022, 00:54:58)
Type 'copyright', 'credits' or 'license' for more information
IPython 8.5.0 -- An enhanced Interactive Python. Type '?' for help.

In [1]: import dpctl.tensor as dpt, numba_dpex

In [2]: x = dpt.ones(10, dtype='i8', device="opencl:gpu")

In [3]: y = dpt.empty_like(x)

In [4]: x.sycl_device == y.sycl_device
Out[4]: True

In [5]: dpt.asnumpy(x), dpt.asnumpy(y)
Out[5]: (array([1, 1, 1, 1, 1, 1, 1, 1, 1, 1]), array([0, 0, 0, 0, 0, 0, 0, 0, 0, 0]))

In [6]: @numba_dpex.kernel
   ...: def foo(X, Y):
   ...:     i = numba_dpex.get_global_id(0)
   ...:     Y[i] = 2*X[i]
   ...:

In [7]: foo[10, numba_dpex.DEFAULT_LOCAL_SIZE](x, y)

In [8]: dpt.asnumpy(x), dpt.asnumpy(y)
Out[8]: (array([1, 1, 1, 1, 1, 1, 1, 1, 1, 1]), array([2, 2, 2, 2, 2, 2, 2, 2, 2, 2]))

[oleksandr-pavlyk]

@diptorupd I think it is clear that this is not an issue with numba_dpex. The issue can be closed.

[oleksandr-pavlyk]

Reopening, I have built a project in C++ which submits the interoperability kernel extracted from this example using both level-zero and opencl backends.

The example works fine in WSL with both backends. Hence, the observed corruption is numba-dpex specific.

The repo can be found at https://github.com/oleksandr-pavlyk/native-reproducer-simulating-numba-dpex-issue-809/