Index: lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h =================================================================== --- lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h +++ lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h @@ -24,6 +24,9 @@ NVPTXTargetStreamer(MCStreamer &S); ~NVPTXTargetStreamer() override; + /// Outputs the list of the DWARF '.file' directives to the streamer. + void outputDwarfFileDirectives(); + /// Record DWARF file directives for later output. /// According to PTX ISA, CUDA Toolkit documentation, 11.5.3. Debugging /// Directives: .file Index: lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp =================================================================== --- lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp +++ lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp @@ -25,6 +25,12 @@ NVPTXTargetStreamer::~NVPTXTargetStreamer() = default; +void NVPTXTargetStreamer::outputDwarfFileDirectives() { + for (const std::string &S : DwarfFiles) + getStreamer().EmitRawText(S.data()); + DwarfFiles.clear(); +} + void NVPTXTargetStreamer::emitDwarfFileDirective(StringRef Directive) { DwarfFiles.emplace_back(Directive); } @@ -82,9 +88,7 @@ OS << "//\t}\n"; if (isDwarfSection(FI, Section)) { // Emit DWARF .file directives in the outermost scope. - for (const std::string &S : DwarfFiles) - getStreamer().EmitRawText(S.data()); - DwarfFiles.clear(); + outputDwarfFileDirectives(); OS << "//\t.section"; Section->PrintSwitchToSection(*getStreamer().getContext().getAsmInfo(), FI->getTargetTriple(), OS, SubSection); Index: lib/Target/NVPTX/NVPTXAsmPrinter.cpp =================================================================== --- lib/Target/NVPTX/NVPTXAsmPrinter.cpp +++ lib/Target/NVPTX/NVPTXAsmPrinter.cpp @@ -16,6 +16,7 @@ #include "InstPrinter/NVPTXInstPrinter.h" #include "MCTargetDesc/NVPTXBaseInfo.h" #include "MCTargetDesc/NVPTXMCAsmInfo.h" +#include "MCTargetDesc/NVPTXTargetStreamer.h" #include "NVPTX.h" #include "NVPTXMCExpr.h" #include "NVPTXMachineFunctionInfo.h" @@ -880,8 +881,22 @@ if (NTM.getDrvInterface() == NVPTX::NVCL) O << ", texmode_independent"; + bool HasFullDebugInfo = false; + for (DICompileUnit *CU : M.debug_compile_units()) { + switch(CU->getEmissionKind()) { + case DICompileUnit::NoDebug: + case DICompileUnit::DebugDirectivesOnly: + break; + case DICompileUnit::LineTablesOnly: + case DICompileUnit::FullDebug: + HasFullDebugInfo = true; + break; + } + if (HasFullDebugInfo) + break; + } // FIXME: remove comment once debug info is properly supported. - if (MMI && MMI->hasDebugInfo()) + if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo) O << "//, debug"; O << "\n"; @@ -938,6 +953,10 @@ if (HasDebugInfo) OutStreamer->EmitRawText("//\t}"); + // Output last DWARF .file directives, if any. + static_cast(OutStreamer->getTargetStreamer()) + ->outputDwarfFileDirectives(); + return ret; //bool Result = AsmPrinter::doFinalization(M); Index: test/DebugInfo/NVPTX/debug-file-loc-only.ll =================================================================== --- /dev/null +++ test/DebugInfo/NVPTX/debug-file-loc-only.ll @@ -0,0 +1,48 @@ +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s + +; // Bitcode int this test case is reduced version of compiled code below: +;extern "C" { +;#line 1 "/source/dir/foo.h" +;__device__ void foo() {} +;#line 2 "/source/dir/bar.cu" +;__device__ void bar() {} +;} + +; CHECK: .target sm_{{[0-9]+$}} + +; CHECK: .visible .func foo() +; CHECK: .loc [[FOO:[0-9]+]] 1 31 +; CHECK: ret; +; CHECK: .visible .func bar() +; CHECK: .loc [[BAR:[0-9]+]] 2 31 +; CHECK: ret; + +define void @foo() !dbg !4 { +bb: + ret void, !dbg !10 +} + +define void @bar() !dbg !7 { +bb: + ret void, !dbg !11 +} + +; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h" +; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu" + +; CHECK-NOT: .section .debug{{.*}} + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!8, !9} + +!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "", isOptimized: false, runtimeVersion: 0, emissionKind: DebugDirectivesOnly, enums: !2) +!1 = !DIFile(filename: "bar.cu", directory: "/source/dir") +!2 = !{} +!4 = distinct !DISubprogram(name: "foo", scope: !5, file: !5, line: 1, type: !6, isLocal: false, isDefinition: true, scopeLine: 1, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2) +!5 = !DIFile(filename: "foo.h", directory: "/source/dir") +!6 = !DISubroutineType(types: !2) +!7 = distinct !DISubprogram(name: "bar", scope: !1, file: !1, line: 2, type: !6, isLocal: false, isDefinition: true, scopeLine: 2, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2) +!8 = !{i32 2, !"Dwarf Version", i32 2} +!9 = !{i32 2, !"Debug Info Version", i32 3} +!10 = !DILocation(line: 1, column: 31, scope: !4) +!11 = !DILocation(line: 2, column: 31, scope: !7)