From: Matt Arsenault <Matthew.Arsenault@amd.com>
Date: Tue, 19 Jul 2016 19:02:01 +0000 (+0000)
Subject: amdgpu: Use right builtn for rsq
X-Git-Tag: llvmorg-4.0.0-rc1~14885
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=633d749da70924c774e1c4357e02d0e65cf95f3a;p=platform%2Fupstream%2Fllvm.git

amdgpu: Use right builtn for rsq

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

llvm-svn: 276009
---

diff --git a/libclc/amdgpu/lib/math/sqrt.cl b/libclc/amdgpu/lib/math/sqrt.cl
index 3e5b17c..395a3f9 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;