Merge pull request numba#9035 from Matt711/main

CUDA: Allow for debuginfo in nvdisasm output
Deezy33-creator · Jun 29, 2023 · 249c8ff · 249c8ff
2 parents b2578ed + 9875576
commit 249c8ff
Show file tree

Hide file tree

Showing 2 changed files with 9 additions and 3 deletions.
diff --git a/numba/cuda/codegen.py b/numba/cuda/codegen.py
@@ -24,7 +24,8 @@ def disassemble_cubin(cubin):
             f.write(cubin)
 
         try:
-            cp = subprocess.run(['nvdisasm', fname], check=True,
+            lineinfo_flag = '-gi'
+            cp = subprocess.run(['nvdisasm', lineinfo_flag, fname], check=True,
                                 stdout=subprocess.PIPE,
                                 stderr=subprocess.PIPE)
         except FileNotFoundError as e:

diff --git a/numba/cuda/tests/cudapy/test_inspect.py b/numba/cuda/tests/cudapy/test_inspect.py
@@ -2,6 +2,7 @@
 
 from io import StringIO
 from numba import cuda, float32, float64, int32, intp
+from numba.cuda.cudadrv.nvvm import NVVM
 from numba.cuda.testing import unittest, CUDATestCase
 from numba.cuda.testing import (skip_on_cudasim, skip_with_nvdisasm,
                                 skip_without_nvdisasm)
@@ -91,13 +92,17 @@ def foo(x, y):
         self.assertIn("foo", asmdict[float64, float64])
 
     def _test_inspect_sass(self, kernel, name, sass):
+        if not NVVM().is_nvvm70:
+            self.skipTest("lineinfo not generated for NVVM 3.4")
         # Ensure function appears in output
         seen_function = False
         for line in sass.split():
             if '.text' in line and name in line:
                 seen_function = True
         self.assertTrue(seen_function)
 
+        self.assertRegex(sass, r'//## File ".*/test_inspect.py", line [0-9]')
+
         # Some instructions common to all supported architectures that should
         # appear in the output
         self.assertIn('S2R', sass)   # Special register to register
@@ -108,7 +113,7 @@ def _test_inspect_sass(self, kernel, name, sass):
     def test_inspect_sass_eager(self):
         sig = (float32[::1], int32[::1])
 
-        @cuda.jit(sig)
+        @cuda.jit(sig, lineinfo=True)
         def add(x, y):
             i = cuda.grid(1)
             if i < len(x):
@@ -118,7 +123,7 @@ def add(x, y):
 
     @skip_without_nvdisasm('nvdisasm needed for inspect_sass()')
     def test_inspect_sass_lazy(self):
-        @cuda.jit
+        @cuda.jit(lineinfo=True)
         def add(x, y):
             i = cuda.grid(1)
             if i < len(x):