Repository: numba/numba

Issue:
<unnamed> (81, 19): parse expected comma after load's type
### Reporting a bug

Im getting an error like pasted below while trying to run the most basic example taken from official WWW:
https://developer.nvidia.com/blog/numba-python-cuda-acceleration/

```
import numpy as np
from numba import vectorize

@vectorize(['float32(float32, float32)'], target='cuda')
def Add(a, b):
  return a + b

N = 100000
A = np.ones(N, dtype=np.float32)
B = np.ones(A.shape, dtype=A.dtype)
C = np.empty_like(A, dtype=A.dtype)

C = Add(A, B)
```
Ubuntu 18.04.5, Numpy 1.19.4, Numba 0.52.0, kernel 5.4.58-050458-generic, CUDA 11.2.0, NVIDIA driver 460.27.04, GeForce GTX 1660

> Traceback (most recent call last):
>   File "GPUtest.py", line 13, in <module>
>     C = Add(A, B)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/dispatcher.py", line 37, in __call__
>     return CUDAUFuncMechanism.call(self.functions, args, kws)
>   File "/usr/local/lib/python3.6/dist-packages/numba/np/ufunc/deviceufunc.py", line 300, in call
>     cr.launch(func, shape[0], stream, devarys)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/dispatcher.py", line 148, in launch
>     func.forall(count, stream=stream)(*args)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 372, in __call__
>     kernel = self.kernel.specialize(*args)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 899, in specialize
>     self._bind, targetoptions)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 811, in __init__
>     self.compile(sigs[0])
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 952, in compile
>     kernel.bind()
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 576, in bind
>     self._func.get()
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 446, in get
>     ptx = self.ptx.get()
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/compiler.py", line 416, in get
>     **self._extra_options)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/cudadrv/nvvm.py", line 548, in llvm_to_ptx
>     ptx = cu.compile(**opts)
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/cudadrv/nvvm.py", line 236, in compile
>     self._try_error(err, 'Failed to compile\n')
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/cudadrv/nvvm.py", line 254, in _try_error
>     self.driver.check_error(err, "%s\n%s" % (msg, self.get_log()))
>   File "/usr/local/lib/python3.6/dist-packages/numba/cuda/cudadrv/nvvm.py", line 144, in check_error
>     raise exc
> numba.cuda.cudadrv.error.NvvmError: Failed to compile
> 
> <unnamed> (81, 19): parse expected comma after load's type
> NVVM_ERROR_COMPILATION

Is some default behavior changed and broke the decorator?

Discussion:

Patch:
diff --git a/numba/cuda/cudadrv/nvvm.py b/numba/cuda/cudadrv/nvvm.py
index 32085511f..a6f6f2af6 100644
--- a/numba/cuda/cudadrv/nvvm.py
+++ b/numba/cuda/cudadrv/nvvm.py
@@ -101,6 +101,11 @@ class NVVM(object):
 
         # nvvmResult nvvmGetProgramLog(nvvmProgram cu, char *buffer)
         'nvvmGetProgramLog': (nvvm_result, nvvm_program, c_char_p),
+
+        # nvvmResult nvvmIRVersion (int* majorIR, int* minorIR, int* majorDbg,
+        #                           int* minorDbg )
+        'nvvmIRVersion': (nvvm_result, POINTER(c_int), POINTER(c_int),
+                          POINTER(c_int), POINTER(c_int)),
     }
 
     # Singleton reference
@@ -127,6 +132,20 @@ class NVVM(object):
 
         return cls.__INSTANCE
 
+    def __init__(self):
+        ir_versions = self.get_ir_version()
+        self._majorIR = ir_versions[0]
+        self._minorIR = ir_versions[1]
+        self._majorDbg = ir_versions[2]
+        self._minorDbg = ir_versions[3]
+
+    @property
+    def is_nvvm70(self):
+        # NVVM70 uses NVVM IR version 1.6. See the documentation for
+        # nvvmAddModuleToProgram in
+        # https://docs.nvidia.com/cuda/libnvvm-api/group__compilation.html
+        return (self._majorIR, self._minorIR) >= (1, 6)
+
     def get_version(self):
         major = c_int()
         minor = c_int()
@@ -134,6 +153,16 @@ class NVVM(object):
         self.check_error(err, 'Failed to get version.')
         return major.value, minor.value
 
+    def get_ir_version(self):
+        majorIR = c_int()
+        minorIR = c_int()
+        majorDbg = c_int()
+        minorDbg = c_int()
+        err = self.nvvmIRVersion(byref(majorIR), byref(minorIR),
+                                 byref(majorDbg), byref(minorDbg))
+        self.check_error(err, 'Failed to get IR version.')
+        return majorIR.value, minorIR.value, majorDbg.value, minorDbg.value
+
     def check_error(self, error, msg, exit=False):
         if error:
             exc = NvvmError(msg, RESULT_CODE_NAMES[error])
@@ -399,15 +428,26 @@ class LibDevice(object):
         return self.bc
 
 
-ir_numba_atomic_cas = """
+ir_numba_cas_hack = """
 define internal {T} @___numba_atomic_{T}_cas_hack({T}* %ptr, {T} %cmp, {T} %val) alwaysinline {{
     %out = cmpxchg volatile {T}* %ptr, {T} %cmp, {T} %val monotonic
     ret {T} %out
 }}
 """ # noqa: E501
 
+cas_nvvm70 = """
+    %cas_success = cmpxchg volatile {Ti}* %iptr, {Ti} %old, {Ti} %new monotonic monotonic
+    %cas = extractvalue {{ {Ti}, i1 }} %cas_success, 0
+""" # noqa: E501
+
+
+cas_nvvm34 = """
+    %cas = cmpxchg volatile {Ti}* %iptr, {Ti} %old, {Ti} %new monotonic
+""" # noqa: E501
+
+
 # Translation of code from CUDA Programming Guide v6.5, section B.12
-ir_numba_atomic_binary = """
+ir_numba_atomic_binary_template = """
 define internal {T} @___numba_atomic_{T}_{FUNC}({T}* %ptr, {T} %val) alwaysinline {{
 entry:
     %iptr = bitcast {T}* %ptr to {Ti}*
@@ -419,7 +459,7 @@ attempt:
     %dold = bitcast {Ti} %old to {T}
     %dnew = {OP} {T} %dold, %val
     %new = bitcast {T} %dnew to {Ti}
-    %cas = cmpxchg volatile {Ti}* %iptr, {Ti} %old, {Ti} %new monotonic
+    {CAS}
     %repeat = icmp ne {Ti} %cas, %old
     br i1 %repeat, label %attempt, label %done
 
@@ -429,10 +469,10 @@ done:
 }}
 """ # noqa: E501
 
-ir_numba_atomic_inc = """
-define internal {T} @___numba_atomic_{Tu}_inc({T}* %ptr, {T} %val) alwaysinline {{
+ir_numba_atomic_inc_template = """
+define internal {T} @___numba_atomic_{Tu}_inc({T}* %iptr, {T} %val) alwaysinline {{
 entry:
-    %old2 = load volatile {T}, {T}* %ptr
+    %old2 = load volatile {T}, {T}* %iptr
     br label %attempt
 
 attempt:
@@ -440,7 +480,7 @@ attempt:
     %bndchk = icmp ult {T} %old, %val
     %inc = add {T} %old, 1
     %new = select i1 %bndchk, {T} %inc, {T} 0
-    %cas = cmpxchg volatile {T}* %ptr, {T} %old, {T} %new monotonic
+    {CAS}
     %repeat = icmp ne {T} %cas, %old
     br i1 %repeat, label %attempt, label %done
 
@@ -449,10 +489,10 @@ done:
 }}
 """ # noqa: E501
 
-ir_numba_atomic_dec = """
-define internal {T} @___numba_atomic_{Tu}_dec({T}* %ptr, {T} %val) alwaysinline {{
+ir_numba_atomic_dec_template = """
+define internal {T} @___numba_atomic_{Tu}_dec({T}* %iptr, {T} %val) alwaysinline {{
 entry:
-    %old2 = load volatile {T}, {T}* %ptr
+    %old2 = load volatile {T}, {T}* %iptr
     br label %attempt
 
 attempt:
@@ -460,7 +500,7 @@ attempt:
     %dec = add {T} %old, -1
     %bndchk = icmp ult {T} %dec, %val
     %new = select i1 %bndchk, {T} %dec, {T} %val
-    %cas = cmpxchg volatile {T}* %ptr, {T} %old, {T} %new monotonic
+    {CAS}
     %repeat = icmp ne {T} %cas, %old
     br i1 %repeat, label %attempt, label %done
 
@@ -469,7 +509,7 @@ done:
 }}
 """ # noqa: E501
 
-ir_numba_atomic_minmax = """
+ir_numba_atomic_minmax_template = """
 define internal {T} @___numba_atomic_{T}_{NAN}{FUNC}({T}* %ptr, {T} %val) alwaysinline {{
 entry:
     %ptrval = load volatile {T}, {T}* %ptr
@@ -488,10 +528,10 @@ lt_check:
 
 attempt:
     ; Attempt to swap in the value
-    %iold = bitcast {T} %dold to {Ti}
+    %old = bitcast {T} %dold to {Ti}
     %iptr = bitcast {T}* %ptr to {Ti}*
-    %ival = bitcast {T} %val to {Ti}
-    %cas = cmpxchg volatile {Ti}* %iptr, {Ti} %iold, {Ti} %ival monotonic
+    %new = bitcast {T} %val to {Ti}
+    {CAS}
     %dcas = bitcast {Ti} %cas to {T}
     br label %lt_check
 
@@ -501,6 +541,33 @@ done:
 """ # noqa: E501
 
 
+def ir_cas(Ti):
+    if NVVM().is_nvvm70:
+        return cas_nvvm70.format(Ti=Ti)
+    else:
+        return cas_nvvm34.format(Ti=Ti)
+
+
+def ir_numba_atomic_binary(T, Ti, OP, FUNC):
+    params = dict(T=T, Ti=Ti, OP=OP, FUNC=FUNC, CAS=ir_cas(Ti))
+    return ir_numba_atomic_binary_template.format(**params)
+
+
+def ir_numba_atomic_minmax(T, Ti, NAN, OP, PTR_OR_VAL, FUNC):
+    params = dict(T=T, Ti=Ti, NAN=NAN, OP=OP, PTR_OR_VAL=PTR_OR_VAL,
+                  FUNC=FUNC, CAS=ir_cas(Ti))
+
+    return ir_numba_atomic_minmax_template.format(**params)
+
+
+def ir_numba_atomic_inc(T, Tu):
+    return ir_numba_atomic_inc_template.format(T=T, Tu=Tu, CAS=ir_cas(T))
+
+
+def ir_numba_atomic_dec(T, Tu):
+    return ir_numba_atomic_dec_template.format(T=T, Tu=Tu, CAS=ir_cas(T))
+
+
 def _replace_datalayout(llvmir):
     """
     Find the line containing the datalayout and replace it
@@ -515,60 +582,57 @@ def _replace_datalayout(llvmir):
 
 
 def llvm_replace(llvmir):
-    #New LLVM generate a shorthand for datalayout that NVVM does not know
-    llvmir = _replace_datalayout(llvmir)
-
-    # Replace with our cmpxchg and atomic implementations because LLVM 3.5 has
-    # a new semantic for cmpxchg.
     replacements = [
         ('declare double @___numba_atomic_double_add(double*, double)',
-         ir_numba_atomic_binary.format(T='double', Ti='i64', OP='fadd',
-                                       FUNC='add')),
+         ir_numba_atomic_binary(T='double', Ti='i64', OP='fadd', FUNC='add')),
         ('declare float @___numba_atomic_float_sub(float*, float)',
-         ir_numba_atomic_binary.format(T='float', Ti='i32', OP='fsub',
-                                       FUNC='sub')),
+         ir_numba_atomic_binary(T='float', Ti='i32', OP='fsub', FUNC='sub')),
         ('declare double @___numba_atomic_double_sub(double*, double)',
-         ir_numba_atomic_binary.format(T='double', Ti='i64', OP='fsub',
-                                       FUNC='sub')),
+         ir_numba_atomic_binary(T='double', Ti='i64', OP='fsub', FUNC='sub')),
         ('declare i64 @___numba_atomic_u64_inc(i64*, i64)',
-         ir_numba_atomic_inc.format(T='i64', Tu='u64')),
+         ir_numba_atomic_inc(T='i64', Tu='u64')),
         ('declare i64 @___numba_atomic_u64_dec(i64*, i64)',
-         ir_numba_atomic_dec.format(T='i64', Tu='u64')),
-        ('declare i32 @___numba_atomic_i32_cas_hack(i32*, i32, i32)',
-         ir_numba_atomic_cas.format(T='i32')),
-        ('declare i64 @___numba_atomic_i64_cas_hack(i64*, i64, i64)',
-         ir_numba_atomic_cas.format(T='i64')),
+         ir_numba_atomic_dec(T='i64', Tu='u64')),
         ('declare float @___numba_atomic_float_max(float*, float)',
-         ir_numba_atomic_minmax.format(T='float', Ti='i32', NAN='',
-                                       OP='nnan olt', PTR_OR_VAL='ptr',
-                                       FUNC='max')),
+         ir_numba_atomic_minmax(T='float', Ti='i32', NAN='', OP='nnan olt',
+                                PTR_OR_VAL='ptr', FUNC='max')),
         ('declare double @___numba_atomic_double_max(double*, double)',
-         ir_numba_atomic_minmax.format(T='double', Ti='i64', NAN='',
-                                       OP='nnan olt', PTR_OR_VAL='ptr',
-                                       FUNC='max')),
+         ir_numba_atomic_minmax(T='double', Ti='i64', NAN='', OP='nnan olt',
+                                PTR_OR_VAL='ptr', FUNC='max')),
         ('declare float @___numba_atomic_float_min(float*, float)',
-         ir_numba_atomic_minmax.format(T='float', Ti='i32', NAN='',
-                                       OP='nnan ogt', PTR_OR_VAL='ptr',
-                                       FUNC='min')),
+         ir_numba_atomic_minmax(T='float', Ti='i32', NAN='', OP='nnan ogt',
+                                PTR_OR_VAL='ptr', FUNC='min')),
         ('declare double @___numba_atomic_double_min(double*, double)',
-         ir_numba_atomic_minmax.format(T='double', Ti='i64', NAN='',
-                                       OP='nnan ogt', PTR_OR_VAL='ptr',
-                                       FUNC='min')),
+         ir_numba_atomic_minmax(T='double', Ti='i64', NAN='', OP='nnan ogt',
+                                PTR_OR_VAL='ptr', FUNC='min')),
         ('declare float @___numba_atomic_float_nanmax(float*, float)',
-         ir_numba_atomic_minmax.format(T='float', Ti='i32', NAN='nan',
-                                       OP='ult', PTR_OR_VAL='', FUNC='max')),
+         ir_numba_atomic_minmax(T='float', Ti='i32', NAN='nan', OP='ult',
+                                PTR_OR_VAL='', FUNC='max')),
         ('declare double @___numba_atomic_double_nanmax(double*, double)',
-         ir_numba_atomic_minmax.format(T='double', Ti='i64', NAN='nan',
-                                       OP='ult', PTR_OR_VAL='', FUNC='max')),
+         ir_numba_atomic_minmax(T='double', Ti='i64', NAN='nan', OP='ult',
+                                PTR_OR_VAL='', FUNC='max')),
         ('declare float @___numba_atomic_float_nanmin(float*, float)',
-         ir_numba_atomic_minmax.format(T='float', Ti='i32', NAN='nan',
-                                       OP='ugt', PTR_OR_VAL='', FUNC='min')),
+         ir_numba_atomic_minmax(T='float', Ti='i32', NAN='nan', OP='ugt',
+                                PTR_OR_VAL='', FUNC='min')),
         ('declare double @___numba_atomic_double_nanmin(double*, double)',
-         ir_numba_atomic_minmax.format(T='double', Ti='i64', NAN='nan',
-                                       OP='ugt', PTR_OR_VAL='', FUNC='min')),
+         ir_numba_atomic_minmax(T='double', Ti='i64', NAN='nan', OP='ugt',
+                                PTR_OR_VAL='', FUNC='min')),
         ('immarg', '')
     ]
 
+    if not NVVM().is_nvvm70:
+        # Replace with our cmpxchg implementation because LLVM 3.5 has a new
+        # semantic for cmpxchg.
+        replacements += [
+            ('declare i32 @___numba_atomic_i32_cas_hack(i32*, i32, i32)',
+             ir_numba_cas_hack.format(T='i32')),
+            ('declare i64 @___numba_atomic_i64_cas_hack(i64*, i64, i64)',
+             ir_numba_cas_hack.format(T='i64'))
+        ]
+        # Newer LLVMs generate a shorthand for datalayout that NVVM34 does not
+        # know
+        llvmir = _replace_datalayout(llvmir)
+
     for decl, fn in replacements:
         llvmir = llvmir.replace(decl, fn)
 
@@ -578,7 +642,10 @@ def llvm_replace(llvmir):
     # pass to NVVM.
     llvmir = llvmir.replace('llvm.numba_nvvm.atomic', 'llvm.nvvm.atomic')
 
-    llvmir = llvm39_to_34_ir(llvmir)
+    if NVVM().is_nvvm70:
+        llvmir = llvm100_to_70_ir(llvmir)
+    else:
+        llvmir = llvm100_to_34_ir(llvmir)
 
     return llvmir
 
@@ -650,10 +717,51 @@ re_unsupported_keywords = re.compile(r"\b(local_unnamed_addr|writeonly)\b")
 
 re_parenthesized_list = re.compile(r"\((.*)\)")
 
+re_spflags = re.compile(r"spFlags: (.*),")
+
+spflagmap = {
+    'DISPFlagDefinition': 'isDefinition',
+    'DISPFlagOptimized': 'isOptimized',
+}
+
 
-def llvm39_to_34_ir(ir):
+def llvm100_to_70_ir(ir):
     """
-    Convert LLVM 3.9 IR for LLVM 3.4.
+    Convert LLVM 10.0 IR for LLVM 7.0.
+    """
+    buf = []
+    for line in ir.splitlines():
+        if line.startswith('attributes #'):
+            # Remove function attributes unsupported by LLVM 7.0
+            m = re_attributes_def.match(line)
+            attrs = m.group(1).split()
+            attrs = ' '.join(a for a in attrs if a != 'willreturn')
+            line = line.replace(m.group(1), attrs)
+
+        if '!DISubprogram' in line:
+            # Replace the DISPFlags (LLVM 10.0) with main subprogram DIFlags
+            # (LLVM 7.0). Example:
+            #
+            #     spflags: DISPFlagDefinition | DISPFlagOptimized
+            #
+            # becomes:
+            #
+            #     isDefinition: true, isOptimized: true
+            m = re_spflags.search(line)
+            flags = m.group(1).split(' | ')
+            new_flags = ", ".join([ '%s: true' % spflagmap[f] for f in flags ])
+            start_of_line = line[:m.span()[0]]
+            end_of_line = line[m.span()[1] - 1:]
+            line = start_of_line + new_flags + end_of_line
+
+        buf.append(line)
+
+    return '\n'.join(buf)
+
+
+def llvm100_to_34_ir(ir):
+    """
+    Convert LLVM 10.0 IR for LLVM 3.4.
     """
     def parse_out_leading_type(s):
         par_level = 0
@@ -796,9 +904,16 @@ def set_cuda_kernel(lfunc):
     nmd = m.get_or_insert_named_metadata('nvvm.annotations')
     nmd.add(md)
 
-    # set nvvm ir version
+    # Set NVVM IR version
     i32 = ir.IntType(32)
-    md_ver = m.add_metadata([i32(1), i32(2), i32(2), i32(0)])
+    if NVVM().is_nvvm70:
+        # NVVM IR 1.6, DWARF 3.0
+        ir_versions = [i32(1), i32(6), i32(3), i32(0)]
+    else:
+        # NVVM IR 1.1, DWARF 2.0
+        ir_versions = [i32(1), i32(2), i32(2), i32(0)]
+
+    md_ver = m.add_metadata(ir_versions)
     m.add_named_metadata('nvvmir.version', md_ver)
 
 
diff --git a/numba/cuda/cudaimpl.py b/numba/cuda/cudaimpl.py
index 6ae42281f..332e3f02b 100644
--- a/numba/cuda/cudaimpl.py
+++ b/numba/cuda/cudaimpl.py
@@ -774,9 +774,7 @@ def ptx_atomic_cas_tuple(context, builder, sig, args):
     if aryty.dtype in (cuda.cudadecl.integer_numba_types):
         lmod = builder.module
         bitwidth = aryty.dtype.bitwidth
-        return builder.call(nvvmutils.declare_atomic_cas_int(lmod,
-                                                             bitwidth),
-                            (ptr, old, val))
+        return nvvmutils.atomic_cmpxchg(builder, lmod, bitwidth, ptr, old, val)
     else:
         raise TypeError('Unimplemented atomic compare_and_swap '
                         'with %s array' % dtype)
diff --git a/numba/cuda/nvvmutils.py b/numba/cuda/nvvmutils.py
index ae7bedf5c..ba984067f 100644
--- a/numba/cuda/nvvmutils.py
+++ b/numba/cuda/nvvmutils.py
@@ -12,6 +12,15 @@ def declare_atomic_cas_int(lmod, isize):
                              lc.Type.int(isize)))
     return lmod.get_or_insert_function(fnty, fname)
 
+
+def atomic_cmpxchg(builder, lmod, isize, ptr, cmp, val):
+    if nvvm.NVVM().is_nvvm70:
+        out = builder.cmpxchg(ptr, cmp, val, 'monotonic', 'monotonic')
+        return builder.extract_value(out, 0)
+    else:
+        return builder.call(declare_atomic_cas_int(lmod, isize),
+                            (ptr, cmp, val))
+
 # For atomic intrinsics, "numba_nvvm" prevents LLVM 9 onwards auto-upgrading
 # them into atomicrmw instructions that are not recognized by NVVM. It is
 # replaced with "nvvm" in llvm_to_ptx later, after the module has been parsed
diff --git a/numba/cuda/target.py b/numba/cuda/target.py
index 73f9417e2..fcb11f751 100644
--- a/numba/cuda/target.py
+++ b/numba/cuda/target.py
@@ -61,7 +61,13 @@ VALID_CHARS = re.compile(r'[^a-z0-9]', re.I)
 class CUDATargetContext(BaseContext):
     implement_powi_as_math_call = True
     strict_alignment = True
-    DIBuilder = debuginfo.NvvmDIBuilder
+
+    @property
+    def DIBuilder(self):
+        if nvvm.NVVM().is_nvvm70:
+            return debuginfo.DIBuilder
+        else:
+            return debuginfo.NvvmDIBuilder
 
     @property
     def enable_boundscheck(self):
@@ -177,13 +183,18 @@ class CUDATargetContext(BaseContext):
                 # Use atomic cmpxchg to prevent rewriting the error status
                 # Only the first error is recorded
 
-                casfnty = lc.Type.function(old.type, [gv_exc.type, old.type,
-                                                      old.type])
-
-                cas_hack = "___numba_atomic_i32_cas_hack"
-                casfn = wrapper_module.add_function(casfnty, name=cas_hack)
-                xchg = builder.call(casfn, [gv_exc, old, status.code])
-                changed = builder.icmp(ICMP_EQ, xchg, old)
+                if nvvm.NVVM().is_nvvm70:
+                    xchg = builder.cmpxchg(gv_exc, old, status.code,
+                                           'monotonic', 'monotonic')
+                    changed = builder.extract_value(xchg, 1)
+                else:
+                    casfnty = lc.Type.function(old.type, [gv_exc.type, old.type,
+                                                          old.type])
+
+                    cas_hack = "___numba_atomic_i32_cas_hack"
+                    casfn = wrapper_module.add_function(casfnty, name=cas_hack)
+                    xchg = builder.call(casfn, [gv_exc, old, status.code])
+                    changed = builder.icmp(ICMP_EQ, xchg, old)
 
                 # If the xchange is successful, save the thread ID.
                 sreg = nvvmutils.SRegBuilder(builder)


Test patch:
diff --git a/numba/cuda/tests/cudadrv/test_ir_patch.py b/numba/cuda/tests/cudadrv/test_ir_patch.py
index a91045918..29ed43422 100644
--- a/numba/cuda/tests/cudadrv/test_ir_patch.py
+++ b/numba/cuda/tests/cudadrv/test_ir_patch.py
@@ -6,9 +6,9 @@ from numba.cuda.testing import skip_on_cudasim
 class TestIRPatch(unittest.TestCase):
     def patch(self, ir):
         # Import here to avoid error in CUDASIM
-        from numba.cuda.cudadrv.nvvm import llvm39_to_34_ir
+        from numba.cuda.cudadrv.nvvm import llvm100_to_34_ir
 
-        return llvm39_to_34_ir(ir)
+        return llvm100_to_34_ir(ir)
 
     def test_load_rewrite(self):
         text = "%myload = not really"
diff --git a/numba/cuda/tests/cudadrv/test_nvvm_driver.py b/numba/cuda/tests/cudadrv/test_nvvm_driver.py
index eee882faa..b6f614a31 100644
--- a/numba/cuda/tests/cudadrv/test_nvvm_driver.py
+++ b/numba/cuda/tests/cudadrv/test_nvvm_driver.py
@@ -4,7 +4,7 @@ from numba.cuda.cudadrv.nvvm import (llvm_to_ptx, set_cuda_kernel,
                                      get_supported_ccs)
 from ctypes import c_size_t, c_uint64, sizeof
 from numba.cuda.testing import unittest
-from numba.cuda.cudadrv.nvvm import LibDevice, NvvmError
+from numba.cuda.cudadrv.nvvm import LibDevice, NvvmError, NVVM
 from numba.cuda.testing import skip_on_cudasim
 
 is64bit = sizeof(c_size_t) == sizeof(c_uint64)
@@ -12,14 +12,16 @@ is64bit = sizeof(c_size_t) == sizeof(c_uint64)
 
 @skip_on_cudasim('NVVM Driver unsupported in the simulator')
 class TestNvvmDriver(unittest.TestCase):
-    def get_ptx(self):
-        if is64bit:
-            return gpu64
+    def get_nvvmir(self):
+        if NVVM().is_nvvm70:
+            metadata = metadata_nvvm70
         else:
-            return gpu32
+            metadata = metadata_nvvm34
+
+        return nvvmir_generic + metadata
 
     def test_nvvm_compile_simple(self):
-        nvvmir = self.get_ptx()
+        nvvmir = self.get_nvvmir()
         ptx = llvm_to_ptx(nvvmir).decode('utf8')
         self.assertTrue('simple' in ptx)
         self.assertTrue('ave' in ptx)
@@ -41,8 +43,8 @@ class TestNvvmDriver(unittest.TestCase):
             self.assertTrue('.address_size 32' in ptx)
 
     def _test_nvvm_support(self, arch):
-        nvvmir = self.get_ptx()
         compute_xx = 'compute_{0}{1}'.format(*arch)
+        nvvmir = self.get_nvvmir()
         ptx = llvm_to_ptx(nvvmir, arch=compute_xx, ftz=1, prec_sqrt=0,
                           prec_div=0).decode('utf8')
         self.assertIn(".target sm_{0}{1}".format(*arch), ptx)
@@ -99,7 +101,7 @@ class TestLibDevice(unittest.TestCase):
         self._libdevice_load('compute_52', 'compute_50')
 
 
-gpu64 = '''
+nvvmir_generic = '''\
 target triple="nvptx64-"
 target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
@@ -129,46 +131,23 @@ declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() nounwind readnone
 declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() nounwind readnone
 
 declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() nounwind readnone
-
-!nvvm.annotations = !{!1}
-!1 = metadata !{void (i32*)* @simple, metadata !"kernel", i32 1}
 '''  # noqa: E501
 
-gpu32 = '''
-target triple="nvptx-"
-target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
-define i32 @ave(i32 %a, i32 %b) {
-entry:
-%add = add nsw i32 %a, %b
-%div = sdiv i32 %add, 2
-ret i32 %div
-}
+metadata_nvvm70 = '''
+!nvvmir.version = !{!1}
+!1 = !{i32 1, i32 6, i32 3, i32 0}
 
-define void @simple(i32* %data) {
-entry:
-%0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
-%1 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-%mul = mul i32 %0, %1
-%2 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-%add = add i32 %mul, %2
-%call = call i32 @ave(i32 %add, i32 %add)
-%idxprom = sext i32 %add to i64
-%arrayidx = getelementptr inbounds i32, i32* %data, i64 %idxprom
-store i32 %call, i32* %arrayidx, align 4
-ret void
-}
-
-declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() nounwind readnone
-
-declare i32 @llvm.nvvm.read.ptx.sreg.ntid.x() nounwind readnone
+!nvvm.annotations = !{!2}
+!2 = !{void (i32*)* @simple, !"kernel", i32 1}
+'''  # noqa: E501
 
-declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() nounwind readnone
 
+metadata_nvvm34 = '''
 !nvvm.annotations = !{!1}
 !1 = metadata !{void (i32*)* @simple, metadata !"kernel", i32 1}
+'''
 
-'''  # noqa: E501
 
 if __name__ == '__main__':
     unittest.main()
diff --git a/numba/cuda/tests/cudadrv/test_runtime.py b/numba/cuda/tests/cudadrv/test_runtime.py
index 5a14ce399..967a4b625 100644
--- a/numba/cuda/tests/cudadrv/test_runtime.py
+++ b/numba/cuda/tests/cudadrv/test_runtime.py
@@ -23,7 +23,7 @@ class TestRuntime(unittest.TestCase):
             supported_versions = (-1, -1),
         else:
             supported_versions = ((9, 0), (9, 1), (9, 2), (10, 0),
-                                  (10, 1), (10, 2), (11, 0), (11, 1))
+                                  (10, 1), (10, 2), (11, 0), (11, 1), (11, 2))
         self.assertIn(runtime.get_version(), supported_versions)
 
 
diff --git a/numba/cuda/tests/cudapy/test_constmem.py b/numba/cuda/tests/cudapy/test_constmem.py
index ae3f00fbd..f32700aa0 100644
--- a/numba/cuda/tests/cudapy/test_constmem.py
+++ b/numba/cuda/tests/cudapy/test_constmem.py
@@ -134,7 +134,11 @@ class TestCudaConstantMemory(CUDATestCase):
         self.assertTrue(np.all(A == CONST3D))
 
         if not ENABLE_CUDASIM:
-            if cuda.runtime.get_version() in ((8, 0), (9, 0), (9, 1)):
+            # CUDA 9.2 - 11.1 use two f32 loads to load the complex. CUDA < 9.2
+            # and > 11.1 use a vector of 2x f32. The root cause of these
+            # codegen differences is not known, but must be accounted for in
+            # this test.
+            if cuda.runtime.get_version() in ((8, 0), (9, 0), (9, 1), (11, 2)):
                 complex_load = 'ld.const.v2.f32'
                 description = 'Load the complex as a vector of 2x f32'
             else:
@@ -164,12 +168,14 @@ class TestCudaConstantMemory(CUDATestCase):
         B = np.zeros(2, dtype=int)
         jcuconst = cuda.jit(cuconstRec).specialize(A, B)
 
-        old_runtime = cuda.runtime.get_version() in ((8, 0), (9, 0), (9, 1))
+        rtver = cuda.runtime.get_version()
+        old_runtime = rtver in ((9, 0), (9, 1))
+        nvvm70_runtime = rtver >= (11, 2)
         windows = sys.platform.startswith('win')
 
         if old_runtime:
             if windows:
-                # for some reason Win64 / Py3 / CUDA 9.1 decides to do two u32
+                # For some reason Win64 / CUDA 9.1 and 9.2 decide to do two u32
                 # loads, and shifts and ors the values to get the float `x`
                 # field, then uses another ld.const.u32 to load the int `y` as
                 # a 32-bit value!
@@ -179,6 +185,16 @@ class TestCudaConstantMemory(CUDATestCase):
                 # Load of the x and y fields fused into a single instruction
                 self.assertIn('ld.const.v2.f64', jcuconst.ptx,
                               'load record fields as vector of 2x f64')
+        elif nvvm70_runtime:
+            if windows:
+                # Two ld.const.u32 as above, but using a bit-field insert to
+                # combine them
+                self.assertIn('ld.const.u32', jcuconst.ptx,
+                              'load record fields as u32')
+            else:
+                # Load of the x and y fields fused into a single instruction
+                self.assertIn('ld.const.v2.u64', jcuconst.ptx,
+                              'load record fields as vector of 2x u64')
         else:
             # In newer toolkits, constant values are all loaded 8 bits at a
             # time. Check that there are enough 8-bit loads for everything to
@@ -224,8 +240,8 @@ class TestCudaConstantMemory(CUDATestCase):
         E = np.zeros(2, dtype=np.float64)
         jcuconst = cuda.jit(cuconstRecAlign).specialize(A, B, C, D, E)
 
-        if rtver >= (10, 2):
-            # Code generation differs slightly in 10.2 onwards - the first
+        if rtver >= (10, 2) and rtver <= (11, 1):
+            # Code generation differs slightly in 10.2 - 11.1 - the first
             # bytes are loaded as individual bytes, so we'll check that
             # ld.const.u8 occurs at least four times (the first three bytes,
             # then the last byte by itself)
diff --git a/numba/cuda/tests/cudapy/test_debuginfo.py b/numba/cuda/tests/cudapy/test_debuginfo.py
index 4300a29ce..aaea8a230 100644
--- a/numba/cuda/tests/cudapy/test_debuginfo.py
+++ b/numba/cuda/tests/cudapy/test_debuginfo.py
@@ -3,6 +3,7 @@ from numba.cuda.testing import skip_on_cudasim
 from numba import cuda
 from numba.core import types
 from numba.cuda.testing import CUDATestCase
+import re
 import unittest
 
 
@@ -17,8 +18,10 @@ class TestCudaDebugInfo(CUDATestCase):
 
     def _check(self, fn, sig, expect):
         asm = self._getasm(fn, sig=sig)
-        assertfn = self.assertIn if expect else self.assertNotIn
-        assertfn('.section .debug_info {', asm, msg=asm)
+        re_section_dbginfo = re.compile(r"\.section\s+\.debug_info\s+{")
+        match = re_section_dbginfo.search(asm)
+        assertfn = self.assertIsNotNone if expect else self.assertIsNone
+        assertfn(match, msg=asm)
 
     def test_no_debuginfo_in_asm(self):
         @cuda.jit(debug=False)
diff --git a/numba/cuda/tests/nocuda/test_nvvm.py b/numba/cuda/tests/nocuda/test_nvvm.py
index 2c7ba0e47..f3b8f3618 100644
--- a/numba/cuda/tests/nocuda/test_nvvm.py
+++ b/numba/cuda/tests/nocuda/test_nvvm.py
@@ -34,13 +34,13 @@ class TestNvvmWithoutCuda(SerialMixin, unittest.TestCase):
         ptx = nvvm.llvm_to_ptx(llvmir)
         self.assertIn("foo", ptx.decode('ascii'))
 
-    def test_nvvm_memset_fixup(self):
+    def test_nvvm_memset_fixup_for_34(self):
         """
         Test llvm.memset changes in llvm7.
         In LLVM7 the alignment parameter can be implicitly provided as
         an attribute to pointer in the first argument.
         """
-        fixed = nvvm.llvm39_to_34_ir(original)
+        fixed = nvvm.llvm100_to_34_ir(original)
         self.assertIn("call void @llvm.memset", fixed)
 
         for ln in fixed.splitlines():
@@ -51,13 +51,13 @@ class TestNvvmWithoutCuda(SerialMixin, unittest.TestCase):
                     r'i32 4, i1 false\)'.replace(' ', r'\s+'),
                 )
 
-    def test_nvvm_memset_fixup_missing_align(self):
+    def test_nvvm_memset_fixup_for_34_missing_align(self):
         """
         We require alignment to be specified as a parameter attribute to the
         dest argument of a memset.
         """
         with self.assertRaises(ValueError) as e:
-            nvvm.llvm39_to_34_ir(missing_align)
+            nvvm.llvm100_to_34_ir(missing_align)
 
         self.assertIn(str(e.exception),
                       "No alignment attribute found on memset dest")