Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

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

Browse source 
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.
This commit is contained in:
Joseph Huber 2026-06-15 07:26:32 -05:00 committed by GitHub
parent 6794e31e89
commit cac31111d3
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
1 changed files with 8 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

9
libclc/clc/lib/amdgpu/math/clc_recip_fast.inc
View file

@ -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