-
Notifications
You must be signed in to change notification settings - Fork 10.8k
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.
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
[NVPTX][DebugInfo] avoid emitting extra .loc directives #84584
Conversation
@llvm/pr-subscribers-debuginfo Author: Alex MacLean (AlexMaclean) ChangesThis change removes an extra, unneeded debug directive emitted in the PTX at the beginning on non-empty functions:
Full diff: https://github.com/llvm/llvm-project/pull/84584.diff 2 Files Affected:
diff --git a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
index 2219d9f6619aaf..ed5df37cbde2c7 100644
--- a/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -489,8 +489,12 @@ void NVPTXAsmPrinter::emitFunctionEntryLabel() {
OutStreamer->emitRawText(StringRef("{\n"));
setAndEmitFunctionVirtualRegisters(*MF);
// Emit initial .loc debug directive for correct relocation symbol data.
- if (MMI && MMI->hasDebugInfo())
- emitInitialRawDwarfLocDirective(*MF);
+ const DISubprogram *SP = MF->getFunction().getSubprogram();
+ if (SP) {
+ assert(SP->getUnit());
+ if (!SP->getUnit()->isDebugDirectivesOnly() && MMI && MMI->hasDebugInfo())
+ emitInitialRawDwarfLocDirective(*MF);
+ }
}
bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {
diff --git a/llvm/test/DebugInfo/NVPTX/no-extra-loc.ll b/llvm/test/DebugInfo/NVPTX/no-extra-loc.ll
new file mode 100644
index 00000000000000..6d3a69beffe875
--- /dev/null
+++ b/llvm/test/DebugInfo/NVPTX/no-extra-loc.ll
@@ -0,0 +1,26 @@
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64-nvidia-cuda | %ptxas-verify %}
+
+
+define i32 @foo(i32 %a, i32 %b) !dbg !3 {
+
+; CHECK: .loc [[FILE:[0-9]+]] 26 0 // extra-lineinfo.cu:26:0
+; CHECK-NOT: .loc [[FILE]] 26 0 // extra-lineinfo.cu:26:0
+; CHECK: .file [[FILE]] "/test/directory/extra-lineinfo.cu"
+
+ %add = add i32 %b, %a, !dbg !6
+ ret i32 %add, !dbg !6
+}
+
+!llvm.dbg.cu = !{!0}
+!nvvm.annotations = !{}
+!llvm.module.flags = !{!2}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "clang", isOptimized: true, runtimeVersion: 0, emissionKind: DebugDirectivesOnly)
+!1 = !DIFile(filename: "extra-lineinfo.cu", directory: "/test/directory/")
+!2 = !{i32 1, !"Debug Info Version", i32 3}
+!3 = distinct !DISubprogram(name: "kernel", linkageName: "foo", scope: !1, file: !1, line: 123, type: !4, scopeLine: 26, spFlags: DISPFlagDefinition | DISPFlagOptimized, unit: !0)
+!4 = !DISubroutineType(types: !5)
+!5 = !{}
+!6 = !DILocation(line: 40, column: 22, scope: !31)
+!31 = distinct !DILexicalBlock(scope: !3, file: !1, line: 3, column: 17)
|
@Artem-B, when you have a minute, could you please take a look at this? Thanks! |
if (MMI && MMI->hasDebugInfo()) | ||
emitInitialRawDwarfLocDirective(*MF); | ||
const DISubprogram *SP = MF->getFunction().getSubprogram(); | ||
if (SP) { |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Nit: the temp variable can be folded into if
:
if (const DISubprogram *SP = MF->getFunction().getSubprogram()) {
...
}
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Fixed
This change removes an extra, unneeded debug directive emitted in the PTX at the beginning on non-empty functions: ```nvptx .visible .func (.param .b32 func_retval0) foo( .param .b32 foo_param_0, .param .b32 foo_param_1 ) { .reg .b32 %r<4>; .loc 1 26 0 <---- unneeded (removed by the PR) $L__func_begin0: .loc 1 26 0 ```
This change removes an extra, unneeded debug directive emitted in the PTX at the beginning on non-empty functions: