Fix empty DW_AT_location for reused loop variables (#811)

Nvbug5886953 exposes an issue that reused loop index variables have a
DWARF location list with corrupted entries.

The root cause is in the base class `Lower.loadvar`, which suppresses
`!dbg` metadata on loads of function argument variables. This causes
NVVM to emit `.loc line 0` between the loops, and `ptxas` builds a
multi-entry location list where the line-0 ranges inserted between
entries corrupt it, producing an empty `DW_AT_location` in the final
DWARF.

This PR overrides `loadvar` in `CUDALower` to skip the suppression. The
suppression is unnecessary because `loadvar` is never called during the
prologue — arg unpacking in `lower_assign` and `storevar` already handle
suppression there.

Also add `test_arg_load_has_dbg_location` regression test.

Fixes nvbug5886953.

Author: jiel-nv

numba_cuda/numba/cuda/lowering.py

@@ -1695,6 +1695,21 @@ def decref(self, typ, val):
 
 
 class CUDALower(Lower):
+    def loadvar(self, name):
+        if name in self._blk_local_varmap and not self._disable_sroa_like_opt:
+            return self._blk_local_varmap[name]
+        ptr = self.getvar(name)
+
+        # Skip the base-class suspend_emission for arg-named loads.
+        # loadvar is never called in the prologue (arg unpacking in
+        # lower_assign and storevar already suppresses there).
+        # On CUDA the missing !dbg emits ".loc line 0"; when a
+        # variable (e.g. a reused loop counter) has multiple
+        # dbg.value entries, ptxas builds a multi-entry location
+        # list and the line-0 ranges inserted between entries
+        # will result in an corrupted DW_AT_location.
+        return self.builder.load(ptr)
+
     def storevar(self, value, name, argidx=None):
         """
         Store the value into the given variable.

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

@@ -1212,6 +1212,31 @@ def foo(output):
             match, f"No non-zero store to 'bar.N' with !dbg !{continue_dbg_id}"
         )
 
+    def test_arg_load_has_dbg_location(self):
+        """Loads of arg-named variables must carry !dbg in the body.
+
+        Reusing `j` across two loops creates multiple dbg.value
+        entries; a missing !dbg on the arg load (".loc line 0")
+        makes ptxas zero the location list and results in corrupted
+        DW_AT_location. Regression test for nvbug 5886953.
+        """
+        sig = (types.float32[:], types.int64)
+
+        @cuda.jit("void(float32[:], int64)", debug=True, opt=False)
+        def foo(arr, n):
+            for j in range(n):
+                arr[0] += j
+            for j in range(n):
+                arr[0] += j
+
+        llvm_ir = foo.inspect_llvm(sig)
+        pat = re.compile(r'load\s.*%"arr"\s*$', re.MULTILINE)
+        match = pat.search(llvm_ir)
+        self.assertIsNone(
+            match,
+            msg="Load of arg 'arr' missing !dbg metadata",
+        )
+
 
 if __name__ == "__main__":
     unittest.main()