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<DIGlobalVariable *, SmallVector<DwarfCompileUnit::GlobalExpr, 1>>
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<NVPTXTargetStreamer *>(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<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",
@@ -442,6 +447,7 @@
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(),