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 | 633d749da70924c774e1c4357e02d0e65cf95f3a (patch) | |
| tree | 1bf758d6a39377795be86c640eb2f2000adfdc66 /libclc | |
| parent | 5246e0b2c25c293e6896e3a6fff63ade00fcb54e (diff) | |
| download | bcm5719-llvm-633d749da70924c774e1c4357e02d0e65cf95f3a.tar.gz bcm5719-llvm-633d749da70924c774e1c4357e02d0e65cf95f3a.zip | |
amdgpu: Use right builtn for rsq
The r600 path has never actually worked sinced double is not implemented
there.
llvm-svn: 276009
Diffstat (limited to 'libclc')
| -rw-r--r-- | libclc/amdgpu/lib/math/sqrt.cl | 7 |
1 files changed, 6 insertions, 1 deletions
diff --git a/libclc/amdgpu/lib/math/sqrt.cl b/libclc/amdgpu/lib/math/sqrt.cl index 3e5b17c5db0..395a3f9fb21 100644 --- a/libclc/amdgpu/lib/math/sqrt.cl +++ b/libclc/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; |
