[clang] [CIR][AMDGPU] Add support for AMDGCN fcmp builtins (PR #198135)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 10 17:17:46 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/198135
>From 6d36436f65b2c1d23d569d9497f71c7e77604ee8 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 1/2] 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 0a7ba0c194400..09a3dde2111b7 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -402,10 +402,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 ca708bca8587b..963b68d2ae2f7 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: @_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);
+}
>From 4f4bdaa280397079276bdd980e9c8242aa9c1339 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Fri, 10 Jul 2026 20:17:26 -0400
Subject: [PATCH 2/2] remove incorrect codegen
---
.../CIR/Dialect/Builder/CIRBaseBuilder.h | 34 ++++++++++++++++++-
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 13 ++-----
2 files changed, 36 insertions(+), 11 deletions(-)
diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
index f8a3aca76f102..92ed11913b354 100644
--- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
+++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h
@@ -596,9 +596,41 @@ class CIRBaseBuilderTy : public mlir::OpBuilder {
}
//===--------------------------------------------------------------------===//
- // Other Instructions
+ // Compare Instructions
//===--------------------------------------------------------------------===//
+ Value *CreateICmp(CmpInst::Predicate P, Value *LHS, Value *RHS,
+ const Twine &Name = "") {
+ if (auto *V = Folder.FoldCmp(P, LHS, RHS))
+ return V;
+ return Insert(new ICmpInst(P, LHS, RHS), Name);
+ }
+
+ // Create a quiet floating-point comparison (i.e. one that raises an FP
+ // exception only in the case where an input is a signaling NaN).
+ // Note that this differs from CreateFCmpS only if IsFPConstrained is true.
+ Value *CreateFCmp(CmpInst::Predicate P, Value *LHS, Value *RHS,
+ const Twine &Name = "", MDNode *FPMathTag = nullptr) {
+ return CreateFCmpHelper(P, LHS, RHS, Name, FPMathTag, {}, false);
+ }
+
+ mlir::Value createCmp(CmpInst::Predicate Pred, mlir::Value lhs, mlir::Value rhs,
+ const Twine &name = "", MDNode *FPMathTag = nullptr) {
+ return CmpInst::isFPPredicate(Pred)
+ ? CreateFCmp(Pred, LHS, RHS, Name, FPMathTag)
+ : CreateICmp(Pred, LHS, RHS, Name);
+ }
+
+private:
+ // Helper routine to create either a signaling or a quiet FP comparison.
+ mlir::Value *CreateFCmpHelper(CmpInst::Predicate P, Value *LHS, Value *RHS,
+ const Twine &Name, MDNode *FPMathTag,
+ FMFSource FMFSource, bool IsSignaling);
+
+ //===--------------------------------------------------------------------===//
+ // Other Instructions
+ //===--------------------------------------------------------------------===//
+public:
mlir::Value createExtractElement(mlir::Location loc, mlir::Value vec,
uint64_t idx) {
mlir::Value idxVal =
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 09a3dde2111b7..c961ecd8d6154 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -394,21 +394,14 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_uicmp:
case AMDGPU::BI__builtin_amdgcn_uicmpl:
case AMDGPU::BI__builtin_amdgcn_sicmp:
- case AMDGPU::BI__builtin_amdgcn_sicmpl: {
+ case AMDGPU::BI__builtin_amdgcn_sicmpl:
+ 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{};
}
- case AMDGPU::BI__builtin_amdgcn_fcmp:
- case AMDGPU::BI__builtin_amdgcn_fcmpf: {
- 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:
case AMDGPU::BI__builtin_amdgcn_classh: {
More information about the cfe-commits
mailing list