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