Skip to content
This repository was archived by the owner on Apr 23, 2020. It is now read-only.

Commit d77902a

Browse files
alexey-bataevalexey-bataev
committed
[DEBUGINFO, NVPTX] Enable support for the debug info on NVPTX target.
Summary: Enable full support for the debug info. Reviewers: echristo Subscribers: jholewinski, aprantl, JDevlieghere, llvm-commits Differential Revision: https://reviews.llvm.org/D46189 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@351846 91177308-0d34-0410-b5e6-96231b3b80d8
1 parent 27047c8 commit d77902a

File tree

8 files changed

+5257
-5261
lines changed

8 files changed

+5257
-5261
lines changed

‎lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp

+4-5
Original file line numberDiff line numberDiff line change
@@ -37,12 +37,11 @@ NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) {
3737
HiddenDeclarationVisibilityAttr = HiddenVisibilityAttr = MCSA_Invalid;
3838
ProtectedVisibilityAttr = MCSA_Invalid;
3939

40-
// FIXME: remove comment once debug info is properly supported.
41-
Data8bitsDirective = "// .b8 ";
40+
Data8bitsDirective = ".b8 ";
4241
Data16bitsDirective = nullptr; // not supported
43-
Data32bitsDirective = "// .b32 ";
44-
Data64bitsDirective = "// .b64 ";
45-
ZeroDirective = "// .b8";
42+
Data32bitsDirective = ".b32 ";
43+
Data64bitsDirective = ".b64 ";
44+
ZeroDirective = ".b8";
4645
AsciiDirective = nullptr; // not supported
4746
AscizDirective = nullptr; // not supported
4847
SupportsQuotedNames = false;

‎lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp

+2-3
Original file line numberDiff line numberDiff line change
@@ -81,18 +81,17 @@ void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
8181
raw_ostream &OS) {
8282
assert(!SubSection && "SubSection is not null!");
8383
const MCObjectFileInfo *FI = getStreamer().getContext().getObjectFileInfo();
84-
// FIXME: remove comment once debug info is properly supported.
8584
// Emit closing brace for DWARF sections only.
8685
if (isDwarfSection(FI, CurSection))
87-
OS << "//\t}\n";
86+
OS << "\t}\n";
8887
if (isDwarfSection(FI, Section)) {
8988
// Emit DWARF .file directives in the outermost scope.
9089
outputDwarfFileDirectives();
9190
OS << "//\t.section";
9291
Section->PrintSwitchToSection(*getStreamer().getContext().getAsmInfo(),
9392
FI->getTargetTriple(), OS, SubSection);
9493
// DWARF sections are enclosed into braces - emit the open one.
95-
OS << "//\t{\n";
94+
OS << "\t{\n";
9695
}
9796
}
9897

‎lib/Target/NVPTX/NVPTXAsmPrinter.cpp

+2-4
Original file line numberDiff line numberDiff line change
@@ -901,9 +901,8 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
901901
if (HasFullDebugInfo)
902902
break;
903903
}
904-
// FIXME: remove comment once debug info is properly supported.
905904
if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
906-
O << "//, debug";
905+
O << ", debug";
907906

908907
O << "\n";
909908

@@ -954,10 +953,9 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) {
954953
clearAnnotationCache(&M);
955954

956955
delete[] gv_array;
957-
// FIXME: remove comment once debug info is properly supported.
958956
// Close the last emitted section
959957
if (HasDebugInfo)
960-
OutStreamer->EmitRawText("//\t}");
958+
OutStreamer->EmitRawText("\t}");
961959

962960
// Output last DWARF .file directives, if any.
963961
static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())

‎test/DebugInfo/NVPTX/cu-range-hole.ll

+140-140
Large diffs are not rendered by default.

‎test/DebugInfo/NVPTX/dbg-declare-alloca.ll

+164-164
Large diffs are not rendered by default.

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

+47-47
Original file line numberDiff line numberDiff line change
@@ -8,7 +8,7 @@
88
;__device__ void bar() {}
99
;}
1010

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

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

3030
; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h"
3131
; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu"
32-
; CHECK: // .section .debug_abbrev
33-
; CHECK-NEXT: // {
34-
; CHECK-NEXT: // .b8 1 // Abbreviation Code
35-
; CHECK-NEXT: // .b8 17 // DW_TAG_compile_unit
36-
; CHECK-NEXT: // .b8 0 // DW_CHILDREN_no
37-
; CHECK-NEXT: // .b8 37 // DW_AT_producer
38-
; CHECK-NEXT: // .b8 8 // DW_FORM_string
39-
; CHECK-NEXT: // .b8 19 // DW_AT_language
40-
; CHECK-NEXT: // .b8 5 // DW_FORM_data2
41-
; CHECK-NEXT: // .b8 3 // DW_AT_name
42-
; CHECK-NEXT: // .b8 8 // DW_FORM_string
43-
; CHECK-NEXT: // .b8 16 // DW_AT_stmt_list
44-
; CHECK-NEXT: // .b8 6 // DW_FORM_data4
45-
; CHECK-NEXT: // .b8 27 // DW_AT_comp_dir
46-
; CHECK-NEXT: // .b8 8 // DW_FORM_string
47-
; CHECK-NEXT: // .b8 17 // DW_AT_low_pc
48-
; CHECK-NEXT: // .b8 1 // DW_FORM_addr
49-
; CHECK-NEXT: // .b8 18 // DW_AT_high_pc
50-
; CHECK-NEXT: // .b8 1 // DW_FORM_addr
51-
; CHECK-NEXT: // .b8 0 // EOM(1)
52-
; CHECK-NEXT: // .b8 0 // EOM(2)
53-
; CHECK-NEXT: // .b8 0 // EOM(3)
54-
; CHECK-NEXT: // }
55-
; CHECK-NEXT: // .section .debug_info
56-
; CHECK-NEXT: // {
57-
; CHECK-NEXT: // .b32 50 // Length of Unit
58-
; CHECK-NEXT: // .b8 2 // DWARF version number
59-
; CHECK-NEXT: // .b8 0
60-
; CHECK-NEXT: // .b32 .debug_abbrev // Offset Into Abbrev. Section
61-
; CHECK-NEXT: // .b8 8 // Address Size (in bytes)
62-
; CHECK-NEXT: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit
63-
; CHECK-NEXT: // .b8 0 // DW_AT_producer
64-
; CHECK-NEXT: // .b8 4 // DW_AT_language
65-
; CHECK-NEXT: // .b8 0
66-
; CHECK-NEXT: // .b8 98,97,114,46,99,117 // DW_AT_name
67-
; CHECK-NEXT: // .b8 0
68-
; CHECK-NEXT: // .b32 .debug_line // DW_AT_stmt_list
69-
; CHECK-NEXT: // .b8 47,115,111,117,114,99,101,47,100,105,114 // DW_AT_comp_dir
70-
; CHECK-NEXT: // .b8 0
71-
; CHECK-NEXT: // .b64 Lfunc_begin0 // DW_AT_low_pc
72-
; CHECK-NEXT: // .b64 Lfunc_end1 // DW_AT_high_pc
73-
; CHECK-NEXT: // }
74-
; CHECK-NEXT: // .section .debug_macinfo
75-
; CHECK-NEXT: // {
76-
; CHECK-NEXT: // .b8 0 // End Of Macro List Mark
77-
; CHECK: // }
32+
; CHECK: .section .debug_abbrev
33+
; CHECK-NEXT: {
34+
; CHECK-NEXT: .b8 1 // Abbreviation Code
35+
; CHECK-NEXT: .b8 17 // DW_TAG_compile_unit
36+
; CHECK-NEXT: .b8 0 // DW_CHILDREN_no
37+
; CHECK-NEXT: .b8 37 // DW_AT_producer
38+
; CHECK-NEXT: .b8 8 // DW_FORM_string
39+
; CHECK-NEXT: .b8 19 // DW_AT_language
40+
; CHECK-NEXT: .b8 5 // DW_FORM_data2
41+
; CHECK-NEXT: .b8 3 // DW_AT_name
42+
; CHECK-NEXT: .b8 8 // DW_FORM_string
43+
; CHECK-NEXT: .b8 16 // DW_AT_stmt_list
44+
; CHECK-NEXT: .b8 6 // DW_FORM_data4
45+
; CHECK-NEXT: .b8 27 // DW_AT_comp_dir
46+
; CHECK-NEXT: .b8 8 // DW_FORM_string
47+
; CHECK-NEXT: .b8 17 // DW_AT_low_pc
48+
; CHECK-NEXT: .b8 1 // DW_FORM_addr
49+
; CHECK-NEXT: .b8 18 // DW_AT_high_pc
50+
; CHECK-NEXT: .b8 1 // DW_FORM_addr
51+
; CHECK-NEXT: .b8 0 // EOM(1)
52+
; CHECK-NEXT: .b8 0 // EOM(2)
53+
; CHECK-NEXT: .b8 0 // EOM(3)
54+
; CHECK-NEXT: }
55+
; CHECK-NEXT: .section .debug_info
56+
; CHECK-NEXT: {
57+
; CHECK-NEXT: .b32 50 // Length of Unit
58+
; CHECK-NEXT: .b8 2 // DWARF version number
59+
; CHECK-NEXT: .b8 0
60+
; CHECK-NEXT: .b32 .debug_abbrev // Offset Into Abbrev. Section
61+
; CHECK-NEXT: .b8 8 // Address Size (in bytes)
62+
; CHECK-NEXT: .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit
63+
; CHECK-NEXT: .b8 0 // DW_AT_producer
64+
; CHECK-NEXT: .b8 4 // DW_AT_language
65+
; CHECK-NEXT: .b8 0
66+
; CHECK-NEXT: .b8 98,97,114,46,99,117 // DW_AT_name
67+
; CHECK-NEXT: .b8 0
68+
; CHECK-NEXT: .b32 .debug_line // DW_AT_stmt_list
69+
; CHECK-NEXT: .b8 47,115,111,117,114,99,101,47,100,105,114 // DW_AT_comp_dir
70+
; CHECK-NEXT: .b8 0
71+
; CHECK-NEXT: .b64 Lfunc_begin0 // DW_AT_low_pc
72+
; CHECK-NEXT: .b64 Lfunc_end1 // DW_AT_high_pc
73+
; CHECK-NEXT: }
74+
; CHECK-NEXT: .section .debug_macinfo
75+
; CHECK-NEXT: {
76+
; CHECK-NEXT: .b8 0 // End Of Macro List Mark
77+
; CHECK: }
7878

7979
!llvm.dbg.cu = !{!0}
8080
!llvm.module.flags = !{!8, !9}

0 commit comments

Comments
 (0)