[DEBUGINFO, NVPTX]Emit last debugging directives.

[ABataev]

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.

[ABataev]

Fixed comment in the test.

[echristo]

LGTM.

lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h

 class NVPTXTargetStreamer : public MCTargetStreamer {
 private:
   SmallVector<std::string, 4> DwarfFiles;
 
 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
   /// (http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file),
   /// The .file directive is allowed only in the outermost scope, i.e., at the
   /// same level as kernel and device function declarations. Also, the order of
   /// the .loc and .file directive does not matter, .file directives may follow
   /// the .loc directives where the file is referenced.

lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp

 
 //
 // NVPTXTargetStreamer Implemenation
 //
 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);
 }
 
 static bool isDwarfSection(const MCObjectFileInfo *FI,
                            const MCSection *Section) {
   // FIXME: the checks for the DWARF sections are very fragile and should be
   // fixed up in a followup patch.
void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
   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";
   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);
     // DWARF sections are enclosed into braces - emit the open one.
     OS << "//\t{\n";
   }
 }
 

lib/Target/NVPTX/NVPTXAsmPrinter.cpp

 // of machine-dependent LLVM code to NVPTX assembly language.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTXAsmPrinter.h"
 #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"
 #include "NVPTXRegisterInfo.h"
 #include "NVPTXSubtarget.h"
 #include "NVPTXTargetMachine.h"
 #include "NVPTXUtilities.h"
 #include "cl_common_defines.h"
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}");
 
+  // Output last DWARF .file directives, if any.
+  static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
+      ->outputDwarfFileDirectives();
+
   return ret;
 
   //bool Result = AsmPrinter::doFinalization(M);
   // Instead of calling the parents doFinalization, we may
   // clone parents doFinalization and customize here.
   // Currently, we if NVISA out the EmitGlobals() in
   // parent's doFinalization, which is too intrusive.
   //

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

+; 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)