[llvm] [NVPTX] Add native `tanh.approx` support for f16/f16x2/bf16/bf16x2 (PR #203257)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Jun 11 06:00:54 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Varad Rahul Kamthe (varadk27)
<details>
<summary>Changes</summary>
Adds NVPTX backend support for the native PTX `tanh.approx` instructions on half-precision and bfloat types:
- `tanh.approx.f16` and `tanh.approx.f16x2` (PTX 7.0+, sm_75+)
- `tanh.approx.bf16` and `tanh.approx.bf16x2` (PTX 7.8+, sm_90+)
Adds a `FTANHInst` TableGen class with the new patterns in NVPTXInstrInfo.td and splits `ISD::FTANH` out of the unconditional `f16/bf16 -> f32` promotion loop in NVPTXISelLowering.cpp, marking it Legal when the target supports it (scalars promote, vectors expand on older targets). Also guards `tanh.approx.f32` behind sm_75 and adds the missing `AddPromotedToType` for bf16.
PTX Spec Reference: https://docs.nvidia.com/cuda/parallel-thread-execution/#half-precision-floating-point-instructions-tanh
---
Full diff: https://github.com/llvm/llvm-project/pull/203257.diff
4 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+21-1)
- (modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+17-4)
- (modified) llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp (+1)
- (modified) llvm/test/CodeGen/NVPTX/tanhf.ll (+104-12)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 17d9f857312d6..83415be870d0d 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -561,6 +561,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;
}
@@ -1027,7 +1028,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
@@ -1040,6 +1041,25 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
}
setOperationAction(ISD::FREM, {MVT::f32, MVT::f64}, Custom);
+ // 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);
+
+ // Scalar f16/bf16: promote to f32 when not natively supported.
+ setFP16OperationAction(ISD::FTANH, MVT::f16, Legal, Promote);
+ setBF16OperationAction(ISD::FTANH, MVT::bf16, Legal, Promote);
+ if (getOperationAction(ISD::FTANH, MVT::bf16) == Promote)
+ AddPromotedToType(ISD::FTANH, MVT::bf16, MVT::f32);
+
+ // Vector v2f16/v2bf16: expand when not natively supported.
+ setFP16OperationAction(ISD::FTANH, MVT::v2f16, Legal, Expand);
+ setBF16OperationAction(ISD::FTANH, MVT::v2bf16, Legal, Expand);
+
setOperationAction(ISD::FABS, {MVT::f32, MVT::f64}, Legal);
setOperationAction(ISD::FABS, MVT::v2f32, Expand);
if (STI.getPTXVersion() >= 65) {
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index ff3e5144c6fbe..4b2541773c2fa 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -1329,10 +1329,23 @@ 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..3f55604a52631 100644
--- a/llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXSubtarget.cpp
@@ -216,6 +216,7 @@ bool NVPTXSubtarget::hasNativeBF16Support(int Opcode) const {
case ISD::SELECT_CC:
case ISD::SETCC:
case ISD::FEXP2:
+ case ISD::FTANH:
case ISD::FCEIL:
case ISD::FFLOOR:
case ISD::FNEARBYINT:
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>)
``````````
</details>
https://github.com/llvm/llvm-project/pull/203257
More information about the llvm-commits
mailing list