1; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
2
3; CHECK: .target sm_20, debug
4
5; CHECK: .visible .func use_dbg_declare()
6; CHECK: .local .align 8 .b8 __local_depot0[8];
7; CHECK: mov.u64 %SPL, __local_depot0;
8; CHECK: add.u64 %rd1, %SP, 0;
9; CHECK: .loc 1 5 3                   // t.c:5:3
10; CHECK: { // callseq 0, 0
11; CHECK: .reg .b32 temp_param_reg;
12; CHECK: .param .b64 param0;
13; CHECK: st.param.b64 [param0+0], %rd1;
14; CHECK: call.uni
15; CHECK: escape_foo,
16; CHECK: (
17; CHECK: param0
18; CHECK: );
19; CHECK: } // callseq 0
20; CHECK: .loc 1 6 1                   // t.c:6:1
21; CHECK: ret;
22; CHECK: }
23
24; CHECK: .file 1 "test{{(/|\\\\)}}t.c"
25
26; CHECK: .section .debug_abbrev
27; CHECK-NEXT: {
28; CHECK-NEXT: .b8 1                                // Abbreviation Code
29; CHECK-NEXT: .b8 17                               // DW_TAG_compile_unit
30; CHECK-NEXT: .b8 1                                // DW_CHILDREN_yes
31; CHECK-NEXT: .b8 37                               // DW_AT_producer
32; CHECK-NEXT: .b8 8                                // DW_FORM_string
33; CHECK-NEXT: .b8 19                               // DW_AT_language
34; CHECK-NEXT: .b8 5                                // DW_FORM_data2
35; CHECK-NEXT: .b8 3                                // DW_AT_name
36; CHECK-NEXT: .b8 8                                // DW_FORM_string
37; CHECK-NEXT: .b8 16                               // DW_AT_stmt_list
38; CHECK-NEXT: .b8 6                                // DW_FORM_data4
39; CHECK-NEXT: .b8 27                               // DW_AT_comp_dir
40; CHECK-NEXT: .b8 8                                // DW_FORM_string
41; CHECK-NEXT: .b8 17                               // DW_AT_low_pc
42; CHECK-NEXT: .b8 1                                // DW_FORM_addr
43; CHECK-NEXT: .b8 18                               // DW_AT_high_pc
44; CHECK-NEXT: .b8 1                                // DW_FORM_addr
45; CHECK-NEXT: .b8 0                                // EOM(1)
46; CHECK-NEXT: .b8 0                                // EOM(2)
47; CHECK-NEXT: .b8 2                                // Abbreviation Code
48; CHECK-NEXT: .b8 46                               // DW_TAG_subprogram
49; CHECK-NEXT: .b8 1                                // DW_CHILDREN_yes
50; CHECK-NEXT: .b8 17                               // DW_AT_low_pc
51; CHECK-NEXT: .b8 1                                // DW_FORM_addr
52; CHECK-NEXT: .b8 18                               // DW_AT_high_pc
53; CHECK-NEXT: .b8 1                                // DW_FORM_addr
54; CHECK-NEXT: .b8 64                               // DW_AT_frame_base
55; CHECK-NEXT: .b8 10                               // DW_FORM_block1
56; CHECK-NEXT: .b8 3                                // DW_AT_name
57; CHECK-NEXT: .b8 8                                // DW_FORM_string
58; CHECK-NEXT: .b8 58                               // DW_AT_decl_file
59; CHECK-NEXT: .b8 11                               // DW_FORM_data1
60; CHECK-NEXT: .b8 59                               // DW_AT_decl_line
61; CHECK-NEXT: .b8 11                               // DW_FORM_data1
62; CHECK-NEXT: .b8 39                               // DW_AT_prototyped
63; CHECK-NEXT: .b8 12                               // DW_FORM_flag
64; CHECK-NEXT: .b8 63                               // DW_AT_external
65; CHECK-NEXT: .b8 12                               // DW_FORM_flag
66; CHECK-NEXT: .b8 0                                // EOM(1)
67; CHECK-NEXT: .b8 0                                // EOM(2)
68; CHECK-NEXT: .b8 3                                // Abbreviation Code
69; CHECK-NEXT: .b8 52                               // DW_TAG_variable
70; CHECK-NEXT: .b8 0                                // DW_CHILDREN_no
71; CHECK-NEXT: .b8 51                               // DW_AT_address_class
72; CHECK-NEXT: .b8 11                               // DW_FORM_data1
73; CHECK-NEXT: .b8 2                                // DW_AT_location
74; CHECK-NEXT: .b8 10                               // DW_FORM_block1
75; CHECK-NEXT: .b8 3                                // DW_AT_name
76; CHECK-NEXT: .b8 8                                // DW_FORM_string
77; CHECK-NEXT: .b8 58                               // DW_AT_decl_file
78; CHECK-NEXT: .b8 11                               // DW_FORM_data1
79; CHECK-NEXT: .b8 59                               // DW_AT_decl_line
80; CHECK-NEXT: .b8 11                               // DW_FORM_data1
81; CHECK-NEXT: .b8 73                               // DW_AT_type
82; CHECK-NEXT: .b8 19                               // DW_FORM_ref4
83; CHECK-NEXT: .b8 0                                // EOM(1)
84; CHECK-NEXT: .b8 0                                // EOM(2)
85; CHECK-NEXT: .b8 4                                // Abbreviation Code
86; CHECK-NEXT: .b8 19                               // DW_TAG_structure_type
87; CHECK-NEXT: .b8 1                                // DW_CHILDREN_yes
88; CHECK-NEXT: .b8 3                                // DW_AT_name
89; CHECK-NEXT: .b8 8                                // DW_FORM_string
90; CHECK-NEXT: .b8 11                               // DW_AT_byte_size
91; CHECK-NEXT: .b8 11                               // DW_FORM_data1
92; CHECK-NEXT: .b8 58                               // DW_AT_decl_file
93; CHECK-NEXT: .b8 11                               // DW_FORM_data1
94; CHECK-NEXT: .b8 59                               // DW_AT_decl_line
95; CHECK-NEXT: .b8 11                               // DW_FORM_data1
96; CHECK-NEXT: .b8 0                                // EOM(1)
97; CHECK-NEXT: .b8 0                                // EOM(2)
98; CHECK-NEXT: .b8 5                                // Abbreviation Code
99; CHECK-NEXT: .b8 13                               // DW_TAG_member
100; CHECK-NEXT: .b8 0                                // DW_CHILDREN_no
101; CHECK-NEXT: .b8 3                                // DW_AT_name
102; CHECK-NEXT: .b8 8                                // DW_FORM_string
103; CHECK-NEXT: .b8 73                               // DW_AT_type
104; CHECK-NEXT: .b8 19                               // DW_FORM_ref4
105; CHECK-NEXT: .b8 58                               // DW_AT_decl_file
106; CHECK-NEXT: .b8 11                               // DW_FORM_data1
107; CHECK-NEXT: .b8 59                               // DW_AT_decl_line
108; CHECK-NEXT: .b8 11                               // DW_FORM_data1
109; CHECK-NEXT: .b8 56                               // DW_AT_data_member_location
110; CHECK-NEXT: .b8 10                               // DW_FORM_block1
111; CHECK-NEXT: .b8 0                                // EOM(1)
112; CHECK-NEXT: .b8 0                                // EOM(2)
113; CHECK-NEXT: .b8 6                                // Abbreviation Code
114; CHECK-NEXT: .b8 36                               // DW_TAG_base_type
115; CHECK-NEXT: .b8 0                                // DW_CHILDREN_no
116; CHECK-NEXT: .b8 3                                // DW_AT_name
117; CHECK-NEXT: .b8 8                                // DW_FORM_string
118; CHECK-NEXT: .b8 62                               // DW_AT_encoding
119; CHECK-NEXT: .b8 11                               // DW_FORM_data1
120; CHECK-NEXT: .b8 11                               // DW_AT_byte_size
121; CHECK-NEXT: .b8 11                               // DW_FORM_data1
122; CHECK-NEXT: .b8 0                                // EOM(1)
123; CHECK-NEXT: .b8 0                                // EOM(2)
124; CHECK-NEXT: .b8 0                                // EOM(3)
125; CHECK-NEXT: }
126; CHECK-NEXT: .section .debug_info
127; CHECK-NEXT: {
128; CHECK-NEXT: .b32 134                             // Length of Unit
129; CHECK-NEXT: .b8 2                                // DWARF version number
130; CHECK-NEXT: .b8 0
131; CHECK-NEXT: .b32 .debug_abbrev                   // Offset Into Abbrev. Section
132; CHECK-NEXT: .b8 8                                // Address Size (in bytes)
133; CHECK-NEXT: .b8 1                                // Abbrev [1] 0xb:0x7f DW_TAG_compile_unit
134; CHECK-NEXT: .b8 99                               // DW_AT_producer
135; CHECK-NEXT: .b8 108
136; CHECK-NEXT: .b8 97
137; CHECK-NEXT: .b8 110
138; CHECK-NEXT: .b8 103
139; CHECK-NEXT: .b8 0
140; CHECK-NEXT: .b8 12                               // DW_AT_language
141; CHECK-NEXT: .b8 0
142; CHECK-NEXT: .b8 116                              // DW_AT_name
143; CHECK-NEXT: .b8 46
144; CHECK-NEXT: .b8 99
145; CHECK-NEXT: .b8 0
146; CHECK-NEXT: .b32 .debug_line                     // DW_AT_stmt_list
147; CHECK-NEXT: .b8 116                              // DW_AT_comp_dir
148; CHECK-NEXT: .b8 101
149; CHECK-NEXT: .b8 115
150; CHECK-NEXT: .b8 116
151; CHECK-NEXT: .b8 0
152; CHECK-NEXT: .b64 Lfunc_begin0                    // DW_AT_low_pc
153; CHECK-NEXT: .b64 Lfunc_end0                      // DW_AT_high_pc
154; CHECK-NEXT: .b8 2                                // Abbrev [2] 0x31:0x3c DW_TAG_subprogram
155; CHECK-NEXT: .b64 Lfunc_begin0                    // DW_AT_low_pc
156; CHECK-NEXT: .b64 Lfunc_end0                      // DW_AT_high_pc
157; CHECK-NEXT: .b8 1                                // DW_AT_frame_base
158; CHECK-NEXT: .b8 156
159; CHECK-NEXT: .b8 117                              // DW_AT_name
160; CHECK-NEXT: .b8 115
161; CHECK-NEXT: .b8 101
162; CHECK-NEXT: .b8 95
163; CHECK-NEXT: .b8 100
164; CHECK-NEXT: .b8 98
165; CHECK-NEXT: .b8 103
166; CHECK-NEXT: .b8 95
167; CHECK-NEXT: .b8 100
168; CHECK-NEXT: .b8 101
169; CHECK-NEXT: .b8 99
170; CHECK-NEXT: .b8 108
171; CHECK-NEXT: .b8 97
172; CHECK-NEXT: .b8 114
173; CHECK-NEXT: .b8 101
174; CHECK-NEXT: .b8 0
175; CHECK-NEXT: .b8 1                                // DW_AT_decl_file
176; CHECK-NEXT: .b8 3                                // DW_AT_decl_line
177; CHECK-NEXT: .b8 1                                // DW_AT_prototyped
178; CHECK-NEXT: .b8 1                                // DW_AT_external
179; CHECK-NEXT: .b8 3                                // Abbrev [3] 0x58:0x14 DW_TAG_variable
180; CHECK-NEXT: .b8 6                                // DW_AT_address_class
181; CHECK-NEXT: .b8 9                                // DW_AT_location
182; CHECK-NEXT: .b8 3
183; CHECK-NEXT: .b64 __local_depot0
184; CHECK-NEXT: .b8 111                              // DW_AT_name
185; CHECK-NEXT: .b8 0
186; CHECK-NEXT: .b8 1                                // DW_AT_decl_file
187; CHECK-NEXT: .b8 4                                // DW_AT_decl_line
188; CHECK-NEXT: .b32 109                             // DW_AT_type
189; CHECK-NEXT: .b8 0                                // End Of Children Mark
190; CHECK-NEXT: .b8 4                                // Abbrev [4] 0x6d:0x15 DW_TAG_structure_type
191; CHECK-NEXT: .b8 70                               // DW_AT_name
192; CHECK-NEXT: .b8 111
193; CHECK-NEXT: .b8 111
194; CHECK-NEXT: .b8 0
195; CHECK-NEXT: .b8 4                                // DW_AT_byte_size
196; CHECK-NEXT: .b8 1                                // DW_AT_decl_file
197; CHECK-NEXT: .b8 1                                // DW_AT_decl_line
198; CHECK-NEXT: .b8 5                                // Abbrev [5] 0x75:0xc DW_TAG_member
199; CHECK-NEXT: .b8 120                              // DW_AT_name
200; CHECK-NEXT: .b8 0
201; CHECK-NEXT: .b32 130                             // DW_AT_type
202; CHECK-NEXT: .b8 1                                // DW_AT_decl_file
203; CHECK-NEXT: .b8 1                                // DW_AT_decl_line
204; CHECK-NEXT: .b8 2                                // DW_AT_data_member_location
205; CHECK-NEXT: .b8 35
206; CHECK-NEXT: .b8 0
207; CHECK-NEXT: .b8 0                                // End Of Children Mark
208; CHECK-NEXT: .b8 6                                // Abbrev [6] 0x82:0x7 DW_TAG_base_type
209; CHECK-NEXT: .b8 105                              // DW_AT_name
210; CHECK-NEXT: .b8 110
211; CHECK-NEXT: .b8 116
212; CHECK-NEXT: .b8 0
213; CHECK-NEXT: .b8 5                                // DW_AT_encoding
214; CHECK-NEXT: .b8 4                                // DW_AT_byte_size
215; CHECK-NEXT: .b8 0                                // End Of Children Mark
216; CHECK-NEXT: }
217
218%struct.Foo = type { i32 }
219
220; Function Attrs: noinline nounwind uwtable
221define void @use_dbg_declare() #0 !dbg !7 {
222entry:
223  %o = alloca %struct.Foo, align 4
224  call void @llvm.dbg.declare(metadata %struct.Foo* %o, metadata !10, metadata !15), !dbg !16
225  call void @escape_foo(%struct.Foo* %o), !dbg !17
226  ret void, !dbg !18
227}
228
229; Function Attrs: nounwind readnone speculatable
230declare void @llvm.dbg.declare(metadata, metadata, metadata) #1
231
232declare void @escape_foo(%struct.Foo*)
233
234attributes #0 = { noinline nounwind uwtable }
235attributes #1 = { nounwind readnone speculatable }
236
237!llvm.dbg.cu = !{!0}
238!llvm.module.flags = !{!3, !4, !5}
239!llvm.ident = !{!6}
240
241!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, enums: !2, nameTableKind: None)
242!1 = !DIFile(filename: "t.c", directory: "test")
243!2 = !{}
244!3 = !{i32 2, !"Dwarf Version", i32 2}
245!4 = !{i32 2, !"Debug Info Version", i32 3}
246!5 = !{i32 1, !"wchar_size", i32 4}
247!6 = !{!"clang"}
248!7 = distinct !DISubprogram(name: "use_dbg_declare", scope: !1, file: !1, line: 3, type: !8, isLocal: false, isDefinition: true, scopeLine: 3, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2)
249!8 = !DISubroutineType(types: !9)
250!9 = !{null}
251!10 = !DILocalVariable(name: "o", scope: !7, file: !1, line: 4, type: !11)
252!11 = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "Foo", file: !1, line: 1, size: 32, elements: !12)
253!12 = !{!13}
254!13 = !DIDerivedType(tag: DW_TAG_member, name: "x", scope: !11, file: !1, line: 1, baseType: !14, size: 32)
255!14 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
256!15 = !DIExpression()
257!16 = !DILocation(line: 4, column: 14, scope: !7)
258!17 = !DILocation(line: 5, column: 3, scope: !7)
259!18 = !DILocation(line: 6, column: 1, scope: !7)
260