llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Rana Pratap Reddy (ranapratap55) <details> <summary>Changes</summary> Upstreaming clangIR PR: https://github.com/llvm/clangir/pull/2065 Support for lowering of `__builtin_amdgcn_logb` and `scalebn` for AMDGPU builtins to clangIR. Followed similar lowering from clang->llvmir: `clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp`. --- Full diff: https://github.com/llvm/llvm-project/pull/191344.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp (+89-10) - (added) clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip (+42) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index b4b0c455904fc..de9f8951823f9 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -18,6 +18,90 @@ using namespace clang; using namespace clang::CIRGen; +using namespace cir; + +static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin( + CIRGenFunction &CGF, const CallExpr *E, llvm::StringRef IntrinsicName, + llvm::StringRef ConstrainedIntrinsicName) { + mlir::Value Src0 = CGF.emitScalarExpr(E->getArg(0)); + mlir::Value Src1 = CGF.emitScalarExpr(E->getArg(1)); + + auto &Builder = CGF.getBuilder(); + + CIRGenFunction::CIRGenFPOptionsRAII FPOptsRAII(CGF, E); + + if (Builder.getIsFPConstrained()) { + return cir::LLVMIntrinsicCallOp::create( + Builder, CGF.getLoc(E->getExprLoc()), + Builder.getStringAttr(ConstrainedIntrinsicName), Src0.getType(), + {Src0, Src1}) + .getResult(); + } + + return cir::LLVMIntrinsicCallOp::create(Builder, CGF.getLoc(E->getExprLoc()), + Builder.getStringAttr(IntrinsicName), + Src0.getType(), {Src0, Src1}) + .getResult(); +} + +static mlir::Value emitLogbBuiltin(CIRGenFunction &CGF, const CallExpr *E, + bool IsFloat) { + auto &Builder = CGF.getBuilder(); + mlir::Location Loc = CGF.getLoc(E->getExprLoc()); + + mlir::Value Src0 = CGF.emitScalarExpr(E->getArg(0)); + mlir::Type SrcTy = Src0.getType(); + mlir::Type Int32Ty = Builder.getSInt32Ty(); + + cir::RecordType FrExpResTy = + Builder.getAnonRecordTy({SrcTy, Int32Ty}, false, false); + + mlir::Value FrExpResult = + cir::LLVMIntrinsicCallOp::create( + Builder, Loc, Builder.getStringAttr("llvm.frexp"), FrExpResTy, {Src0}) + .getResult(); + + mlir::Value Exp = + cir::ExtractMemberOp::create(Builder, Loc, Int32Ty, FrExpResult, 1); + + mlir::Value NegativeOne = + Builder.getConstant(Loc, cir::IntAttr::get(Int32Ty, -1)); + mlir::Value ExpMinus1 = Builder.createAdd(Loc, Exp, NegativeOne); + + mlir::Value SIToFP = cir::CastOp::create( + Builder, Loc, SrcTy, cir::CastKind::int_to_float, ExpMinus1); + + mlir::Value Fabs = cir::FAbsOp::create(Builder, Loc, SrcTy, Src0); + + llvm::APFloat InfVal = + IsFloat ? llvm::APFloat::getInf(llvm::APFloat::IEEEsingle()) + : llvm::APFloat::getInf(llvm::APFloat::IEEEdouble()); + mlir::Value Inf = Builder.getConstant(Loc, cir::FPAttr::get(SrcTy, InfVal)); + + mlir::Value FabsNegInf = + Builder.createCompare(Loc, cir::CmpOpKind::ne, Fabs, Inf); + + mlir::Value Sel = Builder.createSelect(Loc, FabsNegInf, SIToFP, Fabs); + + llvm::APFloat ZeroValue = + IsFloat ? llvm::APFloat::getZero(llvm::APFloat::IEEEsingle()) + : llvm::APFloat::getZero(llvm::APFloat::IEEEdouble()); + mlir::Value Zero = + Builder.getConstant(Loc, cir::FPAttr::get(SrcTy, ZeroValue)); + + mlir::Value SrcEqZero = + Builder.createCompare(Loc, cir::CmpOpKind::eq, Src0, Zero); + + llvm::APFloat NegInfVal = + IsFloat ? llvm::APFloat::getInf(llvm::APFloat::IEEEsingle(), true) + : llvm::APFloat::getInf(llvm::APFloat::IEEEdouble(), true); + mlir::Value NegInf = + Builder.getConstant(Loc, cir::FPAttr::get(SrcTy, NegInfVal)); + + mlir::Value Result = Builder.createSelect(Loc, SrcEqZero, NegInf, Sel); + + return Result; +} std::optional<mlir::Value> CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, @@ -807,20 +891,15 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, return mlir::Value{}; } case Builtin::BIlogbf: - case Builtin::BI__builtin_logbf: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; - } + case Builtin::BI__builtin_logbf: + return emitLogbBuiltin(*this, expr, true); case Builtin::BIscalbnf: case Builtin::BI__builtin_scalbnf: + return emitLogbBuiltin(*this, expr, false); case Builtin::BIscalbn: case Builtin::BI__builtin_scalbn: { - cgm.errorNYI(expr->getSourceRange(), - std::string("unimplemented AMDGPU builtin call: ") + - getContext().BuiltinInfo.getName(builtinId)); - return mlir::Value{}; + return emitBinaryExpMaybeConstrainedFPBuiltin( + *this, expr, "llvm.ldexp", "llvm.experimental.constrained.ldexp"); } default: return std::nullopt; diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip new file mode 100644 index 0000000000000..6d0cfa6bed5c2 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip @@ -0,0 +1,42 @@ +#include "../CodeGenCUDA/Inputs/cuda.h" + +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \ +// RUN: -target-cpu gfx1100 -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 gfx1100 -fcuda-is-device -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s + +// CIR-LABEL: @_Z11test_logbfff +// CIR: cir.call @logbf({{.*}}){{.*}}: (!cir.float{{.*}}) -> !cir.float +// LLVM: define{{.*}} float @_Z11test_logbfff( +// LLVM: call {{.*}}float @logbf(float{{.*}}%{{.*}}) +__device__ float test_logbff(float a) { + return __builtin_logbf(a); +} + +// CIR-LABEL: @_Z11test_logbddd +// CIR: cir.call @logb({{.*}}){{.*}}: (!cir.double{{.*}}) -> !cir.double +// LLVM: define{{.*}} double @_Z11test_logbddd( +// LLVM: call {{.*}}double @logb(double{{.*}}%{{.*}}) +__device__ double test_logbdd(double a) { + return __builtin_logb(a); +} + +// CIR-LABEL: @_Z14test_scalbnffifi +// CIR: cir.call @scalbnf({{.*}}){{.*}}: (!cir.float{{.*}}, !s32i{{.*}}) -> !cir.float +// LLVM: define{{.*}} float @_Z14test_scalbnffifi( +// LLVM: call {{.*}}float @scalbnf(float{{.*}}%{{.*}}, i32{{.*}}%{{.*}}) +__device__ float test_scalbnffi(float a, int b) { + return __builtin_scalbnf(a, b); +} + +// CIR-LABEL: @_Z14test_scalbnfdidi +// CIR: cir.call @scalbn({{.*}}){{.*}}: (!cir.double{{.*}}, !s32i{{.*}}) -> !cir.double +// LLVM: define{{.*}} double @_Z14test_scalbnfdidi( +// LLVM: call {{.*}}double @scalbn(double{{.*}}%{{.*}}, i32{{.*}}%{{.*}}) +__device__ double test_scalbnfdi(double a, int b) { + return __builtin_scalbn(a, b); +} `````````` </details> https://github.com/llvm/llvm-project/pull/191344 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
