Skip to content
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

[BUG] [Metadata] Incorrect kernel argument types when NUMBA_DUMP_LLVM is turned on #135

Open
jiel-nv opened this issue Feb 25, 2025 · 1 comment · May be fixed by #136
Open

[BUG] [Metadata] Incorrect kernel argument types when NUMBA_DUMP_LLVM is turned on #135

jiel-nv opened this issue Feb 25, 2025 · 1 comment · May be fixed by #136
Labels
bug Something isn't working

Comments

@jiel-nv
Copy link
Contributor

jiel-nv commented Feb 25, 2025

Describe the bug
When NUMBA_DUMP_LLVM is turned on, the debug metadata for kernel arguments still contains return value which should not be in the DISubroutineType.

Steps/Code to reproduce bug

$ cat bug.py
from numba import cuda

@cuda.jit("void(int32, int32)", debug=True, opt=False)
def f(x, y):
    z = x + y
$ NUMBA_DUMP_LLVM=1 python bug.py | grep "\!8"
!8 = !{ !5, !6, !7 }
!9 = !DISubroutineType(types: !8)
!8 = !{ !5, !6, !7 }
!9 = !DISubroutineType(types: !8)

Notice that, even after the "post kernel fixup" dump, the return value still shows up in the DISubroutineType, which was fixed by a previous PR (without the NUMBA_DUMP_LLVM turned on)

Expected behavior
If not using the NUMBA_DUMP_LLVM, and simply use the inspect_llvm() method, we do get the expected metadata as following,

$ cat bug.py
from numba import cuda
from numba.core import types

sig = (types.int32, types.int32)
@cuda.jit("void(int32, int32)", debug=True, opt=False)
def f(x, y):
    z = x + y

llvm_ir = f.inspect_llvm(sig)
llvmir = f.inspect_llvm(sig).splitlines()
for line in llvmir:
    print(line)
$ python bug.py | grep "\!8"
!8 = !{ !6, !7 }
!9 = !DISubroutineType(types: !8)

So something wrong happens when turning on NUMBA_DUMP_LLVM.
@jiel-nv jiel-nv added the bug Something isn't working label Feb 25, 2025
@jiel-nv
Copy link
Contributor Author

jiel-nv commented Feb 25, 2025

Another observation, taking the inspect_llvm() version of bug.py, if turning on NUMBA_DUMP_LLVM, the result of inspect_llvm() will also be wrong, i.e. NUMBA_DUMP_LLVM dominates in this case.

$ NUMBA_DUMP_LLVM=0 python bug.py | grep "\!8"
!8 = !{ !6, !7 }
!9 = !DISubroutineType(types: !8)
$ NUMBA_DUMP_LLVM=1 python bug.py | grep "\!8"
!8 = !{ !5, !6, !7 }
!9 = !DISubroutineType(types: !8)
!8 = !{ !5, !6, !7 }
!9 = !DISubroutineType(types: !8)
!8 = !{ !5, !6, !7 }
!9 = !DISubroutineType(types: !8)

jiel-nv added a commit to jiel-nv/numba-cuda that referenced this issue Feb 25, 2025
@jiel-nv
Emit correct kernel arguments debug metadata types when NUMBA_DUMP_LL…
3fe4229 
…VM is on.

After shifting out the return value type from the operands of 'types' field of
the 'DISubroutineType' metadata node, compiler needs to clear the cached string
representation of the corresponding MDValue object if NUMBA_DUMP_LLVM is on.

Because at that moment, the cached string has already been filled with the
tuple before kernel_fixup() which has the return value type and continue being
used when printing out the LLVM IR.

This change fixes NVIDIA#135
@jiel-nv jiel-nv linked a pull request Feb 25, 2025 that will close this issue
Open
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees
No one assigned
Labels
bug Something isn't working
Projects
None yet
Milestone
No milestone
Development

Successfully merging a pull request may close this issue.

Emit correct kernel arguments debug metadata types when NUMBA_DUMP_LLVM is on jiel-nv/numba-cuda
1 participant
@jiel-nv