Fix prologue debug line info pointing to decorator instead of def line

[jiel-nv]

Fixes nvbug5811746.

The root cause is that, when a kernel function has decorator, the code object's co_firstlineno points to the decorator line rather than the def line. During IR lowering, prologue code has no explicit source location, so co_firstlineno is used as the default line number. This caused prologue code to incorrectly reference the decorator line in DWARF debug information.

The problem can be reproduced with the following code example,

@cuda.jit("void(int64[:])", debug=True, opt=False)  # Line 270
def hybrid_loop_kernel(output):                     # Line 271
    idx = cuda.grid(1)                              # Line 272

This PR added _adjust_line_if_prologue() method in lowering.py that redirects any line number less than the def line to the def line. This ensures prologue instructions are associated with the function definition rather than decorators.

Before Fix - .debug_line table:

Address              Line
0x0000000000000000    271   ← def line
0x0000000000001c30    270   ← BUG: decorator line appears here
0x0000000000001f20    272   ← first statement

After Fix - .debug_line table:

Address              Line
0x0000000000000000    271   ← def line
0x0000000000001f20    272   ← first statement (no spurious 270!)

Also added a new test test_prologue_line_number which passes with the fix, fails without.

[copy-pr-bot]

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 Summary

Fixed debug line information for function prologues when decorators are present. Previously, prologue instructions incorrectly pointed to decorator lines instead of the def line because co_firstlineno references the first decorator. The new _adjust_line_if_prologue() helper redirects any line numbers before the actual function definition to the def line itself.

Key changes:

• Added _adjust_line_if_prologue() method in lowering.py:141-149 that returns defn_loc.line when line < defn_loc.line
• Applied the adjustment to three call sites: lower_inst() (line 516), mark_variable() (line 1674), and storevar() (line 1762)
• The change at line 1762 also simplifies conditional logic that previously used argidx to determine line selection
• Added test test_prologue_line_number() that verifies DILocation entries use def line or statement lines, with CHECK-NOT ensuring decorator line is absent

Confidence Score: 5/5

• This PR is safe to merge with minimal risk
• The fix is well-scoped, addresses a specific bug in debug information generation, includes a comprehensive test that validates the fix, and the implementation is sound - the helper method correctly redirects prologue line numbers while preserving actual source line numbers for non-prologue code
• No files require special attention

Important Files Changed

Filename	Overview
numba_cuda/numba/cuda/lowering.py	Adds _adjust_line_if_prologue() method and applies it consistently to fix debug line info for decorated functions
numba_cuda/numba/cuda/tests/cudapy/test_debuginfo.py	Adds comprehensive test verifying prologue uses def line instead of decorator line in debug info

[jiel-nv]

/ok to test 9deb359