[clang] [CIR][AMDGPU] Add support for AMDGCN fcmp builtins (PR #198135)

[Ayokunle Amodu via cfe-commits]

https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/198135

>From 95774c7d7b2aa9ad09f4ae96a6f19d2bbac51f96 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Sat, 16 May 2026 20:47:38 -0400
Subject: [PATCH] add amdgcn fcmp buitins

---
 clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 10 ++++++----
 clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 16 ++++++++++++++++
 2 files changed, 22 insertions(+), 4 deletions(-)

diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 04ab1c29b0d63..1180870f85e33 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -366,10 +366,12 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
   }
   case AMDGPU::BI__builtin_amdgcn_fcmp:
   case AMDGPU::BI__builtin_amdgcn_fcmpf: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
+    mlir::Value src0 = emitScalarExpr(expr->getArg(0));
+    mlir::Value src1 = emitScalarExpr(expr->getArg(1));
+    mlir::Value src2 = emitScalarExpr(expr->getArg(2));
+    return builder.emitIntrinsicCallOp(getLoc(expr->getExprLoc()),
+                                       "amdgcn.fcmp", builder.getUInt64Ty(),
+                                       mlir::ValueRange{src0, src1, src2});
   }
   case AMDGPU::BI__builtin_amdgcn_class:
   case AMDGPU::BI__builtin_amdgcn_classf:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 4a61fde7aa90c..9fca0a1309332 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -71,3 +71,19 @@ __device__ void test_div_fmas_f64(double* out, double a, double b, double c, int
 __device__ void test_ds_swizzle(int* out, int a) {
   *out = __builtin_amdgcn_ds_swizzle(a, 32);
 }
+
+// CIR-LABEL: @_Z13test_fcmp_f32Pyff
+// CIR: cir.call_llvm_intrinsic "amdgcn.fcmp" {{.*}} : (!cir.float, !cir.float, !s32i) -> !u64i
+// LLVM: define{{.*}} void @_Z13test_fcmp_f32Pyff
+// LLVM: call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f32(float %{{.*}}, float %{{.*}}, i32 5)
+__device__ void test_fcmp_f32(unsigned long long* out, float a, float b) {
+  *out = __builtin_amdgcn_fcmpf(a, b, 5);
+}
+
+// CIR-LABEL: @_Z13test_fcmp_f64Pydd
+// CIR: cir.call_llvm_intrinsic "amdgcn.fcmp" {{.*}} : (!cir.double, !cir.double, !s32i) -> !u64i
+// LLVM: define{{.*}} void @_Z13test_fcmp_f64Pydd
+// LLVM: call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f64(double %{{.*}}, double %{{.*}}, i32 6)
+__device__ void test_fcmp_f64(unsigned long long* out, double a, double b) {
+  *out = __builtin_amdgcn_fcmp(a, b, 3+3);
+}