Skip to content

Commit

Permalink
[DEBUGINFO, NVPTX] Enable support for the debug info on NVPTX target.
Browse files Browse the repository at this point in the history 
Enable full support for the debug info. Recommit to fix the emission of
the not required closing brace.

Differential revision: https://reviews.llvm.org/D46189

llvm-svn: 351972
  • Loading branch information
@alexey-bataev
alexey-bataev committed Jan 23, 2019
1 parent 8681a3b commit fe0b356
Show file tree
Hide file tree
Showing 10 changed files with 5,285 additions and 5,261 deletions.
9 changes: 4 additions & 5 deletions llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -37,12 +37,11 @@ NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) {
HiddenDeclarationVisibilityAttr = HiddenVisibilityAttr = MCSA_Invalid;
ProtectedVisibilityAttr = MCSA_Invalid;

// FIXME: remove comment once debug info is properly supported.
Data8bitsDirective = "// .b8 ";
Data8bitsDirective = ".b8 ";
Data16bitsDirective = nullptr; // not supported
Data32bitsDirective = "// .b32 ";
Data64bitsDirective = "// .b64 ";
ZeroDirective = "// .b8";
Data32bitsDirective = ".b32 ";
Data64bitsDirective = ".b64 ";
ZeroDirective = ".b8";
AsciiDirective = nullptr; // not supported
AscizDirective = nullptr; // not supported
SupportsQuotedNames = false;
Expand Down
11 changes: 8 additions & 3 deletions llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,11 @@ void NVPTXTargetStreamer::outputDwarfFileDirectives() {
DwarfFiles.clear();
}

void NVPTXTargetStreamer::closeLastSection() {
if (HasSections)
getStreamer().EmitRawText("\t}");
}

void NVPTXTargetStreamer::emitDwarfFileDirective(StringRef Directive) {
DwarfFiles.emplace_back(Directive);
}
Expand Down Expand Up @@ -81,18 +86,18 @@ void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
raw_ostream &OS) {
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";
OS << "\t}\n";
if (isDwarfSection(FI, Section)) {
// Emit DWARF .file directives in the outermost scope.
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";
OS << "\t{\n";
HasSections = true;
}
}

Expand Down
3 changes: 3 additions & 0 deletions llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,13 +18,16 @@ class MCSection;
class NVPTXTargetStreamer : public MCTargetStreamer {
private:
SmallVector<std::string, 4> DwarfFiles;
bool HasSections = false;

public:
NVPTXTargetStreamer(MCStreamer &S);
~NVPTXTargetStreamer() override;

/// Outputs the list of the DWARF '.file' directives to the streamer.
void outputDwarfFileDirectives();
/// Close last section.
void closeLastSection();

/// Record DWARF file directives for later output.
/// According to PTX ISA, CUDA Toolkit documentation, 11.5.3. Debugging
Expand Down
7 changes: 3 additions & 4 deletions llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -901,9 +901,8 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
if (HasFullDebugInfo)
break;
}
// FIXME: remove comment once debug info is properly supported.
if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
O << "//, debug";
O << ", debug";

O << "\n";

Expand Down Expand Up @@ -954,10 +953,10 @@ 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}");
static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
->closeLastSection();

// Output last DWARF .file directives, if any.
static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
Expand Down
280 changes: 140 additions & 140 deletions llvm/test/DebugInfo/NVPTX/cu-range-hole.ll
View file

Large diffs are not rendered by default.

328 changes: 164 additions & 164 deletions llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll
View file

Large diffs are not rendered by default.

18 changes: 18 additions & 0 deletions llvm/test/DebugInfo/NVPTX/debug-empty.ll
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s

; CHECK: .target sm_{{[0-9]+$}}
; CHECK-NOT: }

!llvm.dbg.cu = !{!0}
!llvm.module.flags = !{!3, !4, !5, !6, !7}
!llvm.ident = !{!8}

!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "clang version 9.0.0 (trunk 351924) (llvm/trunk 351968)", isOptimized: false, runtimeVersion: 0, emissionKind: DebugDirectivesOnly, enums: !2, nameTableKind: None)
!1 = !DIFile(filename: "new.cc", directory: "/test")
!2 = !{}
!3 = !{i32 2, !"Dwarf Version", i32 2}
!4 = !{i32 2, !"Debug Info Version", i32 3}
!5 = !{i32 1, !"wchar_size", i32 4}
!6 = !{i32 4, !"nvvm-reflect-ftz", i32 0}
!7 = !{i32 7, !"PIC Level", i32 2}
!8 = !{!"clang version 9.0.0 (trunk 351924) (llvm/trunk 351968)"}
94 changes: 47 additions & 47 deletions llvm/test/DebugInfo/NVPTX/debug-file-loc.ll
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,7 @@
;__device__ void bar() {}
;}

; CHECK: .target sm_{{[0-9]+}}//, debug
; CHECK: .target sm_{{[0-9]+}}, debug

; CHECK: .visible .func foo()
; CHECK: .loc [[FOO:[0-9]+]] 1 31
Expand All @@ -29,52 +29,52 @@ bb:

; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h"
; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu"
; CHECK: // .section .debug_abbrev
; CHECK-NEXT: // {
; CHECK-NEXT: // .b8 1 // Abbreviation Code
; CHECK-NEXT: // .b8 17 // DW_TAG_compile_unit
; CHECK-NEXT: // .b8 0 // DW_CHILDREN_no
; CHECK-NEXT: // .b8 37 // DW_AT_producer
; CHECK-NEXT: // .b8 8 // DW_FORM_string
; CHECK-NEXT: // .b8 19 // DW_AT_language
; CHECK-NEXT: // .b8 5 // DW_FORM_data2
; CHECK-NEXT: // .b8 3 // DW_AT_name
; CHECK-NEXT: // .b8 8 // DW_FORM_string
; CHECK-NEXT: // .b8 16 // DW_AT_stmt_list
; CHECK-NEXT: // .b8 6 // DW_FORM_data4
; CHECK-NEXT: // .b8 27 // DW_AT_comp_dir
; CHECK-NEXT: // .b8 8 // DW_FORM_string
; CHECK-NEXT: // .b8 17 // DW_AT_low_pc
; CHECK-NEXT: // .b8 1 // DW_FORM_addr
; CHECK-NEXT: // .b8 18 // DW_AT_high_pc
; CHECK-NEXT: // .b8 1 // DW_FORM_addr
; CHECK-NEXT: // .b8 0 // EOM(1)
; CHECK-NEXT: // .b8 0 // EOM(2)
; CHECK-NEXT: // .b8 0 // EOM(3)
; CHECK-NEXT: // }
; CHECK-NEXT: // .section .debug_info
; CHECK-NEXT: // {
; CHECK-NEXT: // .b32 50 // Length of Unit
; CHECK-NEXT: // .b8 2 // DWARF version number
; CHECK-NEXT: // .b8 0
; CHECK-NEXT: // .b32 .debug_abbrev // Offset Into Abbrev. Section
; CHECK-NEXT: // .b8 8 // Address Size (in bytes)
; CHECK-NEXT: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit
; CHECK-NEXT: // .b8 0 // DW_AT_producer
; CHECK-NEXT: // .b8 4 // DW_AT_language
; CHECK-NEXT: // .b8 0
; CHECK-NEXT: // .b8 98,97,114,46,99,117 // DW_AT_name
; CHECK-NEXT: // .b8 0
; CHECK-NEXT: // .b32 .debug_line // DW_AT_stmt_list
; CHECK-NEXT: // .b8 47,115,111,117,114,99,101,47,100,105,114 // DW_AT_comp_dir
; CHECK-NEXT: // .b8 0
; CHECK-NEXT: // .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK-NEXT: // .b64 Lfunc_end1 // DW_AT_high_pc
; CHECK-NEXT: // }
; CHECK-NEXT: // .section .debug_macinfo
; CHECK-NEXT: // {
; CHECK-NEXT: // .b8 0 // End Of Macro List Mark
; CHECK: // }
; CHECK: .section .debug_abbrev
; CHECK-NEXT: {
; CHECK-NEXT: .b8 1 // Abbreviation Code
; CHECK-NEXT: .b8 17 // DW_TAG_compile_unit
; CHECK-NEXT: .b8 0 // DW_CHILDREN_no
; CHECK-NEXT: .b8 37 // DW_AT_producer
; CHECK-NEXT: .b8 8 // DW_FORM_string
; CHECK-NEXT: .b8 19 // DW_AT_language
; CHECK-NEXT: .b8 5 // DW_FORM_data2
; CHECK-NEXT: .b8 3 // DW_AT_name
; CHECK-NEXT: .b8 8 // DW_FORM_string
; CHECK-NEXT: .b8 16 // DW_AT_stmt_list
; CHECK-NEXT: .b8 6 // DW_FORM_data4
; CHECK-NEXT: .b8 27 // DW_AT_comp_dir
; CHECK-NEXT: .b8 8 // DW_FORM_string
; CHECK-NEXT: .b8 17 // DW_AT_low_pc
; CHECK-NEXT: .b8 1 // DW_FORM_addr
; CHECK-NEXT: .b8 18 // DW_AT_high_pc
; CHECK-NEXT: .b8 1 // DW_FORM_addr
; CHECK-NEXT: .b8 0 // EOM(1)
; CHECK-NEXT: .b8 0 // EOM(2)
; CHECK-NEXT: .b8 0 // EOM(3)
; CHECK-NEXT: }
; CHECK-NEXT: .section .debug_info
; CHECK-NEXT: {
; CHECK-NEXT: .b32 50 // Length of Unit
; CHECK-NEXT: .b8 2 // DWARF version number
; CHECK-NEXT: .b8 0
; CHECK-NEXT: .b32 .debug_abbrev // Offset Into Abbrev. Section
; CHECK-NEXT: .b8 8 // Address Size (in bytes)
; CHECK-NEXT: .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit
; CHECK-NEXT: .b8 0 // DW_AT_producer
; CHECK-NEXT: .b8 4 // DW_AT_language
; CHECK-NEXT: .b8 0
; CHECK-NEXT: .b8 98,97,114,46,99,117 // DW_AT_name
; CHECK-NEXT: .b8 0
; CHECK-NEXT: .b32 .debug_line // DW_AT_stmt_list
; CHECK-NEXT: .b8 47,115,111,117,114,99,101,47,100,105,114 // DW_AT_comp_dir
; CHECK-NEXT: .b8 0
; CHECK-NEXT: .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK-NEXT: .b64 Lfunc_end1 // DW_AT_high_pc
; CHECK-NEXT: }
; CHECK-NEXT: .section .debug_macinfo
; CHECK-NEXT: {
; CHECK-NEXT: .b8 0 // End Of Macro List Mark
; CHECK: }

!llvm.dbg.cu = !{!0}
!llvm.module.flags = !{!8, !9}
Expand Down
Loading

0 comments on commit fe0b356

Please sign in to comment.