CUDA: Fix linkage of device functions when compiling for debug

[gmarkall]

Fixes #7159.

[gmarkall]

@esc @stuartarchibald Is this possible for a CUDA buildfarm run to see if it turns up any issues on any of the other platforms please?

[stuartarchibald]

@esc @stuartarchibald Is this possible for a CUDA buildfarm run to see if it turns up any issues on any of the other platforms please?

Buildfarm ID: numba_smoketest_cuda_yaml_80.

[stuartarchibald]

@esc @stuartarchibald Is this possible for a CUDA buildfarm run to see if it turns up any issues on any of the other platforms please?

Buildfarm ID: numba_smoketest_cuda_yaml_80.

Failed on CUDA < 11.2 with Error: Debugging support cannot be enabled when named metadata !llvm.dbg.cu is not defined.

[sklam]

I can get the reproducer in the original issue to work If I turn off the debug option when !llvm.dbg.cu is not in the LLVM IR like below:

diff --git a/numba/cuda/codegen.py b/numba/cuda/codegen.py
index 70dcb10d6..5cfdaf6f1 100644
--- a/numba/cuda/codegen.py
+++ b/numba/cuda/codegen.py
@@ -125,7 +125,7 @@ class CUDACodeLibrary(serialize.ReduceMixin, CodeLibrary):
         else:
             # Otherwise, we compile all modules with NVVM at once because this
             # results in better optimization than separate compilation.
-            ptxes = [nvvm.llvm_to_ptx(irs, **options)]
+            ptxes = [nvvm.llvm_to_ptx(ir, **options) for ir in irs]
 
         # Sometimes the result from NVVM contains trailing whitespace and
         # nulls, which we strip so that the assembly dump looks a little
diff --git a/numba/cuda/cudadrv/nvvm.py b/numba/cuda/cudadrv/nvvm.py
index eb631b703..963022c5c 100644
--- a/numba/cuda/cudadrv/nvvm.py
+++ b/numba/cuda/cudadrv/nvvm.py
@@ -706,6 +706,9 @@ def llvm_to_ptx(llvmir, **opts):
     for mod in llvmir:
         mod = llvm_replace(mod)
         cu.add_module(mod.encode('utf8'))
+
+        if "!llvm.dbg.cu" not in mod:
+            opts.pop("debug", None)
     cu.lazy_add_module(libdevice.get())
 
     return cu.compile(**opts)

however, the test will still fail with:

numba.cuda.cudadrv.driver.LinkerError: [999] Call to cuLinkComplete results in CUDA_ERROR_UNKNOWN
error   : Undefined reference to '_ZN5numba4cuda5tests6cudapy14test_debuginfo17TestCudaDebugInfo29_test_chained_device_function12$3clocals$3e11kernel$2422Eii' in '<cudapy-ptx>'

[gmarkall]

@esc @sklam could this have another buildfarm run please?

[sklam]

BFID: numba_smoketest_cuda_yaml_81

[sklam]

BF passed!

[sklam]

The PR looks good to me.

This PR will disable debuginfo on NVVM version < 7.0 and it is acceptable because the debug support is not ideal in the old (2014) version. We shouldn't spend too much time fixing for the old toolchain given that the CUDA community is likely to be on the latest CTK with the newer NVVM, for which Numba generates better debuginfo now.

[stuartarchibald]

From OOB discussion, this message is going to be updated to include the function name so it's easy for users to find the function(s) causing this warning.

[stuartarchibald]

Thanks for the patch, couple of minor comments else looks good.

[gmarkall]

@stuartarchibald Thanks for the review - I believe comments are addressed / answered now.

[stuartarchibald]

Thanks for the patch and fixes.

[stuartarchibald]

NOTE: I checked out the patch and manually tested the warnings etc in local code on a CTK=11.0. Also checked that the tests are fine under CTK=11.0 and CTK=11.2.

[stuartarchibald]

Buildfarm ID: numba_smoketest_cuda_yaml_82 ran successfully on 68c4e76.