mirror of
https://github.com/RPCS3/llvm-mirror.git
synced 2024-10-21 20:12:56 +02:00
[DEBUGINFO, NVPTX]Emit last debugging directives.
Summary: We may end up with not emitted debug directives at the end of the module emission. Patch fixes this problem emitting those last directives the end of the module emission. Reviewers: echristo Subscribers: jholewinski, llvm-commits Differential Revision: https://reviews.llvm.org/D54320 llvm-svn: 348495
This commit is contained in:
parent
95ee0bc02b
commit
52ecd327f2
@ -25,6 +25,12 @@ 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);
|
||||
}
|
||||
@ -82,9 +88,7 @@ void NVPTXTargetStreamer::changeSection(const MCSection *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);
|
||||
|
@ -24,6 +24,9 @@ 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
|
||||
|
@ -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"
|
||||
@ -938,6 +939,10 @@ 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);
|
||||
|
48
test/DebugInfo/NVPTX/debug-file-loc-only.ll
Normal file
48
test/DebugInfo/NVPTX/debug-file-loc-only.ll
Normal file
@ -0,0 +1,48 @@
|
||||
; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
|
||||
|
||||
; // Bitcode in 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
|
||||
|
||||
; 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)
|
Loading…
Reference in New Issue
Block a user