Repository: numba/numba
Pull request: CUDA: Fix linkage of device functions when compiling for debug
Fixes #7159.

Problem statement:
CUDA regression: Mixing `debug` kwargs between called functions results in linkage error
The following is fine with Numba 0.53.1:

```python
from numba import cuda

@cuda.jit(device=True, debug=True)
def f2(x):
    return x + 1

@cuda.jit(device=True)
def f1(x, y):
    return x - f2(y)

@cuda.jit(debug=True)
def kernel(x, y):
    f1(x, y)

kernel[1, 1](1, 2)
```

However with current master, linking fails:

```
$ python debuglinkage.py 
/home/gmarkall/numbadev/numba/numba/cuda/compiler.py:865: NumbaPerformanceWarning: Grid size (1) < 2 * SM count (144) will likely result in GPU under utilization due to low occupancy.
  warn(NumbaPerformanceWarning(msg))
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 2221, in complete
    driver.cuLinkComplete(self.handle, byref(cubin), byref(size))
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 302, in safe_cuda_api_call
    self._check_error(fname, retcode)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 342, in _check_error
    raise CudaAPIError(retcode, msg)
numba.cuda.cudadrv.driver.CudaAPIError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
  File "/home/gmarkall/numbadev/issues/cuda-extension-bug/debuglinkage.py", line 19, in <module>
    kernel[1, 1](1, 2)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 868, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1001, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1009, in _compile_for_args
    return self.compile(tuple(argtypes))
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 1115, in compile
    kernel.bind()
  File "/home/gmarkall/numbadev/numba/numba/cuda/compiler.py", line 626, in bind
    self._codelibrary.get_cufunc()
  File "/home/gmarkall/numbadev/numba/numba/cuda/codegen.py", line 188, in get_cufunc
    cubin = self.get_cubin(cc=device.compute_capability)
  File "/home/gmarkall/numbadev/numba/numba/cuda/codegen.py", line 165, in get_cubin
    cubin_buf, size = linker.complete()
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 2223, in complete
    raise LinkerError("%s\n%s" % (e, self.error_log))
numba.cuda.cudadrv.driver.LinkerError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN
error   : Undefined reference to '_ZN8__main__6f2$243Ex' in '<cudapy-ptx>'
```

Hints from issue discussion:

Commit message: Give name of debug=True function in warnings

We now emit one warning per function with debug=True when compiling on
NVVM 3.4, with the warning message specifying the function name to help
the user see the source of the warning.

Patch to review:
diff --git a/numba/cuda/codegen.py b/numba/cuda/codegen.py
index 8525f915c27..f89b8fe69b5 100644
--- a/numba/cuda/codegen.py
+++ b/numba/cuda/codegen.py
@@ -1,8 +1,10 @@
 from llvmlite import binding as ll
 from llvmlite import ir
+from warnings import warn
 
 from numba.core import config, serialize
 from numba.core.codegen import Codegen, CodeLibrary
+from numba.core.errors import NumbaInvalidConfigWarning
 from .cudadrv import devices, driver, nvvm
 
 import ctypes
@@ -113,6 +115,17 @@ def _get_ptxes(self, cc=None):
         arch = nvvm.get_arch_option(*cc)
         options = self._nvvm_options.copy()
         options['arch'] = arch
+        if not nvvm.NVVM().is_nvvm70:
+            # Avoid enabling debug for NVVM 3.4 as it has various issues. We
+            # need to warn the user that we're doing this if any of the
+            # functions that they're compiling have `debug=True` set, which we
+            # can determine by checking the NVVM options.
+            for lib in self.linking_libraries:
+                if lib._nvvm_options.get('debug'):
+                    msg = ("debuginfo is not generated for CUDA versions "
+                           f"< 11.2 (debug=True on {lib.name})")
+                    warn(NumbaInvalidConfigWarning(msg))
+            options['debug'] = False
 
         irs = [str(mod) for mod in self.modules]
 
@@ -235,6 +248,17 @@ def modules(self):
         return [self._module] + [mod for lib in self._linking_libraries
                                  for mod in lib.modules]
 
+    @property
+    def linking_libraries(self):
+        # Libraries we link to may link to other libraries, so we recursively
+        # traverse the linking libraries property to build up a list of all
+        # linked libraries.
+        libs = []
+        for lib in self._linking_libraries:
+            libs.extend(lib.linking_libraries)
+            libs.append(lib)
+        return libs
+
     def finalize(self):
         # Unlike the CPUCodeLibrary, we don't invoke the binding layer here -
         # we only adjust the linkage of functions. Global kernels (with
@@ -258,10 +282,13 @@ def finalize(self):
         # because the device functions are in separate modules, we need them to
         # be externally visible.
         for library in self._linking_libraries:
-            for fn in library._module.functions:
-                if (not fn.is_declaration and
-                        not self._nvvm_options.get('debug', False)):
-                    fn.linkage = 'linkonce_odr'
+            for mod in library.modules:
+                for fn in mod.functions:
+                    if not fn.is_declaration:
+                        if self._nvvm_options.get('debug', False):
+                            fn.linkage = 'weak_odr'
+                        else:
+                            fn.linkage = 'linkonce_odr'
 
         self._finalized = True
 
diff --git a/numba/cuda/tests/cudapy/test_compiler.py b/numba/cuda/tests/cudapy/test_compiler.py
index 0d229a103d4..b8ea4a2a715 100644
--- a/numba/cuda/tests/cudapy/test_compiler.py
+++ b/numba/cuda/tests/cudapy/test_compiler.py
@@ -1,6 +1,7 @@
 from math import sqrt
 from numba import cuda, float32, uint32, void
 from numba.cuda import compile_ptx, compile_ptx_for_current_device
+from numba.cuda.cudadrv.nvvm import NVVM
 
 from numba.cuda.testing import skip_on_cudasim, unittest, CUDATestCase
 
@@ -64,6 +65,9 @@ def f(x, y, z, d):
         self.assertIn('sqrt.approx.ftz.f32', ptx)
 
     def check_debug_info(self, ptx):
+        if not NVVM().is_nvvm70:
+            self.skipTest('debuginfo not generated for NVVM 3.4')
+
         # A debug_info section should exist in the PTX. Whitespace varies
         # between CUDA toolkit versions.
         self.assertRegex(ptx, '\\.section\\s+\\.debug_info')
@@ -73,7 +77,12 @@ def check_debug_info(self, ptx):
         self.assertRegex(ptx, '\\.file.*test_compiler.py"')
 
     def test_device_function_with_debug(self):
-        # See Issue #6719
+        # See Issue #6719 - this ensures that compilation with debug succeeds
+        # with CUDA 11.2 / NVVM 7.0 onwards. Previously it failed because NVVM
+        # IR version metadata was not added when compiling device functions,
+        # and NVVM assumed DBG version 1.0 if not specified, which is
+        # incompatible with the 3.0 IR we use. This was specified only for
+        # kernels.
         def f():
             pass
 
diff --git a/numba/cuda/tests/cudapy/test_debuginfo.py b/numba/cuda/tests/cudapy/test_debuginfo.py
index f9cb0180200..338dd04f9b1 100644
--- a/numba/cuda/tests/cudapy/test_debuginfo.py
+++ b/numba/cuda/tests/cudapy/test_debuginfo.py
@@ -3,9 +3,12 @@
 from numba import cuda
 from numba.cuda.cudadrv.nvvm import NVVM
 from numba.core import types
+from numba.core.errors import NumbaInvalidConfigWarning
 from numba.cuda.testing import CUDATestCase
+import itertools
 import re
 import unittest
+import warnings
 
 
 @skip_on_cudasim('Simulator does not produce debug dumps')
@@ -18,6 +21,9 @@ def _getasm(self, fn, sig):
         return fn.inspect_asm(sig)
 
     def _check(self, fn, sig, expect):
+        if not NVVM().is_nvvm70:
+            self.skipTest("debuginfo not generated for NVVM 3.4")
+
         asm = self._getasm(fn, sig=sig)
         re_section_dbginfo = re.compile(r"\.section\s+\.debug_info\s+{")
         match = re_section_dbginfo.search(asm)
@@ -129,6 +135,140 @@ def kernel(arr):
             if i < len(arr):
                 arr[i] = threadid()
 
+    def _test_chained_device_function(self, kernel_debug, f1_debug, f2_debug):
+        @cuda.jit(device=True, debug=f2_debug)
+        def f2(x):
+            return x + 1
+
+        @cuda.jit(device=True, debug=f1_debug)
+        def f1(x, y):
+            return x - f2(y)
+
+        @cuda.jit((types.int32, types.int32), debug=kernel_debug)
+        def kernel(x, y):
+            f1(x, y)
+
+        kernel[1, 1](1, 2)
+
+    def test_chained_device_function(self):
+        # Calling a device function that calls another device function from a
+        # kernel with should succeed regardless of which jit decorators have
+        # debug=True. See Issue #7159.
+
+        debug_opts = itertools.product(*[(True, False)] * 3)
+
+        for kernel_debug, f1_debug, f2_debug in debug_opts:
+            with self.subTest(kernel_debug=kernel_debug,
+                              f1_debug=f1_debug,
+                              f2_debug=f2_debug):
+                self._test_chained_device_function(kernel_debug,
+                                                   f1_debug,
+                                                   f2_debug)
+
+    def _test_chained_device_function_two_calls(self, kernel_debug, f1_debug,
+                                                f2_debug):
+
+        @cuda.jit(device=True, debug=f2_debug)
+        def f2(x):
+            return x + 1
+
+        @cuda.jit(device=True, debug=f1_debug)
+        def f1(x, y):
+            return x - f2(y)
+
+        @cuda.jit(debug=kernel_debug)
+        def kernel(x, y):
+            f1(x, y)
+            f2(x)
+
+        kernel[1, 1](1, 2)
+
+    def test_chained_device_function_two_calls(self):
+        # Calling a device function that calls a leaf device function from a
+        # kernel, and calling the leaf device function from the kernel should
+        # succeed, regardless of which jit decorators have debug=True. See
+        # Issue #7159.
+
+        debug_opts = itertools.product(*[(True, False)] * 3)
+
+        for kernel_debug, f1_debug, f2_debug in debug_opts:
+            with self.subTest(kernel_debug=kernel_debug,
+                              f1_debug=f1_debug,
+                              f2_debug=f2_debug):
+                self._test_chained_device_function_two_calls(kernel_debug,
+                                                             f1_debug,
+                                                             f2_debug)
+
+    def check_warnings(self, warnings, warn_count):
+        if NVVM().is_nvvm70:
+            # We should not warn on NVVM 7.0.
+            self.assertEqual(len(warnings), 0)
+        else:
+            self.assertEqual(len(warnings), warn_count)
+            # Each warning should warn about not generating debug info.
+            for warning in warnings:
+                self.assertIs(warning.category, NumbaInvalidConfigWarning)
+                self.assertIn('debuginfo is not generated',
+                              str(warning.message))
+
+    def test_debug_warning(self):
+        # We don't generate debug info for NVVM 3.4, and warn accordingly. Here
+        # we check that no warnings appear with NVVM 7.0, and that warnings
+        # appear as appropriate with NVVM 3.4
+        debug_opts = itertools.product(*[(True, False)] * 3)
+
+        for kernel_debug, f1_debug, f2_debug in debug_opts:
+            with self.subTest(kernel_debug=kernel_debug,
+                              f1_debug=f1_debug,
+                              f2_debug=f2_debug):
+                with warnings.catch_warnings(record=True) as w:
+                    self._test_chained_device_function_two_calls(kernel_debug,
+                                                                 f1_debug,
+                                                                 f2_debug)
+                warn_count = kernel_debug + f1_debug + f2_debug
+                self.check_warnings(w, warn_count)
+
+    def test_chained_device_three_functions(self):
+        # Like test_chained_device_function, but with enough functions (three)
+        # to ensure that the recursion visits all the way down the call tree
+        # when fixing linkage of functions for debug.
+        def three_device_fns(kernel_debug, leaf_debug):
+            @cuda.jit(device=True, debug=leaf_debug)
+            def f3(x):
+                return x * x
+
+            @cuda.jit(device=True)
+            def f2(x):
+                return f3(x) + 1
+
+            @cuda.jit(device=True)
+            def f1(x, y):
+                return x - f2(y)
+
+            @cuda.jit(debug=kernel_debug)
+            def kernel(x, y):
+                f1(x, y)
+
+            kernel[1, 1](1, 2)
+
+        # Check when debug on the kernel, on the leaf, and not on any function.
+
+        with warnings.catch_warnings(record=True) as w:
+            three_device_fns(kernel_debug=True, leaf_debug=True)
+        self.check_warnings(w, 2)
+
+        with warnings.catch_warnings(record=True) as w:
+            three_device_fns(kernel_debug=True, leaf_debug=False)
+        self.check_warnings(w, 1)
+
+        with warnings.catch_warnings(record=True) as w:
+            three_device_fns(kernel_debug=False, leaf_debug=True)
+        self.check_warnings(w, 1)
+
+        with warnings.catch_warnings(record=True) as w:
+            three_device_fns(kernel_debug=False, leaf_debug=False)
+        self.check_warnings(w, 0)
+
 
 if __name__ == '__main__':
     unittest.main()


Code review comments:
diff_hunk: @@ -113,6 +115,17 @@ def _get_ptxes(self, cc=None):
         arch = nvvm.get_arch_option(*cc)
         options = self._nvvm_options.copy()
         options['arch'] = arch
+        if not nvvm.NVVM().is_nvvm70:
+            # Avoid enabling debug for NVVM 3.4 as it has various issues. We
+            # need to warn the user that we're doing this if any of the
+            # functions that they're compiling have `debug=True` set, which we
+            # can determine by checking the NVVM options.
+            for lib in self.linking_libraries:
+                if lib._nvvm_options.get('debug'):
+                    msg = ("debuginfo is not generated for CUDA versions "
+                           f"< 11.2 (debug=True on {lib.name})")
path: numba/cuda/codegen.py
text: @user1:
```suggestion
                           f"< 11.2 (debug=True on function: {lib.name})")
```