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 <[email protected]> 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 <[email protected]> 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: { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
