sklam · Nov 2, 2021 · Oct 11, 2021 · Oct 13, 2021
diff --git a/docs/source/cuda/cudapysupported.rst b/docs/source/cuda/cudapysupported.rst
@@ -46,7 +46,13 @@ asynchronous operation - in order to ensure that all output is printed after a
 kernel launch, it is necessary to call :func:`numba.cuda.synchronize`. Eliding
 the call to ``synchronize`` is acceptable, but output from a kernel may appear
 during other later driver operations (e.g. subsequent kernel launches, memory
-transfers, etc.), or fail to appear before the program execution completes.
+transfers, etc.), or fail to appear before the program execution completes. Up
+to 32 arguments may be passed to the ``print`` function - if more are passed
+then a format string will be emitted instead and a warning will be produced.
+This is due to a general limitation in CUDA printing, as outlined in the
+`section on limitations in printing
+<https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#limitations>`_
+in the CUDA C++ Programming Guide.
 
 Built-in types
 ===============

diff --git a/numba/cuda/printimpl.py b/numba/cuda/printimpl.py
@@ -1,8 +1,10 @@
 from functools import singledispatch
 from llvmlite import ir
 from numba.core import types, cgutils
+from numba.core.errors import NumbaWarning
 from numba.core.imputils import Registry
 from numba.cuda import nvvmutils
+from warnings import warn
 
 registry = Registry()
 lower = registry.lower
@@ -68,6 +70,12 @@ def print_varargs(context, builder, sig, args):
         values.extend(argvals)
 
     rawfmt = " ".join(formats) + "\n"
+    if len(args) > 32:
+        msg = ('CUDA print() cannot print more than 32 items. '
+               'The raw format string will be emitted by the kernel instead.')
+        warn(msg, NumbaWarning)
+
+        rawfmt = rawfmt.replace('%', '%%')
     fmt = context.insert_string_const_addrspace(builder, rawfmt)
     array = cgutils.make_anonymous_struct(builder, values)
     arrayptr = cgutils.alloca_once_value(builder, array)

diff --git a/numba/cuda/tests/cudapy/test_print.py b/numba/cuda/tests/cudapy/test_print.py
@@ -1,6 +1,10 @@
 from numba import cuda
-from numba.cuda.testing import captured_cuda_stdout, CUDATestCase
+from numba.core.errors import NumbaWarning
+from numba.cuda.testing import (captured_cuda_stdout, CUDATestCase,
+                                skip_on_cudasim)
+import numpy as np
 import unittest
+import warnings
 
 
 def cuhello():
@@ -23,6 +27,13 @@ def printempty():
     print()
 
 
+def print_too_many(r):
+    print(r[0], r[1], r[2], r[3], r[4], r[5], r[6], r[7], r[8], r[9], r[10],
+          r[11], r[12], r[13], r[14], r[15], r[16], r[17], r[18], r[19], r[20],
+          r[21], r[22], r[23], r[24], r[25], r[26], r[27], r[28], r[29], r[30],
+          r[31], r[32])
+
+
 class TestPrint(CUDATestCase):
 
     def test_cuhello(self):
@@ -59,6 +70,38 @@ def test_string(self):
         expected = ['%d hop! 999\n' % i for i in range(3)]
         self.assertEqual(lines, expected)
 
+    @skip_on_cudasim('cudasim can print unlimited output')
+    def test_too_many_args(self):
+        # Tests that we emit the format string and warn when there are more
+        # than 32 arguments, in common with CUDA C/C++ printf - this is due to
+        # a limitation in CUDA vprintf, see:
+        #    https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#limitations
+
+        cufunc = cuda.jit(print_too_many)
+        r = np.arange(33)
+        with captured_cuda_stdout() as stdout:
+            with warnings.catch_warnings(record=True) as w:
+                cufunc[1, 1](r)
+
+        # Check that the format string was printed instead of formatted garbage
+        expected_fmt_string = ' '.join(['%lld' for _ in range(33)])
+        self.assertIn(expected_fmt_string, stdout.getvalue())
+
+        # Check for the expected warning about formatting more than 32 items
+        found = False
+
+        for warning in w:
+            warnobj = warning.message
+            if isinstance(warnobj, NumbaWarning):
+                expected = ('CUDA print() cannot print more than 32 items. '
+                            'The raw format string will be emitted by the '
+                            'kernel instead.')
+                if warnobj.msg == expected:
+                    found = True
+
+        if not found:
+            self.fail('Expected a warning for printing more than 32 items')
+
 
 if __name__ == '__main__':
     unittest.main()