1 ; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s 2 3 ; // Bitcode int this test case is reduced version of compiled code below: 4 ;extern "C" { 5 ;#line 1 "/source/dir/foo.h" 6 ;__device__ void foo() {} 7 ;#line 2 "/source/dir/bar.cu" 8 ;__device__ void bar() {} 9 ;} 10 11 ; CHECK: .target sm_{{[0-9]+}}//, debug 12 13 ; CHECK: .visible .func foo() 14 ; CHECK: .loc [[FOO:[0-9]+]] 1 31 15 ; CHECK: ret; 16 ; CHECK: .visible .func bar() 17 ; CHECK: .loc [[BAR:[0-9]+]] 2 31 18 ; CHECK: ret; 19 20 define void @foo() !dbg !4 { 21 bb: 22 ret void, !dbg !10 23 } 24 25 define void @bar() !dbg !7 { 26 bb: 27 ret void, !dbg !11 28 } 29 30 ; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h" 31 ; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu" 32 ; CHECK: // .section .debug_abbrev 33 ; CHECK: // { 34 ; CHECK: // .b8 1 // Abbreviation Code 35 ; CHECK: // .b8 17 // DW_TAG_compile_unit 36 ; CHECK: // .b8 0 // DW_CHILDREN_no 37 ; CHECK: // .b8 37 // DW_AT_producer 38 ; CHECK: // .b8 8 // DW_FORM_string 39 ; CHECK: // .b8 19 // DW_AT_language 40 ; CHECK: // .b8 5 // DW_FORM_data2 41 ; CHECK: // .b8 3 // DW_AT_name 42 ; CHECK: // .b8 8 // DW_FORM_string 43 ; CHECK: // .b8 16 // DW_AT_stmt_list 44 ; CHECK: // .b8 6 // DW_FORM_data4 45 ; CHECK: // .b8 27 // DW_AT_comp_dir 46 ; CHECK: // .b8 8 // DW_FORM_string 47 ; CHECK: // .b8 17 // DW_AT_low_pc 48 ; CHECK: // .b8 1 // DW_FORM_addr 49 ; CHECK: // .b8 18 // DW_AT_high_pc 50 ; CHECK: // .b8 1 // DW_FORM_addr 51 ; CHECK: // .b8 0 // EOM(1) 52 ; CHECK: // .b8 0 // EOM(2) 53 ; CHECK: // .b8 0 // EOM(3) 54 ; CHECK: // } 55 ; CHECK: // .section .debug_info 56 ; CHECK: // { 57 ; CHECK: // .b32 50 // Length of Unit 58 ; CHECK: // .b8 2 // DWARF version number 59 ; CHECK: // .b8 0 60 ; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section 61 ; CHECK: // .b8 8 // Address Size (in bytes) 62 ; CHECK: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit 63 ; CHECK: // .b8 0 // DW_AT_producer 64 ; CHECK: // .b8 4 // DW_AT_language 65 ; CHECK: // .b8 0 66 ; CHECK: // .b8 98 // DW_AT_name 67 ; CHECK: // .b8 97 68 ; CHECK: // .b8 114 69 ; CHECK: // .b8 46 70 ; CHECK: // .b8 99 71 ; CHECK: // .b8 117 72 ; CHECK: // .b8 0 73 ; CHECK: // .b32 .debug_line // DW_AT_stmt_list 74 ; CHECK: // .b8 47 // DW_AT_comp_dir 75 ; CHECK: // .b8 115 76 ; CHECK: // .b8 111 77 ; CHECK: // .b8 117 78 ; CHECK: // .b8 114 79 ; CHECK: // .b8 99 80 ; CHECK: // .b8 101 81 ; CHECK: // .b8 47 82 ; CHECK: // .b8 100 83 ; CHECK: // .b8 105 84 ; CHECK: // .b8 114 85 ; CHECK: // .b8 0 86 ; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc 87 ; CHECK: // .b64 Lfunc_end1 // DW_AT_high_pc 88 ; CHECK: // } 89 ; CHECK: // .section .debug_macinfo 90 ; CHECK: // { 91 ; CHECK: // .b8 0 // End Of Macro List Mark 92 ; CHECK: // } 93 94 !llvm.dbg.cu = !{!0} 95 !llvm.module.flags = !{!8, !9} 96 97 !0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2) 98 !1 = !DIFile(filename: "bar.cu", directory: "/source/dir") 99 !2 = !{} 100 !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) 101 !5 = !DIFile(filename: "foo.h", directory: "/source/dir") 102 !6 = !DISubroutineType(types: !2) 103 !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) 104 !8 = !{i32 2, !"Dwarf Version", i32 2} 105 !9 = !{i32 2, !"Debug Info Version", i32 3} 106 !10 = !DILocation(line: 1, column: 31, scope: !4) 107 !11 = !DILocation(line: 2, column: 31, scope: !7) 108
