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

[tra]

Summary

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

[dblaikie]

Needs a test case, presumably?

[jlebar]

Oh cool, there are no tests for this asm printer at all?

[jholewinski]

We do not have any tests for debug support, since it was never properly implemented. We have plenty of tests for other parts of the NVPTX AsmPrinter. Or are you referring to something else?

As for the review, the change looks good, but a test should be included.

[jlebar]

In D17086#348839, @jholewinski wrote:

We do not have any tests for debug support, since it was never properly implemented. We have plenty of tests for other parts of the NVPTX AsmPrinter. Or are you referring to something else?

Oh, no, I just fail at grep. For some reason I was looking for non-.ll files.

[tra]

Added test case.

[dblaikie]

Looks good - but please simplify the test a bit further and include the basic C code that produced the IR (I realize the IR is fairly simple, but it's handy to have what the code is modelling, the metadata format we have currently is still fairly verbose/hard to read)

The test could be simplier if it were C rather than C++ (no name mangling) and just void functions taking no parameters, I would imagine.

void foo() {
}
#line 1 other.c
void bar() {
}

Something like that, I would imagine.

Separately, this code seems strange - why does NVPTXAsmPrinter::recordAndEmitFilenames exist? It seems insufficient (what about file names on types? (DW_AT_decl_file on a user-defined type) why doesn't code like this exist for any other target? etc)

[tra]

In D17086#349010, @dblaikie wrote:

I'll simplify the test case.

Separately, this code seems strange - why does NVPTXAsmPrinter::recordAndEmitFilenames exist? It seems insufficient (what about file names on types? (DW_AT_decl_file on a user-defined type) why doesn't code like this exist for any other target? etc)

At the moment NVPTX only supports emitting line info and never produces any other debugging directives or data. I guess recordAndEmitFilenames() was the bare minimum needed to get line info working.

[tra]

Simplified test case and added original source used to create bitcode.

lib/Target/NVPTX/NVPTXAsmPrinter.cpp

for (DISubprogram *SP : DbgFinder.subprograms()) {
     SmallString<128> FullPathName = Dirname;
     if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
       sys::path::append(FullPathName, Filename);
       Filename = FullPathName;
     }
     if (filenameMap.find(Filename) != filenameMap.end())
       continue;
     filenameMap[Filename] = i;
+    OutStreamer->EmitDwarfFileDirective(i, "", Filename);
     ++i;
   }
 }
 
 static bool isEmptyXXStructor(GlobalVariable *GV) {
   if (!GV) return true;
   const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
   if (!InitList) return true;  // Not an array; we don't know how to parse.

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

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