[llvm] [NVPTX] Add native tanh.approx support for f16 and bf16 (PR #207621)
Hao Ren via llvm-commits
llvm-commits at lists.llvm.org
Sun Jul 5 15:42:07 PDT 2026
https://github.com/nvidia-moomoo created https://github.com/llvm/llvm-project/pull/207621
PTX provides tanh.approx for f32, f16, and f16x2 starting with PTX 7.0 and SM75, and for bf16 and bf16x2 starting with PTX 7.8 and SM90.
Teach NVPTX lowering and instruction selection about these availability constraints. Use native operations when available and promote or expand unsupported f16 and bf16 forms through f32.
Extend the code generation test to cover the scalar and packed forms on SM75 and SM90.
>From 48fad283df3e7fd93a70f00b208ed4f679fb9217 Mon Sep 17 00:00:00 2001
From: Varad Rahul Kamthe <133588066+varadk27 at users.noreply.github.com>
Date: Wed, 1 Jul 2026 18:02:01 +0000
Subject: [PATCH] [NVPTX] Add native tanh.approx support for f16 and bf16
PTX provides tanh.approx for f32, f16, and f16x2 starting with PTX 7.0 and SM75, and for bf16 and bf16x2 starting with PTX 7.8 and SM90.
Teach NVPTX lowering and instruction selection about these availability constraints. Use native operations when available and promote or expand unsupported f16 and bf16 forms through f32.
Extend the code generation test to cover the scalar and packed forms on SM75 and SM90.
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 20 +++-
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td | 20 +++-
llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp | 1 +
llvm/test/CodeGen/NVPTX/tanhf.ll | 116 ++++++++++++++++++--
4 files changed, 140 insertions(+), 17 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index fa0b9fe5f7051..13cfa019655bc 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -562,6 +562,7 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
IsOpSupported &= STI.getSmVersion() >= 80 && STI.getPTXVersion() >= 70;
break;
case ISD::FEXP2:
+ case ISD::FTANH:
IsOpSupported &= STI.getSmVersion() >= 75 && STI.getPTXVersion() >= 70;
break;
}
@@ -1028,7 +1029,7 @@ 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::FTANH}) {
+ {ISD::FDIV, ISD::FREM, ISD::FSQRT, ISD::FSIN, ISD::FCOS}) {
setOperationAction(Op, MVT::f16, Promote);
setOperationAction(Op, MVT::f32, Legal);
// only div/rem/sqrt are legal for f64
@@ -1039,6 +1040,23 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setOperationAction(Op, MVT::bf16, Promote);
AddPromotedToType(Op, MVT::bf16, MVT::f32);
}
+ // FTANH support:
+ // - f32 (sm_75+, PTX 7.0+)
+ // - f16/f16x2 (sm_75+, PTX 7.0+)
+ // - bf16/bf16x2 (sm_90+, PTX 7.8+)
+ // When f16/bf16 types aren't supported, they are promoted/expanded to f32.
+ if (STI.getSmVersion() >= 75 && STI.getPTXVersion() >= 70)
+ setOperationAction(ISD::FTANH, MVT::f32, Legal);
+ setOperationAction(ISD::FTANH, MVT::v2f32, Expand);
+ for (MVT VT : {MVT::f16, MVT::v2f16})
+ setFP16OperationAction(ISD::FTANH, VT, Legal,
+ VT.isVector() ? Expand : Promote);
+ for (MVT VT : {MVT::bf16, MVT::v2bf16})
+ setBF16OperationAction(ISD::FTANH, VT, Legal,
+ VT.isVector() ? Expand : Promote);
+ if (getOperationAction(ISD::FTANH, MVT::bf16) == Promote)
+ AddPromotedToType(ISD::FTANH, MVT::bf16, MVT::f32);
+
setOperationAction(ISD::FREM, {MVT::f32, MVT::f64}, Custom);
setOperationAction(ISD::FABS, {MVT::f32, MVT::f64}, Legal);
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 65dd146926f8b..3a8c3e5f0a043 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1329,10 +1329,22 @@ 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>]>;
+class FTANHInst<RegTyInfo t> :
+ BasicNVPTXInst<(outs t.RC:$dst), (ins t.RC:$src),
+ "tanh.approx." # t.PtxType,
+ [(set t.Ty:$dst, (UnaryOpAllowsApproxFn<ftanh> t.Ty:$src))]>;
+
+def TANH_APPROX_f32 : FTANHInst<F32RT>,
+ Requires<[hasPTX<70>, hasSM<75>]>;
+
+let Predicates = [useFP16Math, hasPTX<70>, hasSM<75>] in {
+ def TANH_APPROX_f16 : FTANHInst<F16RT>;
+ def TANH_APPROX_f16x2 : FTANHInst<F16X2RT>;
+}
+let Predicates = [hasPTX<78>, hasSM<90>] in {
+ def TANH_APPROX_bf16 : FTANHInst<BF16RT>;
+ def TANH_APPROX_bf16x2 : FTANHInst<BF16X2RT>;
+}
//-----------------------------------
// Bitwise operations
diff --git a/llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp b/llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp
index bf3c78d3606bf..f418881ec6e39 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp
@@ -222,6 +222,7 @@ bool NVPTXSubtarget::hasNativeBF16Support(int Opcode) const {
case ISD::FRINT:
case ISD::FROUNDEVEN:
case ISD::FTRUNC:
+ case ISD::FTANH:
return getSmVersion() >= 90 && getPTXVersion() >= 78;
// Several BF16 instructions are available on sm_80 only.
case ISD::FMINNUM:
diff --git a/llvm/test/CodeGen/NVPTX/tanhf.ll b/llvm/test/CodeGen/NVPTX/tanhf.ll
index 94ed44c7361ca..d7f6e02b349e4 100644
--- a/llvm/test/CodeGen/NVPTX/tanhf.ll
+++ b/llvm/test/CodeGen/NVPTX/tanhf.ll
@@ -1,16 +1,18 @@
; 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: llc < %s -mcpu=sm_75 -mattr=+ptx70 | FileCheck %s --check-prefixes=CHECK,CHECK-SM75
+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx78 | FileCheck %s --check-prefixes=CHECK,CHECK-SM90
; RUN: %if ptxas-sm_75 && ptxas-isa-7.0 %{ llc < %s -mcpu=sm_75 -mattr=+ptx70 | %ptxas-verify -arch=sm_75 %}
+; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mcpu=sm_90 -mattr=+ptx78 | %ptxas-verify -arch=sm_90 %}
target triple = "nvptx64-nvidia-cuda"
-define float @test1(float %in) local_unnamed_addr {
-; CHECK-LABEL: test1(
+define float @tanh_f32(float %in) local_unnamed_addr {
+; CHECK-LABEL: tanh_f32(
; CHECK: {
; CHECK-NEXT: .reg .b32 %r<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [test1_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [tanh_f32_param_0];
; CHECK-NEXT: tanh.approx.f32 %r2, %r1;
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
@@ -18,23 +20,113 @@ define float @test1(float %in) local_unnamed_addr {
ret float %call
}
-define half @test2(half %in) local_unnamed_addr {
-; CHECK-LABEL: test2(
+define half @tanh_f16(half %in) local_unnamed_addr {
+; CHECK-LABEL: tanh_f16(
; 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: ld.param.b16 %rs1, [tanh_f16_param_0];
+; CHECK-NEXT: tanh.approx.f16 %rs2, %rs1;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-NEXT: ret;
%call = call afn half @llvm.tanh.f16(half %in)
ret half %call
}
+define <2 x half> @tanh_f16x2(<2 x half> %in) local_unnamed_addr {
+; CHECK-LABEL: tanh_f16x2(
+; CHECK: {
+; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r1, [tanh_f16x2_param_0];
+; CHECK-NEXT: tanh.approx.f16x2 %r2, %r1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ret;
+ %call = call afn <2 x half> @llvm.tanh.v2f16(<2 x half> %in)
+ ret <2 x half> %call
+}
+
+define bfloat @tanh_bf16(bfloat %in) local_unnamed_addr {
+; CHECK-SM75-LABEL: tanh_bf16(
+; CHECK-SM75: {
+; CHECK-SM75-NEXT: .reg .pred %p<2>;
+; CHECK-SM75-NEXT: .reg .b32 %r<10>;
+; CHECK-SM75-EMPTY:
+; CHECK-SM75-NEXT: // %bb.0:
+; CHECK-SM75-NEXT: ld.param.b16 %r1, [tanh_bf16_param_0];
+; CHECK-SM75-NEXT: shl.b32 %r2, %r1, 16;
+; CHECK-SM75-NEXT: tanh.approx.f32 %r3, %r2;
+; CHECK-SM75-NEXT: bfe.u32 %r4, %r3, 16, 1;
+; CHECK-SM75-NEXT: add.s32 %r5, %r4, %r3;
+; CHECK-SM75-NEXT: add.s32 %r6, %r5, 32767;
+; CHECK-SM75-NEXT: setp.nan.f32 %p1, %r3, %r3;
+; CHECK-SM75-NEXT: or.b32 %r7, %r3, 4194304;
+; CHECK-SM75-NEXT: selp.b32 %r8, %r7, %r6, %p1;
+; CHECK-SM75-NEXT: shr.u32 %r9, %r8, 16;
+; CHECK-SM75-NEXT: st.param.b16 [func_retval0], %r9;
+; CHECK-SM75-NEXT: ret;
+;
+; CHECK-SM90-LABEL: tanh_bf16(
+; CHECK-SM90: {
+; CHECK-SM90-NEXT: .reg .b16 %rs<3>;
+; CHECK-SM90-EMPTY:
+; CHECK-SM90-NEXT: // %bb.0:
+; CHECK-SM90-NEXT: ld.param.b16 %rs1, [tanh_bf16_param_0];
+; CHECK-SM90-NEXT: tanh.approx.bf16 %rs2, %rs1;
+; CHECK-SM90-NEXT: st.param.b16 [func_retval0], %rs2;
+; CHECK-SM90-NEXT: ret;
+ %call = call afn bfloat @llvm.tanh.bf16(bfloat %in)
+ ret bfloat %call
+}
+
+define <2 x bfloat> @tanh_bf16x2(<2 x bfloat> %in) local_unnamed_addr {
+; CHECK-SM75-LABEL: tanh_bf16x2(
+; CHECK-SM75: {
+; CHECK-SM75-NEXT: .reg .pred %p<3>;
+; CHECK-SM75-NEXT: .reg .b16 %rs<3>;
+; CHECK-SM75-NEXT: .reg .b32 %r<18>;
+; CHECK-SM75-EMPTY:
+; CHECK-SM75-NEXT: // %bb.0:
+; CHECK-SM75-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [tanh_bf16x2_param_0];
+; CHECK-SM75-NEXT: cvt.u32.u16 %r1, %rs2;
+; CHECK-SM75-NEXT: shl.b32 %r2, %r1, 16;
+; CHECK-SM75-NEXT: tanh.approx.f32 %r3, %r2;
+; CHECK-SM75-NEXT: bfe.u32 %r4, %r3, 16, 1;
+; CHECK-SM75-NEXT: add.s32 %r5, %r4, %r3;
+; CHECK-SM75-NEXT: add.s32 %r6, %r5, 32767;
+; CHECK-SM75-NEXT: setp.nan.f32 %p1, %r3, %r3;
+; CHECK-SM75-NEXT: or.b32 %r7, %r3, 4194304;
+; CHECK-SM75-NEXT: selp.b32 %r8, %r7, %r6, %p1;
+; CHECK-SM75-NEXT: cvt.u32.u16 %r9, %rs1;
+; CHECK-SM75-NEXT: shl.b32 %r10, %r9, 16;
+; CHECK-SM75-NEXT: tanh.approx.f32 %r11, %r10;
+; CHECK-SM75-NEXT: bfe.u32 %r12, %r11, 16, 1;
+; CHECK-SM75-NEXT: add.s32 %r13, %r12, %r11;
+; CHECK-SM75-NEXT: add.s32 %r14, %r13, 32767;
+; CHECK-SM75-NEXT: setp.nan.f32 %p2, %r11, %r11;
+; CHECK-SM75-NEXT: or.b32 %r15, %r11, 4194304;
+; CHECK-SM75-NEXT: selp.b32 %r16, %r15, %r14, %p2;
+; CHECK-SM75-NEXT: prmt.b32 %r17, %r16, %r8, 0x7632U;
+; CHECK-SM75-NEXT: st.param.b32 [func_retval0], %r17;
+; CHECK-SM75-NEXT: ret;
+;
+; CHECK-SM90-LABEL: tanh_bf16x2(
+; CHECK-SM90: {
+; CHECK-SM90-NEXT: .reg .b32 %r<3>;
+; CHECK-SM90-EMPTY:
+; CHECK-SM90-NEXT: // %bb.0:
+; CHECK-SM90-NEXT: ld.param.b32 %r1, [tanh_bf16x2_param_0];
+; CHECK-SM90-NEXT: tanh.approx.bf16x2 %r2, %r1;
+; CHECK-SM90-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-SM90-NEXT: ret;
+ %call = call afn <2 x bfloat> @llvm.tanh.v2bf16(<2 x bfloat> %in)
+ ret <2 x bfloat> %call
+}
+
declare float @llvm.tanh.f32(float)
declare half @llvm.tanh.f16(half)
-
+declare <2 x half> @llvm.tanh.v2f16(<2 x half>)
+declare bfloat @llvm.tanh.bf16(bfloat)
+declare <2 x bfloat> @llvm.tanh.v2bf16(<2 x bfloat>)
More information about the llvm-commits
mailing list