Index: include/llvm/MC/MCTargetOptions.h =================================================================== --- include/llvm/MC/MCTargetOptions.h +++ include/llvm/MC/MCTargetOptions.h @@ -58,6 +58,9 @@ /// Preserve Comments in Assembly. bool PreserveAsmComments : 1; + /// Emit debug directives only, disable emission of the debug sections. + bool DebugDirectivesOnly : 1; + int DwarfVersion = 0; std::string ABIName; Index: lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp =================================================================== --- lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp +++ lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp @@ -269,6 +269,10 @@ } void DwarfCompileUnit::initStmtList() { + // Do not emit DWARF sections if it is not required. + if (Asm->TM.Options.MCOptions.DebugDirectivesOnly) + return; + // Define start line table label for each Compile Unit. MCSymbol *LineTableStartSym; const TargetLoweringObjectFile &TLOF = Asm->getObjFileLowering(); Index: lib/CodeGen/AsmPrinter/DwarfDebug.cpp =================================================================== --- lib/CodeGen/AsmPrinter/DwarfDebug.cpp +++ lib/CodeGen/AsmPrinter/DwarfDebug.cpp @@ -624,6 +624,11 @@ M->debug_compile_units_end()); // Tell MMI whether we have debug info. MMI->setDebugInfoAvailability(NumDebugCUs > 0); + + // Do not emit DWARF sections if it is not required. + if (Asm->TM.Options.MCOptions.DebugDirectivesOnly) + return; + SingleCU = NumDebugCUs == 1; DenseMap> GVMap; @@ -813,6 +818,10 @@ if (!MMI->hasDebugInfo()) return; + // Do not emit DWARF sections if it is not required. + if (Asm->TM.Options.MCOptions.DebugDirectivesOnly) + return; + // Finalize the debug info for the module. finalizeModuleInfo(); Index: lib/MC/MCTargetOptions.cpp =================================================================== --- lib/MC/MCTargetOptions.cpp +++ lib/MC/MCTargetOptions.cpp @@ -18,7 +18,7 @@ MCSaveTempLabels(false), MCUseDwarfDirectory(false), MCIncrementalLinkerCompatible(false), MCPIECopyRelocations(false), ShowMCEncoding(false), ShowMCInst(false), AsmVerbose(false), - PreserveAsmComments(true) {} + PreserveAsmComments(true), DebugDirectivesOnly(false) {} StringRef MCTargetOptions::getABIName() const { return ABIName; 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" @@ -876,7 +877,7 @@ O << ", texmode_independent"; // FIXME: remove comment once debug info is properly supported. - if (MMI && MMI->hasDebugInfo()) + if (MMI && MMI->hasDebugInfo() && !NTM.Options.MCOptions.DebugDirectivesOnly) O << "//, debug"; O << "\n"; @@ -933,6 +934,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/Generic/directives-only.ll =================================================================== --- /dev/null +++ test/DebugInfo/Generic/directives-only.ll @@ -0,0 +1,65 @@ +; RUN: llc -filetype=asm -asm-verbose=0 -O0 -debug-directives-only=true < %s | FileCheck %s --check-prefixes=DISABLED,CHECK +; RUN: llc -filetype=asm -asm-verbose=0 -O0 -debug-directives-only=false < %s | FileCheck %s --check-prefixes=ENABLED,CHECK + +; Check that the assembly output properly handles is_stmt changes. And since +; we're testing anyway, check the integrated assembler too. + +; Generated with clang from multiline.c: +; void f1(); +; void f2() { +; f1(); f1(); f1(); +; f1(); f1(); f1(); +; } + + +; CHECK: .file 1 "/tmp/dbginfo{{.*}}multiline.c" +; CHECK: .loc 1 2 0 +; CHECK: .loc 1 3 3 +; CHECK: .loc 1 3 9 +; CHECK: .loc 1 3 15 +; CHECK: .loc 1 4 3 +; CHECK: .loc 1 4 9 +; CHECK: .loc 1 4 15 +; CHECK: .loc 1 5 1 + +; ENABLED: .section .{{debug.*}} +; DISABLED-NOT: .section .{{debug.*}} + +; Function Attrs: nounwind uwtable +define void @f2() #0 !dbg !4 { +entry: + call void (...) @f1(), !dbg !11 + call void (...) @f1(), !dbg !12 + call void (...) @f1(), !dbg !13 + call void (...) @f1(), !dbg !14 + call void (...) @f1(), !dbg !15 + call void (...) @f1(), !dbg !16 + ret void, !dbg !17 +} + +declare void @f1(...) #1 + +attributes #0 = { nounwind uwtable "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #1 = { "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } + +!llvm.dbg.cu = !{!0} +!llvm.module.flags = !{!8, !9} +!llvm.ident = !{!10} + +!0 = distinct !DICompileUnit(language: DW_LANG_C99, producer: "clang version 3.6.0 (trunk 225000) (llvm/trunk 224999)", isOptimized: false, emissionKind: FullDebug, file: !1, enums: !2, retainedTypes: !2, globals: !2, imports: !2) +!1 = !DIFile(filename: "multiline.c", directory: "/tmp/dbginfo") +!2 = !{} +!4 = distinct !DISubprogram(name: "f2", line: 2, isLocal: false, isDefinition: true, isOptimized: false, unit: !0, scopeLine: 2, file: !1, scope: !5, type: !6, variables: !2) +!5 = !DIFile(filename: "multiline.c", directory: "/tmp/dbginfo") +!6 = !DISubroutineType(types: !7) +!7 = !{null} +!8 = !{i32 2, !"Dwarf Version", i32 4} +!9 = !{i32 2, !"Debug Info Version", i32 3} +!10 = !{!"clang version 3.6.0 (trunk 225000) (llvm/trunk 224999)"} +!11 = !DILocation(line: 3, column: 3, scope: !4) +!12 = !DILocation(line: 3, column: 9, scope: !4) +!13 = !DILocation(line: 3, column: 15, scope: !4) +!14 = !DILocation(line: 4, column: 3, scope: !4) +!15 = !DILocation(line: 4, column: 9, scope: !4) +!16 = !DILocation(line: 4, column: 15, scope: !4) +!17 = !DILocation(line: 5, column: 1, scope: !4) Index: test/DebugInfo/NVPTX/debug-file-loc.ll =================================================================== --- test/DebugInfo/NVPTX/debug-file-loc.ll +++ test/DebugInfo/NVPTX/debug-file-loc.ll @@ -1,4 +1,5 @@ -; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s --check-prefixes=CHECK,DEBUG +; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -debug-directives-only=true | FileCheck %s --check-prefixes=CHECK,LINEINFO ; // Bitcode int this test case is reduced version of compiled code below: ;extern "C" { @@ -8,7 +9,8 @@ ;__device__ void bar() {} ;} -; CHECK: .target sm_{{[0-9]+}}//, debug +; DEBUG: .target sm_{{[0-9]+}}//, debug +; LINEINFO: .target sm_{{[0-9]+$}} ; CHECK: .visible .func foo() ; CHECK: .loc [[FOO:[0-9]+]] 1 31 @@ -29,67 +31,70 @@ ; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h" ; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu" -; CHECK: // .section .debug_abbrev -; CHECK: // { -; CHECK: // .b8 1 // Abbreviation Code -; CHECK: // .b8 17 // DW_TAG_compile_unit -; CHECK: // .b8 0 // DW_CHILDREN_no -; CHECK: // .b8 37 // DW_AT_producer -; CHECK: // .b8 8 // DW_FORM_string -; CHECK: // .b8 19 // DW_AT_language -; CHECK: // .b8 5 // DW_FORM_data2 -; CHECK: // .b8 3 // DW_AT_name -; CHECK: // .b8 8 // DW_FORM_string -; CHECK: // .b8 16 // DW_AT_stmt_list -; CHECK: // .b8 6 // DW_FORM_data4 -; CHECK: // .b8 27 // DW_AT_comp_dir -; CHECK: // .b8 8 // DW_FORM_string -; CHECK: // .b8 17 // DW_AT_low_pc -; CHECK: // .b8 1 // DW_FORM_addr -; CHECK: // .b8 18 // DW_AT_high_pc -; CHECK: // .b8 1 // DW_FORM_addr -; CHECK: // .b8 0 // EOM(1) -; CHECK: // .b8 0 // EOM(2) -; CHECK: // .b8 0 // EOM(3) -; CHECK: // } -; CHECK: // .section .debug_info -; CHECK: // { -; CHECK: // .b32 50 // Length of Unit -; CHECK: // .b8 2 // DWARF version number -; CHECK: // .b8 0 -; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section -; CHECK: // .b8 8 // Address Size (in bytes) -; CHECK: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit -; CHECK: // .b8 0 // DW_AT_producer -; CHECK: // .b8 4 // DW_AT_language -; CHECK: // .b8 0 -; CHECK: // .b8 98 // DW_AT_name -; CHECK: // .b8 97 -; CHECK: // .b8 114 -; CHECK: // .b8 46 -; CHECK: // .b8 99 -; CHECK: // .b8 117 -; CHECK: // .b8 0 -; CHECK: // .b32 .debug_line // DW_AT_stmt_list -; CHECK: // .b8 47 // DW_AT_comp_dir -; CHECK: // .b8 115 -; CHECK: // .b8 111 -; CHECK: // .b8 117 -; CHECK: // .b8 114 -; CHECK: // .b8 99 -; CHECK: // .b8 101 -; CHECK: // .b8 47 -; CHECK: // .b8 100 -; CHECK: // .b8 105 -; CHECK: // .b8 114 -; CHECK: // .b8 0 -; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc -; CHECK: // .b64 Lfunc_end1 // DW_AT_high_pc -; CHECK: // } -; CHECK: // .section .debug_macinfo -; CHECK: // { -; CHECK: // .b8 0 // End Of Macro List Mark -; CHECK: // } + +; LINEINFO-NOT: .section .debug{{.*}} + +; DEBUG: // .section .debug_abbrev +; DEBUG: // { +; DEBUG: // .b8 1 // Abbreviation Code +; DEBUG: // .b8 17 // DW_TAG_compile_unit +; DEBUG: // .b8 0 // DW_CHILDREN_no +; DEBUG: // .b8 37 // DW_AT_producer +; DEBUG: // .b8 8 // DW_FORM_string +; DEBUG: // .b8 19 // DW_AT_language +; DEBUG: // .b8 5 // DW_FORM_data2 +; DEBUG: // .b8 3 // DW_AT_name +; DEBUG: // .b8 8 // DW_FORM_string +; DEBUG: // .b8 16 // DW_AT_stmt_list +; DEBUG: // .b8 6 // DW_FORM_data4 +; DEBUG: // .b8 27 // DW_AT_comp_dir +; DEBUG: // .b8 8 // DW_FORM_string +; DEBUG: // .b8 17 // DW_AT_low_pc +; DEBUG: // .b8 1 // DW_FORM_addr +; DEBUG: // .b8 18 // DW_AT_high_pc +; DEBUG: // .b8 1 // DW_FORM_addr +; DEBUG: // .b8 0 // EOM(1) +; DEBUG: // .b8 0 // EOM(2) +; DEBUG: // .b8 0 // EOM(3) +; DEBUG: // } +; DEBUG: // .section .debug_info +; DEBUG: // { +; DEBUG: // .b32 50 // Length of Unit +; DEBUG: // .b8 2 // DWARF version number +; DEBUG: // .b8 0 +; DEBUG: // .b32 .debug_abbrev // Offset Into Abbrev. Section +; DEBUG: // .b8 8 // Address Size (in bytes) +; DEBUG: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit +; DEBUG: // .b8 0 // DW_AT_producer +; DEBUG: // .b8 4 // DW_AT_language +; DEBUG: // .b8 0 +; DEBUG: // .b8 98 // DW_AT_name +; DEBUG: // .b8 97 +; DEBUG: // .b8 114 +; DEBUG: // .b8 46 +; DEBUG: // .b8 99 +; DEBUG: // .b8 117 +; DEBUG: // .b8 0 +; DEBUG: // .b32 .debug_line // DW_AT_stmt_list +; DEBUG: // .b8 47 // DW_AT_comp_dir +; DEBUG: // .b8 115 +; DEBUG: // .b8 111 +; DEBUG: // .b8 117 +; DEBUG: // .b8 114 +; DEBUG: // .b8 99 +; DEBUG: // .b8 101 +; DEBUG: // .b8 47 +; DEBUG: // .b8 100 +; DEBUG: // .b8 105 +; DEBUG: // .b8 114 +; DEBUG: // .b8 0 +; DEBUG: // .b64 Lfunc_begin0 // DW_AT_low_pc +; DEBUG: // .b64 Lfunc_end1 // DW_AT_high_pc +; DEBUG: // } +; DEBUG: // .section .debug_macinfo +; DEBUG: // { +; DEBUG: // .b8 0 // End Of Macro List Mark +; DEBUG: // } !llvm.dbg.cu = !{!0} !llvm.module.flags = !{!8, !9} Index: tools/llc/llc.cpp =================================================================== --- tools/llc/llc.cpp +++ tools/llc/llc.cpp @@ -80,6 +80,11 @@ cl::desc("Preserve Comments in outputted assembly"), cl::init(true)); +static cl::opt DebugDirectivesOnly( + "debug-directives-only", cl::Hidden, + cl::desc("Emit only debug directives, do not emit DWARF sections."), + cl::init(false)); + // Determine optimization level. static cl::opt OptLevel("O", @@ -442,6 +447,7 @@ Options.MCOptions.PreserveAsmComments = PreserveComments; Options.MCOptions.IASSearchPaths = IncludeDirs; Options.MCOptions.SplitDwarfFile = SplitDwarfFile; + Options.MCOptions.DebugDirectivesOnly = DebugDirectivesOnly; std::unique_ptr Target(TheTarget->createTargetMachine( TheTriple.getTriple(), CPUStr, FeaturesStr, Options, getRelocModel(),