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] Kernel fixup does not adjust debug metadata arguments #95

Open
gmarkall opened this issue Dec 19, 2024 · 0 comments · May be fixed by #97
Open

[BUG] Kernel fixup does not adjust debug metadata arguments #95

gmarkall opened this issue Dec 19, 2024 · 0 comments · May be fixed by #97
Labels
bug Something isn't working

Comments

@gmarkall
Copy link
Collaborator

The kernel fixup added in #50 removes the return value placeholder from the argument list, but does not remove it from the debug metadata. Thus, a kernel like:

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

will have the following prototype in LLVM IR:

define void @"k"(i32 %"arg.x", i32 %"arg.y") !dbg !10

(function name demangled for clarity)

but the debug metadata is:

!3 = !DIBasicType(encoding: DW_ATE_unsigned, name: "i8", size: 8)
!4 = !DIDerivedType(baseType: !3, size: 64, tag: DW_TAG_pointer_type)
!5 = !DIDerivedType(baseType: !4, size: 64, tag: DW_TAG_pointer_type)
!6 = !DIBasicType(encoding: DW_ATE_signed, name: "int32", size: 32)
!7 = !DIBasicType(encoding: DW_ATE_signed, name: "int32", size: 32)
!8 = !{ !5, !6, !7 }
!9 = !DISubroutineType(types: !8)
!10 = distinct !DISubprogram(file: !1, isDefinition: true, isLocal: false, isOptimized: true,
                             line: 5, linkageName: "k", name: "k", scope: !1, scopeLine: 5, type: !9, unit: !2)

The issue being that the arguments in !8 are !{ !5, !6, !7 }. !5 is the placeholder for a return value (the equivalent of void**), which should have been removed to leave only !{ !6, !7 }.
@gmarkall gmarkall added the bug Something isn't working label Dec 19, 2024
@gmarkall gmarkall linked a pull request Dec 20, 2024 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.

Fixup debug metadata in kernel fixup gmarkall/numba-cuda
1 participant
@gmarkall