[clang] [CIR][AMDGPU] Adds amdgcn logb and scalebn builtins (PR #191344)
Rana Pratap Reddy via cfe-commits
cfe-commits at lists.llvm.org
Thu Apr 9 22:57:35 PDT 2026
https://github.com/ranapratap55 created https://github.com/llvm/llvm-project/pull/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`.
>From 9909223f83058968341dd6ecdd8b2b9e002ed6e2 Mon Sep 17 00:00:00 2001
From: ranapratap55 <RanaPratapReddy.Nimmakayala at amd.com>
Date: Fri, 10 Apr 2026 09:52:22 +0530
Subject: [PATCH] [CIR][AMDGPU] Adds amdgcn logb and scalebn builtins
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 99 +++++++++++++++++--
.../builtins-amdgcn-logb-scalbn.hip | 42 ++++++++
2 files changed, 131 insertions(+), 10 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip
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);
+}
More information about the cfe-commits
mailing list