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