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