OSchip
/
llvm-project
mirror of https://github.com/THU-DSP-LAB/llvm-project.git
13
0
Fork 2
Code Issues Packages Projects Releases Wiki Activity

amdgpu: Use right builtn for rsq 

The r600 path has never actually worked sinced double is not implemented
there.

llvm-svn: 276009
This commit is contained in:
Matt Arsenault 2016-07-19 19:02:01 +00:00
parent 5246e0b2c2
commit 633d749da7
1 changed files with 6 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

7
libclc/amdgpu/lib/math/sqrt.cl
View File

@ -30,6 +30,11 @@ _CLC_DEFINE_UNARY_BUILTIN(float, sqrt, __clc_sqrt, float)
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#ifdef __AMDGCN__
#define __clc_builtin_rsq __builtin_amdgcn_rsq
#else
#define __clc_builtin_rsq __builtin_r600_recipsqrt_ieee
#endif
_CLC_OVERLOAD _CLC_DEF double sqrt(double x) {
@ -38,7 +43,7 @@ _CLC_OVERLOAD _CLC_DEF double sqrt(double x) {
unsigned exp1 = vcc ? 0xffffff80 : 0;
double v01 = ldexp(x, exp0);
double v23 = __builtin_amdgpu_rsq(v01);
double v23 = __clc_builtin_rsq(v01);
double v45 = v01 * v23;
v23 = v23 * 0.5;