[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 16 23:11:51 PDT 2026
https://github.com/ranapratap55 updated https://github.com/llvm/llvm-project/pull/191344
>From 95b516e26a6d61d20d0eccdbf5e86752f98fa875 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 1/3] [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);
+}
>From 218e7478662f739f6841fdc011a0f0a78ef71b80 Mon Sep 17 00:00:00 2001
From: ranapratap55 <RanaPratapReddy.Nimmakayala at amd.com>
Date: Tue, 14 Apr 2026 10:17:39 +0530
Subject: [PATCH 2/3] [CIR][AMDGPU] Fix FltSemantics, naming convention, and
CIR APIs
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 113 ++++++++----------
.../builtins-amdgcn-logb-scalbn.hip | 12 ++
2 files changed, 64 insertions(+), 61 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index de9f8951823f9..fad1932b6593c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -21,86 +21,75 @@ 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));
+ 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());
- auto &Builder = CGF.getBuilder();
+ CIRGenBuilderTy &builder = cgf.getBuilder();
- CIRGenFunction::CIRGenFPOptionsRAII FPOptsRAII(CGF, E);
+ 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();
+ if (builder.getIsFPConstrained()) {
+ return builder.emitIntrinsicCallOp(loc, constrainedIntrinsicName,
+ src0.getType(),
+ mlir::ValueRange{src0, src1});
}
- return cir::LLVMIntrinsicCallOp::create(Builder, CGF.getLoc(E->getExprLoc()),
- Builder.getStringAttr(IntrinsicName),
- Src0.getType(), {Src0, Src1})
- .getResult();
+ return builder.emitIntrinsicCallOp(loc, intrinsicName, src0.getType(),
+ mlir::ValueRange{src0, src1});
}
-static mlir::Value emitLogbBuiltin(CIRGenFunction &CGF, const CallExpr *E,
- bool IsFloat) {
- auto &Builder = CGF.getBuilder();
- mlir::Location Loc = CGF.getLoc(E->getExprLoc());
+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();
+ 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);
+ 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 frExpResult = builder.emitIntrinsicCallOp(
+ loc, "frexp", frExpResTy, mlir::ValueRange{src0});
- mlir::Value Exp =
- cir::ExtractMemberOp::create(Builder, Loc, Int32Ty, FrExpResult, 1);
+ 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 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 siToFp = cir::CastOp::create(
+ builder, loc, srcTy, cir::CastKind::int_to_float, expMinus1);
- mlir::Value Fabs = cir::FAbsOp::create(Builder, Loc, SrcTy, Src0);
+ 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));
+ 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 fabsNegInf =
+ builder.createCompare(loc, cir::CmpOpKind::ne, fabs, inf);
- mlir::Value Sel = Builder.createSelect(Loc, FabsNegInf, SIToFP, Fabs);
+ 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));
+ 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);
+ 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));
+ llvm::APFloat negInfVal = llvm::APFloat::getInf(fSem, true);
+ mlir::Value negInf =
+ builder.getConstant(loc, cir::FPAttr::get(srcTy, negInfVal));
- mlir::Value Result = Builder.createSelect(Loc, SrcEqZero, NegInf, Sel);
+ mlir::Value res = builder.createSelect(loc, srcEqZero, negInf, sel);
- return Result;
+ return res;
}
std::optional<mlir::Value>
@@ -892,14 +881,16 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
}
case Builtin::BIlogbf:
case Builtin::BI__builtin_logbf:
- return emitLogbBuiltin(*this, expr, true);
+ 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:
- return emitLogbBuiltin(*this, expr, false);
case Builtin::BIscalbn:
case Builtin::BI__builtin_scalbn: {
return emitBinaryExpMaybeConstrainedFPBuiltin(
- *this, expr, "llvm.ldexp", "llvm.experimental.constrained.ldexp");
+ *this, expr, "ldexp", "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
index 6d0cfa6bed5c2..e6cbe223e6580 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip
@@ -9,10 +9,16 @@
// 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=OGCG --input-file=%t.ll %s
+
// CIR-LABEL: @_Z11test_logbfff
// CIR: cir.call @logbf({{.*}}){{.*}}: (!cir.float{{.*}}) -> !cir.float
// LLVM: define{{.*}} float @_Z11test_logbfff(
// LLVM: call {{.*}}float @logbf(float{{.*}}%{{.*}})
+// OGCG: define{{.*}} float @_Z11test_logbfff(
+// OGCG: call { float, i32 } @llvm.frexp.f32.i32(float %{{.*}})
__device__ float test_logbff(float a) {
return __builtin_logbf(a);
}
@@ -21,6 +27,8 @@ __device__ float test_logbff(float a) {
// CIR: cir.call @logb({{.*}}){{.*}}: (!cir.double{{.*}}) -> !cir.double
// LLVM: define{{.*}} double @_Z11test_logbddd(
// LLVM: call {{.*}}double @logb(double{{.*}}%{{.*}})
+// OGCG: define{{.*}} double @_Z11test_logbddd(
+// OGCG: call { double, i32 } @llvm.frexp.f64.i32(double %{{.*}})
__device__ double test_logbdd(double a) {
return __builtin_logb(a);
}
@@ -29,6 +37,8 @@ __device__ double test_logbdd(double a) {
// CIR: cir.call @scalbnf({{.*}}){{.*}}: (!cir.float{{.*}}, !s32i{{.*}}) -> !cir.float
// LLVM: define{{.*}} float @_Z14test_scalbnffifi(
// LLVM: call {{.*}}float @scalbnf(float{{.*}}%{{.*}}, i32{{.*}}%{{.*}})
+// OGCG: define{{.*}} float @_Z14test_scalbnffifi(
+// OGCG: call {{.*}}float @llvm.ldexp.f32.i32(float %{{.*}}, i32 %{{.*}})
__device__ float test_scalbnffi(float a, int b) {
return __builtin_scalbnf(a, b);
}
@@ -37,6 +47,8 @@ __device__ float test_scalbnffi(float a, int b) {
// CIR: cir.call @scalbn({{.*}}){{.*}}: (!cir.double{{.*}}, !s32i{{.*}}) -> !cir.double
// LLVM: define{{.*}} double @_Z14test_scalbnfdidi(
// LLVM: call {{.*}}double @scalbn(double{{.*}}%{{.*}}, i32{{.*}}%{{.*}})
+// OGCG: define{{.*}} double @_Z14test_scalbnfdidi(
+// OGCG: call {{.*}}double @llvm.ldexp.f64.i32(double %{{.*}}, i32 %{{.*}})
__device__ double test_scalbnfdi(double a, int b) {
return __builtin_scalbn(a, b);
}
>From 02bee9a92b0cabd619ccb769240301eb4e7cbbdb Mon Sep 17 00:00:00 2001
From: ranapratap55 <RanaPratapReddy.Nimmakayala at amd.com>
Date: Fri, 17 Apr 2026 11:40:36 +0530
Subject: [PATCH 3/3] [CIR][AMDGPU] Fix constrained FP and library calls path
---
clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp | 27 ++++++++++++++++++-
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 5 ++--
clang/lib/CIR/CodeGen/TargetInfo.cpp | 2 ++
clang/lib/CIR/CodeGen/TargetInfo.h | 3 +++
.../builtins-amdgcn-logb-scalbn.hip | 26 +++++++-----------
5 files changed, 42 insertions(+), 21 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltin.cpp
index 1a6aeef73fb79..f1d199cc1505e 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));
@@ -2287,7 +2311,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 fad1932b6593c..459643e15eda1 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -32,9 +32,8 @@ static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(
CIRGenFunction::CIRGenFPOptionsRAII fpOptsRAII(cgf, e);
if (builder.getIsFPConstrained()) {
- return builder.emitIntrinsicCallOp(loc, constrainedIntrinsicName,
- src0.getType(),
- mlir::ValueRange{src0, src1});
+ cgf.cgm.errorNYI(e->getSourceRange(),
+ "constrained FP intrinsic support is NYI.");
}
return builder.emitIntrinsicCallOp(loc, intrinsicName, src0.getType(),
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
index e6cbe223e6580..e4091c3610530 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-logb-scalbn.hip
@@ -11,44 +11,36 @@
// 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=OGCG --input-file=%t.ll %s
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
// CIR-LABEL: @_Z11test_logbfff
-// CIR: cir.call @logbf({{.*}}){{.*}}: (!cir.float{{.*}}) -> !cir.float
+// CIR: cir.call_llvm_intrinsic "frexp" {{.*}} : (!cir.float) -> {{.*}}
// LLVM: define{{.*}} float @_Z11test_logbfff(
-// LLVM: call {{.*}}float @logbf(float{{.*}}%{{.*}})
-// OGCG: define{{.*}} float @_Z11test_logbfff(
-// OGCG: call { float, i32 } @llvm.frexp.f32.i32(float %{{.*}})
+// 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 @logb({{.*}}){{.*}}: (!cir.double{{.*}}) -> !cir.double
+// CIR: cir.call_llvm_intrinsic "frexp" {{.*}} : (!cir.double) -> {{.*}}
// LLVM: define{{.*}} double @_Z11test_logbddd(
-// LLVM: call {{.*}}double @logb(double{{.*}}%{{.*}})
-// OGCG: define{{.*}} double @_Z11test_logbddd(
-// OGCG: call { double, i32 } @llvm.frexp.f64.i32(double %{{.*}})
+// 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 @scalbnf({{.*}}){{.*}}: (!cir.float{{.*}}, !s32i{{.*}}) -> !cir.float
+// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.float, !s32i) -> !cir.float
// LLVM: define{{.*}} float @_Z14test_scalbnffifi(
-// LLVM: call {{.*}}float @scalbnf(float{{.*}}%{{.*}}, i32{{.*}}%{{.*}})
-// OGCG: define{{.*}} float @_Z14test_scalbnffifi(
-// OGCG: call {{.*}}float @llvm.ldexp.f32.i32(float %{{.*}}, i32 %{{.*}})
+// 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 @scalbn({{.*}}){{.*}}: (!cir.double{{.*}}, !s32i{{.*}}) -> !cir.double
+// CIR: cir.call_llvm_intrinsic "ldexp" {{.*}} : (!cir.double, !s32i) -> !cir.double
// LLVM: define{{.*}} double @_Z14test_scalbnfdidi(
-// LLVM: call {{.*}}double @scalbn(double{{.*}}%{{.*}}, i32{{.*}}%{{.*}})
-// OGCG: define{{.*}} double @_Z14test_scalbnfdidi(
-// OGCG: call {{.*}}double @llvm.ldexp.f64.i32(double %{{.*}}, i32 %{{.*}})
+// 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