Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly

Appearance settings

[AMDGPU][clang] provide device implementation for __builtin_logb and … #129347

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
May 19, 2025
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
35 changes: 32 additions & 3 deletions 35 clang/lib/CodeGen/CGBuiltin.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,33 @@ using namespace clang;
using namespace CodeGen;
using namespace llvm;

/// Some builtins do not have library implementation on some targets and
/// are instead emitted as LLVM IRs by some target builtin emitters.
/// FIXME: Remove this when library support is added
static bool shouldEmitBuiltinAsIR(unsigned BuiltinID,
const Builtin::Context &BI,
const CodeGenFunction &CGF) {
if (!CGF.CGM.getLangOpts().MathErrno &&
CGF.CurFPFeatures.getExceptionMode() ==
LangOptions::FPExceptionModeKind::FPE_Ignore &&
!CGF.CGM.getTargetCodeGenInfo().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 Value *EmitTargetArchBuiltinExpr(CodeGenFunction *CGF,
unsigned BuiltinID, const CallExpr *E,
ReturnValueSlot ReturnValue,
Expand Down Expand Up @@ -2414,7 +2441,7 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
// disabled.
// Math intrinsics are generated only when math-errno is disabled. Any pragmas
// or attributes that affect math-errno should prevent or allow math
// intrincs to be generated. Intrinsics are generated:
// intrinsics to be generated. Intrinsics are generated:
// 1- In fast math mode, unless math-errno is overriden
// via '#pragma float_control(precise, on)', or via an
// 'attribute__((optnone))'.
Expand Down Expand Up @@ -5999,13 +6026,15 @@ RValue CodeGenFunction::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))
const auto &BI = getContext().BuiltinInfo;
if (!shouldEmitBuiltinAsIR(BuiltinID, BI, *this) &&
BI.isLibFunction(BuiltinID))
return emitLibraryCall(*this, FD, E,
CGM.getBuiltinLibFunction(FD, BuiltinID));

// If this is a predefined lib function (e.g. malloc), emit the call
// using exactly the normal call path.
if (getContext().BuiltinInfo.isPredefinedLibFunction(BuiltinID))
if (BI.isPredefinedLibFunction(BuiltinID))
return emitLibraryCall(*this, FD, E, CGM.getRawFunctionPointer(FD));

// Check that a call to a target specific builtin has the correct target
Expand Down
72 changes: 72 additions & 0 deletions 72 clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,27 @@ using namespace CodeGen;
using namespace llvm;

namespace {

// Has second type mangled argument.
static Value *
emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E,
Intrinsic::ID IntrinsicID,
Intrinsic::ID ConstrainedIntrinsicID) {
llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));

CodeGenFunction::CGFPOptionsRAII FPOptsRAII(CGF, E);
if (CGF.Builder.getIsFPConstrained()) {
Function *F = CGF.CGM.getIntrinsic(ConstrainedIntrinsicID,
{Src0->getType(), Src1->getType()});
return CGF.Builder.CreateConstrainedFPCall(F, {Src0, Src1});
}

Function *F =
CGF.CGM.getIntrinsic(IntrinsicID, {Src0->getType(), Src1->getType()});
return CGF.Builder.CreateCall(F, {Src0, Src1});
}

// If \p E is not null pointer, insert address space cast to match return
// type of \p E if necessary.
Value *EmitAMDGPUDispatchPtr(CodeGenFunction &CGF,
Expand Down Expand Up @@ -1142,6 +1163,57 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_s_prefetch_data:
return emitBuiltinWithOneOverloadedType<2>(
*this, E, Intrinsic::amdgcn_s_prefetch_data);
case Builtin::BIlogbf:
case Builtin::BI__builtin_logbf: {
Value *Src0 = EmitScalarExpr(E->getArg(0));
Function *FrExpFunc = CGM.getIntrinsic(
Intrinsic::frexp, {Src0->getType(), Builder.getInt32Ty()});
CallInst *FrExp = Builder.CreateCall(FrExpFunc, Src0);
Value *Exp = Builder.CreateExtractValue(FrExp, 1);
Value *Add = Builder.CreateAdd(
Exp, ConstantInt::getSigned(Exp->getType(), -1), "", false, true);
Value *SIToFP = Builder.CreateSIToFP(Add, Builder.getFloatTy());
Value *Fabs =
emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::fabs);
Value *FCmpONE = Builder.CreateFCmpONE(
Fabs, ConstantFP::getInfinity(Builder.getFloatTy()));
Value *Sel1 = Builder.CreateSelect(FCmpONE, SIToFP, Fabs);
Value *FCmpOEQ =
Builder.CreateFCmpOEQ(Src0, ConstantFP::getZero(Builder.getFloatTy()));
Value *Sel2 = Builder.CreateSelect(
FCmpOEQ,
ConstantFP::getInfinity(Builder.getFloatTy(), /*Negative=*/true), Sel1);
return Sel2;
}
case Builtin::BIlogb:
case Builtin::BI__builtin_logb: {
choikwa marked this conversation as resolved.
Show resolved Hide resolved
Value *Src0 = EmitScalarExpr(E->getArg(0));
Function *FrExpFunc = CGM.getIntrinsic(
Intrinsic::frexp, {Src0->getType(), Builder.getInt32Ty()});
CallInst *FrExp = Builder.CreateCall(FrExpFunc, Src0);
Value *Exp = Builder.CreateExtractValue(FrExp, 1);
Value *Add = Builder.CreateAdd(
Exp, ConstantInt::getSigned(Exp->getType(), -1), "", false, true);
Value *SIToFP = Builder.CreateSIToFP(Add, Builder.getDoubleTy());
Value *Fabs =
emitBuiltinWithOneOverloadedType<1>(*this, E, Intrinsic::fabs);
Value *FCmpONE = Builder.CreateFCmpONE(
Fabs, ConstantFP::getInfinity(Builder.getDoubleTy()));
Value *Sel1 = Builder.CreateSelect(FCmpONE, SIToFP, Fabs);
Value *FCmpOEQ =
Builder.CreateFCmpOEQ(Src0, ConstantFP::getZero(Builder.getDoubleTy()));
Value *Sel2 = Builder.CreateSelect(
FCmpOEQ,
ConstantFP::getInfinity(Builder.getDoubleTy(), /*Negative=*/true),
Sel1);
return Sel2;
}
case Builtin::BIscalbnf:
case Builtin::BI__builtin_scalbnf:
case Builtin::BIscalbn:
case Builtin::BI__builtin_scalbn:
return emitBinaryExpMaybeConstrainedFPBuiltin(
*this, E, Intrinsic::ldexp, Intrinsic::experimental_constrained_ldexp);
default:
return nullptr;
}
Expand Down
4 changes: 4 additions & 0 deletions 4 clang/lib/CodeGen/TargetInfo.h
Original file line number Diff line number Diff line change
Expand Up @@ -71,6 +71,10 @@ class TargetCodeGenInfo {
return *SwiftInfo;
}

/// supportsLibCall - Query to whether or not target supports all
/// lib calls.
virtual bool supportsLibCall() const { return true; }

/// setTargetAttributes - Provides a convenient hook to handle extra
/// target-specific attributes for the given global.
virtual void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
Expand Down
1 change: 1 addition & 0 deletions 1 clang/lib/CodeGen/Targets/AMDGPU.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -302,6 +302,7 @@ class AMDGPUTargetCodeGenInfo : public TargetCodeGenInfo {
AMDGPUTargetCodeGenInfo(CodeGenTypes &CGT)
: TargetCodeGenInfo(std::make_unique<AMDGPUABIInfo>(CGT)) {}

bool supportsLibCall() const override { return false; }
void setFunctionDeclAttributes(const FunctionDecl *FD, llvm::Function *F,
CodeGenModule &CGM) const;

Expand Down
Loading
Morty Proxy This is a proxified and sanitized view of the page, visit original site.