Repository: numba/numba
Pull request: CUDA: Add support for passing tuples and namedtuples to kernels
Fixes #4405.

This adds support for tuples as per @heyer2's suggestion, and namedtuples with @atait's suggestion. Tests for passing each, and an empty tuple are added.

This passes tests for me locally with the hardware target and cudasim.

<!--

Thanks for wanting to contribute to Numba :)

First, if you need some help or want to chat to the core developers, please
visit https://gitter.im/numba/numba for real time chat or post to the Numba
mailing list https://groups.google.com/a/continuum.io/forum/#!forum/numba-users.

Here's some guidelines to help the review process go smoothly.

0. Please write a description in this text box of the changes that are being
   made.

1. Please ensure that you have written units tests for the changes made/features
   added.

2. If you are closing an issue please use one of the automatic closing words as
   noted here: https://help.github.com/articles/closing-issues-using-keywords/

3. If your pull request is not ready for review but you want to make use of the
   continuous integration testing facilities please label it with `[WIP]` and
   then remove the label when you'd like it to be reviewed.

4. Once review has taken place please do not add features or make changes out of
   the scope of those requested by the reviewer (doing this just add delays as
   already reviewed code ends up having to be re-reviewed/it is hard to tell
   what is new etc!). Further, please do not rebase your branch on master/force
   push/rewrite history, doing any of these causes the context of any comments
   made by reviewers to be lost. If conflicts occur against master they should
   be resolved by merging master into the branch used for making the pull
   request.

Many thanks in advance for your cooperation!

-->


Problem statement:
CUDA functions can be compiled but not called with Tuple and UniTuple arguments
As the title suggests, functions with `Tuple` or  `UniTuple` arguments can be compiled just fine, but are unable to actually be called without raising an exception. Here's an example using `UniTuple`:

```python
import numpy as np
from numba import cuda, types

n_array = 5
@cuda.jit(types.void(types.UniTuple(types.float32[:], n_array), types.float32[:]))
def test(arrays, out):
	i = cuda.threadIdx.x
	tmp = 0
	for arr in arrays:
		tmp += arr[i]
	out[i] = tmp

def main():
	n_elements = 3
	arrays_in = [k * np.ones(n_elements, np.float32) for k in range(n_array)]
	array_out = np.empty(n_elements, np.float32)
	test[1, n_elements](arrays_in, array_out)
	print("output:", array_out)

if __name__=="__main__": main()
```
I have no real knowledge about the internals of Numba,
but I found adding the following lines to `_prepare_args` of the `CUDAKernel` class:


```python
...
elif isinstance(ty, types.UniTuple) or isinstance(ty, types.Tuple):
    assert len(ty)==len(val)
    for t, v in zip(ty, val):
       self._prepare_args(t, v, stream, retr, kernelargs)
...
```

seems to fix the minimal example I provided.
I am using the latest version of Numba.

Hints from issue discussion:
Thanks for the report, I can reproduce. Seems like this could work, one of the core devs will take a look.
Would be wonderful if this was fixed.
@gmarkall : Can you take a look at this? (in your backlog of other CUDA things you are checking out...)  It does seem like a straightforward thing to fix since we unpack tuples, so there is no memory allocated data structure to worry about on the GPU.
Whoops, just noticed this mention - assigning to myself.
This fix is a lifesaver. 
- It enables otherwise impossible patterns for variable number of arguments (including but not limited to @heyer2's example)
- It makes code significantly more readable and modular

Supporting `namedtuple` is also possible with this approach
```python
elif isinstance(ty, (types.UniTuple, types.Tuple, types.NamedUniTuple, types.NamedTuple)):
    ...
```

Commit message: CUDA: Add support for passing tuples and namedtuples to kernels

Fixes #4405.

Patch to review:
diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py
index cbcc0364b56..b55ee24c618 100644
--- a/numba/cuda/compiler.py
+++ b/numba/cuda/compiler.py
@@ -756,6 +756,12 @@ def _prepare_args(self, ty, val, stream, retr, kernelargs):
             devrec = wrap_arg(val).to_device(retr, stream)
             kernelargs.append(devrec)
 
+        elif isinstance(ty, (types.UniTuple, types.Tuple, types.NamedUniTuple,
+                             types.NamedTuple)):
+            assert len(ty) == len(val)
+            for t, v in zip(ty, val):
+                self._prepare_args(t, v, stream, retr, kernelargs)
+
         else:
             raise NotImplementedError(ty, val)
 
diff --git a/numba/cuda/tests/cudapy/test_array_args.py b/numba/cuda/tests/cudapy/test_array_args.py
index a60649e2d64..0f1e3c90f7d 100644
--- a/numba/cuda/tests/cudapy/test_array_args.py
+++ b/numba/cuda/tests/cudapy/test_array_args.py
@@ -1,4 +1,5 @@
 import numpy as np
+from collections import namedtuple
 
 from numba import cuda
 from numba.cuda.testing import unittest, CUDATestCase
@@ -21,6 +22,83 @@ def kernel(x, y):
         kernel[10, 1](x, y)
         self.assertTrue(np.all(x == y))
 
+    def test_unituple(self):
+        @cuda.jit
+        def f(r, x):
+            r[0] = x[0]
+            r[1] = x[1]
+            r[2] = x[2]
+
+        x = (1, 2, 3)
+        r = np.zeros(len(x), dtype=np.int64)
+        f[1, 1](r, x)
+
+        for i in range(len(x)):
+            self.assertEqual(r[i], x[i])
+
+    def test_tuple(self):
+        @cuda.jit
+        def f(r1, r2, x):
+            r1[0] = x[0]
+            r1[1] = x[1]
+            r1[2] = x[2]
+            r2[0] = x[3]
+            r2[1] = x[4]
+            r2[2] = x[5]
+
+        x = (1, 2, 3, 4.5, 5.5, 6.5)
+        r1 = np.zeros(len(x) // 2, dtype=np.int64)
+        r2 = np.zeros(len(x) // 2, dtype=np.float64)
+        f[1, 1](r1, r2, x)
+
+        for i in range(len(r1)):
+            self.assertEqual(r1[i], x[i])
+
+        for i in range(len(r2)):
+            self.assertEqual(r2[i], x[i + len(r1)])
+
+    def test_namedunituple(self):
+        @cuda.jit
+        def f(r, x):
+            r[0] = x.x
+            r[1] = x.y
+
+        Point = namedtuple('Point', ('x', 'y'))
+        x = Point(1, 2)
+        r = np.zeros(len(x), dtype=np.int64)
+        f[1, 1](r, x)
+
+        self.assertEqual(r[0], x.x)
+        self.assertEqual(r[1], x.y)
+
+    def test_namedtuple(self):
+        @cuda.jit
+        def f(r1, r2, x):
+            r1[0] = x.x
+            r1[1] = x.y
+            r2[0] = x.r
+
+        Point = namedtuple('Point', ('x', 'y', 'r'))
+        x = Point(1, 2, 2.236)
+        r1 = np.zeros(2, dtype=np.int64)
+        r2 = np.zeros(1, dtype=np.float64)
+        f[1, 1](r1, r2, x)
+
+        self.assertEqual(r1[0], x.x)
+        self.assertEqual(r1[1], x.y)
+        self.assertEqual(r2[0], x.r)
+
+    def test_empty_tuple(self):
+        @cuda.jit
+        def f(r, x):
+            r[0] = len(x)
+
+        x = tuple()
+        r = np.ones(1, dtype=np.int64)
+        f[1, 1](r, x)
+
+        self.assertEqual(r[0], 0)
+
 
 if __name__ == '__main__':
     unittest.main()


Code review comments:
diff_hunk: @@ -756,6 +756,12 @@ def _prepare_args(self, ty, val, stream, retr, kernelargs):
             devrec = wrap_arg(val).to_device(retr, stream)
             kernelargs.append(devrec)
 
+        elif isinstance(ty, (types.UniTuple, types.Tuple, types.NamedUniTuple,
+                             types.NamedTuple)):
path: numba/cuda/compiler.py
text: @user1:
Can this just check against `types.BaseTuple`? 

@author:
Looks like it can :-)