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