[AMDGPU] Add support for `v_tanh_bf16` on gfx1250 #147425

shiltian · 2025-07-07T23:44:43Z

Co-authored-by: Mekhanoshin, Stanislav [email protected]

shiltian · 2025-07-07T23:45:01Z

[AMDGPU] Add support for v_tanh_bf16 on gfx1250 #147425 👈 (View in Graphite)
main

This stack of pull requests is managed by Graphite. Learn more about stacking.

llvmbot · 2025-07-07T23:45:18Z

@llvm/pr-subscribers-clang
@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-mc

Author: Shilei Tian (shiltian)

Changes

Co-authored-by: Mekhanoshin, Stanislav <[email protected]>

Patch is 71.25 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/147425.diff

24 Files Affected:

(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+4)
(modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+10)
(modified) llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h (+3)
(modified) llvm/lib/Target/AMDGPU/SIInstrInfo.td (+1)
(modified) llvm/lib/Target/AMDGPU/VOP1Instructions.td (+5)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll (+81)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s (+45)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s (+48)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s (+56)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s (+60)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s (+12)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s (+45)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s (+48)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s (+56)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s (+60)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s (+20)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt (+63)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt (+59)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt (+16)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt (+64)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt (+60)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt (+20)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f7a9b65854696..dcbf436c85abf 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -588,6 +588,10 @@ def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;
 def int_amdgcn_ds_append : AMDGPUDSAppendConsumedIntrinsic;
 def int_amdgcn_ds_consume : AMDGPUDSAppendConsumedIntrinsic;
 
+def int_amdgcn_tanh : DefaultAttrsIntrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
+>;
+
 def int_amdgcn_cvt_pk_f16_fp8 : DefaultAttrsIntrinsic<
   [llvm_v2f16_ty], [llvm_i16_ty], [IntrNoMem, IntrSpeculatable]
 >, ClangBuiltin<"__builtin_amdgcn_cvt_pk_f16_fp8">;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 1a1c32fba9d18..1e88fc241a402 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -535,6 +535,12 @@ def FeatureRealTrue16Insts : SubtargetFeature<"real-true16",
   "Use true 16-bit registers"
 >;
 
+def FeatureBF16TransInsts : SubtargetFeature<"bf16-trans-insts",
+  "HasBF16TransInsts",
+  "true",
+  "Has bf16 transcendental instructions"
+>;
+
 def FeatureBF16ConversionInsts : SubtargetFeature<"bf16-cvt-insts",
   "HasBF16ConversionInsts",
   "true",
@@ -1946,6 +1952,7 @@ def FeatureISAVersion12_50 : FeatureSet<
    FeatureDPPSrc1SGPR,
    FeatureBitOp3Insts,
    FeatureTransposeLoadF4F6Insts,
+   FeatureBF16TransInsts,
    FeatureBF16ConversionInsts,
    FeatureCvtPkF16F32Inst,
    FeatureMinimum3Maximum3PKF16,
@@ -2413,6 +2420,9 @@ def UseFakeTrue16Insts : True16PredicateClass<"Subtarget->hasTrue16BitInsts() &&
   // FIXME When we default to RealTrue16 instead of Fake, change the line as follows.
   // AssemblerPredicate<(all_of FeatureTrue16BitInsts, (not FeatureRealTrue16Insts))>;
 
+def HasBF16TransInsts : Predicate<"Subtarget->hasBF16TransInsts()">,
+  AssemblerPredicate<(all_of FeatureBF16TransInsts)>;
+
 def HasBF16ConversionInsts : Predicate<"Subtarget->hasBF16ConversionInsts()">,
   AssemblerPredicate<(all_of FeatureBF16ConversionInsts)>;
 
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h
index 7c24f428d78e4..1e44be8e47201 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h
@@ -59,6 +59,7 @@ class AMDGPUSubtarget {
   bool HasCvtPkF16F32Inst = false;
   bool HasF32ToF16BF16ConversionSRInsts = false;
   bool EnableRealTrue16Insts = false;
+  bool HasBF16TransInsts = false;
   bool HasBF16ConversionInsts = false;
   bool HasMadMixInsts = false;
   bool HasMadMacF32Insts = false;
@@ -202,6 +203,8 @@ class AMDGPUSubtarget {
   // supported and the support for fake True16 instructions is removed.
   bool useRealTrue16Insts() const;
 
+  bool hasBF16TransInsts() const { return HasBF16TransInsts; }
+
   bool hasBF16ConversionInsts() const {
     return HasBF16ConversionInsts;
   }
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index 6d6c2af7ce490..c9d3b85f36da4 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2743,6 +2743,7 @@ def VOP_F16_F16 : VOPProfile<[f16, f16, untyped, untyped]>;
 def VOP_F16_I16 : VOPProfile <[f16, i16, untyped, untyped]>;
 def VOP_I16_F16 : VOPProfile <[i16, f16, untyped, untyped]>;
 def VOP_I16_I16 : VOPProfile <[i16, i16, untyped, untyped]>;
+def VOP_BF16_BF16 : VOPProfile<[bf16, bf16, untyped, untyped]>;
 
 def VOP_F16_F16_F16 : VOPProfile <[f16, f16, f16, untyped]>;
 def VOP_F16_F16_I16 : VOPProfile <[f16, f16, i16, untyped]>;
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index 0dacd9df71305..7159d5ba33e5a 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -522,6 +522,10 @@ defm V_LOG_F16 : VOP1Inst_t16 <"v_log_f16", VOP_F16_F16, AMDGPUlogf16>;
 defm V_EXP_F16 : VOP1Inst_t16 <"v_exp_f16", VOP_F16_F16, AMDGPUexpf16>;
 defm V_SIN_F16 : VOP1Inst_t16 <"v_sin_f16", VOP_F16_F16, AMDGPUsin>;
 defm V_COS_F16 : VOP1Inst_t16 <"v_cos_f16", VOP_F16_F16, AMDGPUcos>;
+
+let SubtargetPredicate = HasBF16TransInsts in {
+defm V_TANH_BF16 : VOP1Inst_t16 <"v_tanh_bf16", VOP_BF16_BF16, int_amdgcn_tanh>;
+}
 } // End TRANS = 1, SchedRW = [WriteTrans32]
 defm V_FREXP_MANT_F16 : VOP1Inst_t16 <"v_frexp_mant_f16", VOP_F16_F16, int_amdgcn_frexp_mant>;
 defm V_FREXP_EXP_I16_F16 : VOP1Inst_t16_with_profiles <"v_frexp_exp_i16_f16",
@@ -1099,6 +1103,7 @@ defm V_CVT_NORM_U16_F16      : VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x064>;
 defm V_CVT_F16_F32           : VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x00a>;
 defm V_CVT_F32_F16           : VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x00b>;
 
+defm V_TANH_BF16             : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
 defm V_CVT_F32_BF16          : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
 defm V_CVT_PK_F16_FP8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
 defm V_CVT_PK_F16_BF8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
new file mode 100644
index 0000000000000..126212ed347e0
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
@@ -0,0 +1,81 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=SDAG-REAL16 %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=SDAG-FAKE16 %s
+; xUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=GISEL-REAL16 %s
+; xUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=GISEL-FAKE16 %s
+
+; FIXME: GlobalISel does not work with bf16
+
+declare bfloat @llvm.amdgcn.tanh.bf16(bfloat) #0
+
+define amdgpu_kernel void @tanh_bf16(ptr addrspace(1) %out, bfloat %src) #1 {
+; SDAG-REAL16-LABEL: tanh_bf16:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b96 s[0:2], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_tanh_bf16_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: tanh_bf16:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b96 s[0:2], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_tanh_bf16_e32 v0, s2
+; SDAG-FAKE16-NEXT:    global_store_b16 v1, v0, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+  %tanh = call bfloat @llvm.amdgcn.tanh.bf16(bfloat %src) #0
+  store bfloat %tanh, ptr addrspace(1) %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @tanh_bf16_constant_4(ptr addrspace(1) %out) #1 {
+; SDAG-REAL16-LABEL: tanh_bf16_constant_4:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_tanh_bf16_e32 v0.l, 4.0
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: tanh_bf16_constant_4:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_tanh_bf16_e32 v0, 4.0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    global_store_b16 v1, v0, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+  %tanh = call bfloat @llvm.amdgcn.tanh.bf16(bfloat 4.0) #0
+  store bfloat %tanh, ptr addrspace(1) %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @tanh_bf16_constant_100(ptr addrspace(1) %out) #1 {
+; SDAG-REAL16-LABEL: tanh_bf16_constant_100:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_tanh_bf16_e32 v0.l, 0x42c8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: tanh_bf16_constant_100:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_tanh_bf16_e32 v0, 0x42c8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    global_store_b16 v1, v0, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+  %tanh = call bfloat @llvm.amdgcn.tanh.bf16(bfloat 100.0) #0
+  store bfloat %tanh, ptr addrspace(1) %out, align 2
+  ret void
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index 7b07c84d56680..d986b2ff6ee13 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -1,6 +1,51 @@
 // NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --version 5
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -show-encoding %s | FileCheck --check-prefix=GFX1250 %s
 
+v_tanh_bf16 v5, v1
+// GFX1250: v_tanh_bf16_e32 v5, v1                  ; encoding: [0x01,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, v127
+// GFX1250: v_tanh_bf16_e32 v5, v127                ; encoding: [0x7f,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, s1
+// GFX1250: v_tanh_bf16_e32 v5, s1                  ; encoding: [0x01,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, s105
+// GFX1250: v_tanh_bf16_e32 v5, s105                ; encoding: [0x69,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_lo
+// GFX1250: v_tanh_bf16_e32 v5, vcc_lo              ; encoding: [0x6a,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_hi
+// GFX1250: v_tanh_bf16_e32 v5, vcc_hi              ; encoding: [0x6b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, ttmp15
+// GFX1250: v_tanh_bf16_e32 v5, ttmp15              ; encoding: [0x7b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, m0
+// GFX1250: v_tanh_bf16_e32 v5, m0                  ; encoding: [0x7d,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_lo
+// GFX1250: v_tanh_bf16_e32 v5, exec_lo             ; encoding: [0x7e,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_hi
+// GFX1250: v_tanh_bf16_e32 v5, exec_hi             ; encoding: [0x7f,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, null
+// GFX1250: v_tanh_bf16_e32 v5, null                ; encoding: [0x7c,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, -1
+// GFX1250: v_tanh_bf16_e32 v5, -1                  ; encoding: [0xc1,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, 0.5
+// GFX1250: v_tanh_bf16_e32 v5, 0.5                 ; encoding: [0xf0,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, src_scc
+// GFX1250: v_tanh_bf16_e32 v5, src_scc             ; encoding: [0xfd,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v127, 0x8000
+// GFX1250: v_tanh_bf16_e32 v127, 0x8000            ; encoding: [0xff,0x94,0xfe,0x7e,0x00,0x80,0x00,0x00]
+
 v_cvt_f32_bf16 v5, v1
 // GFX1250: v_cvt_f32_bf16_e32 v5, v1               ; encoding: [0x01,0xe5,0x0a,0x7e]
 
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
index 30c62c957874d..0b1d0a3342b2f 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
@@ -1,6 +1,54 @@
 // NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --version 5
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | FileCheck --check-prefix=GFX1250 %s
 
+v_tanh_bf16 v5, v1
+// GFX1250: v_tanh_bf16_e32 v5, v1                  ; encoding: [0x01,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, v127
+// GFX1250: v_tanh_bf16_e32 v5, v127                ; encoding: [0x7f,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, s1
+// GFX1250: v_tanh_bf16_e32 v5, s1                  ; encoding: [0x01,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, s105
+// GFX1250: v_tanh_bf16_e32 v5, s105                ; encoding: [0x69,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_lo
+// GFX1250: v_tanh_bf16_e32 v5, vcc_lo              ; encoding: [0x6a,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_hi
+// GFX1250: v_tanh_bf16_e32 v5, vcc_hi              ; encoding: [0x6b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, ttmp15
+// GFX1250: v_tanh_bf16_e32 v5, ttmp15              ; encoding: [0x7b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, m0
+// GFX1250: v_tanh_bf16_e32 v5, m0                  ; encoding: [0x7d,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_lo
+// GFX1250: v_tanh_bf16_e32 v5, exec_lo             ; encoding: [0x7e,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_hi
+// GFX1250: v_tanh_bf16_e32 v5, exec_hi             ; encoding: [0x7f,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, null
+// GFX1250: v_tanh_bf16_e32 v5, null                ; encoding: [0x7c,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, -1
+// GFX1250: v_tanh_bf16_e32 v5, -1                  ; encoding: [0xc1,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, 0.5
+// GFX1250: v_tanh_bf16_e32 v5, 0.5                 ; encoding: [0xf0,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, src_scc
+// GFX1250: v_tanh_bf16_e32 v5, src_scc             ; encoding: [0xfd,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v127, 0x8000
+// GFX1250: v_tanh_bf16_e32 v127, 0x8000            ; encoding: [0xff,0x94,0xfe,0x7e,0x00,0x80,0x00,0x00]
+
+v_tanh_bf16 v5.h, v1.h
+// GFX1250: v_tanh_bf16_e32 v5.h, v1.h              ; encoding: [0x81,0x95,0x0a,0x7f]
+
 v_cvt_f32_bf16 v5, v1
 // GFX1250: v_cvt_f32_bf16_e32 v5, v1               ; encoding: [0x01,0xe5,0x0a,0x7e]
 
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
index e53812bb3fd04..4e5754f3961c1 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
@@ -2,6 +2,62 @@
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -show-encoding %s | FileCheck --check-prefixes=GFX1250 %s
 // RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -mattr=-real-true16 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX12-ERR --implicit-check-not=error: --strict-whitespace %s
 
+v_tanh_bf16 v5, v1 quad_perm:[3,2,1,0]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x1b,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 quad_perm:[0,1,2,3]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_half_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x41,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x01,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x0f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x11,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x1f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_ror:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x21,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_ror:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x2f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x50,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x5f,0x01,0x01]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x60,0x09,0x13]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 bound_ctrl:0 fi:1
+// GFX1250: v_tanh_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding: [0xfa,0x94,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
 v_cvt_f32_bf16 v5, v1 quad_perm:[3,2,1,0]
 // GFX1250: v_cvt_f32_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xe4,0x0a,0x7e,0x01,0x1b,0x00,0xff]
 // GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
index bd767d14fab5f..a6787254ae60f 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
@@ -2,6 +2,66 @@
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | FileCheck --check-prefixes=GFX1250 %s
 // RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -mattr=+real-true16 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX12-ERR --implicit-check-not=error: --strict-whitespace %s
 
+v_tanh_bf16 v5, v1 quad_perm:[3,2,1,0]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x1b,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 quad_perm:[0,1,2,3]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_half_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x41,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x01,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x0f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x11,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xf...
[truncated]

llvmbot · 2025-07-07T23:45:19Z

@llvm/pr-subscribers-llvm-ir

Author: Shilei Tian (shiltian)

Changes

Co-authored-by: Mekhanoshin, Stanislav <[email protected]>

Patch is 71.25 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/147425.diff

24 Files Affected:

(modified) llvm/include/llvm/IR/IntrinsicsAMDGPU.td (+4)
(modified) llvm/lib/Target/AMDGPU/AMDGPU.td (+10)
(modified) llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h (+3)
(modified) llvm/lib/Target/AMDGPU/SIInstrInfo.td (+1)
(modified) llvm/lib/Target/AMDGPU/VOP1Instructions.td (+5)
(added) llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll (+81)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s (+45)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s (+48)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s (+56)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s (+60)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s (+12)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s (+45)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s (+48)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s (+56)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s (+60)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s (+16)
(modified) llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s (+20)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt (+63)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt (+59)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt (+16)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt (+64)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt (+60)
(modified) llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt (+20)

diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f7a9b65854696..dcbf436c85abf 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -588,6 +588,10 @@ def int_amdgcn_ds_ordered_swap : AMDGPUDSOrderedIntrinsic;
 def int_amdgcn_ds_append : AMDGPUDSAppendConsumedIntrinsic;
 def int_amdgcn_ds_consume : AMDGPUDSAppendConsumedIntrinsic;
 
+def int_amdgcn_tanh : DefaultAttrsIntrinsic<
+  [llvm_anyfloat_ty], [LLVMMatchType<0>], [IntrNoMem, IntrSpeculatable]
+>;
+
 def int_amdgcn_cvt_pk_f16_fp8 : DefaultAttrsIntrinsic<
   [llvm_v2f16_ty], [llvm_i16_ty], [IntrNoMem, IntrSpeculatable]
 >, ClangBuiltin<"__builtin_amdgcn_cvt_pk_f16_fp8">;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPU.td b/llvm/lib/Target/AMDGPU/AMDGPU.td
index 1a1c32fba9d18..1e88fc241a402 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPU.td
+++ b/llvm/lib/Target/AMDGPU/AMDGPU.td
@@ -535,6 +535,12 @@ def FeatureRealTrue16Insts : SubtargetFeature<"real-true16",
   "Use true 16-bit registers"
 >;
 
+def FeatureBF16TransInsts : SubtargetFeature<"bf16-trans-insts",
+  "HasBF16TransInsts",
+  "true",
+  "Has bf16 transcendental instructions"
+>;
+
 def FeatureBF16ConversionInsts : SubtargetFeature<"bf16-cvt-insts",
   "HasBF16ConversionInsts",
   "true",
@@ -1946,6 +1952,7 @@ def FeatureISAVersion12_50 : FeatureSet<
    FeatureDPPSrc1SGPR,
    FeatureBitOp3Insts,
    FeatureTransposeLoadF4F6Insts,
+   FeatureBF16TransInsts,
    FeatureBF16ConversionInsts,
    FeatureCvtPkF16F32Inst,
    FeatureMinimum3Maximum3PKF16,
@@ -2413,6 +2420,9 @@ def UseFakeTrue16Insts : True16PredicateClass<"Subtarget->hasTrue16BitInsts() &&
   // FIXME When we default to RealTrue16 instead of Fake, change the line as follows.
   // AssemblerPredicate<(all_of FeatureTrue16BitInsts, (not FeatureRealTrue16Insts))>;
 
+def HasBF16TransInsts : Predicate<"Subtarget->hasBF16TransInsts()">,
+  AssemblerPredicate<(all_of FeatureBF16TransInsts)>;
+
 def HasBF16ConversionInsts : Predicate<"Subtarget->hasBF16ConversionInsts()">,
   AssemblerPredicate<(all_of FeatureBF16ConversionInsts)>;
 
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h
index 7c24f428d78e4..1e44be8e47201 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h
@@ -59,6 +59,7 @@ class AMDGPUSubtarget {
   bool HasCvtPkF16F32Inst = false;
   bool HasF32ToF16BF16ConversionSRInsts = false;
   bool EnableRealTrue16Insts = false;
+  bool HasBF16TransInsts = false;
   bool HasBF16ConversionInsts = false;
   bool HasMadMixInsts = false;
   bool HasMadMacF32Insts = false;
@@ -202,6 +203,8 @@ class AMDGPUSubtarget {
   // supported and the support for fake True16 instructions is removed.
   bool useRealTrue16Insts() const;
 
+  bool hasBF16TransInsts() const { return HasBF16TransInsts; }
+
   bool hasBF16ConversionInsts() const {
     return HasBF16ConversionInsts;
   }
diff --git a/llvm/lib/Target/AMDGPU/SIInstrInfo.td b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
index 6d6c2af7ce490..c9d3b85f36da4 100644
--- a/llvm/lib/Target/AMDGPU/SIInstrInfo.td
+++ b/llvm/lib/Target/AMDGPU/SIInstrInfo.td
@@ -2743,6 +2743,7 @@ def VOP_F16_F16 : VOPProfile<[f16, f16, untyped, untyped]>;
 def VOP_F16_I16 : VOPProfile <[f16, i16, untyped, untyped]>;
 def VOP_I16_F16 : VOPProfile <[i16, f16, untyped, untyped]>;
 def VOP_I16_I16 : VOPProfile <[i16, i16, untyped, untyped]>;
+def VOP_BF16_BF16 : VOPProfile<[bf16, bf16, untyped, untyped]>;
 
 def VOP_F16_F16_F16 : VOPProfile <[f16, f16, f16, untyped]>;
 def VOP_F16_F16_I16 : VOPProfile <[f16, f16, i16, untyped]>;
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index 0dacd9df71305..7159d5ba33e5a 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -522,6 +522,10 @@ defm V_LOG_F16 : VOP1Inst_t16 <"v_log_f16", VOP_F16_F16, AMDGPUlogf16>;
 defm V_EXP_F16 : VOP1Inst_t16 <"v_exp_f16", VOP_F16_F16, AMDGPUexpf16>;
 defm V_SIN_F16 : VOP1Inst_t16 <"v_sin_f16", VOP_F16_F16, AMDGPUsin>;
 defm V_COS_F16 : VOP1Inst_t16 <"v_cos_f16", VOP_F16_F16, AMDGPUcos>;
+
+let SubtargetPredicate = HasBF16TransInsts in {
+defm V_TANH_BF16 : VOP1Inst_t16 <"v_tanh_bf16", VOP_BF16_BF16, int_amdgcn_tanh>;
+}
 } // End TRANS = 1, SchedRW = [WriteTrans32]
 defm V_FREXP_MANT_F16 : VOP1Inst_t16 <"v_frexp_mant_f16", VOP_F16_F16, int_amdgcn_frexp_mant>;
 defm V_FREXP_EXP_I16_F16 : VOP1Inst_t16_with_profiles <"v_frexp_exp_i16_f16",
@@ -1099,6 +1103,7 @@ defm V_CVT_NORM_U16_F16      : VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x064>;
 defm V_CVT_F16_F32           : VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x00a>;
 defm V_CVT_F32_F16           : VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x00b>;
 
+defm V_TANH_BF16             : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x04a>;
 defm V_CVT_F32_BF16          : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x072, "v_cvt_f32_bf16", "V_CVT_F32_BF16_gfx1250">;
 defm V_CVT_PK_F16_FP8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x075>;
 defm V_CVT_PK_F16_BF8        : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x076>;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
new file mode 100644
index 0000000000000..126212ed347e0
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll
@@ -0,0 +1,81 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=SDAG-REAL16 %s
+; RUN: llc -global-isel=0 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=SDAG-FAKE16 %s
+; xUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=GISEL-REAL16 %s
+; xUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=GISEL-FAKE16 %s
+
+; FIXME: GlobalISel does not work with bf16
+
+declare bfloat @llvm.amdgcn.tanh.bf16(bfloat) #0
+
+define amdgpu_kernel void @tanh_bf16(ptr addrspace(1) %out, bfloat %src) #1 {
+; SDAG-REAL16-LABEL: tanh_bf16:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b96 s[0:2], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    v_tanh_bf16_e32 v0.l, s2
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: tanh_bf16:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b96 s[0:2], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    v_tanh_bf16_e32 v0, s2
+; SDAG-FAKE16-NEXT:    global_store_b16 v1, v0, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+  %tanh = call bfloat @llvm.amdgcn.tanh.bf16(bfloat %src) #0
+  store bfloat %tanh, ptr addrspace(1) %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @tanh_bf16_constant_4(ptr addrspace(1) %out) #1 {
+; SDAG-REAL16-LABEL: tanh_bf16_constant_4:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_tanh_bf16_e32 v0.l, 4.0
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: tanh_bf16_constant_4:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_tanh_bf16_e32 v0, 4.0
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    global_store_b16 v1, v0, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+  %tanh = call bfloat @llvm.amdgcn.tanh.bf16(bfloat 4.0) #0
+  store bfloat %tanh, ptr addrspace(1) %out, align 2
+  ret void
+}
+
+define amdgpu_kernel void @tanh_bf16_constant_100(ptr addrspace(1) %out) #1 {
+; SDAG-REAL16-LABEL: tanh_bf16_constant_100:
+; SDAG-REAL16:       ; %bb.0:
+; SDAG-REAL16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-REAL16-NEXT:    v_tanh_bf16_e32 v0.l, 0x42c8
+; SDAG-REAL16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-REAL16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-REAL16-NEXT:    flat_store_b16 v1, v0, s[0:1]
+; SDAG-REAL16-NEXT:    s_endpgm
+;
+; SDAG-FAKE16-LABEL: tanh_bf16_constant_100:
+; SDAG-FAKE16:       ; %bb.0:
+; SDAG-FAKE16-NEXT:    s_load_b64 s[0:1], s[4:5], 0x0
+; SDAG-FAKE16-NEXT:    v_tanh_bf16_e32 v0, 0x42c8
+; SDAG-FAKE16-NEXT:    v_mov_b32_e32 v1, 0
+; SDAG-FAKE16-NEXT:    s_wait_kmcnt 0x0
+; SDAG-FAKE16-NEXT:    global_store_b16 v1, v0, s[0:1]
+; SDAG-FAKE16-NEXT:    s_endpgm
+  %tanh = call bfloat @llvm.amdgcn.tanh.bf16(bfloat 100.0) #0
+  store bfloat %tanh, ptr addrspace(1) %out, align 2
+  ret void
+}
+
+attributes #0 = { nounwind readnone }
+attributes #1 = { nounwind }
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index 7b07c84d56680..d986b2ff6ee13 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -1,6 +1,51 @@
 // NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --version 5
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -show-encoding %s | FileCheck --check-prefix=GFX1250 %s
 
+v_tanh_bf16 v5, v1
+// GFX1250: v_tanh_bf16_e32 v5, v1                  ; encoding: [0x01,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, v127
+// GFX1250: v_tanh_bf16_e32 v5, v127                ; encoding: [0x7f,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, s1
+// GFX1250: v_tanh_bf16_e32 v5, s1                  ; encoding: [0x01,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, s105
+// GFX1250: v_tanh_bf16_e32 v5, s105                ; encoding: [0x69,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_lo
+// GFX1250: v_tanh_bf16_e32 v5, vcc_lo              ; encoding: [0x6a,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_hi
+// GFX1250: v_tanh_bf16_e32 v5, vcc_hi              ; encoding: [0x6b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, ttmp15
+// GFX1250: v_tanh_bf16_e32 v5, ttmp15              ; encoding: [0x7b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, m0
+// GFX1250: v_tanh_bf16_e32 v5, m0                  ; encoding: [0x7d,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_lo
+// GFX1250: v_tanh_bf16_e32 v5, exec_lo             ; encoding: [0x7e,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_hi
+// GFX1250: v_tanh_bf16_e32 v5, exec_hi             ; encoding: [0x7f,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, null
+// GFX1250: v_tanh_bf16_e32 v5, null                ; encoding: [0x7c,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, -1
+// GFX1250: v_tanh_bf16_e32 v5, -1                  ; encoding: [0xc1,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, 0.5
+// GFX1250: v_tanh_bf16_e32 v5, 0.5                 ; encoding: [0xf0,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, src_scc
+// GFX1250: v_tanh_bf16_e32 v5, src_scc             ; encoding: [0xfd,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v127, 0x8000
+// GFX1250: v_tanh_bf16_e32 v127, 0x8000            ; encoding: [0xff,0x94,0xfe,0x7e,0x00,0x80,0x00,0x00]
+
 v_cvt_f32_bf16 v5, v1
 // GFX1250: v_cvt_f32_bf16_e32 v5, v1               ; encoding: [0x01,0xe5,0x0a,0x7e]
 
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
index 30c62c957874d..0b1d0a3342b2f 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
@@ -1,6 +1,54 @@
 // NOTE: Assertions have been autogenerated by utils/update_mc_test_checks.py UTC_ARGS: --version 5
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | FileCheck --check-prefix=GFX1250 %s
 
+v_tanh_bf16 v5, v1
+// GFX1250: v_tanh_bf16_e32 v5, v1                  ; encoding: [0x01,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, v127
+// GFX1250: v_tanh_bf16_e32 v5, v127                ; encoding: [0x7f,0x95,0x0a,0x7e]
+
+v_tanh_bf16 v5, s1
+// GFX1250: v_tanh_bf16_e32 v5, s1                  ; encoding: [0x01,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, s105
+// GFX1250: v_tanh_bf16_e32 v5, s105                ; encoding: [0x69,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_lo
+// GFX1250: v_tanh_bf16_e32 v5, vcc_lo              ; encoding: [0x6a,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, vcc_hi
+// GFX1250: v_tanh_bf16_e32 v5, vcc_hi              ; encoding: [0x6b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, ttmp15
+// GFX1250: v_tanh_bf16_e32 v5, ttmp15              ; encoding: [0x7b,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, m0
+// GFX1250: v_tanh_bf16_e32 v5, m0                  ; encoding: [0x7d,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_lo
+// GFX1250: v_tanh_bf16_e32 v5, exec_lo             ; encoding: [0x7e,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, exec_hi
+// GFX1250: v_tanh_bf16_e32 v5, exec_hi             ; encoding: [0x7f,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, null
+// GFX1250: v_tanh_bf16_e32 v5, null                ; encoding: [0x7c,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, -1
+// GFX1250: v_tanh_bf16_e32 v5, -1                  ; encoding: [0xc1,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, 0.5
+// GFX1250: v_tanh_bf16_e32 v5, 0.5                 ; encoding: [0xf0,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v5, src_scc
+// GFX1250: v_tanh_bf16_e32 v5, src_scc             ; encoding: [0xfd,0x94,0x0a,0x7e]
+
+v_tanh_bf16 v127, 0x8000
+// GFX1250: v_tanh_bf16_e32 v127, 0x8000            ; encoding: [0xff,0x94,0xfe,0x7e,0x00,0x80,0x00,0x00]
+
+v_tanh_bf16 v5.h, v1.h
+// GFX1250: v_tanh_bf16_e32 v5.h, v1.h              ; encoding: [0x81,0x95,0x0a,0x7f]
+
 v_cvt_f32_bf16 v5, v1
 // GFX1250: v_cvt_f32_bf16_e32 v5, v1               ; encoding: [0x01,0xe5,0x0a,0x7e]
 
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
index e53812bb3fd04..4e5754f3961c1 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
@@ -2,6 +2,62 @@
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -show-encoding %s | FileCheck --check-prefixes=GFX1250 %s
 // RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -mattr=-real-true16 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX12-ERR --implicit-check-not=error: --strict-whitespace %s
 
+v_tanh_bf16 v5, v1 quad_perm:[3,2,1,0]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x1b,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 quad_perm:[0,1,2,3]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_half_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x41,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x01,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x0f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x11,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x1f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_ror:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_ror:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x21,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_ror:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_ror:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x2f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_share:0 row_mask:0xf bank_mask:0xf
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_share:0 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x50,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_share:15 row_mask:0x0 bank_mask:0x1 ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x5f,0x01,0x01]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 fi:0
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_xmask:0 row_mask:0x1 bank_mask:0x3 bound_ctrl:1 ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x60,0x09,0x13]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 bound_ctrl:0 fi:1
+// GFX1250: v_tanh_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3 bank_mask:0x0 fi:1 ; encoding: [0xfa,0x94,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
 v_cvt_f32_bf16 v5, v1 quad_perm:[3,2,1,0]
 // GFX1250: v_cvt_f32_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0xe4,0x0a,0x7e,0x01,0x1b,0x00,0xff]
 // GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
index bd767d14fab5f..a6787254ae60f 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
@@ -2,6 +2,66 @@
 // RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -show-encoding %s | FileCheck --check-prefixes=GFX1250 %s
 // RUN: not llvm-mc -triple=amdgcn -mcpu=gfx1200 -mattr=+real-true16 -show-encoding %s 2>&1 | FileCheck --check-prefix=GFX12-ERR --implicit-check-not=error: --strict-whitespace %s
 
+v_tanh_bf16 v5, v1 quad_perm:[3,2,1,0]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x1b,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 quad_perm:[0,1,2,3]
+// GFX1250: v_tanh_bf16_dpp v5, v1 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x40,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_half_mirror
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_half_mirror row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x41,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x01,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shl:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shl:15 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x0f,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:1
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:1 row_mask:0xf bank_mask:0xf ; encoding: [0xfa,0x94,0x0a,0x7e,0x01,0x11,0x01,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_tanh_bf16 v5, v1 row_shr:15
+// GFX1250: v_tanh_bf16_dpp v5, v1 row_shr:15 row_mask:0xf bank_mask:0xf ; encoding: [0xf...
[truncated]

arsenm · 2025-07-08T01:05:01Z

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

    case Intrinsic::amdgcn_frexp_mant:
    case Intrinsic::amdgcn_fdot2:
    case Intrinsic::amdgcn_trig_preop:
+    case Intrinsic::amdgcn_tanh:


This isn't tested

It was but not by this instruction, since this intrinsic is used for those different types. I'll remove them then?

I don't follow, this shouldn't be hard to test? e.g. fcanonicalize-elimination.ll, you should just need a canonicalize call after these which should fold out

We can't do it at this moment because the support for v_cvt_pk_bf16_f32 on gfx1250 has not been upstreamed. Without it llvm.canonicalize.bf16 doesn't work.

I'll add the test to downstream though.

llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

shiltian · 2025-07-08T05:11:37Z

For some reason there is a crash in llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll.

LLVM ERROR: Cannot select: t28: ch = store<(store (s16) into %ir.out.load, addrspace 1)> t0, t27, t30, undef:i64
  t27: i16 = bitcast t21
    t21: bf16 = llvm.amdgcn.tanh TargetConstant:i64<3419>, t19
      t19: bf16 = bitcast t18
        t18: i16 = truncate t17
          t17: i32,ch = load<(dereferenceable invariant load (s32) from %ir.src.kernarg.offset.align.down, align 8, addrspace 4)> t0, t41, undef:i64
            t41: i64 = bitcast t40
              t40: v2i32 = BUILD_VECTOR t39, t37
                t39: i32 = or t36, Constant:i32<8>
                  t36: i32 = extract_vector_elt t33, Constant:i32<0>

                t37: i32 = extract_vector_elt t33, Constant:i32<1>
                  t33: v2i32 = bitcast t14

  t30: i64 = bitcast t29
    t29: v2i32,ch = load<(dereferenceable invariant load (s64) from %ir.out.kernarg.offset1, align 16, addrspace 4)> t0, t14, undef:i64
      t14: i64 = AssertAlign<16> t32
        t32: i64,ch = CopyFromReg t0, Register:i64 %3
In function: tanh_bf16

Not sure why the store to s16 can't be selected. Need to figure it out.

github-actions · 2025-07-11T17:19:23Z

⚠️ undef deprecator found issues in your code. ⚠️

You can test this locally with the following command:

git diff -U0 --pickaxe-regex -S '([^a-zA-Z0-9#_-]undef[^a-zA-Z0-9_-]|UndefValue::get)' 'HEAD~1' HEAD llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp llvm/lib/Target/AMDGPU/AMDGPUSubtarget.h llvm/lib/Target/AMDGPU/SIISelLowering.cpp llvm/lib/TargetParser/TargetParser.cpp

The following files introduce new uses of undef:

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll

Undef is now deprecated and should only be used in the rare cases where no replacement is possible. For example, a load of uninitialized memory yields undef. You should use poison values for placeholders instead.

In tests, avoid using undef and having tests that trigger undefined behavior. If you need an operand with some unimportant value, you can add a new argument to the function and use that instead.

For example, this is considered a bad practice:

define void @fn() {
  ...
  br i1 undef, ...
}

Please use the following instead:

define void @fn(i1 %cond) {
  ...
  br i1 %cond, ...
}

Please refer to the Undefined Behavior Manual for more information.

shiltian · 2025-07-11T18:08:15Z

This PR is messed up at this moment.

llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt

llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt

shiltian · 2025-07-14T18:52:18Z

FWIW, the crash still exists.

rampitec

LGTM modulo the crash fix.

Co-authored-by: Mekhanoshin, Stanislav <[email protected]>

shiltian · 2025-07-14T19:29:48Z

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.tanh.ll

+; xUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=+real-true16 %s -o - | FileCheck -check-prefix=GISEL-REAL16 %s
+; xUN: llc -global-isel=1 -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1250 -mattr=-real-true16 %s -o - | FileCheck -check-prefix=GISEL-FAKE16 %s
+
+; FIXME: t16 doesn't work at the moment because the store of s16 under t16 mode fails to select.


@rampitec This is the reason it crashes. I've marked the t16 test as xUN. This needs to be reenabled back when the new flat store stuff is upstreamed.

shiltian requested review from changpeng and rampitec July 7, 2025 23:44

llvmbot added backend:AMDGPU llvm:mc Machine (object) code llvm:ir labels Jul 7, 2025

shiltian force-pushed the users/shiltian/v_tanh_bf16 branch from 1b01d0f to 0f941d3 Compare July 7, 2025 23:53

llvmbot added clang Clang issues not falling into any other category clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:codegen IR generation bugs: mangling, exceptions, etc. labels Jul 7, 2025

changpeng approved these changes Jul 8, 2025

View reviewed changes

arsenm reviewed Jul 8, 2025

View reviewed changes

shiltian force-pushed the users/shiltian/v_tanh_bf16 branch from 0f941d3 to 21d18bb Compare July 11, 2025 17:14

shiltian force-pushed the users/shiltian/v_tanh_bf16 branch from 21d18bb to 088b396 Compare July 11, 2025 17:56

shiltian force-pushed the users/shiltian/v_tanh_bf16 branch from 088b396 to 2c36f06 Compare July 14, 2025 17:25

rampitec reviewed Jul 14, 2025

View reviewed changes

llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt Outdated Show resolved Hide resolved

llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt Outdated Show resolved Hide resolved

shiltian force-pushed the users/shiltian/v_tanh_bf16 branch from 2c36f06 to f75f85e Compare July 14, 2025 18:51

rampitec approved these changes Jul 14, 2025

View reviewed changes

[AMDGPU] Add support for v_tanh_bf16 on gfx1250

81d3ce5

Co-authored-by: Mekhanoshin, Stanislav <[email protected]>

shiltian force-pushed the users/shiltian/v_tanh_bf16 branch from f75f85e to 81d3ce5 Compare July 14, 2025 19:28

shiltian commented Jul 14, 2025

View reviewed changes

shiltian merged commit d7ec80c into main Jul 14, 2025
8 of 9 checks passed

shiltian deleted the users/shiltian/v_tanh_bf16 branch July 14, 2025 20:30

[AMDGPU] Add support for v_tanh_bf16 on gfx1250 #147425

[AMDGPU] Add support for v_tanh_bf16 on gfx1250 #147425

Uh oh!

Conversation

shiltian commented Jul 7, 2025

Uh oh!

shiltian commented Jul 7, 2025

Uh oh!

llvmbot commented Jul 7, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jul 7, 2025

Uh oh!

arsenm Jul 8, 2025

Choose a reason for hiding this comment

Uh oh!

shiltian Jul 8, 2025

Choose a reason for hiding this comment

Uh oh!

arsenm Jul 8, 2025

Choose a reason for hiding this comment

Uh oh!

shiltian Jul 8, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

shiltian commented Jul 8, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions bot commented Jul 11, 2025

Uh oh!

shiltian commented Jul 11, 2025

Uh oh!

Uh oh!

Uh oh!

shiltian commented Jul 14, 2025

Uh oh!

rampitec left a comment

Choose a reason for hiding this comment

Uh oh!

shiltian Jul 14, 2025

Choose a reason for hiding this comment

Uh oh!

rampitec Jul 14, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

[AMDGPU] Add support for `v_tanh_bf16` on gfx1250 #147425

[AMDGPU] Add support for `v_tanh_bf16` on gfx1250 #147425

llvmbot commented Jul 7, 2025 •

edited

Loading

shiltian commented Jul 8, 2025 •

edited

Loading