1// REQUIRES: amdgpu-registered-target
2// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
3// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
4// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
5// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
6
7#pragma OPENCL EXTENSION cl_khr_fp16 : enable
8
9typedef unsigned long ulong;
10
11// CHECK-LABEL: @test_div_fixup_f16
12// CHECK: call half @llvm.amdgcn.div.fixup.f16
13void test_div_fixup_f16(global half* out, half a, half b, half c)
14{
15  *out = __builtin_amdgcn_div_fixuph(a, b, c);
16}
17
18// CHECK-LABEL: @test_rcp_f16
19// CHECK: call half @llvm.amdgcn.rcp.f16
20void test_rcp_f16(global half* out, half a)
21{
22  *out = __builtin_amdgcn_rcph(a);
23}
24
25// CHECK-LABEL: @test_sqrt_f16
26// CHECK: call half @llvm.amdgcn.sqrt.f16
27void test_sqrt_f16(global half* out, half a)
28{
29  *out = __builtin_amdgcn_sqrth(a);
30}
31
32// CHECK-LABEL: @test_rsq_f16
33// CHECK: call half @llvm.amdgcn.rsq.f16
34void test_rsq_f16(global half* out, half a)
35{
36  *out = __builtin_amdgcn_rsqh(a);
37}
38
39// CHECK-LABEL: @test_sin_f16
40// CHECK: call half @llvm.amdgcn.sin.f16
41void test_sin_f16(global half* out, half a)
42{
43  *out = __builtin_amdgcn_sinh(a);
44}
45
46// CHECK-LABEL: @test_cos_f16
47// CHECK: call half @llvm.amdgcn.cos.f16
48void test_cos_f16(global half* out, half a)
49{
50  *out = __builtin_amdgcn_cosh(a);
51}
52
53// CHECK-LABEL: @test_ldexp_f16
54// CHECK: call half @llvm.amdgcn.ldexp.f16
55void test_ldexp_f16(global half* out, half a, int b)
56{
57  *out = __builtin_amdgcn_ldexph(a, b);
58}
59
60// CHECK-LABEL: @test_frexp_mant_f16
61// CHECK: call half @llvm.amdgcn.frexp.mant.f16
62void test_frexp_mant_f16(global half* out, half a)
63{
64  *out = __builtin_amdgcn_frexp_manth(a);
65}
66
67// CHECK-LABEL: @test_frexp_exp_f16
68// CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16
69void test_frexp_exp_f16(global short* out, half a)
70{
71  *out = __builtin_amdgcn_frexp_exph(a);
72}
73
74// CHECK-LABEL: @test_fract_f16
75// CHECK: call half @llvm.amdgcn.fract.f16
76void test_fract_f16(global half* out, half a)
77{
78  *out = __builtin_amdgcn_fracth(a);
79}
80
81// CHECK-LABEL: @test_class_f16
82// CHECK: call i1 @llvm.amdgcn.class.f16
83void test_class_f16(global half* out, half a, int b)
84{
85  *out = __builtin_amdgcn_classh(a, b);
86}
87
88// CHECK-LABEL: @test_s_memrealtime
89// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
90void test_s_memrealtime(global ulong* out)
91{
92  *out = __builtin_amdgcn_s_memrealtime();
93}
94
95// CHECK-LABEL: @test_s_dcache_wb()
96// CHECK: call void @llvm.amdgcn.s.dcache.wb()
97void test_s_dcache_wb()
98{
99  __builtin_amdgcn_s_dcache_wb();
100}
101
102// CHECK-LABEL: @test_mov_dpp
103// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 undef, i32 %src, i32 0, i32 0, i32 0, i1 false)
104void test_mov_dpp(global int* out, int src)
105{
106  *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
107}
108
109// CHECK-LABEL: @test_update_dpp
110// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
111void test_update_dpp(global int* out, int arg1, int arg2)
112{
113  *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
114}
115
116// CHECK-LABEL: @test_ds_fadd
117// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
118void test_ds_faddf(local float *out, float src) {
119  *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false);
120}
121
122// CHECK-LABEL: @test_ds_fmin
123// CHECK: call float @llvm.amdgcn.ds.fmin.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
124void test_ds_fminf(local float *out, float src) {
125  *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
126}
127
128// CHECK-LABEL: @test_ds_fmax
129// CHECK: call float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
130void test_ds_fmaxf(local float *out, float src) {
131  *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
132}
133