Revert "[DEBUGINFO, NVPTX]DO not emit ',debug' option if no debug info or only debug directives are requested."

This reverts commit r345972. Need to update the description + possibly
to update the patch itself after discussion with Eric Christofer.

llvm-svn: 346508
This commit is contained in:
Alexey Bataev 2018-11-09 16:22:35 +00:00
parent 60087383d9
commit 336616d4e5
4 changed files with 4 additions and 78 deletions

View File

@ -25,12 +25,6 @@ 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);
}
@ -88,7 +82,9 @@ void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
OS << "//\t}\n";
if (isDwarfSection(FI, Section)) {
// Emit DWARF .file directives in the outermost scope.
outputDwarfFileDirectives();
for (const std::string &S : DwarfFiles)
getStreamer().EmitRawText(S.data());
DwarfFiles.clear();
OS << "//\t.section";
Section->PrintSwitchToSection(*getStreamer().getContext().getAsmInfo(),
FI->getTargetTriple(), OS, SubSection);

View File

@ -24,9 +24,6 @@ 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

View File

@ -16,7 +16,6 @@
#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"
@ -881,22 +880,8 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
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() && HasFullDebugInfo)
if (MMI && MMI->hasDebugInfo())
O << "//, debug";
O << "\n";
@ -953,10 +938,6 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) {
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);

View File

@ -1,48 +0,0 @@
; 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)