diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-07-19 19:02:01 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2016-07-19 19:02:01 +0000 |
| commit | a3d6cec243bbe0446e8500cfd6a241c3357aca41 (patch) | |
| tree | b580bf83f9d32b48ba9b0846a519313a35c5a7f3 | |
| parent | d6e912ee32f0a4b7a1c24b8965c970c023de1360 (diff) | |
amdgpu: Use right builtn for rsq
The r600 path has never actually worked sinced double is not implemented
there.
git-svn-id: https://llvm.org/svn/llvm-project/libclc/trunk@276009 91177308-0d34-0410-b5e6-96231b3b80d8
| -rw-r--r-- | amdgpu/lib/math/sqrt.cl | 7 |
1 files changed, 6 insertions, 1 deletions
diff --git a/amdgpu/lib/math/sqrt.cl b/amdgpu/lib/math/sqrt.cl index 3e5b17c..395a3f9 100644 --- a/amdgpu/lib/math/sqrt.cl +++ b/amdgpu/lib/math/sqrt.cl @@ -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; |