Skip to content

Fix kernel return type in DISubroutineType debug metadata #745

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

Jump to bottom
Merged
gmarkall merged 2 commits into from
Jan 23, 2026
Merged

Fix kernel return type in DISubroutineType debug metadata #745

gmarkall merged 2 commits into from
Jan 23, 2026

Conversation

@jiel-nv
Copy link
Contributor

@jiel-nv jiel-nv commented Jan 23, 2026

Fixes nvbug#5804997

CUDA kernels were incorrectly showing the first parameter type as the return
type in DWARF debug info:

DW_AT_type (0x00000176 "float32")

This was caused by kernel_fixup() generating incorrect LLVM IR debug metadata DISubroutineType. The types tuple was:

!7 = !{ !6, !6 } // Wrong: first param type becomes return type

After this fix, generates null as the first element to represent void return type:

!7 = !{ null, !6, !6 } // Correct: null means void return

The resulting DWARF omits the DW_AT_type entry for the subprogram, which is the standard way to represent void return type in DWARF.

Also updated the existing debug info tests to reflect this change.

@copy-pr-bot
Copy link

copy-pr-bot bot commented Jan 23, 2026

Auto-sync is disabled for ready for review pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

@greptile-apps
Copy link
Contributor

greptile-apps bot commented Jan 23, 2026 •
edited
Loading

Greptile Summary

  • Fixes CUDA kernel debug metadata where DWARF debug info incorrectly showed the first parameter type as the kernel return type instead of void
  • Updates kernel_fixup() function in numba_cuda/numba/cuda/compiler.py to properly prepend null to DISubroutineType metadata tuple to represent void return types
  • Modifies test expectations in numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py to match the corrected metadata format with null as first element

Important Files Changed

Filename Overview
numba_cuda/numba/cuda/compiler.py Modified kernel_fixup() to prepend null to DISubroutineType types tuple instead of removing first element, fixing void return type representation
numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py Updated regex pattern to expect null as first element in DISubroutineType metadata tuple matching corrected compiler behavior

Confidence score: 5/5

  • This PR is safe to merge with minimal risk as it fixes a specific debug metadata bug without affecting runtime behavior
  • Score reflects targeted fix with clear problem/solution mapping, comprehensive test updates, and adherence to LLVM/DWARF standards
  • No files require special attention as both changes are straightforward and well-coordinated

@jiel-nv
Copy link
Contributor Author

jiel-nv commented Jan 23, 2026

/ok to test f6f0d5e

@jiel-nv
Copy link
Contributor Author

jiel-nv commented Jan 23, 2026

/ok to test af2c371

@gmarkall gmarkall added the 4 - Waiting on CI Waiting for a CI run to finish successfully label Jan 23, 2026
@gmarkall gmarkall enabled auto-merge (squash) January 23, 2026 20:44
@gmarkall gmarkall merged commit 1c39fb8 into NVIDIA:main Jan 23, 2026
105 checks passed
gmarkall added a commit to gmarkall/numba-cuda that referenced this pull request Jan 27, 2026
@gmarkall
Bump version to 0.25.0
cf01534 
- Add Python 3.14 to the wheel publishing matrix (NVIDIA#750)
- feat: swap out internal device array usage with `StridedMemoryView` (NVIDIA#703)
- Fix max block size computation in `forall` (NVIDIA#744)
- Fix prologue debug line info pointing to decorator instead of def line (NVIDIA#746)
- Fix kernel return type in DISubroutineType debug metadata (NVIDIA#745)
- Fix missing line info in Jupyter notebooks (NVIDIA#742)
- Fix: Pass correct flags to linker when debugging in the presence of LTOIR code (NVIDIA#698)
- chore(deps): add cuda-pathfinder to pixi deps (NVIDIA#741)
- fix: enable flake8-bugbear lints and fix found problems (NVIDIA#708)
- fix: Fix race condition in CUDA Simulator (NVIDIA#690)
- ci: run tests in parallel (NVIDIA#740)
- feat: users can pass `shared_memory_carveout` to @cuda.jit (NVIDIA#642)
- Fix compatibility with NumPy 2.4: np.trapz and np.in1d removed (NVIDIA#739)
- Pass the -numba-debug flag to libnvvm (NVIDIA#681)
- ci: remove rapids containers from conda ci (NVIDIA#737)
- Use `pathfinder` for dynamic libraries (NVIDIA#308)
- CI: Add CUDA 13.1 testing support (NVIDIA#705)
- Adding `pixi run test` and `pixi run test-par` support (NVIDIA#724)
- Disable per-PR nvmath tests + follow same test practice (NVIDIA#723)
- chore(deps): regenerate pixi lockfile (NVIDIA#722)
- Fix DISubprogram line number to point to function definition line (NVIDIA#695)
- revert: chore(dev): build pixi using rattler (NVIDIA#713) (NVIDIA#719)
- [feat] Initial version of the Numba CUDA GDB pretty-printer (NVIDIA#692)
- chore(dev): build pixi using rattler (NVIDIA#713)
- build(deps): bump the actions-monthly group across 1 directory with 8 updates (NVIDIA#704)
@gmarkall gmarkall mentioned this pull request Jan 27, 2026
Merged
kkraus14 pushed a commit that referenced this pull request Jan 28, 2026
@gmarkall
Bump version to 0.25.0 (#752)
bfa805a 
- Add Python 3.14 to the wheel publishing matrix (#750)
- feat: swap out internal device array usage with `StridedMemoryView`
(#703)
- Fix max block size computation in `forall` (#744)
- Fix prologue debug line info pointing to decorator instead of def line
(#746)
- Fix kernel return type in DISubroutineType debug metadata (#745)
- Fix missing line info in Jupyter notebooks (#742)
- Fix: Pass correct flags to linker when debugging in the presence of
LTOIR code (#698)
- chore(deps): add cuda-pathfinder to pixi deps (#741)
- fix: enable flake8-bugbear lints and fix found problems (#708)
- fix: Fix race condition in CUDA Simulator (#690)
- ci: run tests in parallel (#740)
- feat: users can pass `shared_memory_carveout` to @cuda.jit (#642)
- Fix compatibility with NumPy 2.4: np.trapz and np.in1d removed (#739)
- Pass the -numba-debug flag to libnvvm (#681)
- ci: remove rapids containers from conda ci (#737)
- Use `pathfinder` for dynamic libraries (#308)
- CI: Add CUDA 13.1 testing support (#705)
- Adding `pixi run test` and `pixi run test-par` support (#724)
- Disable per-PR nvmath tests + follow same test practice (#723)
- chore(deps): regenerate pixi lockfile (#722)
- Fix DISubprogram line number to point to function definition line
(#695)
- revert: chore(dev): build pixi using rattler (#713) (#719)
- [feat] Initial version of the Numba CUDA GDB pretty-printer (#692)
- chore(dev): build pixi using rattler (#713)
- build(deps): bump the actions-monthly group across 1 directory with 8
updates (#704)

<!--

Thank you for contributing to numba-cuda :)

Here are some guidelines to help the review process go smoothly.

1. Please write a description in this text box of the changes that are
being
   made.

2. Please ensure that you have written units tests for the changes
made/features
   added.

3. If you are closing an issue please use one of the automatic closing
words as
noted here:
https://help.github.com/articles/closing-issues-using-keywords/

4. If your pull request is not ready for review but you want to make use
of the
continuous integration testing facilities please label it with `[WIP]`.

5. If your pull request is ready to be reviewed without requiring
additional
work on top of it, then remove the `[WIP]` label (if present) and
replace
it with `[REVIEW]`. If assistance is required to complete the
functionality,
for example when the C/C++ code of a feature is complete but Python
bindings
are still required, then add the label `[HELP-REQ]` so that others can
triage
and assist. The additional changes then can be implemented on top of the
same PR. If the assistance is done by members of the rapidsAI team, then
no
additional actions are required by the creator of the original PR for
this,
otherwise the original author of the PR needs to give permission to the
person(s) assisting to commit to their personal fork of the project. If
that
doesn't happen then a new PR based on the code of the original PR can be
opened by the person assisting, which then will be the PR that will be
   merged.

6. Once all work has been done and review has taken place please do not
add
features or make changes out of the scope of those requested by the
reviewer
(doing this just add delays as already reviewed code ends up having to
be
re-reviewed/it is hard to tell what is new etc!). Further, please do not
rebase your branch on main/force push/rewrite history, doing any of
these
   causes the context of any comments made by reviewers to be lost. If
   conflicts occur against main they should be resolved by merging main
   into the branch used for making the pull request.

Many thanks in advance for your cooperation!

-->
@isVoid isVoid mentioned this pull request Feb 11, 2026
Open
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@gmarkall gmarkall gmarkall approved these changes

Assignees

No one assigned

Labels

4 - Waiting on CI Waiting for a CI run to finish successfully

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@jiel-nv @gmarkall