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 6 void 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 14 void 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 22 void 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 30 void 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() 38 void 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() 47 void 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]] 61 void 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
