[DEBUGINFO, NVPTX] Disable emission of ',debug' option if only debug directives are allowed.

[ABataev]

Summary

If the output of debug directives only is requested, we should drop
emission of ',debug' option from the target directive. Required for
supporting of nvprof profiler.

[ABataev]

Updated and reworked to latest changes

[ABataev]

Ping

[echristo]

Do you mean basically just line tables or something else? (Minus the need in normal dwarf for a minimal cu to associate the line table with). Also, how does this handle inlining etc?

[ABataev]

In D46061#1224055, @echristo wrote:

Do you mean basically just line tables or something else? (Minus the need in normal dwarf for a minimal cu to associate the line table with). Also, how does this handle inlining etc?

When we have full debug info, we need to add debug option to the target directive. When we do not emit debug info or emit just debug directives (i.e. only lineinfo), this options should not be added to the target directive. Plus we need to render all the debug directives (.line|.file) to be able to get the correct data from the profiler.
The inlining is handled only in terms of files/location debug directives, nothing else.

[ABataev]

Ping!

[ABataev]

Removed an option, used emission kind for the check of the full debug info.

[probinson]

LGTM

test/DebugInfo/NVPTX/debug-file-loc-only.ll

3 ↗

(On Diff #172164)

Typo: "Bitcode in this test..." (not "int").

[dblaikie]

Looks like Eric still had some questions here? Or at least hadn't commented
about whether the answers/changes addressed his concerns?

[probinson]

In D46061#1285431, @dblaikie wrote:

Looks like Eric still had some questions here? Or at least hadn't commented
about whether the answers/changes addressed his concerns?

My bad. It looked answered to me, but yes should have waited on his confirmation.

[echristo]

Some inline comments - I'm still not happy with some of how this patch is and would like to see some changes and elaborations of how we split things out. Mostly it's bikeshed naming things, but the current state is a bit more confusing than without.

Thanks!

-eric

llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

888–890 ↗

(On Diff #172350)

This seems to conflict with what you said above:

"When we have full debug info, we need to add debug option to the target directive. When we do not emit debug info or emit just debug directives (i.e. only lineinfo),"

as LineTablesOnly is line tables. If we need to split line tables into "things that handle inlining and things that don't" we should handle that discrepancy, but DebugDirectivesOnly isn't very helpful as a descriptor of what the compiler should be doing and paying attention to.

956–958 ↗

(On Diff #172350)

What's going on here?

[ABataev]

Update after revert.

[ABataev]

Removed extra functionality from the patch.

[ABataev]

Fixed test comment.

[ABataev]

llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

888–890 ↗

(On Diff #172350)

Yes, because I reworked the patch for the frontend. LineTablesOnly is considered as normal debug info option now. Plus, even before, when we saw -line-table-only option for NVPTX target, we still emitted DICompileUnit::DebugDirectivesOnly debug info kind. So, this is ok even for the previous description.

956–958 ↗

(On Diff #172350)

I split this patch into 2 different patches.

include/llvm/MC/MCTargetOptions.h

public:
   bool MCPIECopyRelocations : 1;
   bool ShowMCEncoding : 1;
   bool ShowMCInst : 1;
   bool AsmVerbose : 1;
 
   /// 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;
   std::string SplitDwarfFile;
 
   /// Additional paths to search for `.include` directives when using the
   /// integrated assembler.
   std::vector<std::string> IASSearchPaths;

lib/CodeGen/AsmPrinter/DwarfCompileUnit.cpp

if (CURanges.empty() || !SameAsPrevCU ||
     CURanges.push_back(Range);
     return;
   }
 
   CURanges.back().setEnd(Range.getEnd());
 }
 
 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();
   if (DD->useSectionsAsReferences()) {
     LineTableStartSym = TLOF.getDwarfLineSection()->getBeginSymbol();
   } else {
     LineTableStartSym =
         Asm->OutStreamer->getDwarfLineTableSymbol(getUniqueID());

lib/CodeGen/AsmPrinter/DwarfDebug.cpp

if (DisableDebugInfoPrinting)
     return;
 
   const Module *M = MMI->getModule();
 
   unsigned NumDebugCUs = std::distance(M->debug_compile_units_begin(),
                                        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<DIGlobalVariable *, SmallVector<DwarfCompileUnit::GlobalExpr, 1>>
       GVMap;
   for (const GlobalVariable &Global : M->globals()) {
     SmallVector<DIGlobalVariableExpression *, 1> GVs;
     Global.getDebugInfo(GVs);
     for (auto *GVE : GVs)
       GVMap[GVE->getVariable()].push_back({&Global, GVE->getExpression()});
void DwarfDebug::endModule() {
   assert(CurMI == nullptr);
 
   // If we aren't actually generating debug info (check beginModule -
   // conditionalized on !DisableDebugInfoPrinting and the presence of the
   // llvm.dbg.cu metadata node)
   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();
 
   emitDebugStr();
 
   if (useSplitDwarf())
     emitDebugLocDWO();
   else

lib/MC/MCTargetOptions.cpp

 using namespace llvm;
 
 MCTargetOptions::MCTargetOptions()
     : SanitizeAddress(false), MCRelaxAll(false), MCNoExecStack(false),
       MCFatalWarnings(false), MCNoWarn(false), MCNoDeprecatedWarn(false),
       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;
 }

lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h

 class NVPTXTargetStreamer : public MCTargetStreamer {
 private:
   SmallVector<std::string, 4> DwarfFiles;
 
 public:
   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
   /// (http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file),
   /// The .file directive is allowed only in the outermost scope, i.e., at the
   /// same level as kernel and device function declarations. Also, the order of
   /// the .loc and .file directive does not matter, .file directives may follow
   /// the .loc directives where the file is referenced.

lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp

 
 //
 // NVPTXTargetStreamer Implemenation
 //
 NVPTXTargetStreamer::NVPTXTargetStreamer(MCStreamer &S) : MCTargetStreamer(S) {}
 
 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);
 }
 
 static bool isDwarfSection(const MCObjectFileInfo *FI,
                            const MCSection *Section) {
   // FIXME: the checks for the DWARF sections are very fragile and should be
   // fixed up in a followup patch.
void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
   assert(!SubSection && "SubSection is not null!");
   const MCObjectFileInfo *FI = getStreamer().getContext().getObjectFileInfo();
   // FIXME: remove comment once debug info is properly supported.
   // Emit closing brace for DWARF sections only.
   if (isDwarfSection(FI, CurSection))
     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);
     // DWARF sections are enclosed into braces - emit the open one.
     OS << "//\t{\n";
   }
 }

lib/Target/NVPTX/NVPTXAsmPrinter.cpp

 // of machine-dependent LLVM code to NVPTX assembly language.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTXAsmPrinter.h"
 #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"
 #include "NVPTXRegisterInfo.h"
 #include "NVPTXSubtarget.h"
 #include "NVPTXTargetMachine.h"
 #include "NVPTXUtilities.h"
 #include "cl_common_defines.h"
void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
   O << ".target ";
   O << STI.getTargetName();
 
   const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
   if (NTM.getDrvInterface() == NVPTX::NVCL)
     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";
 
   O << ".address_size ";
   if (NTM.is64Bit())
     O << "64";
   else
bool NVPTXAsmPrinter::doFinalization(Module &M) {
   clearAnnotationCache(&M);
 
   delete[] gv_array;
   // FIXME: remove comment once debug info is properly supported.
   // Close the last emitted section
   if (HasDebugInfo)
     OutStreamer->EmitRawText("//\t}");
 
+  // Output last DWARF .file directives, if any.
+  static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
+      ->outputDwarfFileDirectives();
+
   return ret;
 
   //bool Result = AsmPrinter::doFinalization(M);
   // Instead of calling the parents doFinalization, we may
   // clone parents doFinalization and customize here.
   // Currently, we if NVISA out the EmitGlobals() in
   // parent's doFinalization, which is too intrusive.
   //

test/DebugInfo/Generic/directives-only.ll

+; 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)

test/DebugInfo/NVPTX/debug-file-loc.ll

-; 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" {
 ;#line 1 "/source/dir/foo.h"
 ;__device__ void foo() {}
 ;#line 2 "/source/dir/bar.cu"
 ;__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
 ; 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: // .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}
 
 !0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, 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, variables: !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, variables: !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)

tools/llc/llc.cpp

 NoIntegratedAssembler("no-integrated-as", cl::Hidden,
                       cl::desc("Disable integrated assembler"));
 
 static cl::opt<bool>
     PreserveComments("preserve-as-comments", cl::Hidden,
                      cl::desc("Preserve Comments in outputted assembly"),
                      cl::init(true));
 
+static cl::opt<bool> 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<char>
 OptLevel("O",
          cl::desc("Optimization level. [-O0, -O1, -O2, or -O3] "
                   "(default = '-O2')"),
          cl::Prefix,
          cl::ZeroOrMore,
          cl::init(' '));
static int compileModule(char **argv, LLVMContext &Context) {
   TargetOptions Options = InitTargetOptionsFromCodeGenFlags();
   Options.DisableIntegratedAS = NoIntegratedAssembler;
   Options.MCOptions.ShowMCEncoding = ShowMCEncoding;
   Options.MCOptions.MCUseDwarfDirectory = EnableDwarfDirectory;
   Options.MCOptions.AsmVerbose = AsmVerbose;
   Options.MCOptions.PreserveAsmComments = PreserveComments;
   Options.MCOptions.IASSearchPaths = IncludeDirs;
   Options.MCOptions.SplitDwarfFile = SplitDwarfFile;
+  Options.MCOptions.DebugDirectivesOnly = DebugDirectivesOnly;
 
   std::unique_ptr<TargetMachine> Target(TheTarget->createTargetMachine(
       TheTriple.getTriple(), CPUStr, FeaturesStr, Options, getRelocModel(),
       getCodeModel(), OLvl));
 
   assert(Target && "Could not allocate target machine!");
 
   // If we don't have a module then just exit now. We do this down