Clang Project

clang_source_code/test/CodeGenOpenCL/builtins-amdgcn.cl
1// REQUIRES: amdgpu-registered-target
2// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
3
4#pragma OPENCL EXTENSION cl_khr_fp64 : enable
5
6typedef unsigned long ulong;
7typedef unsigned int uint;
8
9// CHECK-LABEL: @test_div_scale_f64
10// CHECK: call { double, i1 } @llvm.amdgcn.div.scale.f64(double %a, double %b, i1 true)
11// CHECK-DAG: [[FLAG:%.+]] = extractvalue { double, i1 } %{{.+}}, 1
12// CHECK-DAG: [[VAL:%.+]] = extractvalue { double, i1 } %{{.+}}, 0
13// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i32
14// CHECK: store i32 [[FLAGEXT]]
15void test_div_scale_f64(global double* out, global int* flagout, double a, double b)
16{
17  bool flag;
18  *out = __builtin_amdgcn_div_scale(a, b, true, &flag);
19  *flagout = flag;
20}
21
22// CHECK-LABEL: @test_div_scale_f32(
23// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true)
24// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1
25// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0
26// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8
27// CHECK: store i8 [[FLAGEXT]]
28void test_div_scale_f32(global float* out, global bool* flagout, float a, float b)
29{
30  bool flag;
31  *out = __builtin_amdgcn_div_scalef(a, b, true, &flag);
32  *flagout = flag;
33}
34
35// CHECK-LABEL: @test_div_scale_f32_global_ptr(
36// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true)
37// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1
38// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0
39// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8
40// CHECK: store i8 [[FLAGEXT]]
41void test_div_scale_f32_global_ptr(global float* out, global int* flagout, float a, float b, global bool* flag)
42{
43  *out = __builtin_amdgcn_div_scalef(a, b, true, flag);
44}
45
46// CHECK-LABEL: @test_div_scale_f32_generic_ptr(
47// CHECK: call { float, i1 } @llvm.amdgcn.div.scale.f32(float %a, float %b, i1 true)
48// CHECK-DAG: [[FLAG:%.+]] = extractvalue { float, i1 } %{{.+}}, 1
49// CHECK-DAG: [[VAL:%.+]] = extractvalue { float, i1 } %{{.+}}, 0
50// CHECK: [[FLAGEXT:%.+]] = zext i1 [[FLAG]] to i8
51// CHECK: store i8 [[FLAGEXT]]
52void test_div_scale_f32_generic_ptr(global float* out, global int* flagout, float a, float b, global bool* flag_arg)
53{
54  generic bool* flag = flag_arg;
55  *out = __builtin_amdgcn_div_scalef(a, b, true, flag);
56}
57
58// CHECK-LABEL: @test_div_fmas_f32
59// CHECK: call float @llvm.amdgcn.div.fmas.f32
60void test_div_fmas_f32(global float* out, float a, float b, float c, int d)
61{
62  *out = __builtin_amdgcn_div_fmasf(a, b, c, d);
63}
64
65// CHECK-LABEL: @test_div_fmas_f64
66// CHECK: call double @llvm.amdgcn.div.fmas.f64
67void test_div_fmas_f64(global double* out, double a, double b, double c, int d)
68{
69  *out = __builtin_amdgcn_div_fmas(a, b, c, d);
70}
71
72// CHECK-LABEL: @test_div_fixup_f32
73// CHECK: call float @llvm.amdgcn.div.fixup.f32
74void test_div_fixup_f32(global float* out, float a, float b, float c)
75{
76  *out = __builtin_amdgcn_div_fixupf(a, b, c);
77}
78
79// CHECK-LABEL: @test_div_fixup_f64
80// CHECK: call double @llvm.amdgcn.div.fixup.f64
81void test_div_fixup_f64(global double* out, double a, double b, double c)
82{
83  *out = __builtin_amdgcn_div_fixup(a, b, c);
84}
85
86// CHECK-LABEL: @test_trig_preop_f32
87// CHECK: call float @llvm.amdgcn.trig.preop.f32
88void test_trig_preop_f32(global float* out, float a, int b)
89{
90  *out = __builtin_amdgcn_trig_preopf(a, b);
91}
92
93// CHECK-LABEL: @test_trig_preop_f64
94// CHECK: call double @llvm.amdgcn.trig.preop.f64
95void test_trig_preop_f64(global double* out, double a, int b)
96{
97  *out = __builtin_amdgcn_trig_preop(a, b);
98}
99
100// CHECK-LABEL: @test_rcp_f32
101// CHECK: call float @llvm.amdgcn.rcp.f32
102void test_rcp_f32(global float* out, float a)
103{
104  *out = __builtin_amdgcn_rcpf(a);
105}
106
107// CHECK-LABEL: @test_rcp_f64
108// CHECK: call double @llvm.amdgcn.rcp.f64
109void test_rcp_f64(global double* out, double a)
110{
111  *out = __builtin_amdgcn_rcp(a);
112}
113
114// CHECK-LABEL: @test_rsq_f32
115// CHECK: call float @llvm.amdgcn.rsq.f32
116void test_rsq_f32(global float* out, float a)
117{
118  *out = __builtin_amdgcn_rsqf(a);
119}
120
121// CHECK-LABEL: @test_rsq_f64
122// CHECK: call double @llvm.amdgcn.rsq.f64
123void test_rsq_f64(global double* out, double a)
124{
125  *out = __builtin_amdgcn_rsq(a);
126}
127
128// CHECK-LABEL: @test_rsq_clamp_f32
129// CHECK: call float @llvm.amdgcn.rsq.clamp.f32
130void test_rsq_clamp_f32(global float* out, float a)
131{
132  *out = __builtin_amdgcn_rsq_clampf(a);
133}
134
135// CHECK-LABEL: @test_rsq_clamp_f64
136// CHECK: call double @llvm.amdgcn.rsq.clamp.f64
137void test_rsq_clamp_f64(global double* out, double a)
138{
139  *out = __builtin_amdgcn_rsq_clamp(a);
140}
141
142// CHECK-LABEL: @test_sin_f32
143// CHECK: call float @llvm.amdgcn.sin.f32
144void test_sin_f32(global float* out, float a)
145{
146  *out = __builtin_amdgcn_sinf(a);
147}
148
149// CHECK-LABEL: @test_cos_f32
150// CHECK: call float @llvm.amdgcn.cos.f32
151void test_cos_f32(global float* out, float a)
152{
153  *out = __builtin_amdgcn_cosf(a);
154}
155
156// CHECK-LABEL: @test_log_clamp_f32
157// CHECK: call float @llvm.amdgcn.log.clamp.f32
158void test_log_clamp_f32(global float* out, float a)
159{
160  *out = __builtin_amdgcn_log_clampf(a);
161}
162
163// CHECK-LABEL: @test_ldexp_f32
164// CHECK: call float @llvm.amdgcn.ldexp.f32
165void test_ldexp_f32(global float* out, float a, int b)
166{
167  *out = __builtin_amdgcn_ldexpf(a, b);
168}
169
170// CHECK-LABEL: @test_ldexp_f64
171// CHECK: call double @llvm.amdgcn.ldexp.f64
172void test_ldexp_f64(global double* out, double a, int b)
173{
174  *out = __builtin_amdgcn_ldexp(a, b);
175}
176
177// CHECK-LABEL: @test_frexp_mant_f32
178// CHECK: call float @llvm.amdgcn.frexp.mant.f32
179void test_frexp_mant_f32(global float* out, float a)
180{
181  *out = __builtin_amdgcn_frexp_mantf(a);
182}
183
184// CHECK-LABEL: @test_frexp_mant_f64
185// CHECK: call double @llvm.amdgcn.frexp.mant.f64
186void test_frexp_mant_f64(global double* out, double a)
187{
188  *out = __builtin_amdgcn_frexp_mant(a);
189}
190
191// CHECK-LABEL: @test_frexp_exp_f32
192// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f32
193void test_frexp_exp_f32(global int* out, float a)
194{
195  *out = __builtin_amdgcn_frexp_expf(a);
196}
197
198// CHECK-LABEL: @test_frexp_exp_f64
199// CHECK: call i32 @llvm.amdgcn.frexp.exp.i32.f64
200void test_frexp_exp_f64(global int* out, double a)
201{
202  *out = __builtin_amdgcn_frexp_exp(a);
203}
204
205// CHECK-LABEL: @test_fract_f32
206// CHECK: call float @llvm.amdgcn.fract.f32
207void test_fract_f32(global int* out, float a)
208{
209  *out = __builtin_amdgcn_fractf(a);
210}
211
212// CHECK-LABEL: @test_fract_f64
213// CHECK: call double @llvm.amdgcn.fract.f64
214void test_fract_f64(global int* out, double a)
215{
216  *out = __builtin_amdgcn_fract(a);
217}
218
219// CHECK-LABEL: @test_lerp
220// CHECK: call i32 @llvm.amdgcn.lerp
221void test_lerp(global int* out, int a, int b, int c)
222{
223  *out = __builtin_amdgcn_lerp(a, b, c);
224}
225
226// CHECK-LABEL: @test_sicmp_i32
227// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32)
228void test_sicmp_i32(global ulong* out, int a, int b)
229{
230  *out = __builtin_amdgcn_sicmp(a, b, 32);
231}
232
233// CHECK-LABEL: @test_uicmp_i32
234// CHECK: call i64 @llvm.amdgcn.icmp.i32(i32 %a, i32 %b, i32 32)
235void test_uicmp_i32(global ulong* out, uint a, uint b)
236{
237  *out = __builtin_amdgcn_uicmp(a, b, 32);
238}
239
240// CHECK-LABEL: @test_sicmp_i64
241// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 38)
242void test_sicmp_i64(global ulong* out, long a, long b)
243{
244  *out = __builtin_amdgcn_sicmpl(a, b, 39-1);
245}
246
247// CHECK-LABEL: @test_uicmp_i64
248// CHECK: call i64 @llvm.amdgcn.icmp.i64(i64 %a, i64 %b, i32 35)
249void test_uicmp_i64(global ulong* out, ulong a, ulong b)
250{
251  *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
252}
253
254// CHECK-LABEL: @test_ds_swizzle
255// CHECK: call i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32)
256void test_ds_swizzle(global int* out, int a)
257{
258  *out = __builtin_amdgcn_ds_swizzle(a, 32);
259}
260
261// CHECK-LABEL: @test_ds_permute
262// CHECK: call i32 @llvm.amdgcn.ds.permute(i32 %a, i32 %b)
263void test_ds_permute(global int* out, int a, int b)
264{
265  out[0] = __builtin_amdgcn_ds_permute(a, b);
266}
267
268// CHECK-LABEL: @test_ds_bpermute
269// CHECK: call i32 @llvm.amdgcn.ds.bpermute(i32 %a, i32 %b)
270void test_ds_bpermute(global int* out, int a, int b)
271{
272  *out = __builtin_amdgcn_ds_bpermute(a, b);
273}
274
275// CHECK-LABEL: @test_readfirstlane
276// CHECK: call i32 @llvm.amdgcn.readfirstlane(i32 %a)
277void test_readfirstlane(global int* out, int a)
278{
279  *out = __builtin_amdgcn_readfirstlane(a);
280}
281
282// CHECK-LABEL: @test_readlane
283// CHECK: call i32 @llvm.amdgcn.readlane(i32 %a, i32 %b)
284void test_readlane(global int* out, int a, int b)
285{
286  *out = __builtin_amdgcn_readlane(a, b);
287}
288
289// CHECK-LABEL: @test_fcmp_f32
290// CHECK: call i64 @llvm.amdgcn.fcmp.f32(float %a, float %b, i32 5)
291void test_fcmp_f32(global ulong* out, float a, float b)
292{
293  *out = __builtin_amdgcn_fcmpf(a, b, 5);
294}
295
296// CHECK-LABEL: @test_fcmp_f64
297// CHECK: call i64 @llvm.amdgcn.fcmp.f64(double %a, double %b, i32 6)
298void test_fcmp_f64(global ulong* out, double a, double b)
299{
300  *out = __builtin_amdgcn_fcmp(a, b, 3+3);
301}
302
303// CHECK-LABEL: @test_class_f32
304// CHECK: call i1 @llvm.amdgcn.class.f32
305void test_class_f32(global float* out, float a, int b)
306{
307  *out = __builtin_amdgcn_classf(a, b);
308}
309
310// CHECK-LABEL: @test_class_f64
311// CHECK: call i1 @llvm.amdgcn.class.f64
312void test_class_f64(global double* out, double a, int b)
313{
314  *out = __builtin_amdgcn_class(a, b);
315}
316
317// CHECK-LABEL: @test_buffer_wbinvl1
318// CHECK: call void @llvm.amdgcn.buffer.wbinvl1(
319void test_buffer_wbinvl1()
320{
321  __builtin_amdgcn_buffer_wbinvl1();
322}
323
324// CHECK-LABEL: @test_s_dcache_inv
325// CHECK: call void @llvm.amdgcn.s.dcache.inv(
326void test_s_dcache_inv()
327{
328  __builtin_amdgcn_s_dcache_inv();
329}
330
331// CHECK-LABEL: @test_s_waitcnt
332// CHECK: call void @llvm.amdgcn.s.waitcnt(
333void test_s_waitcnt()
334{
335  __builtin_amdgcn_s_waitcnt(0);
336}
337
338// CHECK-LABEL: @test_s_sendmsg
339// CHECK: call void @llvm.amdgcn.s.sendmsg(
340void test_s_sendmsg()
341{
342  __builtin_amdgcn_s_sendmsg(1, 0);
343}
344
345// CHECK-LABEL: @test_s_sendmsg_var
346// CHECK: call void @llvm.amdgcn.s.sendmsg(
347void test_s_sendmsg_var(int in)
348{
349  __builtin_amdgcn_s_sendmsg(1, in);
350}
351
352// CHECK-LABEL: @test_s_sendmsghalt
353// CHECK: call void @llvm.amdgcn.s.sendmsghalt(
354void test_s_sendmsghalt()
355{
356  __builtin_amdgcn_s_sendmsghalt(1, 0);
357}
358
359// CHECK-LABEL: @test_s_sendmsghalt
360// CHECK: call void @llvm.amdgcn.s.sendmsghalt(
361void test_s_sendmsghalt_var(int in)
362{
363  __builtin_amdgcn_s_sendmsghalt(1, in);
364}
365
366// CHECK-LABEL: @test_s_barrier
367// CHECK: call void @llvm.amdgcn.s.barrier(
368void test_s_barrier()
369{
370  __builtin_amdgcn_s_barrier();
371}
372
373// CHECK-LABEL: @test_wave_barrier
374// CHECK: call void @llvm.amdgcn.wave.barrier(
375void test_wave_barrier()
376{
377  __builtin_amdgcn_wave_barrier();
378}
379
380// CHECK-LABEL: @test_s_memtime
381// CHECK: call i64 @llvm.amdgcn.s.memtime()
382void test_s_memtime(global ulong* out)
383{
384  *out = __builtin_amdgcn_s_memtime();
385}
386
387// CHECK-LABEL: @test_s_sleep
388// CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
389// CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
390void test_s_sleep()
391{
392  __builtin_amdgcn_s_sleep(1);
393  __builtin_amdgcn_s_sleep(15);
394}
395
396// CHECK-LABEL: @test_s_incperflevel
397// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 1)
398// CHECK: call void @llvm.amdgcn.s.incperflevel(i32 15)
399void test_s_incperflevel()
400{
401  __builtin_amdgcn_s_incperflevel(1);
402  __builtin_amdgcn_s_incperflevel(15);
403}
404
405// CHECK-LABEL: @test_s_decperflevel
406// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 1)
407// CHECK: call void @llvm.amdgcn.s.decperflevel(i32 15)
408void test_s_decperflevel()
409{
410  __builtin_amdgcn_s_decperflevel(1);
411  __builtin_amdgcn_s_decperflevel(15);
412}
413
414// CHECK-LABEL: @test_cubeid(
415// CHECK: call float @llvm.amdgcn.cubeid(float %a, float %b, float %c)
416void test_cubeid(global float* out, float a, float b, float c) {
417  *out = __builtin_amdgcn_cubeid(a, b, c);
418}
419
420// CHECK-LABEL: @test_cubesc(
421// CHECK: call float @llvm.amdgcn.cubesc(float %a, float %b, float %c)
422void test_cubesc(global float* out, float a, float b, float c) {
423  *out = __builtin_amdgcn_cubesc(a, b, c);
424}
425
426// CHECK-LABEL: @test_cubetc(
427// CHECK: call float @llvm.amdgcn.cubetc(float %a, float %b, float %c)
428void test_cubetc(global float* out, float a, float b, float c) {
429  *out = __builtin_amdgcn_cubetc(a, b, c);
430}
431
432// CHECK-LABEL: @test_cubema(
433// CHECK: call float @llvm.amdgcn.cubema(float %a, float %b, float %c)
434void test_cubema(global float* out, float a, float b, float c) {
435  *out = __builtin_amdgcn_cubema(a, b, c);
436}
437
438// CHECK-LABEL: @test_read_exec(
439// CHECK: call i64 @llvm.read_register.i64(metadata ![[$EXEC:[0-9]+]]) #[[$READ_EXEC_ATTRS:[0-9]+]]
440void test_read_exec(global ulong* out) {
441  *out = __builtin_amdgcn_read_exec();
442}
443
444// CHECK: declare i64 @llvm.read_register.i64(metadata) #[[$NOUNWIND_READONLY:[0-9]+]]
445
446// CHECK-LABEL: @test_read_exec_lo(
447// CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_LO:[0-9]+]]) #[[$READ_EXEC_ATTRS]]
448void test_read_exec_lo(global uint* out) {
449  *out = __builtin_amdgcn_read_exec_lo();
450}
451
452// CHECK-LABEL: @test_read_exec_hi(
453// CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_HI:[0-9]+]]) #[[$READ_EXEC_ATTRS]]
454void test_read_exec_hi(global uint* out) {
455  *out = __builtin_amdgcn_read_exec_hi();
456}
457
458// CHECK-LABEL: @test_dispatch_ptr
459// CHECK: call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
460void test_dispatch_ptr(__constant unsigned char ** out)
461{
462  *out = __builtin_amdgcn_dispatch_ptr();
463}
464
465// CHECK-LABEL: @test_queue_ptr
466// CHECK: call i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
467void test_queue_ptr(__constant unsigned char ** out)
468{
469  *out = __builtin_amdgcn_queue_ptr();
470}
471
472// CHECK-LABEL: @test_kernarg_segment_ptr
473// CHECK: call i8 addrspace(4)* @llvm.amdgcn.kernarg.segment.ptr()
474void test_kernarg_segment_ptr(__constant unsigned char ** out)
475{
476  *out = __builtin_amdgcn_kernarg_segment_ptr();
477}
478
479// CHECK-LABEL: @test_implicitarg_ptr
480// CHECK: call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
481void test_implicitarg_ptr(__constant unsigned char ** out)
482{
483  *out = __builtin_amdgcn_implicitarg_ptr();
484}
485
486// CHECK-LABEL: @test_get_group_id(
487// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
488// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
489// CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()
490void test_get_group_id(int d, global int *out)
491{
492 switch (d) {
493 case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;
494 case 1: *out = __builtin_amdgcn_workgroup_id_y(); break;
495 case 2: *out = __builtin_amdgcn_workgroup_id_z(); break;
496 default: *out = 0;
497 }
498}
499
500// CHECK-LABEL: @test_s_getreg(
501// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 0)
502// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 1)
503// CHECK: tail call i32 @llvm.amdgcn.s.getreg(i32 65535)
504void test_s_getreg(volatile global uint *out)
505{
506  *out = __builtin_amdgcn_s_getreg(0);
507  *out = __builtin_amdgcn_s_getreg(1);
508  *out = __builtin_amdgcn_s_getreg(65535);
509}
510
511// CHECK-LABEL: @test_get_local_id(
512// CHECK: tail call i32 @llvm.amdgcn.workitem.id.x(), !range [[$WI_RANGE:![0-9]*]]
513// CHECK: tail call i32 @llvm.amdgcn.workitem.id.y(), !range [[$WI_RANGE]]
514// CHECK: tail call i32 @llvm.amdgcn.workitem.id.z(), !range [[$WI_RANGE]]
515void test_get_local_id(int d, global int *out)
516{
517 switch (d) {
518 case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
519 case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
520 case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
521 default: *out = 0;
522 }
523}
524
525// CHECK-LABEL: @test_fmed3_f32
526// CHECK: call float @llvm.amdgcn.fmed3.f32(
527void test_fmed3_f32(global float* out, float a, float b, float c)
528{
529  *out = __builtin_amdgcn_fmed3f(a, b, c);
530}
531
532// CHECK-LABEL: @test_s_getpc
533// CHECK: call i64 @llvm.amdgcn.s.getpc()
534void test_s_getpc(global ulong* out)
535{
536  *out = __builtin_amdgcn_s_getpc();
537}
538
539// CHECK-LABEL: @test_ds_append_lds(
540// CHECK: call i32 @llvm.amdgcn.ds.append.p3i32(i32 addrspace(3)* %ptr, i1 false)
541kernel void test_ds_append_lds(global int* out, local int* ptr) {
542  *out = __builtin_amdgcn_ds_append(ptr);
543}
544
545// CHECK-LABEL: @test_ds_consume_lds(
546// CHECK: call i32 @llvm.amdgcn.ds.consume.p3i32(i32 addrspace(3)* %ptr, i1 false)
547kernel void test_ds_consume_lds(global int* out, local int* ptr) {
548  *out = __builtin_amdgcn_ds_consume(ptr);
549}
550
551// CHECK-DAG: [[$WI_RANGE]] = !{i32 0, i32 1024}
552// CHECK-DAG: attributes #[[$NOUNWIND_READONLY:[0-9]+]] = { nounwind readonly }
553// CHECK-DAG: attributes #[[$READ_EXEC_ATTRS]] = { convergent }
554// CHECK-DAG: ![[$EXEC]] = !{!"exec"}
555// CHECK-DAG: ![[$EXEC_LO]] = !{!"exec_lo"}
556// CHECK-DAG: ![[$EXEC_HI]] = !{!"exec_hi"}
557
CodeThinker, Enjoy your reading code