Skip to content

Fix DISubprogram line number to point to function definition line#695

Merged
gmarkall merged 2 commits into
NVIDIA:mainfrom
jiel-nv:decl-line
Jan 15, 2026
Merged

Fix DISubprogram line number to point to function definition line#695
gmarkall merged 2 commits into
NVIDIA:mainfrom
jiel-nv:decl-line

Conversation

@jiel-nv

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

Copy link
Copy Markdown
Contributor

Fixes incorrect DW_AT_decl_line where the line number was first_stmt_line - 1 instead of the actual def line, causing off-by-one errors when comments exist between the def line and first statement.

  • Add get_func_def_lineno() in firstlinefinder.py to return the actual 'def' line
  • Update lowering.py to use get_func_def_lineno() directly without offset calculation
  • Add test_DISubprogram_def_line_number covering single/multi-line decorators with/without comments

This change resolves nvbug#5723821.

- Add get_func_def_lineno() in firstlinefinder.py to return the actual 'def' line
- Update lowering.py to use get_func_def_lineno() directly without offset calculation
- Add test_DISubprogram_def_line_number covering single/multi-line decorators with/without comments

Fixes incorrect DW_AT_decl_line where the line number was first_stmt_line - 1
instead of the actual def line, causing off-by-one errors when comments exist
between the def line and first statement.
@copy-pr-bot

copy-pr-bot Bot commented Jan 6, 2026

Copy link
Copy Markdown

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

greptile-apps Bot commented Jan 6, 2026

Copy link
Copy Markdown
Contributor

Greptile Summary

This PR fixes an off-by-one error in debug information where DW_AT_decl_line was incorrectly pointing to first_stmt_line - 1 instead of the actual function definition line.

Changes:

  • Added get_func_def_lineno() function in firstlinefinder.py that returns the actual def line number from the AST
  • Modified FindDefFirstLine class to capture def_lineno (line 31, 51)
  • Updated lowering.py to use get_func_def_lineno() directly without offset calculation (removed the -1 adjustment)
  • Added comprehensive test test_DISubprogram_def_line_number covering single/multi-line decorators with/without comments

The fix ensures that debug information correctly points to the function definition line regardless of comments between the def line and the first statement.

Confidence Score: 5/5

  • This PR is safe to merge with minimal risk
  • The change is straightforward and well-tested. It fixes a specific bug in debug information generation by using the correct line number from the AST. The new function mirrors the existing get_func_body_first_lineno() pattern, comprehensive tests cover multiple scenarios (single/multi-line decorators, with/without comments), and the change is isolated to debug information generation without affecting runtime behavior
  • No files require special attention

Important Files Changed

Filename Overview
numba_cuda/numba/cuda/misc/firstlinefinder.py Adds get_func_def_lineno() function to return the actual 'def' line number from AST, and stores def_lineno in FindDefFirstLine class
numba_cuda/numba/cuda/lowering.py Switches from get_func_body_first_lineno() to get_func_def_lineno() and removes the -1 offset calculation
numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py Adds comprehensive test coverage for DISubprogram line numbers with various decorator and comment configurations

@greptile-apps

greptile-apps Bot commented Jan 6, 2026

Copy link
Copy Markdown
Contributor

Greptile's behavior is changing!

From now on, if a review finishes with no comments, we will not post an additional "statistics" comment to confirm that our review found nothing to comment on. However, you can confirm that we reviewed your changes in the status check section.

This feature can be toggled off in your Code Review Settings by deselecting "Create a status check for each PR".

@jiel-nv

jiel-nv commented Jan 6, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test 6fe1b2f

@jiel-nv

jiel-nv commented Jan 6, 2026

Copy link
Copy Markdown
Contributor Author

/ok to test 8286bc0

@jiel-nv jiel-nv added the 3 - Ready for Review Ready for review by team label Jan 6, 2026
@jiel-nv
jiel-nv requested a review from gmarkall January 6, 2026 22:05
@jiel-nv jiel-nv added the 4 - Waiting on reviewer Waiting for reviewer to respond to author label Jan 15, 2026
@gmarkall

Copy link
Copy Markdown
Contributor

Apologies for the delay in reviewing. Whilst testing it, I discovered and investigated #721, which I think needs fixing at some point but could be done separately from this PR. I think this PR is otherwise correct and makes an improvement.

@gmarkall gmarkall added 5 - Ready to merge Testing and reviews complete, ready to merge and removed 3 - Ready for Review Ready for review by team 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Jan 15, 2026
@gmarkall
gmarkall merged commit b706197 into NVIDIA:main Jan 15, 2026
130 checks passed
gmarkall added a commit to gmarkall/numba-cuda that referenced this pull request Jan 27, 2026
- 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
kkraus14 pushed a commit that referenced this pull request Jan 28, 2026
- 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!

-->
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

5 - Ready to merge Testing and reviews complete, ready to merge

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants