[libclc] cac3111 - [libclc] Use unchecked division for f64 AMDGPU reciprocal (#203809)

[via cfe-commits]

Author: Joseph Huber
Date: 2026-06-15T07:26:32-05:00
New Revision: cac31111d37cd27b222aa2ddb11e2238db27017f

URL: https://github.com/llvm/llvm-project/commit/cac31111d37cd27b222aa2ddb11e2238db27017f
DIFF: https://github.com/llvm/llvm-project/commit/cac31111d37cd27b222aa2ddb11e2238db27017f.diff

LOG: [libclc] Use unchecked division for f64 AMDGPU reciprocal (#203809)

Summary:
This matches what the AMD device libraries does. We can avoid extra
steps by only performing the two steps of the Newton-Raphson
approximation of 1 / x. The exceptional cases should not appear in these
math functions, this is local to AMDGPU, and I verified they are bitwise
identical to the AMD math functions with parity in performance now.

Added: 
    

Modified: 
    libclc/clc/lib/amdgpu/math/clc_recip_fast.inc

Removed: 
    


################################################################################
diff  --git a/libclc/clc/lib/amdgpu/math/clc_recip_fast.inc b/libclc/clc/lib/amdgpu/math/clc_recip_fast.inc
index 9d635cc700442..e19ec82d7566d 100644
--- a/libclc/clc/lib/amdgpu/math/clc_recip_fast.inc
+++ b/libclc/clc/lib/amdgpu/math/clc_recip_fast.inc
@@ -6,10 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 
-// On AMDGPU the "fast" reciprocal is the hardware v_rcp_f32 approximation,
+// On AMDGPU the "fast" reciprocal is the hardware v_rcp approximation.
 _CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_recip_fast(__CLC_GENTYPE x) {
 #if defined(__CLC_SCALAR) && __CLC_FPSIZE == 32
   return __builtin_amdgcn_rcpf(x);
+#elif defined(__CLC_SCALAR) && __CLC_FPSIZE == 64
+  // Hardware v_rcp_f64 seed refined with two Newton-Raphson iterations. This
+  // computes 1.0 / x without the full IEEE scaling and subnormal fixups.
+  __CLC_GENTYPE r = __builtin_amdgcn_rcp(x);
+  r = __builtin_fma(__builtin_fma(-x, r, 1.0), r, r);
+  r = __builtin_fma(__builtin_fma(-x, r, 1.0), r, r);
+  return r;
 #else
   return ((__CLC_GENTYPE)1.0) / x;
 #endif