raise warning for inputs that has non-LTOIR objects

Author: isVoid

numba_cuda/numba/cuda/codegen.py

@@ -8,7 +8,7 @@
 import os
 import subprocess
 import tempfile
-
+from warnings import warn
 
 CUDA_TRIPLE = 'nvptx64-nvidia-cuda'
 
@@ -207,11 +207,20 @@ def get_cubin(self, cc=None):
                 lto=self._lto
             )
             self._link_all(linker, cc)
-            ptx = linker.get_linked_ptx().decode('utf-8')
 
-            print(("ASSEMBLY (AFTER LTO) %s" % self._name).center(80, '-'))
-            print(ptx)
-            print('=' * 80)
+            try:
+                ptx = linker.get_linked_ptx().decode('utf-8')
+
+                print(("ASSEMBLY (AFTER LTO) %s" % self._name).center(80, '-'))
+                print(ptx)
+                print('=' * 80)
+            except driver.LinkerError as e:
+                if linkererr_cause := getattr(e, "__cause__", None):
+                    if "-ptx requires that all inputs have LTOIR" in str(linkererr_cause):
+                        warn(
+                            "Linker input contains non-LTOIR objects, nvjitlink "
+                            "cannot generate LTO-ed PTX."
+                        )
 
         linker = driver.Linker.new(
             max_registers=self._max_registers,

numba_cuda/numba/cuda/tests/cudadrv/test_nvjitlink.py

@@ -7,6 +7,7 @@
 import os
 import io
 import contextlib
+import warnings
 
 from numba.cuda import get_current_device
 from numba import cuda
@@ -201,6 +202,38 @@ def kernel(result):
         config.DUMP_ASSEMBLY = False
 
 
+    def test_nvjitlink_jit_with_linkable_code_lto_dump_assembly_warn(self):
+        files = [
+            test_device_functions_a,
+            test_device_functions_cubin,
+            test_device_functions_fatbin,
+            test_device_functions_o,
+            test_device_functions_ptx,
+        ]
+
+        config.DUMP_ASSEMBLY = True
+        
+        for file in files:
+            with self.subTest(file=file):
+                with warnings.catch_warnings(record=True) as w:
+                    with contextlib.redirect_stdout(None): # suppress other PTX
+                        sig = "uint32(uint32, uint32)"
+                        add_from_numba = cuda.declare_device("add_from_numba", sig)
+
+                        @cuda.jit(link=[file], lto=True)
+                        def kernel(result):
+                            result[0] = add_from_numba(1, 2)
+
+                        result = cuda.device_array(1)
+                        kernel[1, 1](result)
+                        assert result[0] == 3
+
+                assert len(w) == 1
+                self.assertIn("cannot generate LTO-ed PTX", str(w[0].message))
+
+
+        config.DUMP_ASSEMBLY = False
+
     def test_nvjitlink_jit_with_invalid_linkable_code(self):
         with open(test_device_functions_cubin, "rb") as f:
             content = f.read()