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; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits