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

[Artem Belevich via llvm-commits]

tra updated this revision to Diff 47531.
tra added a comment.

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


http://reviews.llvm.org/D17086

Files:
  lib/Target/NVPTX/NVPTXAsmPrinter.cpp
  test/CodeGen/NVPTX/debug-file-loc.ll

Index: test/CodeGen/NVPTX/debug-file-loc.ll
===================================================================
--- /dev/null
+++ 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)
Index: lib/Target/NVPTX/NVPTXAsmPrinter.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -798,6 +798,7 @@
     if (filenameMap.find(Filename) != filenameMap.end())
       continue;
     filenameMap[Filename] = i;
+    OutStreamer->EmitDwarfFileDirective(i, "", Filename);
     ++i;
   }
 }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D17086.47531.patch
Type: text/x-patch
Size: 2148 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20160210/34bf3987/attachment.bin>