Improve debug value range coverage

numba_cuda/numba/cuda/lowering.py

@@ -1698,33 +1698,71 @@ def storevar(self, value, name, argidx=None):
             self.context.enable_debuginfo
             # Conditions used to elide stores in parent method
             and self.store_var_needed(name)
-            # No emission of debuginfo for internal names
-            and not name.startswith("$")
         ):
-            # Emit debug value for user variable
             fetype = self.typeof(name)
             lltype = self.context.get_value_type(fetype)
             int_type = (llvm_ir.IntType,)
             real_type = llvm_ir.FloatType, llvm_ir.DoubleType
             if isinstance(lltype, int_type + real_type):
-                src_name = name.split(".")[0]
-                if src_name in self.poly_var_typ_map:
-                    # Do not emit debug value on polymorphic type var
-                    return
-                # Emit debug value for scalar variable
                 sizeof = self.context.get_abi_sizeof(lltype)
                 datamodel = self.context.data_model_manager[fetype]
                 line = self.loc.line if argidx is None else self.defn_loc.line
-                self.debuginfo.update_variable(
-                    self.builder,
-                    value,
-                    name,
-                    lltype,
-                    sizeof,
-                    line,
-                    datamodel,
-                    argidx,
-                )
+                if not name.startswith("$"):
+                    # Emit debug value for user variable
+                    src_name = name.split(".")[0]
+                    if src_name not in self.poly_var_typ_map:
+                        # Insert the llvm.dbg.value intrinsic call
+                        self.debuginfo.update_variable(
+                            self.builder,
+                            value,
+                            src_name,
+                            lltype,
+                            sizeof,
+                            line,
+                            datamodel,
+                            argidx,
+                        )
+                elif isinstance(value, llvm_ir.LoadInstr):
+                    # Emit debug value for user variable that falls out of the
+                    # coverage of dbg.value range per basic block
+                    ld_name = value.operands[0].name
+                    if not ld_name.startswith(("$", ".")):
+                        src_name = ld_name.split(".")[0]
+                        if (
+                            src_name not in self.poly_var_typ_map
+                            # Not yet covered by the dbg.value range
+                            and src_name not in self.dbg_val_names
+                        ):
+                            for index, item in enumerate(self.fnargs):
+                                if item.name == src_name:
+                                    argidx = index + 1
+                                    break
+                            # Insert the llvm.dbg.value intrinsic call
+                            self.debuginfo.update_variable(
+                                self.builder,
+                                value,
+                                src_name,
+                                lltype,
+                                sizeof,
+                                line,
+                                datamodel,
+                                argidx,
+                            )
+
+    def pre_block(self, block):
+        super().pre_block(block)
+
+        # dbg.value range covered names
+        self.dbg_val_names = set()
+
+        if self.context.enable_debuginfo and self._disable_sroa_like_opt:
+            for x in block.find_insts(ir.Assign):
+                if x.target.name.startswith("$"):
+                    continue
+                ssa_name = x.target.name
+                src_name = ssa_name.split(".")[0]
+                if src_name not in self.dbg_val_names:
+                    self.dbg_val_names.add(src_name)
 
     def pre_lower(self):
         """

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

@@ -348,6 +348,34 @@ def f(x, y):
         match = re.compile(pat2).search(llvm_ir)
         self.assertIsNotNone(match, msg=llvm_ir)
 
+    def test_llvm_dbg_value_range(self):
+        sig = (types.int64,)
+
+        @cuda.jit("void(int64,)", debug=True, opt=False)
+        def foo(x):
+            """
+            CHECK: store i1 true, i1* %"second.1"
+            CHECK: call void @"llvm.dbg.value"
+            CHECK: store i1 true, i1* %"second.2"
+            CHECK: call void @"llvm.dbg.value"
+
+            CHECK: %[[VAL_1:.*]] = load i1, i1* %"second.2"
+            CHECK: %[[VAL_2:.*]] = load i1, i1* %[[VAL_3:.*]]
+            CHECK: store i1 %[[VAL_1]], i1* %[[VAL_3]]
+            CHECK: call void @"llvm.dbg.value"(metadata i1 %[[VAL_1]], metadata ![[VAL_4:[0-9]+]]
+
+            CHECK: ![[VAL_4]] = !DILocalVariable{{.+}}name: "second"
+            """
+            if x > 0:
+                second = x > 10
+            else:
+                second = True
+            if second:
+                pass
+
+        ir = foo.inspect_llvm()[sig]
+        self.assertFileCheckMatches(ir, foo.__doc__)
+
     def test_no_user_var_alias(self):
         sig = (types.int32, types.int32)