Repository: numba/numba
Pull request: CUDA: Support for multiple signatures
When a list of signatures is passed to the `@cuda.jit` decorator, the dispatcher is compiled for each signature. This does not result in a specialized kernel, even if the list contained just one signature, though compilation will be disabled for the dispatcher.

This ensures that:

- An exact match is required when a list of signatures is provided (matching the dispatch behaviour of the CPU target)
- An exact match is required, but not checked for, when a single signature is provided (matching the current CUDA target behaviour).
- Other functionality that depends on whether a kernel is specalialized (getting shared memory per block, for example) retains its existing behaviour for specialized kernels, such as returning the shared memory per block directly instead of in a dict for each implementation. For dispatchers where multiple signatures were provided, a dict should be returned by these functions.

Tests for multiple signatures are ported from the CPU multiple signature tests in `numba.tests.test_dispatcher`. Note that there is a difference between expectations for kernels and device functions when ambiguity cannot be resolved for an overload - launching a kernel with ambiguous overloads is not permitted, but compiling a call to a device function with ambiguous overloads is permitted - this feels a little odd but is the way Numba currently seems to work for dispatch vs. function calls, and is recorded separately as Issue #8307.

Fixes #3226.

Problem statement:
CUDA JIT errors when multiple signatures provided
With Numba's CPU JIT, one is able to supply multiple signatures (e.g. single and double precision floating values) for a function. However when trying the same with Numba's CUDA JIT, this results in an error. Limiting the CUDA JIT to a single signature fixes the issue. Code snippet demonstrating this below.

<details>
<summary>Example:</summary>

```python
In [1]: from numba import jit

In [2]: from numba.cuda import jit as cuda_jit

In [3]: @jit(["f4(f4, f4)", "f8(f8, f8)"])
   ...: def f(a, b):
   ...:     return a + b
   ...: 
   ...: 

In [4]: @cuda_jit(["f4(f4, f4)", "f8(f8, f8)"])
   ...: def g(a, b):
   ...:     return a + b
   ...: 
   ...: 
---------------------------------------------------------------------------
ValueError                                Traceback (most recent call last)
<ipython-input-4-5a86fdfdf6d9> in <module>()
----> 1 @cuda_jit(["f4(f4, f4)", "f8(f8, f8)"])
      2 def g(a, b):
      3     return a + b

~/.conda/envs/numba/lib/python3.6/site-packages/numba/cuda/compiler.py in __call__(self, *args)
    700         Specialize and invoke this kernel with *args*.
    701         '''
--> 702         kernel = self.specialize(*args)
    703         cfg = kernel[self.griddim, self.blockdim, self.stream, self.sharedmem]
    704         cfg(*args)

~/.conda/envs/numba/lib/python3.6/site-packages/numba/cuda/compiler.py in specialize(self, *args)
    710         '''
    711         argtypes = tuple(
--> 712             [self.typingctx.resolve_argument_type(a) for a in args])
    713         kernel = self.compile(argtypes)
    714         return kernel

~/.conda/envs/numba/lib/python3.6/site-packages/numba/cuda/compiler.py in <listcomp>(.0)
    710         '''
    711         argtypes = tuple(
--> 712             [self.typingctx.resolve_argument_type(a) for a in args])
    713         kernel = self.compile(argtypes)
    714         return kernel

~/.conda/envs/numba/lib/python3.6/site-packages/numba/typing/context.py in resolve_argument_type(self, val)
    288         ValueError is raised for unsupported types.
    289         """
--> 290         return typeof(val, Purpose.argument)
    291 
    292     def resolve_value_type(self, val):

~/.conda/envs/numba/lib/python3.6/site-packages/numba/typing/typeof.py in typeof(val, purpose)
     31         msg = errors.termcolor.errmsg(
     32             "cannot determine Numba type of %r") % (type(val),)
---> 33         raise ValueError(msg)
     34     return ty
     35 

ValueError: cannot determine Numba type of <class 'function'>
```

</details>

Hints from issue discussion:
Thanks for the report. It seems that this is the case, lists of signatures are not handled, we should either raise with a more specific error message or just fix the problem!

It's also worth noting from your example that `@cuda.jit` functions are GPU kernel functions and should be `void` [by definition](http://numba.pydata.org/numba-doc/dev/cuda/kernels.html#kernel-declaration).

Also, might be worth looking at `@vectorize` and `@guvectorize`, both of which support the kwarg `target='cuda'` (and multiple signatures). Here's a quick example of both ways:

```python
from numba.cuda import jit as cuda_jit
from numba import cuda, vectorize
import numpy as np

@cuda_jit("void(f8[:], f8)")
def g(a, b):
    tx = cuda.threadIdx.x
    ty = cuda.blockIdx.x
    bw = cuda.blockDim.x
    pos = tx + ty * bw
    if pos < a.size:
        a[pos] += b

@vectorize(['f4(f4, f4)', 'f8(f8, f8)'], target='cuda')
def vec_g(a, b):
    return a + b

x = np.arange(64.)
y = 17

g[1, 64](x, y)
print(x)

x = np.arange(64.)
print(vec_g(x, y))
```
It appears that there may be cases when kernels with signatures specified don't catch type errors as they should do, e.g.: https://github.com/numba/numba/pull/5388#issuecomment-603226418 - when looking into adding support for multiple signatures, type checking should also be checked / considered.
It would be **really** cool if this behavior of `cuda.jit` was (explicitly) documented ;) I have just burned a couple of hours on this, thinking that I made an utterly stupid mistake, until I ended up digging into the sources and issues ... In the meantime, some more useful exceptions could also help. When passing a list of strings, `cuda.jit` currently complains: `[...] is not a callable object` ... which kind of makes sense but is inconsistent with `numba.jit` and does not tell the actual story: Only one signature allowed.

Just out of interest: What are the odds (or technical feasibility) for `cuda.jit` to support multiple signatures just like `numba.jit`?
> Just out of interest: What are the odds (or technical feasibility) for `cuda.jit` to support multiple signatures just like `numba.jit`?

Hopefully much better than they were when this issue was first reported, now that the CUDA Dispatcher uses a lot more of the infrastructure that the CPU target does for dispatch. I'll look into this and see what the gaps are.
Looks like this is pretty doable, so I'm adding it to the 0.57 milestone. With the changes in my branch so far: https://github.com/gmarkall/numba/tree/issue-3226 (I just have one commit, https://github.com/gmarkall/numba/commit/b2f5a097a7701ac087379abc3f9670c52d876752, in it at the moment) the following:

```python
from numba import cuda
import numpy as np

# Define a kernel with multiple signatures

sigs = [
    'void(int32[::1], int32[::1])',
    'void(float32[::1], float32[::1])'
]


@cuda.jit(sigs)
def f(r, x):
    r[0] = x[0] * 2 + 1


# Create arguments for supported and unsupported types

x_int32 = np.ones(1, dtype=np.int32)
x_int64 = np.ones(1, dtype=np.int64)
x_float32 = np.ones(1, dtype=np.float32)
x_float64 = np.ones(1, dtype=np.float64)

r_int32 = np.zeros_like(x_int32)
r_int64 = np.zeros_like(x_int64)
r_float32 = np.zeros_like(x_float32)
r_float64 = np.zeros_like(x_float64)


# Execute with supported types and print results
f[1, 1](r_int32, x_int32)
f[1, 1](r_float32, x_float32)

print(r_int32)
print(r_float32)


# Demonstrate that unsupported types are rejected

def unsupported_call(r, x):
    try:
        f[1, 1](r, x)
    except TypeError as te:
        print(f'Got TypeError: {te.args[0]}')


unsupported_call(r_int64, x_int64)
unsupported_call(r_float64, x_float64)
``` 

produces:

```
[3]
[3.]
Got TypeError: No matching definition for argument type(s) array(int64, 1d, C), array(int64, 1d, C)
Got TypeError: No matching definition for argument type(s) array(float64, 1d, C), array(float64, 1d, C)
```

which is what I would hope to get. 

This work still needs:

- [ ] ~~The example above formalizing into a test case~~ (decided to just use the relevant tests from the CPU dispatcher)
- [X] Other relevant test cases from the CPU dispatcher tests porting across
- [ ] Documentation updating
- [ ] Checking that device functions also behave correctly with multiple signatures

(**edited**  to make the to-do list a checklist to track progress)

Commit message: Update cuda.jit documentation for multiple signatures

Patch to review:
diff --git a/docs/source/reference/jit-compilation.rst b/docs/source/reference/jit-compilation.rst
index ac67a593b14..06f7a787568 100644
--- a/docs/source/reference/jit-compilation.rst
+++ b/docs/source/reference/jit-compilation.rst
@@ -1,12 +1,11 @@
 Just-in-Time compilation
 ========================
 
+.. _jit-decorator:
 
 JIT functions
 -------------
 
-.. _jit-decorator:
-
 .. decorator:: numba.jit(signature=None, nopython=False, nogil=False, cache=False, forceobj=False, parallel=False, error_model='python', fastmath=False, locals={}, boundscheck=False)
 
    Compile the decorated function on-the-fly to produce efficient machine
diff --git a/numba/cuda/decorators.py b/numba/cuda/decorators.py
index 6703a6feacc..5a34df46918 100644
--- a/numba/cuda/decorators.py
+++ b/numba/cuda/decorators.py
@@ -14,21 +14,22 @@
 def jit(func_or_sig=None, device=False, inline=False, link=[], debug=None,
         opt=True, cache=False, **kws):
     """
-    JIT compile a python function conforming to the CUDA Python specification.
-    If a signature is supplied, then a function is returned that takes a
-    function to compile.
+    JIT compile a Python function for CUDA GPUs.
 
-    :param func_or_sig: A function to JIT compile, or a signature of a function
-       to compile. If a function is supplied, then a
-       :class:`numba.cuda.compiler.AutoJitCUDAKernel` is returned. If a
-       signature is supplied, then a function is returned. The returned
-       function accepts another function, which it will compile and then return
-       a :class:`numba.cuda.compiler.AutoJitCUDAKernel`.
+    :param func_or_sig: A function to JIT compile, or *signatures* of a
+       function to compile. If a function is supplied, then a
+       :class:`Dispatcher <numba.cuda.dispatcher.CUDADispatcher>` is returned.
+       Otherwise, ``func_or_sig`` may be a signature or a list of signatures,
+       and a function is returned. The returned function accepts another
+       function, which it will compile and then return a :class:`Dispatcher
+       <numba.cuda.dispatcher.CUDADispatcher>`. See :ref:`jit-decorator` for
+       more information about passing signatures.
 
        .. note:: A kernel cannot have any return value.
     :param device: Indicates whether this is a device function.
     :type device: bool
-    :param link: A list of files containing PTX source to link with the function
+    :param link: A list of files containing PTX or CUDA C/C++ source to link
+       with the function
     :type link: list
     :param debug: If True, check for exceptions thrown when executing the
        kernel. Since this degrades performance, this should only be used for
@@ -83,16 +84,20 @@ def jit(func_or_sig=None, device=False, inline=False, link=[], debug=None,
         raise ValueError("link keyword invalid for device function")
 
     if sigutils.is_signature(func_or_sig):
+        signatures = [func_or_sig]
+        specialized = True
+    elif isinstance(func_or_sig, list):
+        signatures = func_or_sig
+        specialized = False
+    else:
+        signatures = None
+
+    if signatures:
         if config.ENABLE_CUDASIM:
             def jitwrapper(func):
                 return FakeCUDAKernel(func, device=device, fastmath=fastmath)
             return jitwrapper
 
-        argtypes, restype = sigutils.normalize_signature(func_or_sig)
-
-        if restype and not device and restype != types.void:
-            raise TypeError("CUDA kernel must have void return type.")
-
         def _jit(func):
             targetoptions = kws.copy()
             targetoptions['debug'] = debug
@@ -107,14 +112,20 @@ def _jit(func):
             if cache:
                 disp.enable_caching()
 
-            if device:
-                from numba.core import typeinfer
-                with typeinfer.register_dispatcher(disp):
-                    disp.compile_device(argtypes)
-            else:
-                disp.compile(argtypes)
+            for sig in signatures:
+                argtypes, restype = sigutils.normalize_signature(sig)
+
+                if restype and not device and restype != types.void:
+                    raise TypeError("CUDA kernel must have void return type.")
+
+                if device:
+                    from numba.core import typeinfer
+                    with typeinfer.register_dispatcher(disp):
+                        disp.compile_device(argtypes)
+                else:
+                    disp.compile(argtypes)
 
-            disp._specialized = True
+            disp._specialized = specialized
             disp.disable_compile()
 
             return disp
diff --git a/numba/cuda/tests/cudapy/test_dispatcher.py b/numba/cuda/tests/cudapy/test_dispatcher.py
index efbe4291292..ee050131af3 100644
--- a/numba/cuda/tests/cudapy/test_dispatcher.py
+++ b/numba/cuda/tests/cudapy/test_dispatcher.py
@@ -2,6 +2,7 @@
 import threading
 
 from numba import cuda, float32, float64, int32, int64, void
+from numba.core.errors import TypingError
 from numba.cuda.testing import skip_on_cudasim, unittest, CUDATestCase
 import math
 
@@ -213,6 +214,171 @@ def wrapper():
             t.join()
         self.assertFalse(errors)
 
+    def test_explicit_signatures(self):
+        sigs = ["(int64[::1], int64, int64)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+        self.assertIn("No matching definition", str(cm.exception))
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float64[::1], float32, float32)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Fail to resolve ambiguity between the two best overloads
+        # (Also deliberate float64[::1] for the first argument in all cases)
+        f = cuda.jit(["(float64[::1], float32,float64)",
+                      "(float64[::1], float64,float32)",
+                      "(float64[::1], int64,int64)"])(add_kernel)
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.float64)
+            f[1, 1](r, 1.0, 2.0)
+
+        # The two best matches are output in the error message, as well
+        # as the actual argument types.
+        self.assertRegexpMatches(
+            str(cm.exception),
+            r"Ambiguous overloading for <function add_kernel [^>]*> "
+            r"\(array\(float64, 1d, C\), float64, float64\):\n"
+            r"\(array\(float64, 1d, C\), float32, float64\) -> none\n"
+            r"\(array\(float64, 1d, C\), float64, float32\) -> none"
+        )
+        # The integer signature is not part of the best matches
+        self.assertNotIn("int64", str(cm.exception))
+
+    @unittest.expectedFailure
+    def test_explicit_signatures_unsafe(self):
+        # These tests are from test_explicit_signatures, but have to be xfail
+        # at present because _prepare_args in the CUDA target cannot handle
+        # unsafe conversions of arguments.
+        f = cuda.jit("(int64[::1], int64,int64)")(add_kernel)
+        r = np.zeros(1, dtype=np.int64)
+
+        # Approximate match (unsafe conversion)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 3)
+        self.assertEqual(len(f.overloads), 1, f.overloads)
+
+        sigs = ["(int64[::1], int64,int64)", "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+        r = np.zeros(1, dtype=np.float64)
+        # Approximate match (int32 -> float64 is a safe conversion)
+        f[1, 1](r, np.int32(1), 2.5)
+        self.assertPreciseEqual(r[0], 3.5)
+
+    def add_device_usecase(self, sigs):
+        # Generate a kernel that calls the add device function compiled with a
+        # given set of signatures
+        add_device = cuda.jit(sigs, device=True)(add)
+
+        @cuda.jit
+        def f(r, x, y):
+            r[0] = add_device(x, y)
+
+        return f
+
+    def test_explicit_signatures_device(self):
+        # Tests similar to test_explicit_signatures, but on a device function
+        # instead of a kernel
+        sigs = ["(int64, int64)", "(float64, float64)"]
+        f = self.add_device_usecase(sigs)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypingError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+
+        msg = str(cm.exception)
+        self.assertIn("Invalid use of type", msg)
+        self.assertIn("with parameters (complex128, complex128)", msg)
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float32, float32)", "(float64, float64)"]
+        f = self.add_device_usecase(sigs)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Ambiguity between the two best overloads resolves. This is somewhat
+        # surprising given that ambiguity is not permitted for dispatching
+        # overloads when launching a kernel, but seems to be the general
+        # behaviour of Numba (See Issue #8307:
+        # https://github.com/numba/numba/issues/8307).
+        sigs = ["(float32,float64)", "(float64,float32)", "(int64,int64)"]
+        f = self.add_device_usecase(sigs)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+    def test_explicit_signatures_device_unsafe(self):
+        # These tests are from test_explicit_signatures. The device function
+        # variant of these tests can succeed on CUDA because the compilation
+        # can handle unsafe casting (c.f. test_explicit_signatures_unsafe which
+        # has to xfail due to _prepare_args not supporting unsafe casting).
+        sigs = ["(int64, int64)"]
+        f = self.add_device_usecase(sigs)
+
+        # Approximate match (unsafe conversion)
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 3)
+        self.assertEqual(len(f.overloads), 1, f.overloads)
+
+        sigs = ["(int64, int64)", "(float64, float64)"]
+        f = self.add_device_usecase(sigs)
+
+        # Approximate match (int32 -> float64 is a safe conversion)
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.int32(1), 2.5)
+        self.assertPreciseEqual(r[0], 3.5)
+
     def test_get_regs_per_thread_unspecialized(self):
         # A kernel where the register usage per thread is likely to differ
         # between different specializations


Code review comments:
diff_hunk: @@ -213,6 +214,171 @@ def wrapper():
             t.join()
         self.assertFalse(errors)
 
+    def test_explicit_signatures(self):
+        sigs = ["(int64[::1], int64, int64)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+        self.assertIn("No matching definition", str(cm.exception))
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float64[::1], float32, float32)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Fail to resolve ambiguity between the two best overloads
+        # (Also deliberate float64[::1] for the first argument in all cases)
+        f = cuda.jit(["(float64[::1], float32,float64)",
+                      "(float64[::1], float64,float32)",
+                      "(float64[::1], int64,int64)"])(add_kernel)
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.float64)
+            f[1, 1](r, 1.0, 2.0)
+
+        # The two best matches are output in the error message, as well
+        # as the actual argument types.
+        self.assertRegexpMatches(
+            str(cm.exception),
+            r"Ambiguous overloading for <function add_kernel [^>]*> "
+            r"\(array\(float64, 1d, C\), float64, float64\):\n"
+            r"\(array\(float64, 1d, C\), float32, float64\) -> none\n"
+            r"\(array\(float64, 1d, C\), float64, float32\) -> none"
+        )
+        # The integer signature is not part of the best matches
+        self.assertNotIn("int64", str(cm.exception))
+
+    @unittest.expectedFailure
+    def test_explicit_signatures_unsafe(self):
+        # These tests are from test_explicit_signatures, but have to be xfail
+        # at present because _prepare_args in the CUDA target cannot handle
+        # unsafe conversions of arguments.
+        f = cuda.jit("(int64[::1], int64,int64)")(add_kernel)
path: numba/cuda/tests/cudapy/test_dispatcher.py
text: @user1:
```suggestion
        f = cuda.jit("(int64[::1], int64, int64)")(add_kernel)
```

diff_hunk: @@ -213,6 +214,171 @@ def wrapper():
             t.join()
         self.assertFalse(errors)
 
+    def test_explicit_signatures(self):
+        sigs = ["(int64[::1], int64, int64)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+        self.assertIn("No matching definition", str(cm.exception))
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float64[::1], float32, float32)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Fail to resolve ambiguity between the two best overloads
+        # (Also deliberate float64[::1] for the first argument in all cases)
+        f = cuda.jit(["(float64[::1], float32,float64)",
+                      "(float64[::1], float64,float32)",
+                      "(float64[::1], int64,int64)"])(add_kernel)
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.float64)
+            f[1, 1](r, 1.0, 2.0)
+
+        # The two best matches are output in the error message, as well
+        # as the actual argument types.
+        self.assertRegexpMatches(
+            str(cm.exception),
+            r"Ambiguous overloading for <function add_kernel [^>]*> "
+            r"\(array\(float64, 1d, C\), float64, float64\):\n"
+            r"\(array\(float64, 1d, C\), float32, float64\) -> none\n"
+            r"\(array\(float64, 1d, C\), float64, float32\) -> none"
+        )
+        # The integer signature is not part of the best matches
+        self.assertNotIn("int64", str(cm.exception))
+
+    @unittest.expectedFailure
+    def test_explicit_signatures_unsafe(self):
+        # These tests are from test_explicit_signatures, but have to be xfail
+        # at present because _prepare_args in the CUDA target cannot handle
+        # unsafe conversions of arguments.
+        f = cuda.jit("(int64[::1], int64,int64)")(add_kernel)
+        r = np.zeros(1, dtype=np.int64)
+
+        # Approximate match (unsafe conversion)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 3)
+        self.assertEqual(len(f.overloads), 1, f.overloads)
+
+        sigs = ["(int64[::1], int64,int64)", "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+        r = np.zeros(1, dtype=np.float64)
+        # Approximate match (int32 -> float64 is a safe conversion)
+        f[1, 1](r, np.int32(1), 2.5)
+        self.assertPreciseEqual(r[0], 3.5)
+
+    def add_device_usecase(self, sigs):
+        # Generate a kernel that calls the add device function compiled with a
+        # given set of signatures
+        add_device = cuda.jit(sigs, device=True)(add)
+
+        @cuda.jit
+        def f(r, x, y):
+            r[0] = add_device(x, y)
+
+        return f
+
+    def test_explicit_signatures_device(self):
+        # Tests similar to test_explicit_signatures, but on a device function
+        # instead of a kernel
+        sigs = ["(int64, int64)", "(float64, float64)"]
+        f = self.add_device_usecase(sigs)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypingError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+
+        msg = str(cm.exception)
+        self.assertIn("Invalid use of type", msg)
+        self.assertIn("with parameters (complex128, complex128)", msg)
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float32, float32)", "(float64, float64)"]
+        f = self.add_device_usecase(sigs)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Ambiguity between the two best overloads resolves. This is somewhat
+        # surprising given that ambiguity is not permitted for dispatching
+        # overloads when launching a kernel, but seems to be the general
+        # behaviour of Numba (See Issue #8307:
+        # https://github.com/numba/numba/issues/8307).
+        sigs = ["(float32,float64)", "(float64,float32)", "(int64,int64)"]
path: numba/cuda/tests/cudapy/test_dispatcher.py
text: @user1:
```suggestion
        sigs = ["(float32, float64)", "(float64, float32)", "(int64, int64)"]
```

diff_hunk: @@ -83,16 +84,20 @@ def jit(func_or_sig=None, device=False, inline=False, link=[], debug=None,
         raise ValueError("link keyword invalid for device function")
 
     if sigutils.is_signature(func_or_sig):
+        signatures = [func_or_sig]
+        specialized = True
+    elif isinstance(func_or_sig, list):
+        signatures = func_or_sig
+        specialized = False
+    else:
+        signatures = None
+
+    if signatures:
path: numba/cuda/decorators.py
text: @user1:
```suggestion
    if signatures is not None:
```

diff_hunk: @@ -213,6 +214,171 @@ def wrapper():
             t.join()
         self.assertFalse(errors)
 
+    def test_explicit_signatures(self):
+        sigs = ["(int64[::1], int64, int64)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+        self.assertIn("No matching definition", str(cm.exception))
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float64[::1], float32, float32)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Fail to resolve ambiguity between the two best overloads
+        # (Also deliberate float64[::1] for the first argument in all cases)
+        f = cuda.jit(["(float64[::1], float32,float64)",
+                      "(float64[::1], float64,float32)",
+                      "(float64[::1], int64,int64)"])(add_kernel)
path: numba/cuda/tests/cudapy/test_dispatcher.py
text: @user1:
```suggestion
        f = cuda.jit(["(float64[::1], float32, float64)",
                      "(float64[::1], float64, float32)",
                      "(float64[::1], int64, int64)"])(add_kernel)
```

diff_hunk: @@ -213,6 +214,171 @@ def wrapper():
             t.join()
         self.assertFalse(errors)
 
+    def test_explicit_signatures(self):
+        sigs = ["(int64[::1], int64, int64)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        # Exact signature matches
+        r = np.zeros(1, dtype=np.int64)
+        f[1, 1](r, 1, 2)
+        self.assertPreciseEqual(r[0], 3)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 4.0)
+
+        # No conversion
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.complex128)
+            f[1, 1](r, 1j, 1j)
+        self.assertIn("No matching definition", str(cm.exception))
+        self.assertEqual(len(f.overloads), 2, f.overloads)
+
+        # A more interesting one...
+        # (Note that the type of r is deliberately float64 in both cases so
+        # that dispatch is differentiated on the types of x and y only, to
+        # closely preserve the intent of the original test from
+        # numba.tests.test_dispatcher)
+        sigs = ["(float64[::1], float32, float32)",
+                "(float64[::1], float64, float64)"]
+        f = cuda.jit(sigs)(add_kernel)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, np.float32(1), np.float32(2**-25))
+        self.assertPreciseEqual(r[0], 1.0)
+
+        r = np.zeros(1, dtype=np.float64)
+        f[1, 1](r, 1, 2**-25)
+        self.assertPreciseEqual(r[0], 1.0000000298023224)
+
+        # Fail to resolve ambiguity between the two best overloads
+        # (Also deliberate float64[::1] for the first argument in all cases)
+        f = cuda.jit(["(float64[::1], float32,float64)",
+                      "(float64[::1], float64,float32)",
+                      "(float64[::1], int64,int64)"])(add_kernel)
+        with self.assertRaises(TypeError) as cm:
+            r = np.zeros(1, dtype=np.float64)
+            f[1, 1](r, 1.0, 2.0)
+
+        # The two best matches are output in the error message, as well
+        # as the actual argument types.
+        self.assertRegexpMatches(
+            str(cm.exception),
+            r"Ambiguous overloading for <function add_kernel [^>]*> "
+            r"\(array\(float64, 1d, C\), float64, float64\):\n"
+            r"\(array\(float64, 1d, C\), float32, float64\) -> none\n"
+            r"\(array\(float64, 1d, C\), float64, float32\) -> none"
+        )
+        # The integer signature is not part of the best matches
+        self.assertNotIn("int64", str(cm.exception))
+
+    @unittest.expectedFailure
+    def test_explicit_signatures_unsafe(self):
+        # These tests are from test_explicit_signatures, but have to be xfail
+        # at present because _prepare_args in the CUDA target cannot handle
+        # unsafe conversions of arguments.
+        f = cuda.jit("(int64[::1], int64,int64)")(add_kernel)
+        r = np.zeros(1, dtype=np.int64)
+
+        # Approximate match (unsafe conversion)
+        f[1, 1](r, 1.5, 2.5)
+        self.assertPreciseEqual(r[0], 3)
+        self.assertEqual(len(f.overloads), 1, f.overloads)
+
+        sigs = ["(int64[::1], int64,int64)", "(float64[::1], float64, float64)"]
path: numba/cuda/tests/cudapy/test_dispatcher.py
text: @user1:
```suggestion
        sigs = ["(int64[::1], int64, int64)", "(float64[::1], float64, float64)"]
```