[libclc] 35/58: amdgpu: Use right builtn for rsq
Andreas Boll
aboll-guest at moszumanska.debian.org
Thu Oct 6 08:16:32 UTC 2016
This is an automated email from the git hooks/post-receive script.
aboll-guest pushed a commit to branch master
in repository libclc.
commit a3d6cec243bbe0446e8500cfd6a241c3357aca41
Author: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Tue Jul 19 19:02:01 2016 +0000
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
---
amdgpu/lib/math/sqrt.cl | 7 ++++++-
1 file changed, 6 insertions(+), 1 deletion(-)
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;
--
Alioth's /usr/local/bin/git-commit-notice on /srv/git.debian.org/git/pkg-opencl/libclc.git
More information about the Pkg-opencl-commits
mailing list