Use dbg.declare for scalar kernel parameters (#828)

## Summary
- Describe scalar kernel parameters with `dbg.declare` on their stack
slots to avoid unstable parameter-space locations in downstream DWARF
(boolean parameters remain on `dbg.value` to avoid an NVVM crash).
- Avoid emitting additional `dbg.value` entries for function arguments
once declared.
- Add a regression test for scalar formal parameters.

Author: cpcloud

numba_cuda/numba/cuda/debuginfo.py

@@ -866,8 +866,27 @@ def mark_variable(
             int_type = (ir.IntType,)
             real_type = ir.FloatType, ir.DoubleType
             if isinstance(lltype, int_type + real_type):
-                # Start with scalar variable, swtiching llvm.dbg.declare
-                # to llvm.dbg.value
+                # For scalar locals we use llvm.dbg.value instead of
+                # llvm.dbg.declare, but for scalar *arguments* we still want a
+                # stable stack location so they don't get encoded as absolute
+                # parameter-space addresses in downstream DWARF.
+                if argidx is not None:
+                    # NVVM has been observed to crash on some boolean-parameter
+                    # debug.declare patterns - use dbg.value for these instead.
+                    if datamodel is not None and isinstance(
+                        datamodel.fe_type, types.Boolean
+                    ):
+                        return
+                    return super().mark_variable(
+                        builder,
+                        allocavalue,
+                        name,
+                        lltype,
+                        size,
+                        line,
+                        datamodel,
+                        argidx,
+                    )
                 return
             else:
                 return super().mark_variable(

numba_cuda/numba/cuda/lowering.py

@@ -1722,6 +1722,12 @@ def loadvar(self, name):
                 src_name not in self.dbg_val_names
                 and src_name not in self.poly_var_typ_map
             ):
+                # Function arguments are declared once in the prologue. Emitting
+                # additional dbg.value entries for them can cause ptxas to
+                # describe them via parameter-space addresses (DW_OP_addr),
+                # which is both unstable and confusing for debuggers.
+                if src_name in self.fndesc.args:
+                    return val
                 fetype = self.typeof(name)
                 lltype = self.context.get_value_type(fetype)
                 int_type = (llvm_ir.IntType,)
@@ -1730,10 +1736,6 @@ def loadvar(self, name):
                     sizeof = self.context.get_abi_sizeof(lltype)
                     datamodel = self.context.data_model_manager[fetype]
                     line = self._adjust_line_if_prologue(self.loc.line)
-                    if src_name in self.fndesc.args:
-                        argidx = self.fndesc.args.index(src_name) + 1
-                    else:
-                        argidx = None
                     self.debuginfo.update_variable(
                         self.builder,
                         val,
@@ -1742,7 +1744,7 @@ def loadvar(self, name):
                         sizeof,
                         line,
                         datamodel,
-                        argidx,
+                        argidx=None,
                     )
                     self.dbg_val_names.add(src_name)
         return val
@@ -1817,6 +1819,23 @@ def storevar(self, value, name, argidx=None):
                     # Emit debug value for user variable
                     src_name = name.split(".")[0]
                     if src_name not in self.poly_var_typ_map:
+                        # Function arguments are described via dbg.declare on
+                        # their stack slots in the prologue.
+                        if argidx is not None:
+                            # Boolean parameters are kept on dbg.value to avoid
+                            # NVVM crashes with dbg.declare.
+                            if isinstance(fetype, types.Boolean):
+                                self.debuginfo.update_variable(
+                                    self.builder,
+                                    value,
+                                    src_name,
+                                    lltype,
+                                    sizeof,
+                                    line,
+                                    datamodel,
+                                    argidx,
+                                )
+                            return
                         # Insert the llvm.dbg.value intrinsic call
                         self.debuginfo.update_variable(
                             self.builder,
@@ -1839,9 +1858,8 @@ def storevar(self, value, name, argidx=None):
                             # Not yet covered by the dbg.value range
                             and src_name not in self.dbg_val_names
                         ):
-                            # Use fndesc.args to get correct argidx for func args
                             if src_name in self.fndesc.args:
-                                argidx = self.fndesc.args.index(src_name) + 1
+                                return
                             # Insert the llvm.dbg.value intrinsic call
                             self.debuginfo.update_variable(
                                 self.builder,
@@ -1851,7 +1869,7 @@ def storevar(self, value, name, argidx=None):
                                 sizeof,
                                 line,
                                 datamodel,
-                                argidx,
+                                argidx=None,
                             )
 
     def pre_block(self, block):

numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py

@@ -114,6 +114,37 @@ def f(x, y):
         match = re.compile(pat).search(llvm_ir)
         self.assertIsNotNone(match, msg=llvm_ir)
 
+    def test_scalar_kernel_parameters_use_dbg_declare(self):
+        sig = (types.int32, types.int32, types.int32, types.int32[::1])
+
+        @cuda.jit(sig, debug=True, opt=False)
+        def f(depth1, depth2, depth3, out):
+            acc = 0
+            for i in range(depth1):
+                for j in range(depth2):
+                    for k in range(depth3):
+                        acc += i + j + k
+            out[0] = acc
+
+        llvm_ir = f.inspect_llvm(sig)
+
+        # Scalar arguments should be described via dbg.declare on their stack
+        # slots (stable), rather than only via dbg.value (which ptxas may encode
+        # as parameter-space DW_OP_addr locations).
+        for argno, name in enumerate(("depth1", "depth2", "depth3"), start=1):
+            md_pat = rf"^!(\d+)\s+=\s+!DILocalVariable\([^)]*arg:\s*{argno}[^)]*name:\s*\"{name}\""
+            md_match = re.compile(md_pat, re.MULTILINE).search(llvm_ir)
+            self.assertIsNotNone(md_match, msg=llvm_ir)
+            md_id = md_match.group(1)
+
+            declare_pat = (
+                r"call void @\"llvm\.dbg\.declare\"\("
+                r"[\s\S]*?"
+                rf"metadata\s+!{md_id}\b"
+            )
+            declare_match = re.compile(declare_pat).search(llvm_ir)
+            self.assertIsNotNone(declare_match, msg=llvm_ir)
+
     def test_grid_group_type(self):
         sig = (types.int32,)
 
@@ -343,13 +374,23 @@ def f(x, y):
             z4 = True  # noqa: F841
 
         llvm_ir = f.inspect_llvm(sig)
-        # Verify the call to llvm.dbg.declare is replaced by llvm.dbg.value
-        pat1 = r'call void @"llvm.dbg.declare"'
-        match = re.compile(pat1).search(llvm_ir)
-        self.assertIsNone(match, msg=llvm_ir)
-        pat2 = r'call void @"llvm.dbg.value"'
-        match = re.compile(pat2).search(llvm_ir)
-        self.assertIsNotNone(match, msg=llvm_ir)
+        # Scalar locals should be described via llvm.dbg.value, not
+        # llvm.dbg.declare (formal parameters may still use dbg.declare).
+        for name in ("z1", "z2", "z3", "z4"):
+            md_pat = rf'^!(\d+)\s+=\s+!DILocalVariable\(.*arg:\s*0,.*name:\s+"{name}"'
+            md_match = re.compile(md_pat, re.MULTILINE).search(llvm_ir)
+            self.assertIsNotNone(md_match, msg=llvm_ir)
+            md_id = md_match.group(1)
+
+            value_pat = rf'call void @"llvm\.dbg\.value"\(metadata [^,]+, metadata !{md_id}\b'
+            value_match = re.compile(value_pat).search(llvm_ir)
+            self.assertIsNotNone(value_match, msg=llvm_ir)
+
+            declare_pat = (
+                rf'call void @"llvm\.dbg\.declare"\([^)]*metadata !{md_id}\b'
+            )
+            declare_match = re.compile(declare_pat).search(llvm_ir)
+            self.assertIsNone(declare_match, msg=llvm_ir)
 
     def test_llvm_dbg_value_range(self):
         sig = (types.int64,)
@@ -383,10 +424,9 @@ def test_llvm_dbg_value_loadvar_coverage(self):
         @cuda.jit("void(int32[:], int32)", debug=True, opt=False)
         def foo(arr, scalar):
             """
-            CHECK: call void @"llvm.dbg.value"(metadata i32 %"scalar"
+            CHECK: call void @"llvm.dbg.declare"(metadata i32* %"scalar.1", metadata ![[SC:[0-9]+]]
 
             CHECK: load i32, i32* %"scalar.1"
-            CHECK: call void @"llvm.dbg.value"(metadata i32 %"{{[^"]+}}", metadata ![[SC:[0-9]+]]
 
             CHECK: ![[SC]] = !DILocalVariable{{.+}}name: "scalar"
             """