[libclc] [libclc] Improve performance and precision of reciprocal functions (PR #203805)

Matt Arsenault via cfe-commits cfe-commits at lists.llvm.org
Sun Jun 14 23:03:15 PDT 2026


================
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// On AMDGPU the "fast" reciprocal is the hardware v_rcp_f32 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);
----------------
arsenm wrote:

I thought I was trying to avoid using this. Don't you get the raw instruction with afn + reciprocal pattern with all denormal modes?

https://github.com/llvm/llvm-project/pull/203805


More information about the cfe-commits mailing list