[clang] af4649f - [CIR][AMDGPU] Adds amdgcn logb and scalebn builtins (#191344)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Apr 29 22:33:58 PDT 2026
Author: Rana Pratap Reddy
Date: 2026-04-30T11:03:53+05:30
New Revision: af4649fe1c92022365043ccb48d217403e54de14
URL: https://github.com/llvm/llvm-project/commit/af4649fe1c92022365043ccb48d217403e54de14
DIFF: https://github.com/llvm/llvm-project/commit/af4649fe1c92022365043ccb48d217403e54de14.diff
LOG: [CIR][AMDGPU] Adds amdgcn logb and scalebn builtins (#191344)
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`.
Added:
clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip
Modified:
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
clang/lib/CIR/CodeGen/TargetInfo.cpp
clang/lib/CIR/CodeGen/TargetInfo.h
Removed:
################################################################################
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<mlir::Value>
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<AMDGPUABIInfo>(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<cir::FuncOp>(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);
+}
More information about the cfe-commits
mailing list