numba_dpex's kernel disregards strides of input array

[oleksandr-pavlyk]

import dpctl
import numba_dpex as dppy
import numpy as np
import dpnp

import dpctl.tensor as dpt


@dppy.kernel
def data_parallel_sum(a, b, c):
    """
    Vector addition using the ``kernel`` decorator.
    """
    i = dppy.get_global_id(0)
    c[i] = a[i] + b[i]


def driver(a, b, c, global_size):
    print("A : ", (a.shape, a.strides))
    print("B : ", (b.shape, b.strides))
    data_parallel_sum[dppy.Range(global_size)](a, b, c)
    print("C ", (c.shape, c.strides))
    # assert np.allclose(dpt.asnumpy(c), dpt.asnumpy(a) + dpt.asnumpy(b))


global_size = 10
N = global_size
print("N", N)

a = dpnp.random.random(N)  # C-contig
b = dpnp.random.random(N)  # C-contig
out = dpnp.zeros((2 * N,), dtype=a.dtype)  # C-contig
c = out[::2]  # Strided, every other element

driver(a, b, c, global_size)

print(c)
print(out)

Executing this script outputs

N 10
Using device ...
    Name            Intel(R) UHD Graphics [0x9bca]
    Driver version  1.1.20678
    Vendor          Intel(R) Corporation
    Profile         FULL_PROFILE
    Filter string   level_zero:gpu:0
A :  ((10,), (1,))
B :  ((10,), (1,))
C  ((10,), (2,))
Done...
[0.83166494 1.27176889 0.61416447 0.57984313 1.04820722 0.
 0.         0.         0.         0.        ]
[0.83166494 0.50883122 1.27176889 1.08111028 0.61416447 0.77052362
 0.57984313 1.55592601 1.04820722 1.12209132 0.         0.
 0.         0.         0.         0.         0.         0.
 0.         0.        ]

The expected result should have been

[0.83166494 0.50883122 1.27176889 1.08111028 0.61416447 0.77052362
 0.57984313 1.55592601 1.04820722 1.12209132]
[0.83166494 0.       0.50883122  0.       1.27176889  0.        1.08111028  0.        
0.61416447 0.        0.77052362  0.       0.57984313  0.        1.55592601  0.    
1.04820722  0.       1.12209132    0.  ]

UPDATE : Updated the reproducer to work with latest numba-dpex.

[PokhodenkoSA]

This is a limitation of how numba-dppy supports dpnp features. I will qualify it as an enhancement for future work.

[PokhodenkoSA]

I think for preventing this situation numba-dppy should raise exception if receives array w/ strides.

[chudur-budur]

A better minimal reproducer:

@dpex.kernel
def change_values(x):
    i = dpex.get_global_id(0)
    x[i] = -3

N = 10

out = dpnp.arange(0, 20, dtype=dpnp.int64)
b = out[::2]

print(f"out = {out}, out.shape = {out.shape}, out.strides = {out.strides}")
print(f"b = {b}, b.shape = {b.shape}, b.strides = {b.strides}")

change_values[dpex.Range(N)](b)

print(f"out = {out}, out.shape = {out.shape}, out.strides = {out.strides}")
print(f"b = {b}, b.shape = {b.shape}, b.strides = {b.strides}")

output:

out = [ 0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 16 17 18 19], out.shape = (20,), out.strides = (1,)
b = [ 0  2  4  6  8 10 12 14 16 18], b.shape = (10,), b.strides = (2,)
out = [          -3           -3 -12884901886  -8589934593   4294967295
           -3           -3 -12884901881  -8589934593   4294967295
           -3           -3           12           13           14
           15           16           17           18           19], out.shape = (20,), out.strides = (1,)
b = [          -3 -12884901886   4294967295           -3  -8589934593
           -3           12           14           16           18], b.shape = (10,), b.strides = (2,)

Expected output:

out = [ 0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 16 17 18 19], out.shape = (20,), out.strides = (1,)
b = [ 0  2  4  6  8 10 12 14 16 18], b.shape = (10,), b.strides = (2,)
out = [ -3 1 -3 3 -3 5 -3 7 -3 9 -3 11 -3 13 -3 15 -3 17 -3 19], out.shape = (20,), out.strides = (1,)
b = [-3 -3 -3 -3 -3 -3 -3 -3 -3 -3], b.shape = (10,), b.strides = (2,)

[chudur-budur]

This is related to

IntelPython/dpnp#1649

[chudur-budur]

Resolved by #1271

[chudur-budur]

#1271 works only in 1-dim case.

If the array dimension is more than 1, the indexing doesn't skip the right amount of rows across the axis-0.

Reopening.

[chudur-budur]

Resolved by #1285

[diptorupd]

Please close issues when the associated PR is merged. Reopening the ticket.