| 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 | |
|---|
| 6 | typedef unsigned long ulong; |
|---|
| 7 | |
|---|
| 8 | // CHECK-LABEL: @test_div_fixup_f16 |
|---|
| 9 | // CHECK: call half @llvm.amdgcn.div.fixup.f16 |
|---|
| 10 | void 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 |
|---|
| 17 | void 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 |
|---|
| 24 | void 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 |
|---|
| 31 | void 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 |
|---|
| 38 | void 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 |
|---|
| 45 | void 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 |
|---|
| 52 | void 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 |
|---|
| 59 | void 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 |
|---|
| 66 | void 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 |
|---|
| 73 | void 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() |
|---|
| 80 | void test_s_memrealtime(global ulong* out) |
|---|
| 81 | { |
|---|
| 82 | *out = __builtin_amdgcn_s_memrealtime(); |
|---|
| 83 | } |
|---|
| 84 | |
|---|
| 85 | // CHECK-LABEL: @test_s_dcache_wb() |
|---|
| 86 | // CHECK: call void @llvm.amdgcn.s.dcache.wb() |
|---|
| 87 | void test_s_dcache_wb() |
|---|
| 88 | { |
|---|
| 89 | __builtin_amdgcn_s_dcache_wb(); |
|---|
| 90 | } |
|---|
| 91 | |
|---|
| 92 | // CHECK-LABEL: @test_mov_dpp |
|---|
| 93 | // CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 undef, i32 %src, i32 0, i32 0, i32 0, i1 false) |
|---|
| 94 | void test_mov_dpp(global int* out, int src) |
|---|
| 95 | { |
|---|
| 96 | *out = __builtin_amdgcn_mov_dpp(src, 0, 0, 0, false); |
|---|
| 97 | } |
|---|
| 98 | |
|---|
| 99 | // CHECK-LABEL: @test_update_dpp |
|---|
| 100 | // CHECK: call i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false) |
|---|
| 101 | void test_update_dpp(global int* out, int arg1, int arg2) |
|---|
| 102 | { |
|---|
| 103 | *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false); |
|---|
| 104 | } |
|---|
| 105 | |
|---|
| 106 | // CHECK-LABEL: @test_ds_fadd |
|---|
| 107 | // CHECK: call float @llvm.amdgcn.ds.fadd(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) |
|---|
| 108 | void test_ds_faddf(local float *out, float src) { |
|---|
| 109 | *out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, false); |
|---|
| 110 | } |
|---|
| 111 | |
|---|
| 112 | // CHECK-LABEL: @test_ds_fmin |
|---|
| 113 | // CHECK: call float @llvm.amdgcn.ds.fmin(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) |
|---|
| 114 | void test_ds_fminf(local float *out, float src) { |
|---|
| 115 | *out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, false); |
|---|
| 116 | } |
|---|
| 117 | |
|---|
| 118 | // CHECK-LABEL: @test_ds_fmax |
|---|
| 119 | // CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* %out, float %src, i32 0, i32 0, i1 false) |
|---|
| 120 | void test_ds_fmaxf(local float *out, float src) { |
|---|
| 121 | *out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false); |
|---|
| 122 | } |
|---|
| 123 | |
|---|