[NVPTX] emit .file directives for files referenced by subprograms.

.. so .loc directives referring to those files work correctly.

Differential Revision: http://reviews.llvm.org/D17086

git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@260557 91177308-0d34-0410-b5e6-96231b3b80d8

lib/Target/NVPTX/NVPTXAsmPrinter.cpp

@@ -798,6 +798,7 @@ void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) {
     if (filenameMap.find(Filename) != filenameMap.end())
       continue;
     filenameMap[Filename] = i;
+    OutStreamer->EmitDwarfFileDirective(i, "", Filename);
     ++i;
   }
 }

test/CodeGen/NVPTX/debug-file-loc.ll

@@ -0,0 +1,44 @@
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
+
+; // Bitcode int 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: .file   1 "/source/dir/bar.cu"
+; CHECK: .file   2 "/source/dir/foo.h"
+
+; CHECK-LABEL: @foo
+define void @foo() !dbg !4 {
+bb:
+  ret void, !dbg !10
+}
+; CHECK: .loc 2 1
+; CHECK: ret
+
+; CHECK-LABEL: @bar
+define void @bar() !dbg !7 {
+bb:
+  ret void, !dbg !11
+}
+; CHECK: .loc 1 2
+; CHECK: ret
+
+!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: 2, enums: !2, subprograms: !3)
+!1 = !DIFile(filename: "bar.cu", directory: "/source/dir")
+!2 = !{}
+!3 = !{!4, !7}
+!4 = distinct !DISubprogram(name: "foo", scope: !5, file: !5, line: 1, type: !6, isLocal: false, isDefinition: true, scopeLine: 1, flags: DIFlagPrototyped, isOptimized: false, variables: !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, variables: !2)
+!8 = !{i32 2, !"Dwarf Version", i32 4}
+!9 = !{i32 2, !"Debug Info Version", i32 3}
+!10 = !DILocation(line: 1, column: 31, scope: !4)
+!11 = !DILocation(line: 2, column: 31, scope: !7)