llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Ayokunle Amodu (ayokunle321) <details> <summary>Changes</summary> Adds codegen for the following AMDGPU floating point comparison builtins: - __builtin_amdgcn_fcmp (double) - __builtin_amdgcn_fcmpf (float) These are lowered to the corresponding `llvm.amdgcn.fcmp` intrinsic, returning a 64-bit lane mask. --- Full diff: https://github.com/llvm/llvm-project/pull/198135.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+6-4) - (modified) clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip (+16) ``````````diff 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); +} `````````` </details> https://github.com/llvm/llvm-project/pull/198135 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
