[libclc] r276009 - amdgpu: Use right builtn for rsq

[Matt Arsenault via cfe-commits]

Author: arsenm
Date: Tue Jul 19 14:02:01 2016
New Revision: 276009

URL: http://llvm.org/viewvc/llvm-project?rev=276009&view=rev
Log:
amdgpu: Use right builtn for rsq

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

Modified:
    libclc/trunk/amdgpu/lib/math/sqrt.cl

Modified: libclc/trunk/amdgpu/lib/math/sqrt.cl
URL: http://llvm.org/viewvc/llvm-project/libclc/trunk/amdgpu/lib/math/sqrt.cl?rev=276009&r1=276008&r2=276009&view=diff
==============================================================================
--- libclc/trunk/amdgpu/lib/math/sqrt.cl (original)
+++ libclc/trunk/amdgpu/lib/math/sqrt.cl Tue Jul 19 14:02:01 2016
@@ -30,6 +30,11 @@ _CLC_DEFINE_UNARY_BUILTIN(float, sqrt, _
 
 #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(doubl
   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;