Skip to content

[LLVM][NVPTX] Upstream tanh intrinsic for libdevice #149596

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.

Sign up for GitHub

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 4 commits into from
Jul 24, 2025
Merged

[LLVM][NVPTX] Upstream tanh intrinsic for libdevice #149596

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
52fadde
initial commit
mjulian31 Jul 17, 2025
484adfe
fix comment
mjulian31 Jul 19, 2025
bfda682
simplify test
mjulian31 Jul 21, 2025
12ed616
add ptx70 and sm75 predicates, update test
mjulian31 Jul 23, 2025
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
7 changes: 7 additions & 0 deletions llvm/lib/IR/AutoUpgrade.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1450,6 +1450,7 @@ static bool upgradeIntrinsicFunction1(Function *F, Function *&NewFn,
.Case("popc.ll", true)
.Case("h2f", true)
.Case("swap.lo.hi.b64", true)
.Case("tanh.approx.f32", true)
.Default(false);

if (Expand) {
Expand Down Expand Up @@ -2543,6 +2544,12 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
MDNode *MD = MDNode::get(Builder.getContext(), {});
LD->setMetadata(LLVMContext::MD_invariant_load, MD);
return LD;
} else if (Name == "tanh.approx.f32") {
// nvvm.tanh.approx.f32 -> afn llvm.tanh.f32
FastMathFlags FMF;
FMF.setApproxFunc();
Rep = Builder.CreateUnaryIntrinsic(Intrinsic::tanh, CI->getArgOperand(0),
FMF);
} else if (Name == "barrier0" || Name == "barrier.n" || Name == "bar.sync") {
Value *Arg =
Name.ends_with('0') ? Builder.getInt32(0) : CI->getArgOperand(0);
Expand Down
7 changes: 5 additions & 2 deletions llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -952,10 +952,13 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
// promoted to f32. v2f16 is expanded to f16, which is then promoted
// to f32.
for (const auto &Op :
{ISD::FDIV, ISD::FREM, ISD::FSQRT, ISD::FSIN, ISD::FCOS}) {
{ISD::FDIV, ISD::FREM, ISD::FSQRT, ISD::FSIN, ISD::FCOS, ISD::FTANH}) {
setOperationAction(Op, MVT::f16, Promote);
setOperationAction(Op, MVT::f32, Legal);
setOperationAction(Op, MVT::f64, Legal);
// only div/rem/sqrt are legal for f64
if (Op == ISD::FDIV || Op == ISD::FREM || Op == ISD::FSQRT) {
setOperationAction(Op, MVT::f64, Legal);
Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this change tested? What will happen for the cases where this is not legal?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

In testing TANH, I noticed that we don't actually support f64 for these, despite marking them legal. See here for an example: https://godbolt.org/z/qjPhMc9nf. I figured I should fix this while I am changing this for TANH anyways.

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I see, I was more wondering if there is a unit test that covers this change. My guess is that the compiler will now just crash in a slightly different way as a result of this change, but I suppose it's better to have a more accurate operation action table regardless.

}
setOperationAction(Op, {MVT::v2f16, MVT::v2bf16, MVT::v2f32}, Expand);
setOperationAction(Op, MVT::bf16, Promote);
AddPromotedToType(Op, MVT::bf16, MVT::f32);
Expand Down
6 changes: 5 additions & 1 deletion llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1234,7 +1234,7 @@ defm FMA_F32 : FMA<F32RT, allow_ftz = true>;
defm FMA_F32x2 : FMA<F32X2RT, allow_ftz = true, preds = [hasF32x2Instructions]>;
defm FMA_F64 : FMA<F64RT, allow_ftz = false>;

// sin/cos
// sin/cos/tanh

class UnaryOpAllowsApproxFn<SDPatternOperator operator>
: PatFrag<(ops node:$A),
Expand All @@ -1250,6 +1250,10 @@ def COS_APPROX_f32 :
BasicFlagsNVPTXInst<(outs B32:$dst), (ins B32:$src), (ins FTZFlag:$ftz),
"cos.approx$ftz.f32",
[(set f32:$dst, (UnaryOpAllowsApproxFn<fcos> f32:$src))]>;
def TANH_APPROX_f32 :
BasicNVPTXInst<(outs B32:$dst), (ins B32:$src), "tanh.approx.f32",
[(set f32:$dst, (UnaryOpAllowsApproxFn<ftanh> f32:$src))]>,
Requires<[hasPTX<70>, hasSM<75>]>;

//-----------------------------------
// Bitwise operations
Expand Down
9 changes: 9 additions & 0 deletions llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,6 +17,8 @@ declare float @llvm.nvvm.fabs.f(float)
declare float @llvm.nvvm.fabs.ftz.f(float)
declare double @llvm.nvvm.fabs.d(double)

declare float @llvm.nvvm.tanh.approx.f32(float)

declare i16 @llvm.nvvm.max.s(i16, i16)
declare i32 @llvm.nvvm.max.i(i32, i32)
declare i64 @llvm.nvvm.max.ll(i64, i64)
Expand Down Expand Up @@ -138,6 +140,13 @@ define void @fabs(float %a, double %b) {
ret void
}

; CHECK-LABEL: @tanh
define void @tanh(float %a) {
; CHECK: call afn float @llvm.tanh.f32(float %a)
%r1 = call float @llvm.nvvm.tanh.approx.f32(float %a)
ret void
}

; CHECK-LABEL: @min_max
define void @min_max(i16 %a1, i16 %a2, i32 %b1, i32 %b2, i64 %c1, i64 %c2) {
; CHECK: [[maxs:%[a-zA-Z0-9.]+]] = icmp sge i16 %a1, %a2
Expand Down
40 changes: 40 additions & 0 deletions llvm/test/CodeGen/NVPTX/tanhf.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mcpu=sm_75 -mattr=+ptx70 | FileCheck %s
; RUN: %if ptxas-11.0 %{ llc < %s -mcpu=sm_75 -mattr=+ptx70 | %ptxas-verify -arch=sm_75 %}

target triple = "nvptx64-nvidia-cuda"

define float @test1(float %in) local_unnamed_addr {
; CHECK-LABEL: test1(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b32 %r1, [test1_param_0];
; CHECK-NEXT: tanh.approx.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%call = call afn float @llvm.tanh.f32(float %in)
ret float %call
}

define half @test2(half %in) local_unnamed_addr {
; CHECK-LABEL: test2(
; CHECK: {
; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [test2_param_0];
; CHECK-NEXT: cvt.f32.f16 %r1, %rs1;
; CHECK-NEXT: tanh.approx.f32 %r2, %r1;
; CHECK-NEXT: cvt.rn.f16.f32 %rs2, %r2;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-NEXT: ret;
%call = call afn half @llvm.tanh.f16(half %in)
ret half %call
}

declare float @llvm.tanh.f32(float)
declare half @llvm.tanh.f16(half)