CUDA: Add support for passing tuples to kernels via vectorize

Feature request

numba version used: 0.53.0.dev0+357.g4d3d2673c.dirty

I would like to pass tuples to CUDA kernels by using vectorize similar to #6343. Here is the code and subsequent error

from numba import cuda, vectorize
import numpy as np
import pkg_resources
# working based on https://github.com/numba/numba/pull/6343

print('numba version used: ' + str(pkg_resources.get_distribution("numba").version))



@cuda.jit
def f(r, x):
    r[0] = x[0][0]
    r[1] = x[1]
    r[2] = x[2][0]

z = np.ones(1, dtype=np.int64)
x = (1 * z, 10, (4, 3))
r = np.zeros(len(x), dtype=np.int64)
# Use the syntax: kernel_function[blockspergrid, threadsperblock](arg0, arg1, ..., argn)
f[1, 1](r, x)
print(r)



def f2(r2, x2): 
    out = 0
    # if the next two lines are commented out, the error is TypeError: argument #2 has invalid type of <class 'tuple'>
    for i in range(len(x2)):
        out += x2[i]
    return out

z2 = np.arange(1024, dtype=np.int64)
list_tuple = []
for n in range(1024):
    list_tuple.append((n, n+1, n+2))
x2 = (1 * z2, list_tuple)
r2 = np.zeros((len(z2),len(x2)), dtype=np.int64)

v_cuda = vectorize(['int64(int64, int64)'], target = 'cuda')(f2)  
out = v_cuda(r2, x2)

Error is

  File "D:\Documents\Cristi\Tools\Python\GitHub\numba\6343.py", line 39, in <module>
    v_cuda = vectorize(['int64(int64, int64)'],)(f2)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\np\ufunc\decorators.py", line 120, in wrap
    vec.add(sig)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\np\ufunc\deviceufunc.py", line 399, in add
    corefn, return_type = self._compile_core(devfnsig)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\vectorizers.py", line 15, in _compile_core
    cudevfn = cuda.jit(sig, device=True, inline=True)(self.pyfunc)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\decorators.py", line 134, in device_jit
    debug=debug)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 294, in compile_device
    return DeviceFunction(pyfunc, return_type, args, inline=True, debug=False)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 325, in __init__
    debug=self.debug, inline=self.inline)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 48, in compile_cuda
    locals={})

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 603, in compile_extra
    return pipeline.compile_extra(func)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 353, in compile_extra
    return self._compile_bytecode()

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 415, in _compile_bytecode
    return self._compile_core()

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 395, in _compile_core
    raise e

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 386, in _compile_core
    pm.run(self.state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 339, in run
    raise patched_exception

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 330, in run
    self._runPass(idx, pass_inst, state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_lock.py", line 32, in _acquire_compile_lock
    return func(*args, **kwargs)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 289, in _runPass
    mutated |= check(pss.run_pass, internal_state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 262, in check
    mangled = func(compiler_state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typed_passes.py", line 100, in run_pass
    raise_errors=self._raise_errors)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typed_passes.py", line 72, in type_inference_stage
    infer.propagate(raise_errors=raise_errors)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typeinfer.py", line 1071, in propagate
    raise errors[0]

TypingError: No implementation of function Function(<built-in function len>) found for signature:
 
len(int64)
 
There are 16 candidate implementations:
    - Of which 16 did not match due to:
    Overload of function 'len': File: <numerous>: Line N/A.
      With argument(s): '(int64)':
     No match.

During: resolving callee type: Function(<built-in function len>)
During: typing of call at D:\Documents\Cristi\Tools\Python\GitHub\numba\6343.py (28)

Thanks for the request. Just to check, you'd like to be able to pass a tuple as an argument to a @vectorize function, ideally targetting the CUDA backend?

@vectorize

Yes, exactly.
My understanding, based at least for my needs, is that the degree of difficulty to add new features is ranked something like:

vectorize(cuda)
guvectorize(cuda)
njit(cuda)
vectorize(parallel)
guvectorize(parallel)
vectorize(cpu)
guvectorize(cpu)

Most features are added at the "bottom" (since it is easier to do small steps improvements) and then make way to introduce them for the ones on the "top".
For using some easier stuff, I found @vectorize cuda, is really powerful, but if I want to run some more complex stuff, I really hate to fall back to guvectorize(cpu) or use other modules e.g. numpy.
The reason I request the above, was that I was hoping it could trickle up the difficulty ranking.
Maybe it would also be helpful to make in the docs, a simplified euler diagram of which features do not work with each "decorator". So that in the future, it will be a unified decorator (hopefully vectorize) ... to rule them all :)

@vectorize

Yes, exactly.
My understanding, based at least for my needs, is that the degree of difficulty to add new features is ranked something like:

vectorize(cuda)
guvectorize(cuda)
njit(cuda)
vectorize(parallel)
guvectorize(parallel)
vectorize(cpu)
guvectorize(cpu)

Most features are added at the "bottom" (since it is easier to do small steps improvements) and then make way to introduce them for the ones on the "top".

I'm not really sure what this means. The difficulty in adding new features is not hugely connected to the decorators involved, it's often more about how well a feature maps to a given target and what the constraints are in the implementation, e.g. @vectorize has to produce a valid NumPy ufunc.

For using some easier stuff, I found @vectorize cuda, is really powerful, but if I want to run some more complex stuff, I really hate to fall back to guvectorize(cpu) or use other modules e.g. numpy.

You may wish to look at CuPy, which provides a NumPy interface for CUDA GPUs and can share "arrays" with Numba via https://numba.readthedocs.io/en/stable/cuda/cuda_array_interface.html

The reason I request the above, was that I was hoping it could trickle up the difficulty ranking.
Maybe it would also be helpful to make in the docs, a simplified euler diagram of which features do not work with each "decorator". So that in the future, it will be a unified decorator (hopefully vectorize) ... to rule them all :)

These decorators do different things to provide flexibility of expression for users. For each hardware target, @vectorize exists to permit writing ufuncs, @guvectorize exists to permit writing generalised ufuncs and @jit exists to provide a more general compilation mechanism.

From what I have experienced, vectorize has teh least amount of features.
I could use CuPy or tensorflow, but I can not use the numba's vectorize implementation.

From what I have experienced, vectorize has teh least amount of features.
I could use CuPy or tensorflow, but I can not use the numba's vectorize implementation.

@vectorize is implementing this specification: https://numpy.org/doc/stable/reference/ufuncs.html out of interest, what specific features are you looking for from that specification?

Support for passing tuples, and later on lists, dictionaries, on the cuda target
Also matmul and minverse (currently not available in numpy), would be welcomed also

Support for passing tuples, and later on lists, dictionaries, on the cuda target
Also matmul and minverse (currently not available in numpy), would be welcomed also

I think some of this is outside of the specification but agree that accepting more array-like types could be useful.

It seems that regular NumPy ufuncs accept tuples:

import numpy as np

print(np.cos((1, 2, 3)))
# array([ 0.54030231, -0.41614684, -0.9899925 ])

So I think this is a reasonable feature request for the CUDA target.

A starter patch:

diff --git a/numba/np/ufunc/deviceufunc.py b/numba/np/ufunc/deviceufunc.py
index 583d84056..83070af2b 100644
--- a/numba/np/ufunc/deviceufunc.py
+++ b/numba/np/ufunc/deviceufunc.py
@@ -90,16 +90,17 @@ class UFuncMechanism(object):
         Get all arguments in array form
         """
         for i, arg in enumerate(self.args):
-            if isinstance(arg, np.ndarray):
-                self.arrays[i] = arg
-            elif self.is_device_array(arg):
+            if self.is_device_array(arg):
                 self.arrays[i] = self.as_device_array(arg)
             elif isinstance(arg, (int, float, complex, np.number)):
                 # Is scalar
                 self.scalarpos.append(i)
             else:
-                raise TypeError("argument #%d has invalid type of %s" \
-                % (i + 1, type(arg) ))
+                array = np.asarray(arg)
+                if array.dtype == np.object:
+                    fmt = "argument #%d has invalid type of %s"
+                    raise TypeError(fmt % (i + 1, type(arg) ))
+                self.arrays[i] = array
 
     def _fill_argtypes(self):
         """
@@ -107,7 +108,7 @@ class UFuncMechanism(object):
         """
         for i, ary in enumerate(self.arrays):
             if ary is not None:
-                self.argtypes[i] = ary.dtype
+                self.argtypes[i] = np.asarray(ary).dtype
 
     def _resolve_signature(self):
         """Resolve signature.

which enables the following to work:

from numba import vectorize, float64
import math


@vectorize([float64(float64)], target='cuda')
def cuda_cos(x):
    return math.cos(x)


print(cuda_cos((1.0, 2.0, 3.0)))

but causes one test fail:

======================================================================
ERROR: test_auto_transfer (numba.cuda.tests.cudapy.test_vectorize.TestCUDAVectorize)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/cudapy/test_vectorize.py", line 163, in test_auto_transfer
    y = cuda_ufunc(x, dx).copy_to_host()
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 37, in __call__
    return CUDAUFuncMechanism.call(self.functions, args, kws)
  File "/home/gmarkall/numbadev/numba/numba/np/ufunc/deviceufunc.py", line 254, in call
    args = cr.get_arguments()
  File "/home/gmarkall/numbadev/numba/numba/np/ufunc/deviceufunc.py", line 204, in get_arguments
    self._resolve_signature()
  File "/home/gmarkall/numbadev/numba/numba/np/ufunc/deviceufunc.py", line 144, in _resolve_signature
    raise TypeError("No matching version.  GPU ufunc requires array "
TypeError: No matching version.  GPU ufunc requires array arguments to have the exact types.  This behaves like regular ufunc with casting='no'.

----------------------------------------------------------------------

I've added this to the 0.54 milestone, to remind me to take a look at a proper implementation ahead of 0.54.

@gmarkall

While looking into starter patch by @gmarkall

The test_auto_transfer is failing due the following change :-
Without it all tests pass without problem

@@ -107,7 +108,7 @@ class UFuncMechanism(object):
         """
         for i, ary in enumerate(self.arrays):
             if ary is not None:
-                self.argtypes[i] = ary.dtype
+                self.argtypes[i] = np.asarray(ary).dtype

Since ary would always be either np.ndarray or DeviceNDArray type, keeping ary.dtype should be fine? (correct me if I missed something )

Ideally, the above change should not break anything, but for some reason, if ary is of DeviceNDArray type then np.asarray(ary).dtype gives the wrong dtype.
For DeviceNDArray ary, np.asarray(ary) will always be a float64 ndarray irrespective of original dtype

>>> x = np.arange(n, dtype=np.int32)
>>> dx = cuda.to_device(x)
>>> print(dx.dtype)
>>> print(np.asarray(dx).dtype)
int32
float64

After digging more, looks like problem might be in DeviceNDArray.__array__() method

numba/numba/cuda/cudadrv/devicearray.py

Lines 482 to 486 in 00ad127

    
               def __array__(self, dtype=None): 
        
                   """ 
        
                   :return: an `numpy.ndarray`, so copies to the host. 
        
                   """ 
        
                   return self.copy_to_host().__array__(dtype)

In above, passing even None (dtype) to __array__ of ndarray will return new array with default dtype (float64)
ndarray.__array__ docs - https://numpy.org/doc/stable/reference/generated/numpy.ndarray.__array__.html

>>> import numpy as np
>>> a = np.arange(10, dtype=np.int16)   
>>> b = a.__array__(np.int16)          # no copy
>>> c = a.__array__(None)              # copy
>>> d = a.__array__()                  # no copy
>>> print(f"{a.dtype=}")
>>> print(f"{b.dtype=}")
>>> print(f"{c.dtype=}")
>>> print(f"{d.dtype=}")
a.dtype=dtype('int16')
b.dtype=dtype('int16')
c.dtype=dtype('float64')
d.dtype=dtype('int16')

Changing __array__(dtype) to either __array__(self.dtype) or __array__() should fix this.

@ashutoshvarma

My local testing indicates that PR #6698 fixes the issue that comes up when running tests with the starter patch. Many thanks for tracing through this and finding fix, @ashutoshvarma!

Since ary would always be either np.ndarray or DeviceNDArray type, keeping ary.dtype should be fine? (correct me if I missed something )

I'd have to go back and check, but I think that ary ended up being a tuple if the user passes a tuple to the ufunc, as in #6563 (comment) (if that's not the case, then somehow I found that we'd get something that didn't have a dtype popping up here, but don't recall how it occurred).

CUDA: Support passing tuples to ufuncs

Now that #6698 is in (the fix for #6697) this can be fixed - #6878 contains the fix.

Mayne I am missing something, but my desired initial output is not working

Please reopen the issue

numba version used: 0.54.0.dev0+540.gd4d09f378
[ 1 10  4]
Found 1 CUDA devices
id 0      b'GeForce GT 1030'                              [SUPPORTED]
                      Compute Capability: 6.1
                           PCI Device ID: 0
                              PCI Bus ID: 16
                                Watchdog: Enabled
                            Compute Mode: WDDM
             FP32/FP64 Performance Ratio: 32
Summary:
	1/1 devices are supported
Traceback (most recent call last):

  File "D:\Documents\Cristi\Tools\Python\GitHub\numba\6563.py", line 40, in <module>
    v_cuda = vectorize(['int64(int64, int64)'],)(f2)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\np\ufunc\decorators.py", line 125, in wrap
    vec.add(sig)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\np\ufunc\deviceufunc.py", line 397, in add
    corefn, return_type = self._compile_core(devfnsig)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\vectorizers.py", line 15, in _compile_core
    cudevfn = cuda.jit(sig, device=True, inline=True)(self.pyfunc)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\decorators.py", line 100, in device_jit
    debug=debug)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 396, in compile_device
    return DeviceFunction(pyfunc, return_type, args, inline=True, debug=False)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 426, in __init__
    debug=self.debug, inline=self.inline)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 165, in compile_cuda
    pipeline_class=CUDACompiler)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 675, in compile_extra
    return pipeline.compile_extra(func)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 419, in compile_extra
    return self._compile_bytecode()

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 483, in _compile_bytecode
    return self._compile_core()

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 462, in _compile_core
    raise e

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 453, in _compile_core
    pm.run(self.state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 339, in run
    raise patched_exception

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 330, in run
    self._runPass(idx, pass_inst, state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 289, in _runPass
    mutated |= check(pss.run_pass, internal_state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 262, in check
    mangled = func(compiler_state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typed_passes.py", line 112, in run_pass
    raise_errors=self._raise_errors)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typed_passes.py", line 83, in type_inference_stage
    errs = infer.propagate(raise_errors=raise_errors)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typeinfer.py", line 1071, in propagate
    raise errors[0]

TypingError: No implementation of function Function(<built-in function len>) found for signature:
 
len(int64)
 
There are 6 candidate implementations:
  - Of which 6 did not match due to:
  Overload of function 'len': File: <numerous>: Line N/A.
    With argument(s): '(int64)':
   No match.

During: resolving callee type: Function(<built-in function len>)
During: typing of call at D:\Documents\Cristi\Tools\Python\GitHub\numba\6563.py (29)



runfile('D:/Documents/Cristi/Tools/Python/GitHub/numba/6563.py', wdir='D:/Docu
ments/Cristi/Tools/Python/GitHub/numba')
numba version used: 0.54.0.dev0+540.gd4d09f378
[ 1 10  4]
Traceback (most recent call last):

  File "D:\Documents\Cristi\Tools\Python\GitHub\numba\6563.py", line 39, in <module>
    v_cuda = vectorize(['int64(int64, int64)'],)(f2)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\np\ufunc\decorators.py", line 125, in wrap
    vec.add(sig)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\np\ufunc\deviceufunc.py", line 397, in add
    corefn, return_type = self._compile_core(devfnsig)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\vectorizers.py", line 15, in _compile_core
    cudevfn = cuda.jit(sig, device=True, inline=True)(self.pyfunc)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\decorators.py", line 100, in device_jit
    debug=debug)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 396, in compile_device
    return DeviceFunction(pyfunc, return_type, args, inline=True, debug=False)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 426, in __init__
    debug=self.debug, inline=self.inline)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\cuda\compiler.py", line 165, in compile_cuda
    pipeline_class=CUDACompiler)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 675, in compile_extra
    return pipeline.compile_extra(func)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 419, in compile_extra
    return self._compile_bytecode()

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 483, in _compile_bytecode
    return self._compile_core()

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 462, in _compile_core
    raise e

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler.py", line 453, in _compile_core
    pm.run(self.state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 339, in run
    raise patched_exception

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 330, in run
    self._runPass(idx, pass_inst, state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_lock.py", line 35, in _acquire_compile_lock
    return func(*args, **kwargs)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 289, in _runPass
    mutated |= check(pss.run_pass, internal_state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\compiler_machinery.py", line 262, in check
    mangled = func(compiler_state)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typed_passes.py", line 112, in run_pass
    raise_errors=self._raise_errors)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typed_passes.py", line 83, in type_inference_stage
    errs = infer.propagate(raise_errors=raise_errors)

  File "D:\Kit\PortableApps\Miniconda_numba_dev\lib\site-packages\numba\core\typeinfer.py", line 1071, in propagate
    raise errors[0]

TypingError: No implementation of function Function(<built-in function len>) found for signature:
 
len(int64)
 
There are 12 candidate implementations:
  - Of which 12 did not match due to:
  Overload of function 'len': File: <numerous>: Line N/A.
    With argument(s): '(int64)':
   No match.

During: resolving callee type: Function(<built-in function len>)
During: typing of call at D:\Documents\Cristi\Tools\Python\GitHub\numba\6563.py (28)

I have a couple of queries about your initial example:

v_cuda = vectorize(['int64(int64, int64)'], target = 'cuda')(f2)

makes the argument x2 of f2 an int, but f2 tries to take its length:

    for i in range(len(x2)):

Secondly, x2 is a tuple of an array and a list of tuples:

x2 = (1 * z2, list_tuple)

Is the example as you intended? If so, can you explain a bit more about why either of the above should work? (perhaps I'm missing something?)

For the first part:
I was trying to do a loop over the input data with vectorize, since I found guvectorize too slow, but it should be fix now in #6889.
If it is not possible at the moment to do operations with a for loop using vectorize, the incident can remain closed.

For the second part:
I was trying to use smth like an 2D array equivalent to pass to vectorize, in order to use the for loop, since again, I found guvectorize too slow.

For the first part:
I was trying to do a loop over the input data with vectorize, since I found guvectorize too slow, but it should be fix now in #6889.

There's no change in performance with #6889 - it only tells the user if they're using guvectorize in an inefficient way. Once it's merged, you might find that your guvectorize case produces a warning - if so, guvectorize may not be appropriate for the use case, or the way in which the problem is mapped to a guvectorize function may need to be changed.

If it is not possible at the moment to do operations with a for loop using vectorize, the incident can remain closed.

You can use for loops in @vectorize functions - the issue in the above code is that it treats a scalar as if it is something that has length. If you replace len(x2) with something else (e.g. an integer), then this would remove that problem.

For the second part:
I was trying to use smth like an 2D array equivalent to pass to vectorize, in order to use the for loop, since again, I found guvectorize too slow.

The problem here is the list - you can't pass a list to a vectorized function, only tuples and arrays.

For further discussion related to the above code, it would be good to create a thread on the Numba Discourse forums: https://numba.discourse.group

Navigation Menu

Repository navigation

CUDA: Add support for passing tuples to kernels via vectorize #6563

Feature request

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

	def __array__(self, dtype=None):
	"""
	:return: an `numpy.ndarray`, so copies to the host.
	"""
	return self.copy_to_host().__array__(dtype)

CUDA: Add support for passing tuples to kernels via vectorize #6563

Description

Feature request

Activity

stuartarchibald commented on Dec 11, 2020

cristipurdel commented on Dec 11, 2020

stuartarchibald commented on Dec 14, 2020

cristipurdel commented on Dec 14, 2020

stuartarchibald commented on Dec 14, 2020

cristipurdel commented on Dec 15, 2020

stuartarchibald commented on Dec 15, 2020

gmarkall commented on Jan 29, 2021

gmarkall commented on Jan 29, 2021

gmarkall commented on Jan 29, 2021

ashutoshvarma commented on Feb 5, 2021

gmarkall commented on Feb 8, 2021

gmarkall commented on Mar 31, 2021

cristipurdel commented on May 14, 2021

gmarkall commented on May 14, 2021

cristipurdel commented on May 15, 2021

gmarkall commented on May 17, 2021

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions