[clang] da13805 - [CIR][AMDGPU] Add support for AMDGCN rcp builtins (#197447)

[via cfe-commits]

Author: Ayokunle Amodu
Date: 2026-07-13T06:21:15-04:00
New Revision: da138053e78c5af7b83f49674b4e190e1803102d

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

LOG: [CIR][AMDGPU] Add support for AMDGCN rcp builtins (#197447)

Adds codegen for the following AMDGCN reciprocal builtins:

- __builtin_amdgcn_rcp (double)
- __builtin_amdgcn_rcpf (float)
- __builtin_amdgcn_rcph (half)
- __builtin_amdgcn_rcp_bf16 (bfloat16)

These are lowered to the corresponding `llvm.amdgcn.rcp` intrinsic.

Added: 
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip

Modified: 
    clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
    clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
    clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0a7ba0c194400..5736eb70bc123 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -251,10 +251,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   case AMDGPU::BI__builtin_amdgcn_rcpf:
   case AMDGPU::BI__builtin_amdgcn_rcph:
   case AMDGPU::BI__builtin_amdgcn_rcp_bf16: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.rcp").getValue();
   }
   case AMDGPU::BI__builtin_amdgcn_sqrt:
   case AMDGPU::BI__builtin_amdgcn_sqrtf:

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
new file mode 100644
index 0000000000000..64fa56b45e39a
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-gfx1250.hip
@@ -0,0 +1,26 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN:            -target-cpu gfx1250 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN:            -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
+// RUN:            -target-cpu gfx1250 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU builtins
+//===----------------------------------------------------------------------===//
+
+#define __device__ __attribute__((device))
+
+// CIR-LABEL: @_Z13test_rcp_bf16PDF16bDF16b
+// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.bf16) -> !cir.bf16
+// LLVM: define{{.*}} void @_Z13test_rcp_bf16PDF16bDF16b
+// LLVM: call{{.*}} bfloat @llvm.amdgcn.rcp.bf16(bfloat %{{.*}})
+__device__ void test_rcp_bf16(__bf16* out, __bf16 a) {
+  *out = __builtin_amdgcn_rcp_bf16(a);
+}

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
index c21359ff7a88b..3ee9f1664f61a 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-vi-f16.hip
@@ -63,3 +63,11 @@
 __device__ void test_div_fixup_f16(_Float16* out, _Float16 a, _Float16 b, _Float16 c) {
   *out = __builtin_amdgcn_div_fixuph(a, b, c);
 }
+
+// CIR-LABEL: @_Z12test_rcp_f16PDF16_DF16_
+// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.f16) -> !cir.f16
+// LLVM: define{{.*}} void @_Z12test_rcp_f16PDF16_DF16_
+// LLVM: call{{.*}} half @llvm.amdgcn.rcp.f16(half %{{.*}})
+__device__ void test_rcp_f16(_Float16* out, _Float16 a) {
+  *out = __builtin_amdgcn_rcph(a);
+}

diff  --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index ca708bca8587b..9e279d66de1fa 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -111,3 +111,19 @@ __device__ void test_readfirstlane(int* out, int a) {
 __device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) {
   *out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr();
 }
+
+// CIR-LABEL: @_Z12test_rcp_f32Pff
+// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.float) -> !cir.float
+// LLVM: define{{.*}} void @_Z12test_rcp_f32Pff
+// LLVM: call{{.*}} float @llvm.amdgcn.rcp.f32(float %{{.*}})
+__device__ void test_rcp_f32(float* out, float a) {
+  *out = __builtin_amdgcn_rcpf(a);
+}
+
+// CIR-LABEL: @_Z12test_rcp_f64Pdd
+// CIR: cir.call_llvm_intrinsic "amdgcn.rcp" {{.*}} : (!cir.double) -> !cir.double
+// LLVM: define{{.*}} void @_Z12test_rcp_f64Pdd
+// LLVM: call{{.*}} double @llvm.amdgcn.rcp.f64(double %{{.*}})
+__device__ void test_rcp_f64(double* out, double a) {
+  *out = __builtin_amdgcn_rcp(a);
+}