Skip to content

Improve debug info for kernel arguments #242

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
gmarkall merged 3 commits into from
May 20, 2025
Merged

Improve debug info for kernel arguments #242

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
1a79dfb
remove prefix on arguments
jiel-nv Apr 24, 2025
993e244
formatting
jiel-nv May 6, 2025
663e23a
formatting indent
jiel-nv May 6, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
27 changes: 27 additions & 0 deletions numba_cuda/numba/cuda/debuginfo.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,33 @@ def _var_type(self, lltype, size, datamodel=None):
# For other cases, use upstream Numba implementation
return super()._var_type(lltype, size, datamodel=datamodel)

def _di_subroutine_type(self, line, function, argmap):
# The function call conv needs encoding.
llfunc = function
md = []

# Create metadata type for return value
if len(llfunc.args) > 0:
lltype = llfunc.args[0].type
Copy link
Contributor

@gmarkall gmarkall May 20, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Does omitting the replacement of . with $ compared to the original implementation not cause issues for GDB? I note the comment in the original version:

name = llarg.name.replace('.', '$')    # for gdb to work correctly

I thought that a dot in a name might be problematic for it? (c.f. the same syntax for traversing the members of a struct or union)

Copy link
Contributor Author

@jiel-nv jiel-nv May 20, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That . versioned name used to be written to the metadata node as the value of 'name' of DILocalVariable. But currently the . versioned variable nodes are merged to the unversioned / user symbol named metadata node. Therefore, they are no longer to be present in DWARF entry.
And when represented as member of a struct or union, their names in the metadata are also unversioned names with corresponding types, so there won't be issues omitting this original replacement.
In short, neither . nor $ will show up in debugger.

size = self.cgctx.get_abi_sizeof(lltype)
mdtype = self._var_type(lltype, size, datamodel=None)
md.append(mdtype)

# Create metadata type for arguments
for idx, (name, nbtype) in enumerate(argmap.items()):
datamodel = self.cgctx.data_model_manager[nbtype]
lltype = self.cgctx.get_value_type(nbtype)
size = self.cgctx.get_abi_sizeof(lltype)
mdtype = self._var_type(lltype, size, datamodel=datamodel)
md.append(mdtype)

return self.module.add_debug_info(
"DISubroutineType",
{
"types": self.module.add_metadata(md),
},
)

def mark_variable(
self,
builder,
Expand Down
11 changes: 10 additions & 1 deletion numba_cuda/numba/cuda/target.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -290,7 +290,16 @@ def get_ufunc_info(self, ufunc_key):


class CUDACallConv(MinimalCallConv):
pass
def decorate_function(self, fn, args, fe_argtypes, noalias=False):
"""
Set names and attributes of function arguments.
"""
assert not noalias
arginfo = self._get_arg_packer(fe_argtypes)
# Do not prefix "arg." on argument name, so that nvvm compiler
Copy link
Contributor

@gmarkall gmarkall May 20, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think that it's OK not to prefix arg. onto the name - as far as I can tell it's a convention that was added that has not become load-bearing in any way - I don't think a duplicate name will appear inside a function body.

# can track debug info of argument more accurately
arginfo.assign_names(self.get_arguments(fn), args)
fn.args[0].name = ".ret"


class CUDACABICallConv(BaseCallConv):
Expand Down
17 changes: 17 additions & 0 deletions numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -310,6 +310,23 @@ def test_kernel_args_types_dump(self):
with captured_stdout():
self._test_kernel_args_types()

def test_kernel_args_names(self):
sig = (types.int32,)

@cuda.jit("void(int32)", debug=True, opt=False)
def f(x):
z = x # noqa: F841

llvm_ir = f.inspect_llvm(sig)

# Verify argument name is not prefixed with "arg."
pat = r"define void @.*\(i32 %\"x\"\)"
match = re.compile(pat).search(llvm_ir)
self.assertIsNotNone(match, msg=llvm_ir)
pat = r"define void @.*\(i32 %\"arg\.x\"\)"
match = re.compile(pat).search(llvm_ir)
self.assertIsNone(match, msg=llvm_ir)

def test_llvm_dbg_value(self):
sig = (types.int32, types.int32)

Expand Down