sklam · Nov 19, 2021 · May 26, 2021 · Apr 12, 2021 · May 26, 2021 · Aug 11, 2021
diff --git a/docs/source/cuda/cudapysupported.rst b/docs/source/cuda/cudapysupported.rst
@@ -68,6 +68,9 @@ The following built-in types support are inherited from CPU nopython mode.
 
 See :ref:`nopython built-in types <pysupported-builtin-types>`.
 
+There is also some very limited support for character sequences (bytes and
+unicode strings) used in NumPy arrays. Note that this support can only be used
+with CUDA 11.2 onwards.
 
 Built-in functions
 ==================

diff --git a/numba/cuda/compiler.py b/numba/cuda/compiler.py
@@ -15,10 +15,11 @@
 from numba.core.compiler import (CompilerBase, DefaultPassBuilder,
                                  compile_result, Flags, Option)
 from numba.core.compiler_lock import global_compiler_lock
-from numba.core.compiler_machinery import (LoweringPass, PassManager,
-                                           register_pass)
+from numba.core.compiler_machinery import (LoweringPass, AnalysisPass,
+                                           PassManager, register_pass)
 from numba.core.dispatcher import CompilingCounter, OmittedArg
-from numba.core.errors import NumbaInvalidConfigWarning
+from numba.core.errors import (NumbaPerformanceWarning,
+                               NumbaInvalidConfigWarning, TypingError)
 from numba.core.typed_passes import IRLegalization, NativeLowering
 from numba.core.typing.typeof import Purpose, typeof
 from warnings import warn
@@ -29,7 +30,6 @@
 from .errors import missing_launch_config_msg, normalize_kernel_dimensions
 from .api import get_current_device
 from .args import wrap_arg
-from numba.core.errors import NumbaPerformanceWarning
 from .descriptor import cuda_target
 from . import types as cuda_types
 
@@ -103,6 +103,33 @@ def run_pass(self, state):
         return True
 
 
+@register_pass(mutates_CFG=False, analysis_only=True)
+class CUDALegalization(AnalysisPass):
+
+    _name = "cuda_legalization"
+
+    def __init__(self):
+        AnalysisPass.__init__(self)
+
+    def run_pass(self, state):
+        # Early return if NVVM 7
+        from numba.cuda.cudadrv.nvvm import NVVM
+        if NVVM().is_nvvm70:
+            return False
+        # NVVM < 7, need to check for charseq
+        typmap = state.typemap
+        for k, v in typmap.items():
+            if isinstance(v, types.Array):
+                if isinstance(v.dtype, (types.UnicodeCharSeq, types.CharSeq)):
+                    msg = (f"{k} is a char sequence type. This type is not "
+                           "supported with CUDA toolkit versions < 11.2. To "
+                           "use this type, you need to update your CUDA "
+                           "toolkit - try 'conda install cudatoolkit=11' if "
+                           "you are using conda to manage your environment.")
+                    raise TypingError(msg)
+        return False
+
+
 class CUDACompiler(CompilerBase):
     def define_pipelines(self):
         dpb = DefaultPassBuilder
@@ -113,6 +140,7 @@ def define_pipelines(self):
 
         typed_passes = dpb.define_typed_pipeline(self.state)
         pm.passes.extend(typed_passes.passes)
+        pm.add_pass(CUDALegalization, "CUDA legalization")
 
         lowering_passes = self.define_cuda_lowering_pipeline(self.state)
         pm.passes.extend(lowering_passes.passes)

diff --git a/numba/cuda/target.py b/numba/cuda/target.py
@@ -89,7 +89,8 @@ def load_additional_registries(self):
         # side effect of import needed for numba.cpython.*, the builtins
         # registry is updated at import time.
         from numba.cpython import numbers, tupleobj, slicing # noqa: F401
-        from numba.cpython import rangeobj, iterators, unicode # noqa: F401
+        from numba.cpython import rangeobj, iterators # noqa: F401
+        from numba.cpython import unicode, charseq # noqa: F401
         from numba.cpython import cmathimpl
         from numba.np import arrayobj # noqa: F401
         from numba.np import npdatetime # noqa: F401

diff --git a/numba/cuda/tests/cudapy/test_const_string.py b/numba/cuda/tests/cudapy/test_const_string.py
@@ -1,10 +1,14 @@
 import re
-from numba.cuda.testing import unittest, skip_on_cudasim
+import numpy as np
+from numba import config, cuda
+from numba.core.errors import TypingError
+from numba.cuda.cudadrv.nvvm import NVVM
+from numba.cuda.testing import unittest, skip_on_cudasim, CUDATestCase
 from llvmlite import ir
 
 
 @skip_on_cudasim("This is testing CUDA backend code generation")
-class TestCudaConstString(unittest.TestCase):
+class TestConstStringCodegen(unittest.TestCase):
     def test_const_string(self):
         # These imports are incompatible with CUDASIM
         from numba.cuda.descriptor import cuda_target
@@ -52,5 +56,82 @@ def test_const_string(self):
         self.assertEqual(len(matches), 1)
 
 
+# Inspired by the reproducer from Issue #7041.
+class TestConstString(CUDATestCase):
+
+    def skip_on_nvvm34(self):
+        # There is no NVVM on cudasim, but we might as well run this test.
+        if not config.ENABLE_CUDASIM and not NVVM().is_nvvm70:
+            self.skipTest('Character sequences unsupported on NVVM 3.4')
+
+    def skip_on_nvvm70(self):
+        if NVVM().is_nvvm70:
+            self.skipTest('Character sequences are permitted with NVVM 7.0')
+
+    def test_assign_const_unicode_string(self):
+        self.skip_on_nvvm34()
+
+        @cuda.jit
+        def str_assign(arr):
+            i = cuda.grid(1)
+            if i < len(arr):
+                arr[i] = "XYZ"
+
+        n_strings = 8
+        arr = np.zeros(n_strings + 1, dtype="<U12")
+        str_assign[1, n_strings](arr)
+
+        # Expected result, e.g.:
+        #     ['XYZ' 'XYZ' 'XYZ' 'XYZ' 'XYZ' 'XYZ' 'XYZ' 'XYZ' '']
+        expected = np.zeros_like(arr)
+        expected[:-1] = 'XYZ'
+        expected[-1] = ''
+        np.testing.assert_equal(arr, expected)
+
+    def test_assign_const_byte_string(self):
+        self.skip_on_nvvm34()
+
+        @cuda.jit
+        def bytes_assign(arr):
+            i = cuda.grid(1)
+            if i < len(arr):
+                arr[i] = b"XYZ"
+
+        n_strings = 8
+        arr = np.zeros(n_strings + 1, dtype="S12")
+        bytes_assign[1, n_strings](arr)
+
+        # Expected result, e.g.:
+        #     [b'XYZ' b'XYZ' b'XYZ' b'XYZ' b'XYZ' b'XYZ' b'XYZ' b'XYZ' b'']
+        expected = np.zeros_like(arr)
+        expected[:-1] = b'XYZ'
+        expected[-1] = b''
+        np.testing.assert_equal(arr, expected)
+
+    @skip_on_cudasim('Const strings not disallowed on cudasim')
+    def test_const_unicode_string_disallowed_nvvm34(self):
+        self.skip_on_nvvm70()
+
+        @cuda.jit
+        def str_assign(arr):
+            arr[0] = "XYZ"
+
+        arr = np.zeros(1, dtype="<U12")
+        with self.assertRaisesRegex(TypingError, '.*is a char sequence.*'):
+            str_assign[1, 1](arr)
+
+    @skip_on_cudasim('Const strings not disallowed on cudasim')
+    def test_const_byte_string_disallowed_nvvm34(self):
+        self.skip_on_nvvm70()
+
+        @cuda.jit
+        def bytes_assign(arr):
+            arr[0] = b"XYZ"
+
+        arr = np.zeros(1, dtype=b"S12")
+        with self.assertRaisesRegex(TypingError, '.*is a char sequence.*'):
+            bytes_assign[1, 1](arr)
+
+
 if __name__ == '__main__':
     unittest.main()