__cuda_array_interface__

[sklam]

for interop with 3rd party CUDA array.

[codecov-io]

Codecov Report

Merging #2860 into master will decrease coverage by 0.05%.
The diff coverage is 4.76%.

@@            Coverage Diff             @@
##           master    #2860      +/-   ##
==========================================
- Coverage   85.84%   85.79%   -0.06%     
==========================================
  Files         326      327       +1     
  Lines       68311    68422     +111     
  Branches     7721     7729       +8     
==========================================
+ Hits        58643    58703      +60     
- Misses       8428     8481      +53     
+ Partials     1240     1238       -2

[seibert]

I can confirm this works in combination with cupy/cupy#1144.

import cupy

from numba import cuda

@cuda.jit
def add(x, y, out):
    start = cuda.grid(1)
    stride = cuda.gridsize(1)
    for i in range(start, x.shape[0], stride):
        out[i] = x[i] + y[i]

a = cupy.arange(10)
b = a * 2
out = cupy.empty_like(a)

print('out before:', out)

add[1, 32](a, b, out)

print('out after:', out)
print('array types:', type(a), type(b), type(out))

Output:

(cupy_dev) seibert@nvidia1:~/repos/numba_cupy$ python demo.py
out before: [0 0 0 0 0 0 0 0 0 0]
out after: [ 0  3  6  9 12 15 18 21 24 27]
array types: <class 'cupy.core.core.ndarray'> <class 'cupy.core.core.ndarray'> <class 'cupy.core.core.ndarray'>

[seibert]

It looks like support for @vectorize/@guvectorize needs to be done still.

[seibert]

I think you forgot to check in cuda_array_interface.rst :)

[sklam]

Oops, added missing file now

[seibert]

I think we should phrase this as "These objects also can be manually converted into a Numba device array by creating a view of the GPU buffer using the following APIs:"

[seibert]

s/numba/Numba/

[seibert]

Two minor documentation nits, and then this looks good to merge.