llvm.org GIT mirror llvm / 12ed311
[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 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@348495 91177308-0d34-0410-b5e6-96231b3b80d8 Alexey Bataev 1 year, 11 months ago
4 changed file(s) with 63 addition(s) and 3 deletion(s). Raw diff Collapse all Expand all
2323 NVPTXTargetStreamer::NVPTXTargetStreamer(MCStreamer &S) : MCTargetStreamer(S) {}
2424
2525 NVPTXTargetStreamer::~NVPTXTargetStreamer() = default;
26
27 void NVPTXTargetStreamer::outputDwarfFileDirectives() {
28 for (const std::string &S : DwarfFiles)
29 getStreamer().EmitRawText(S.data());
30 DwarfFiles.clear();
31 }
2632
2733 void NVPTXTargetStreamer::emitDwarfFileDirective(StringRef Directive) {
2834 DwarfFiles.emplace_back(Directive);
8187 OS << "//\t}\n";
8288 if (isDwarfSection(FI, Section)) {
8389 // Emit DWARF .file directives in the outermost scope.
84 for (const std::string &S : DwarfFiles)
85 getStreamer().EmitRawText(S.data());
86 DwarfFiles.clear();
90 outputDwarfFileDirectives();
8791 OS << "//\t.section";
8892 Section->PrintSwitchToSection(*getStreamer().getContext().getAsmInfo(),
8993 FI->getTargetTriple(), OS, SubSection);
2323 NVPTXTargetStreamer(MCStreamer &S);
2424 ~NVPTXTargetStreamer() override;
2525
26 /// Outputs the list of the DWARF '.file' directives to the streamer.
27 void outputDwarfFileDirectives();
28
2629 /// Record DWARF file directives for later output.
2730 /// According to PTX ISA, CUDA Toolkit documentation, 11.5.3. Debugging
2831 /// Directives: .file
1515 #include "InstPrinter/NVPTXInstPrinter.h"
1616 #include "MCTargetDesc/NVPTXBaseInfo.h"
1717 #include "MCTargetDesc/NVPTXMCAsmInfo.h"
18 #include "MCTargetDesc/NVPTXTargetStreamer.h"
1819 #include "NVPTX.h"
1920 #include "NVPTXMCExpr.h"
2021 #include "NVPTXMachineFunctionInfo.h"
937938 if (HasDebugInfo)
938939 OutStreamer->EmitRawText("//\t}");
939940
941 // Output last DWARF .file directives, if any.
942 static_cast(OutStreamer->getTargetStreamer())
943 ->outputDwarfFileDirectives();
944
940945 return ret;
941946
942947 //bool Result = AsmPrinter::doFinalization(M);
0 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
1
2 ; // Bitcode in this test case is reduced version of compiled code below:
3 ;extern "C" {
4 ;#line 1 "/source/dir/foo.h"
5 ;__device__ void foo() {}
6 ;#line 2 "/source/dir/bar.cu"
7 ;__device__ void bar() {}
8 ;}
9
10 ; CHECK: .target sm_{{[0-9]+}}//, debug
11
12 ; CHECK: .visible .func foo()
13 ; CHECK: .loc [[FOO:[0-9]+]] 1 31
14 ; CHECK: ret;
15 ; CHECK: .visible .func bar()
16 ; CHECK: .loc [[BAR:[0-9]+]] 2 31
17 ; CHECK: ret;
18
19 define void @foo() !dbg !4 {
20 bb:
21 ret void, !dbg !10
22 }
23
24 define void @bar() !dbg !7 {
25 bb:
26 ret void, !dbg !11
27 }
28
29 ; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h"
30 ; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu"
31
32 ; CHECK-NOT: .section .debug{{.*}}
33
34 !llvm.dbg.cu = !{!0}
35 !llvm.module.flags = !{!8, !9}
36
37 !0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "", isOptimized: false, runtimeVersion: 0, emissionKind: DebugDirectivesOnly, enums: !2)
38 !1 = !DIFile(filename: "bar.cu", directory: "/source/dir")
39 !2 = !{}
40 !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)
41 !5 = !DIFile(filename: "foo.h", directory: "/source/dir")
42 !6 = !DISubroutineType(types: !2)
43 !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)
44 !8 = !{i32 2, !"Dwarf Version", i32 2}
45 !9 = !{i32 2, !"Debug Info Version", i32 3}
46 !10 = !DILocation(line: 1, column: 31, scope: !4)
47 !11 = !DILocation(line: 2, column: 31, scope: !7)