1// REQUIRES: amdgpu-registered-target 2// RUN: %clang_cc1 -triple r600-unknown-unknown -target-cpu cypress -S -emit-llvm -o - %s | FileCheck %s 3 4// CHECK-LABEL: @test_rsq_f32 5// CHECK: call float @llvm.r600.rsq.f32 6void test_rsq_f32(global float* out, float a) 7{ 8 *out = __builtin_amdgpu_rsqf(a); 9} 10 11#if cl_khr_fp64 12// XCHECK-LABEL: @test_rsq_f64 13// XCHECK: call double @llvm.r600.rsq.f64 14void test_rsq_f64(global double* out, double a) 15{ 16 *out = __builtin_amdgpu_rsq(a); 17} 18#endif 19 20// CHECK-LABEL: @test_legacy_ldexp_f32 21// CHECK: call float @llvm.AMDGPU.ldexp.f32 22void test_legacy_ldexp_f32(global float* out, float a, int b) 23{ 24 *out = __builtin_amdgpu_ldexpf(a, b); 25} 26 27#if cl_khr_fp64 28// XCHECK-LABEL: @test_legacy_ldexp_f64 29// XCHECK: call double @llvm.AMDGPU.ldexp.f64 30void test_legacy_ldexp_f64(global double* out, double a, int b) 31{ 32 *out = __builtin_amdgpu_ldexp(a, b); 33} 34#endif 35 36// CHECK-LABEL: @test_implicitarg_ptr 37// CHECK: call i8 addrspace(7)* @llvm.r600.implicitarg.ptr() 38void test_implicitarg_ptr(__attribute__((address_space(7))) unsigned char ** out) 39{ 40 *out = __builtin_r600_implicitarg_ptr(); 41} 42 43// CHECK-LABEL: @test_get_group_id( 44// CHECK: tail call i32 @llvm.r600.read.tgid.x() 45// CHECK: tail call i32 @llvm.r600.read.tgid.y() 46// CHECK: tail call i32 @llvm.r600.read.tgid.z() 47void test_get_group_id(int d, global int *out) 48{ 49 switch (d) { 50 case 0: *out = __builtin_r600_read_tgid_x(); break; 51 case 1: *out = __builtin_r600_read_tgid_y(); break; 52 case 2: *out = __builtin_r600_read_tgid_z(); break; 53 default: *out = 0; 54 } 55} 56 57// CHECK-LABEL: @test_get_local_id( 58// CHECK: tail call i32 @llvm.r600.read.tidig.x(), !range [[WI_RANGE:![0-9]*]] 59// CHECK: tail call i32 @llvm.r600.read.tidig.y(), !range [[WI_RANGE]] 60// CHECK: tail call i32 @llvm.r600.read.tidig.z(), !range [[WI_RANGE]] 61void test_get_local_id(int d, global int *out) 62{ 63 switch (d) { 64 case 0: *out = __builtin_r600_read_tidig_x(); break; 65 case 1: *out = __builtin_r600_read_tidig_y(); break; 66 case 2: *out = __builtin_r600_read_tidig_z(); break; 67 default: *out = 0; 68 } 69} 70 71// CHECK-DAG: [[WI_RANGE]] = !{i32 0, i32 1024} 72