1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -emit-llvm -o - %s | FileCheck %s 3// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s 4 5#pragma OPENCL EXTENSION cl_khr_fp16 : enable 6typedef unsigned int uint; 7typedef unsigned long ulong; 8 9// CHECK-LABEL: @test_fmed3_f16 10// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c) 11void test_fmed3_f16(global half* out, half a, half b, half c) 12{ 13 *out = __builtin_amdgcn_fmed3h(a, b, c); 14} 15 16// CHECK-LABEL: @test_s_memtime 17// CHECK: call i64 @llvm.amdgcn.s.memtime() 18void test_s_memtime(global ulong* out) 19{ 20 *out = __builtin_amdgcn_s_memtime(); 21} 22 23// CHECK-LABEL: @test_groupstaticsize 24// CHECK: call i32 @llvm.amdgcn.groupstaticsize() 25void test_groupstaticsize(global uint* out) 26{ 27 *out = __builtin_amdgcn_groupstaticsize(); 28} 29