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;
10typedef unsigned int  uint;
11
12// CHECK-LABEL: @test_div_fixup_f16
13// CHECK: call half @llvm.amdgcn.div.fixup.f16
14void test_div_fixup_f16(global half* out, half a, half b, half c)
15{
16  *out = __builtin_amdgcn_div_fixuph(a, b, c);
17}
18
19// CHECK-LABEL: @test_rcp_f16
20// CHECK: call half @llvm.amdgcn.rcp.f16
21void test_rcp_f16(global half* out, half a)
22{
23  *out = __builtin_amdgcn_rcph(a);
24}
25
26// CHECK-LABEL: @test_sqrt_f16
27// CHECK: call half @llvm.amdgcn.sqrt.f16
28void test_sqrt_f16(global half* out, half a)
29{
30  *out = __builtin_amdgcn_sqrth(a);
31}
32
33// CHECK-LABEL: @test_rsq_f16
34// CHECK: call half @llvm.amdgcn.rsq.f16
35void test_rsq_f16(global half* out, half a)
36{
37  *out = __builtin_amdgcn_rsqh(a);
38}
39
40// CHECK-LABEL: @test_sin_f16
41// CHECK: call half @llvm.amdgcn.sin.f16
42void test_sin_f16(global half* out, half a)
43{
44  *out = __builtin_amdgcn_sinh(a);
45}
46
47// CHECK-LABEL: @test_cos_f16
48// CHECK: call half @llvm.amdgcn.cos.f16
49void test_cos_f16(global half* out, half a)
50{
51  *out = __builtin_amdgcn_cosh(a);
52}
53
54// CHECK-LABEL: @test_ldexp_f16
55// CHECK: call half @llvm.amdgcn.ldexp.f16
56void test_ldexp_f16(global half* out, half a, int b)
57{
58  *out = __builtin_amdgcn_ldexph(a, b);
59}
60
61// CHECK-LABEL: @test_frexp_mant_f16
62// CHECK: call half @llvm.amdgcn.frexp.mant.f16
63void test_frexp_mant_f16(global half* out, half a)
64{
65  *out = __builtin_amdgcn_frexp_manth(a);
66}
67
68// CHECK-LABEL: @test_frexp_exp_f16
69// CHECK: call i16 @llvm.amdgcn.frexp.exp.i16.f16
70void test_frexp_exp_f16(global short* out, half a)
71{
72  *out = __builtin_amdgcn_frexp_exph(a);
73}
74
75// CHECK-LABEL: @test_fract_f16
76// CHECK: call half @llvm.amdgcn.fract.f16
77void test_fract_f16(global half* out, half a)
78{
79  *out = __builtin_amdgcn_fracth(a);
80}
81
82// CHECK-LABEL: @test_class_f16
83// CHECK: call i1 @llvm.amdgcn.class.f16
84void test_class_f16(global half* out, half a, int b)
85{
86  *out = __builtin_amdgcn_classh(a, b);
87}
88
89// CHECK-LABEL: @test_s_memrealtime
90// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
91void test_s_memrealtime(global ulong* out)
92{
93  *out = __builtin_amdgcn_s_memrealtime();
94}
95
96// CHECK-LABEL: @test_s_dcache_wb()
97// CHECK: call void @llvm.amdgcn.s.dcache.wb()
98void test_s_dcache_wb()
99{
100  __builtin_amdgcn_s_dcache_wb();
101}
102
103// CHECK-LABEL: @test_mov_dpp
104// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 undef, i32 %src, i32 0, i32 0, i32 0, i1 false)
105void test_mov_dpp(global int* out, int src)
106{
107  *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false);
108}
109
110// CHECK-LABEL: @test_update_dpp
111// CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false)
112void test_update_dpp(global int* out, int arg1, int arg2)
113{
114  *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false);
115}
116
117// CHECK-LABEL: @test_ds_fadd
118// CHECK: call float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
119void test_ds_faddf(local float *out, float src) {
120  *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false);
121}
122
123// CHECK-LABEL: @test_ds_fmin
124// CHECK: call float @llvm.amdgcn.ds.fmin.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
125void test_ds_fminf(local float *out, float src) {
126  *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false);
127}
128
129// CHECK-LABEL: @test_ds_fmax
130// CHECK: call float @llvm.amdgcn.ds.fmax.f32(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false)
131void test_ds_fmaxf(local float *out, float src) {
132  *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
133}
134
135// CHECK-LABEL: @test_s_memtime
136// CHECK: call i64 @llvm.amdgcn.s.memtime()
137void test_s_memtime(global ulong* out)
138{
139  *out = __builtin_amdgcn_s_memtime();
140}
141
142// CHECK-LABEL: @test_perm
143// CHECK: call i32 @llvm.amdgcn.perm(i32 %a, i32 %b, i32 %s)
144void test_perm(global uint* out, uint a, uint b, uint s)
145{
146  *out = __builtin_amdgcn_perm(a, b, s);
147}
148
149// CHECK-LABEL: @test_groupstaticsize
150// CHECK: call i32 @llvm.amdgcn.groupstaticsize()
151void test_groupstaticsize(global uint* out)
152{
153  *out = __builtin_amdgcn_groupstaticsize();
154}
155