diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp index 589833c0971b8..fa0d02f9e4eef 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp @@ -33,6 +33,30 @@ using namespace clang; using namespace clang::CIRGen; using namespace llvm; +static bool shouldEmitBuiltinAsIR(unsigned builtinID, + const Builtin::Context &bi, + const CIRGenFunction &cgf) { + if (!cgf.cgm.getLangOpts().MathErrno && + cgf.curFPFeatures.getExceptionMode() == + LangOptions::FPExceptionModeKind::FPE_Ignore && + !cgf.cgm.getTargetCIRGenInfo().supportsLibCall()) { + switch (builtinID) { + default: + return false; + case Builtin::BIlogbf: + case Builtin::BI__builtin_logbf: + case Builtin::BIlogb: + case Builtin::BI__builtin_logb: + case Builtin::BIscalbnf: + case Builtin::BI__builtin_scalbnf: + case Builtin::BIscalbn: + case Builtin::BI__builtin_scalbn: + return true; + } + } + return false; +} + static RValue emitLibraryCall(CIRGenFunction &cgf, const FunctionDecl *fd, const CallExpr *e, mlir::Operation *calleeValue) { CIRGenCallee callee = CIRGenCallee::forDirect(calleeValue, GlobalDecl(fd)); @@ -2386,7 +2410,8 @@ RValue CIRGenFunction::emitBuiltinExpr(const GlobalDecl &gd, unsigned builtinID, // If this is an alias for a lib function (e.g. __builtin_sin), emit // the call using the normal call path, but using the unmangled // version of the function name. - if (getContext().BuiltinInfo.isLibFunction(builtinID)) + if (!shouldEmitBuiltinAsIR(builtinID, getContext().BuiltinInfo, *this) && + getContext().BuiltinInfo.isLibFunction(builtinID)) return emitLibraryCall(*this, fd, e, cgm.getBuiltinLibFunction(fd, builtinID)); diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 186c9d3cf47e7..929cdf8e88789 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -18,6 +18,78 @@ 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)); + mlir::Location loc = cgf.getLoc(e->getExprLoc()); + + CIRGenBuilderTy &builder = cgf.getBuilder(); + + CIRGenFunction::CIRGenFPOptionsRAII fpOptsRAII(cgf, e); + + if (builder.getIsFPConstrained()) { + cgf.cgm.errorNYI(e->getSourceRange(), + "constrained FP intrinsic support is NYI."); + } + + return builder.emitIntrinsicCallOp(loc, intrinsicName, src0.getType(), + mlir::ValueRange{src0, src1}); +} + +static mlir::Value emitLogbBuiltin(CIRGenFunction &cgf, const CallExpr *e, + const llvm::fltSemantics &fSem) { + CIRGenBuilderTy &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 = builder.emitIntrinsicCallOp( + loc, "frexp", frExpResTy, mlir::ValueRange{src0}); + + 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 = llvm::APFloat::getInf(fSem); + 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 = llvm::APFloat::getZero(fSem); + 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 = llvm::APFloat::getInf(fSem, true); + mlir::Value negInf = + builder.getConstant(loc, cir::FPAttr::get(srcTy, negInfVal)); + + mlir::Value res = builder.createSelect(loc, srcEqZero, negInf, sel); + + return res; +} std::optional CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, @@ -837,20 +909,17 @@ 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, llvm::APFloat::IEEEsingle()); + case Builtin::BIlogb: + case Builtin::BI__builtin_logb: + return emitLogbBuiltin(*this, expr, llvm::APFloat::IEEEdouble()); case Builtin::BIscalbnf: case Builtin::BI__builtin_scalbnf: 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, "ldexp", "experimental.constrained.ldexp"); } default: return std::nullopt; diff --git a/clang/lib/CIR/CodeGen/TargetInfo.cpp b/clang/lib/CIR/CodeGen/TargetInfo.cpp index f674299168960..fc939cd9605ab 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.cpp +++ b/clang/lib/CIR/CodeGen/TargetInfo.cpp @@ -56,6 +56,8 @@ class AMDGPUTargetCIRGenInfo : public TargetCIRGenInfo { AMDGPUTargetCIRGenInfo(CIRGenTypes &cgt) : TargetCIRGenInfo(std::make_unique(cgt)) {} + bool supportsLibCall() const override { return false; } + void setTargetAttributes(const clang::Decl *decl, mlir::Operation *global, CIRGenModule &cgm) const override { if (auto func = mlir::dyn_cast(global)) { diff --git a/clang/lib/CIR/CodeGen/TargetInfo.h b/clang/lib/CIR/CodeGen/TargetInfo.h index 5e0103093827b..ecdfb7cb42c0e 100644 --- a/clang/lib/CIR/CodeGen/TargetInfo.h +++ b/clang/lib/CIR/CodeGen/TargetInfo.h @@ -50,6 +50,9 @@ class TargetCIRGenInfo { /// Returns ABI info helper for the target. const ABIInfo &getABIInfo() const { return *info; } + /// Returns true if the target supports math library calls. + virtual bool supportsLibCall() const { return true; } + /// Get target favored AST address space of a global variable for languages /// other than OpenCL and CUDA. /// If \p d is nullptr, returns the default target favored address space 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..e4091c3610530 --- /dev/null +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip @@ -0,0 +1,46 @@ +#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 + +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \ +// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s + +// CIR-LABEL: @_Z11test_logbfff +// CIR: cir.call_llvm_intrinsic "frexp" {{.*}} : (!cir.float) -> {{.*}} +// LLVM: define{{.*}} float @_Z11test_logbfff( +// LLVM: call { float, i32 } @llvm.frexp.f32.i32(float {{.*}}) +__device__ float test_logbff(float a) { + return __builtin_logbf(a); +} + +// CIR-LABEL: @_Z11test_logbddd +// CIR: cir.call_llvm_intrinsic "frexp" {{.*}} : (!cir.double) -> {{.*}} +// LLVM: define{{.*}} double @_Z11test_logbddd( +// LLVM: call { double, i32 } @llvm.frexp.f64.i32(double {{.*}}) +__device__ double test_logbdd(double a) { + return __builtin_logb(a); +} + +// CIR-LABEL: @_Z14test_scalbnffifi +// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.float, !s32i) -> !cir.float +// LLVM: define{{.*}} float @_Z14test_scalbnffifi( +// LLVM: call {{.*}}float @llvm.ldexp.f32.i32(float {{.*}}, i32 {{.*}}) +__device__ float test_scalbnffi(float a, int b) { + return __builtin_scalbnf(a, b); +} + +// CIR-LABEL: @_Z14test_scalbnfdidi +// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.double, !s32i) -> !cir.double +// LLVM: define{{.*}} double @_Z14test_scalbnfdidi( +// LLVM: call {{.*}}double @llvm.ldexp.f64.i32(double {{.*}}, i32 {{.*}}) +__device__ double test_scalbnfdi(double a, int b) { + return __builtin_scalbn(a, b); +}