diff --git a/numba_cuda/numba/cuda/lowering.py b/numba_cuda/numba/cuda/lowering.py index 7edbcbafe..570c50aa5 100644 --- a/numba_cuda/numba/cuda/lowering.py +++ b/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): """ diff --git a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py b/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py index 87fe592bf..7d240299b 100644 --- a/numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py +++ b/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)