https://github.com/shiltian updated
https://github.com/llvm/llvm-project/pull/146302
>From befea46e97c499f3b1ad0e3ac17ecadebc74acc1 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i...@tianshilei.me>
Date: Sun, 29 Jun 2025 23:47:12 -0400
Subject: [PATCH] [AMDGPU] Add support for `v_cvt_f16_fp8` on gfx1250
Co-authored-by: Mekhanoshin, Stanislav <stanislav.mekhanos...@amd.com>
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 +
.../CodeGenOpenCL/builtins-amdgcn-gfx1250.cl | 38 ++++
.../builtins-amdgcn-error-gfx1250-param.cl | 4 +
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 10 +
.../Target/AMDGPU/AMDGPURegisterBankInfo.cpp | 1 +
llvm/lib/Target/AMDGPU/VOP1Instructions.td | 21 ++-
llvm/lib/Target/AMDGPU/VOPInstructions.td | 21 ++-
.../CodeGen/AMDGPU/llvm.amdgcn.cvt.f16.fp8.ll | 171 ++++++++++++++++++
llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s | 12 ++
llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s | 15 ++
.../MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s | 8 +
llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s | 12 ++
.../MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s | 8 +
llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s | 12 ++
.../gfx1250_asm_vop3_from_vop1-fake16.s | 27 +++
.../MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s | 27 +++
.../gfx1250_asm_vop3_from_vop1_dpp16-fake16.s | 20 ++
.../AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s | 24 +++
.../gfx1250_asm_vop3_from_vop1_dpp8-fake16.s | 28 +++
.../AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s | 32 +++-
.../Disassembler/AMDGPU/gfx1250_dasm_vop1.txt | 23 ++-
.../AMDGPU/gfx1250_dasm_vop1_dpp16.txt | 15 +-
.../AMDGPU/gfx1250_dasm_vop1_dpp8.txt | 15 +-
.../AMDGPU/gfx1250_dasm_vop3_from_vop1.txt | 36 ++++
.../gfx1250_dasm_vop3_from_vop1_dpp16.txt | 24 +++
.../gfx1250_dasm_vop3_from_vop1_dpp8.txt | 32 ++++
26 files changed, 615 insertions(+), 22 deletions(-)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 4e28f3bb7ef81..41fe1ebc4c2ce 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -655,6 +655,7 @@ TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8bf16,
"V8yV8y*3", "nc", "gfx
TARGET_BUILTIN(__builtin_amdgcn_s_setprio_inc_wg, "vIs", "n",
"setprio-inc-wg-inst")
+TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts")
TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_bf8, "V2hs", "nc", "gfx1250-insts")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
index 864e301859682..150d4a243f9e2 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx1250.cl
@@ -15,6 +15,44 @@ void test_setprio_inc_wg() {
__builtin_amdgcn_s_setprio_inc_wg(10);
}
+// CHECK-LABEL: @test_cvt_f16_fp8(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
+// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[OUT_ADDR]] to ptr
+// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[A_ADDR]] to ptr
+// CHECK-NEXT: store ptr addrspace(1) [[OUT:%.*]], ptr [[OUT_ADDR_ASCAST]],
align 8
+// CHECK-NEXT: store i32 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP1:%.*]] = call half @llvm.amdgcn.cvt.f16.fp8(i32
[[TMP0]], i32 0)
+// CHECK-NEXT: [[TMP2:%.*]] = load ptr addrspace(1), ptr
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds half, ptr
addrspace(1) [[TMP2]], i64 0
+// CHECK-NEXT: store half [[TMP1]], ptr addrspace(1) [[ARRAYIDX]], align 2
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call half @llvm.amdgcn.cvt.f16.fp8(i32
[[TMP3]], i32 1)
+// CHECK-NEXT: [[TMP5:%.*]] = load ptr addrspace(1), ptr
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds half, ptr
addrspace(1) [[TMP5]], i64 1
+// CHECK-NEXT: store half [[TMP4]], ptr addrspace(1) [[ARRAYIDX1]], align 2
+// CHECK-NEXT: [[TMP6:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP7:%.*]] = call half @llvm.amdgcn.cvt.f16.fp8(i32
[[TMP6]], i32 2)
+// CHECK-NEXT: [[TMP8:%.*]] = load ptr addrspace(1), ptr
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds half, ptr
addrspace(1) [[TMP8]], i64 2
+// CHECK-NEXT: store half [[TMP7]], ptr addrspace(1) [[ARRAYIDX2]], align 2
+// CHECK-NEXT: [[TMP9:%.*]] = load i32, ptr [[A_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP10:%.*]] = call half @llvm.amdgcn.cvt.f16.fp8(i32
[[TMP9]], i32 3)
+// CHECK-NEXT: [[TMP11:%.*]] = load ptr addrspace(1), ptr
[[OUT_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[ARRAYIDX3:%.*]] = getelementptr inbounds half, ptr
addrspace(1) [[TMP11]], i64 3
+// CHECK-NEXT: store half [[TMP10]], ptr addrspace(1) [[ARRAYIDX3]], align 2
+// CHECK-NEXT: ret void
+//
+void test_cvt_f16_fp8(global half* out, int a)
+{
+ out[0] = __builtin_amdgcn_cvt_f16_fp8(a, 0);
+ out[1] = __builtin_amdgcn_cvt_f16_fp8(a, 1);
+ out[2] = __builtin_amdgcn_cvt_f16_fp8(a, 2);
+ out[3] = __builtin_amdgcn_cvt_f16_fp8(a, 3);
+}
+
// CHECK-LABEL: @test_cvt_pk_f16_fp8(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
index b69fcb5f445bc..b8a71a5ba98a6 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-param.cl
@@ -4,3 +4,7 @@
void test_setprio_inc_wg(short a) {
__builtin_amdgcn_s_setprio_inc_wg(a); // expected-error
{{'__builtin_amdgcn_s_setprio_inc_wg' must be a constant integer}}
}
+
+void test__builtin_amdgcn_cvt_f16_fp8(int a, int b) {
+ __builtin_amdgcn_cvt_f16_fp8(a, b); // expected-error
{{'__builtin_amdgcn_cvt_f16_fp8' must be a constant integer}}
+}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index ce37702b91486..63d649f9d38a1 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -3500,6 +3500,16 @@ def int_amdgcn_ashr_pk_u8_i32 :
ClangBuiltin<"__builtin_amdgcn_ashr_pk_u8_i32">,
DefaultAttrsIntrinsic<[llvm_i16_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]>;
+//===----------------------------------------------------------------------===//
+// gfx1250 intrinsics
+//
===----------------------------------------------------------------------===//
+
+// llvm.amdgcn.cvt.f16.fp8 half vdst, int srcA, imm byte_sel [0..3]
+def int_amdgcn_cvt_f16_fp8 : ClangBuiltin<"__builtin_amdgcn_cvt_f16_fp8">,
+ DefaultAttrsIntrinsic<[llvm_half_ty],
+ [llvm_i32_ty, llvm_i32_ty],
+ [IntrNoMem, ImmArg<ArgIndex<1>>]>;
+
//===----------------------------------------------------------------------===//
// Special Intrinsics for backend internal use only. No frontend
// should emit calls to these.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
index 6874657a4ffe7..2cf9c73e3ec81 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
@@ -4596,6 +4596,7 @@ AMDGPURegisterBankInfo::getInstrMapping(const
MachineInstr &MI) const {
case Intrinsic::amdgcn_cvt_sr_bf8_f32:
case Intrinsic::amdgcn_cvt_sr_bf16_f32:
case Intrinsic::amdgcn_cvt_sr_f16_f32:
+ case Intrinsic::amdgcn_cvt_f16_fp8:
case Intrinsic::amdgcn_cvt_scalef32_pk32_fp6_f16:
case Intrinsic::amdgcn_cvt_scalef32_pk32_bf6_f16:
case Intrinsic::amdgcn_cvt_scalef32_pk32_fp6_bf16:
diff --git a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
index d504c8134202d..55e7eb15bd5a0 100644
--- a/llvm/lib/Target/AMDGPU/VOP1Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP1Instructions.td
@@ -687,6 +687,12 @@ class VOPProfile_Base_CVT_F_F8_ByteSel<ValueType DstVT> :
VOPProfile<[DstVT, i32
let HasModifiers = 0;
}
+let IsSingle = 0, HasOpSel = 1, HasModifiers = 1 in {
+def V_CVT_F16_F8_Profile : VOPProfile_Base_CVT_F_F8_ByteSel<f16>;
+def V_CVT_F16_F8_True16_Profile : VOP3_Profile_True16<V_CVT_F16_F8_Profile>;
+def V_CVT_F16_F8_Fake16_Profile : VOP3_Profile_Fake16<V_CVT_F16_F8_Profile>;
+}
+
let SubtargetPredicate = isGFX12Plus, OtherPredicates =
[HasFP8ConversionInsts],
mayRaiseFPException = 0, SchedRW = [WriteFloatCvt] in {
defm V_CVT_F32_FP8_OP_SEL : VOP1Inst<"v_cvt_f32_fp8_op_sel",
VOPProfile_Base_CVT_F_F8_ByteSel<f32>>;
@@ -702,9 +708,10 @@ let SubtargetPredicate = isGFX12Plus, OtherPredicates =
[HasFP8ConversionInsts],
}
}
-class Cvt_F_F8_Pat_ByteSel<SDPatternOperator node, VOP3_Pseudo inst> : GCNPat<
+class Cvt_F_F8_Pat_ByteSel<SDPatternOperator node, VOP3_Pseudo inst, bit
HasOpSel = 0> : GCNPat<
(node i32:$src0, timm:$byte_sel),
- (inst $src0, (as_i32timm $byte_sel))
+ !if(HasOpSel, (inst 0, $src0, (as_i32timm $byte_sel)),
+ (inst $src0, (as_i32timm $byte_sel)))
>;
let SubtargetPredicate = isGFX12Plus, OtherPredicates =
[HasFP8ConversionInsts] in {
@@ -738,6 +745,8 @@ def VOPProfile_CVT_PK_F16_F8_fake16 :
VOP3_Profile_Fake16<VOPProfile_CVT_PK_F16_
let SubtargetPredicate = isGFX1250Plus in {
let mayRaiseFPException = 0, SchedRW = [WriteFloatCvt] in {
+ defm V_CVT_F16_FP8 : VOP1Inst_t16_with_profiles<"v_cvt_f16_fp8",
+ V_CVT_F16_F8_Profile, V_CVT_F16_F8_True16_Profile,
V_CVT_F16_F8_Fake16_Profile>;
defm V_CVT_PK_F16_FP8 : VOP1Inst_t16_with_profiles<"v_cvt_pk_f16_fp8",
VOPProfile_CVT_PK_F16_F8, VOPProfile_CVT_PK_F16_F8_true16,
VOPProfile_CVT_PK_F16_F8_fake16,
int_amdgcn_cvt_pk_f16_fp8>;
@@ -745,6 +754,13 @@ let SubtargetPredicate = isGFX1250Plus in {
VOPProfile_CVT_PK_F16_F8, VOPProfile_CVT_PK_F16_F8_true16,
VOPProfile_CVT_PK_F16_F8_fake16,
int_amdgcn_cvt_pk_f16_bf8>;
}
+
+ let True16Predicate = UseRealTrue16Insts in {
+ def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_fp8, V_CVT_F16_FP8_t16_e64,
1>;
+ }
+ let True16Predicate = UseFakeTrue16Insts in {
+ def : Cvt_F_F8_Pat_ByteSel<int_amdgcn_cvt_f16_fp8,
V_CVT_F16_FP8_fake16_e64, 1>;
+ }
} // End SubtargetPredicate = isGFX1250Plus
let SubtargetPredicate = isGFX10Plus in {
@@ -1082,6 +1098,7 @@ defm V_CVT_F32_F16 :
VOP1_Real_FULL_t16_and_fake16_gfx11_gfx12<0x00b>;
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>;
+defm V_CVT_F16_FP8 : VOP1_Real_FULL_t16_and_fake16_gfx1250<0x077>;
//===----------------------------------------------------------------------===//
// GFX10.
diff --git a/llvm/lib/Target/AMDGPU/VOPInstructions.td
b/llvm/lib/Target/AMDGPU/VOPInstructions.td
index 1e47acb5fde4f..f8edfec1100a2 100644
--- a/llvm/lib/Target/AMDGPU/VOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOPInstructions.td
@@ -348,7 +348,8 @@ class VOP3OpSel_gfx11_gfx12<bits<10> op, VOPProfile p> :
VOP3OpSel_gfx10<op, p>;
class VOP3FP8OpSel_src_bytesel_gfx11_gfx12<bits<10> op, VOPProfile p> :
VOP3e_gfx10<op, p> {
bits<2> byte_sel;
let Inst{11-12} = byte_sel; // NB: bit order is intentionally reversed!
- let Inst{14-13} = 0; // op_sel2/3
+ let Inst{13} = !if(!and(p.HasOpSel, p.HasSrc2), src2_modifiers{2}, 0);
+ let Inst{14} = !if(!and(p.HasOpSel, p.HasDst), src0_modifiers{3}, 0);
}
class VOP3FP8OpSel_dst_bytesel_gfx11_gfx12<bits<10> op, VOPProfile p> :
VOP3e_gfx10<op, p> {
@@ -526,7 +527,7 @@ class VOP3PXe <bits<7> op, VOPProfile MFMAPfl, bit acc_cd =
0> : Enc128, VOP3Pe_
bits<9> scale_src1;
//MFMALdScaleModifierOp transforms 2 bit opsel input to 4 bit value
- //where opsel and opselHi are in 3rd and 4th bit.
+ //where opsel and opselHi are in 3rd and 4th bit.
bits<4> src0_modifiers;
bits<4> src1_modifiers;
@@ -869,14 +870,14 @@ class VOP3_DPPe_Common_Base<bits<10> op, VOPProfile P> :
Enc96 {
let Inst{9} = !if(P.HasSrc1Mods, src1_modifiers{1}, 0);
let Inst{10} = !if(P.HasSrc2Mods, src2_modifiers{1}, 0);
// 16-bit select fields which can be interpreted as OpSel or hi/lo suffix
- let Inst{11} = !if(P.HasOpSel, !if(P.HasSrc0Mods, src0_modifiers{2}, 0),
- !if(P.HasFP8SrcByteSel, byte_sel{1}, ?));
- let Inst{12} = !if(P.HasOpSel, !if(P.HasSrc1Mods, src1_modifiers{2}, 0),
- !if(P.HasFP8SrcByteSel, byte_sel{0}, ?));
- let Inst{13} = !if(P.HasOpSel, !if(P.HasSrc2Mods, src2_modifiers{2}, 0),
- !if(P.HasFP8DstByteSel, byte_sel{0}, ?));
- let Inst{14} = !if(P.HasOpSel, !if(P.HasSrc0Mods, src0_modifiers{3}, 0),
- !if(P.HasFP8DstByteSel, byte_sel{1}, ?));
+ let Inst{11} = !if(P.HasFP8SrcByteSel, byte_sel{1},
+ !if(P.HasOpSel, !if(P.HasSrc0Mods, src0_modifiers{2}, 0), ?));
+ let Inst{12} = !if(P.HasFP8SrcByteSel, byte_sel{0},
+ !if(P.HasOpSel, !if(P.HasSrc1Mods, src1_modifiers{2}, 0), ?));
+ let Inst{13} = !if(P.HasFP8DstByteSel, byte_sel{0},
+ !if(P.HasOpSel, !if(P.HasSrc2Mods, src2_modifiers{2}, 0), ?));
+ let Inst{14} = !if(P.HasFP8DstByteSel, byte_sel{1},
+ !if(P.HasOpSel, !if(P.HasSrc0Mods, src0_modifiers{3}, 0), ?));
let Inst{15} = !if(P.HasClamp, clamp, 0);
let Inst{25-16} = op;
let Inst{31-26} = 0x35;
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.f16.fp8.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.f16.fp8.ll
index 243f6c4d23732..25889ded91681 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.f16.fp8.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.cvt.f16.fp8.ll
@@ -4,6 +4,177 @@
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 %s
-o - | FileCheck -check-prefixes=GFX1250,GFX1250-GISEL-REAL16 %s
; RUN: llc -global-isel=1 -mtriple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 %s
-o - | FileCheck -check-prefixes=GFX1250,GFX1250-GISEL-FAKE16 %s
+declare half @llvm.amdgcn.cvt.f16.bf8(i32, i32)
+declare half @llvm.amdgcn.cvt.f16.fp8(i32, i32)
+declare <2 x half> @llvm.amdgcn.cvt.pk.f16.bf8(i16)
+declare <2 x half> @llvm.amdgcn.cvt.pk.f16.fp8(i16)
+
+define amdgpu_ps float @test_cvt_f16_fp8_byte0(i32 %a) {
+; GFX1250-SDAG-REAL16-LABEL: test_cvt_f16_fp8_byte0:
+; GFX1250-SDAG-REAL16: ; %bb.0:
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f16_fp8_e32 v0.l, v0
+; GFX1250-SDAG-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-SDAG-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-SDAG-FAKE16-LABEL: test_cvt_f16_fp8_byte0:
+; GFX1250-SDAG-FAKE16: ; %bb.0:
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f16_fp8_e32 v0, v0
+; GFX1250-SDAG-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-SDAG-FAKE16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-REAL16-LABEL: test_cvt_f16_fp8_byte0:
+; GFX1250-GISEL-REAL16: ; %bb.0:
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f16_fp8_e32 v0.l, v0
+; GFX1250-GISEL-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-GISEL-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-FAKE16-LABEL: test_cvt_f16_fp8_byte0:
+; GFX1250-GISEL-FAKE16: ; %bb.0:
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f16_fp8_e32 v0, v0
+; GFX1250-GISEL-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-GISEL-FAKE16-NEXT: ; return to shader part epilog
+ %cvt = tail call half @llvm.amdgcn.cvt.f16.fp8(i32 %a, i32 0)
+ %ret = fpext half %cvt to float
+ ret float %ret
+}
+
+define amdgpu_ps float @test_cvt_f16_fp8_byte1(i32 %a) {
+; GFX1250-SDAG-REAL16-LABEL: test_cvt_f16_fp8_byte1:
+; GFX1250-SDAG-REAL16: ; %bb.0:
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:1
+; GFX1250-SDAG-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-SDAG-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-SDAG-FAKE16-LABEL: test_cvt_f16_fp8_byte1:
+; GFX1250-SDAG-FAKE16: ; %bb.0:
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:1
+; GFX1250-SDAG-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-SDAG-FAKE16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-REAL16-LABEL: test_cvt_f16_fp8_byte1:
+; GFX1250-GISEL-REAL16: ; %bb.0:
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:1
+; GFX1250-GISEL-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-GISEL-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-FAKE16-LABEL: test_cvt_f16_fp8_byte1:
+; GFX1250-GISEL-FAKE16: ; %bb.0:
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:1
+; GFX1250-GISEL-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-GISEL-FAKE16-NEXT: ; return to shader part epilog
+ %cvt = tail call half @llvm.amdgcn.cvt.f16.fp8(i32 %a, i32 1)
+ %ret = fpext half %cvt to float
+ ret float %ret
+}
+
+define amdgpu_ps float @test_cvt_f16_fp8_byte2(i32 %a) {
+; GFX1250-SDAG-REAL16-LABEL: test_cvt_f16_fp8_byte2:
+; GFX1250-SDAG-REAL16: ; %bb.0:
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:2
+; GFX1250-SDAG-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-SDAG-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-SDAG-FAKE16-LABEL: test_cvt_f16_fp8_byte2:
+; GFX1250-SDAG-FAKE16: ; %bb.0:
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:2
+; GFX1250-SDAG-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-SDAG-FAKE16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-REAL16-LABEL: test_cvt_f16_fp8_byte2:
+; GFX1250-GISEL-REAL16: ; %bb.0:
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:2
+; GFX1250-GISEL-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-GISEL-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-FAKE16-LABEL: test_cvt_f16_fp8_byte2:
+; GFX1250-GISEL-FAKE16: ; %bb.0:
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:2
+; GFX1250-GISEL-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-GISEL-FAKE16-NEXT: ; return to shader part epilog
+ %cvt = tail call half @llvm.amdgcn.cvt.f16.fp8(i32 %a, i32 2)
+ %ret = fpext half %cvt to float
+ ret float %ret
+}
+
+define amdgpu_ps float @test_cvt_f16_fp8_byte3(i32 %a) {
+; GFX1250-SDAG-REAL16-LABEL: test_cvt_f16_fp8_byte3:
+; GFX1250-SDAG-REAL16: ; %bb.0:
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:3
+; GFX1250-SDAG-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-SDAG-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-SDAG-FAKE16-LABEL: test_cvt_f16_fp8_byte3:
+; GFX1250-SDAG-FAKE16: ; %bb.0:
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:3
+; GFX1250-SDAG-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-SDAG-FAKE16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-REAL16-LABEL: test_cvt_f16_fp8_byte3:
+; GFX1250-GISEL-REAL16: ; %bb.0:
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:3
+; GFX1250-GISEL-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f32_f16_e32 v0, v0.l
+; GFX1250-GISEL-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-FAKE16-LABEL: test_cvt_f16_fp8_byte3:
+; GFX1250-GISEL-FAKE16: ; %bb.0:
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:3
+; GFX1250-GISEL-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f32_f16_e32 v0, v0
+; GFX1250-GISEL-FAKE16-NEXT: ; return to shader part epilog
+ %cvt = tail call half @llvm.amdgcn.cvt.f16.fp8(i32 %a, i32 3)
+ %ret = fpext half %cvt to float
+ ret float %ret
+}
+
+define amdgpu_ps float @test_cvt_f16_fp8_byte3_hi(i32 %a) {
+; GFX1250-SDAG-REAL16-LABEL: test_cvt_f16_fp8_byte3_hi:
+; GFX1250-SDAG-REAL16: ; %bb.0:
+; GFX1250-SDAG-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.h, v0 byte_sel:3
+; GFX1250-SDAG-REAL16-NEXT: v_mov_b16_e32 v0.l, 0
+; GFX1250-SDAG-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-SDAG-FAKE16-LABEL: test_cvt_f16_fp8_byte3_hi:
+; GFX1250-SDAG-FAKE16: ; %bb.0:
+; GFX1250-SDAG-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:3
+; GFX1250-SDAG-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-SDAG-FAKE16-NEXT: v_perm_b32 v0, v0, 0, 0x5040100
+; GFX1250-SDAG-FAKE16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-REAL16-LABEL: test_cvt_f16_fp8_byte3_hi:
+; GFX1250-GISEL-REAL16: ; %bb.0:
+; GFX1250-GISEL-REAL16-NEXT: v_cvt_f16_fp8_e64 v0.l, v0 byte_sel:3
+; GFX1250-GISEL-REAL16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-REAL16-NEXT: v_lshl_or_b32 v0, v0, 16, 0
+; GFX1250-GISEL-REAL16-NEXT: ; return to shader part epilog
+;
+; GFX1250-GISEL-FAKE16-LABEL: test_cvt_f16_fp8_byte3_hi:
+; GFX1250-GISEL-FAKE16: ; %bb.0:
+; GFX1250-GISEL-FAKE16-NEXT: v_cvt_f16_fp8_e64 v0, v0 byte_sel:3
+; GFX1250-GISEL-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX1250-GISEL-FAKE16-NEXT: v_lshl_or_b32 v0, v0, 16, 0
+; GFX1250-GISEL-FAKE16-NEXT: ; return to shader part epilog
+ %cvt = tail call half @llvm.amdgcn.cvt.f16.fp8(i32 %a, i32 3)
+ %ins.0 = insertelement <2 x half> undef, half 0.0, i32 0
+ %ins.1 = insertelement <2 x half> %ins.0, half %cvt, i32 1
+ %ret = bitcast <2 x half> %ins.1 to float
+ ret float %ret
+}
+
define amdgpu_ps float @test_cvt_pk_f16_bf8_v(i16 %a) {
; GFX1250-SDAG-REAL16-LABEL: test_cvt_pk_f16_bf8_v:
; GFX1250-SDAG-REAL16: ; %bb.0:
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
index e62eb6fbb723c..d3b9d403e5088 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1-fake16.s
@@ -46,6 +46,18 @@ v_cvt_f32_bf16 v5, src_scc
v_cvt_f32_bf16 v127, 0x8000
// GFX1250: v_cvt_f32_bf16_e32 v127, 0x8000 ; encoding:
[0xff,0xe4,0xfe,0x7e,0x00,0x80,0x00,0x00]
+v_cvt_f16_fp8 v1, v2
+// GFX1250: v_cvt_f16_fp8_e32 v1, v2 ; encoding:
[0x02,0xef,0x02,0x7e]
+
+v_cvt_f16_fp8 v1, s2
+// GFX1250: v_cvt_f16_fp8_e32 v1, s2 ; encoding:
[0x02,0xee,0x02,0x7e]
+
+v_cvt_f16_fp8 v1, 2
+// GFX1250: v_cvt_f16_fp8_e32 v1, 2 ; encoding:
[0x82,0xee,0x02,0x7e]
+
+v_cvt_f16_fp8 v1, 0x1234
+// GFX1250: v_cvt_f16_fp8_e32 v1, 0x1234 ; encoding:
[0xff,0xee,0x02,0x7e,0x34,0x12,0x00,0x00]
+
v_cvt_pk_f16_bf8 v1, v2
// GFX1250: v_cvt_pk_f16_bf8 v1, v2 ; encoding:
[0x02,0xed,0x02,0x7e]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
index 37f39546ae13d..dd070651e58ca 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1.s
@@ -49,6 +49,21 @@ v_cvt_f32_bf16 v127, 0x8000
v_cvt_f32_bf16 v5, v1.h
// GFX1250: v_cvt_f32_bf16_e32 v5, v1.h ; encoding:
[0x81,0xe5,0x0a,0x7e]
+v_cvt_f16_fp8 v1.l, v2
+// GFX1250: v_cvt_f16_fp8_e32 v1.l, v2 ; encoding:
[0x02,0xef,0x02,0x7e]
+
+v_cvt_f16_fp8 v1.l, s2
+// GFX1250: v_cvt_f16_fp8_e32 v1.l, s2 ; encoding:
[0x02,0xee,0x02,0x7e]
+
+v_cvt_f16_fp8 v1.l, 2
+// GFX1250: v_cvt_f16_fp8_e32 v1.l, 2 ; encoding:
[0x82,0xee,0x02,0x7e]
+
+v_cvt_f16_fp8 v1.l, 0x1234
+// GFX1250: v_cvt_f16_fp8_e32 v1.l, 0x1234 ; encoding:
[0xff,0xee,0x02,0x7e,0x34,0x12,0x00,0x00]
+
+v_cvt_f16_fp8 v1.h, v2
+// GFX1250: v_cvt_f16_fp8_e32 v1.h, v2 ; encoding:
[0x02,0xef,0x02,0x7f]
+
v_cvt_pk_f16_bf8 v1, v2
// GFX1250: v_cvt_pk_f16_bf8 v1, v2 ; encoding:
[0x02,0xed,0x02,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 766992b0761d9..f2751b7aecb49 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16-fake16.s
@@ -58,6 +58,14 @@ v_cvt_f32_bf16 v127, -|v127| row_xmask:15 row_mask:0x3
bank_mask:0x0 bound_ctrl:
// GFX1250: v_cvt_f32_bf16_dpp v127, -|v127| row_xmask:15 row_mask:0x3
bank_mask:0x0 fi:1 ; encoding: [0xfa,0xe4,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_cvt_f16_fp8 v1, v2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_dpp v1, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf fi:1
+// GFX1250: v_cvt_f16_fp8_dpp v1, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x04,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v2 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf fi:1
// GFX1250: v_cvt_pk_f16_bf8_dpp v1, v2 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xec,0x02,0x7e,0x02,0xe4,0x04,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 d674a9ea06843..525963a8c5ba5 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp16.s
@@ -62,6 +62,18 @@ v_cvt_f32_bf16 v5, v1.h quad_perm:[3,2,1,0]
// GFX1250: v_cvt_f32_bf16_dpp v5, v1.h quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xe4,0x0a,0x7e,0x81,0x1b,0x00,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_cvt_f16_fp8 v1.l, v2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_dpp v1.l, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.l, v2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf fi:1
+// GFX1250: v_cvt_f16_fp8_dpp v1.l, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x04,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.h, v2 quad_perm:[0,1,2,3]
+// GFX1250: v_cvt_f16_fp8_dpp v1.h, v2 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xee,0x02,0x7f,0x02,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v2.l quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf fi:1
// GFX1250: v_cvt_pk_f16_bf8_dpp v1, v2.l quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xec,0x02,0x7e,0x02,0xe4,0x04,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
index 481c069de6c9b..1182f4279e159 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8-fake16.s
@@ -14,6 +14,14 @@ v_cvt_f32_bf16 v127, v127 dpp8:[0,0,0,0,0,0,0,0] fi:0
// GFX1250: v_cvt_f32_bf16_dpp v127, v127 dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xe9,0xe4,0xfe,0x7e,0x7f,0x00,0x00,0x00]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_cvt_f16_fp8 v1, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0xe9,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// GFX1250: v_cvt_f16_fp8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ; encoding:
[0xea,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
// GFX1250: v_cvt_pk_f16_bf8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0xea,0xec,0x02,0x7e,0x02,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
index 81a40acc35871..14291a3dea5e1 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop1_dpp8.s
@@ -18,6 +18,18 @@ v_cvt_f32_bf16 v5, v1.h dpp8:[7,6,5,4,3,2,1,0]
// GFX1250: v_cvt_f32_bf16_dpp v5, v1.h dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0xe9,0xe4,0x0a,0x7e,0x81,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_cvt_f16_fp8 v1, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0xe9,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// GFX1250: v_cvt_f16_fp8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ; encoding:
[0xea,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.h, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_dpp v1.h, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0xe9,0xee,0x02,0x7f,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v2.l dpp8:[7,6,5,4,3,2,1,0] fi:1
// GFX1250: v_cvt_pk_f16_bf8_dpp v1, v2.l dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0xea,0xec,0x02,0x7e,0x02,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
index f6c7cf8006508..44e0e3efd965f 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1-fake16.s
@@ -76,6 +76,33 @@ v_cvt_f32_bf16_e64 v5, -1 op_sel:[1]
v_cvt_f32_bf16_e64 v5, src_scc op_sel:[1]
// GFX1250: v_cvt_f32_bf16_e64 v5, src_scc op_sel:[1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0xfd,0x00,0x00,0x00]
+v_cvt_f16_fp8 v150, v2
+// GFX1250: v_cvt_f16_fp8_e64 v150, v2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v150, s2
+// GFX1250: v_cvt_f16_fp8_e64 v150, s2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x00,0x00,0x00]
+
+v_cvt_f16_fp8 v150, 2
+// GFX1250: v_cvt_f16_fp8_e64 v150, 2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x82,0x00,0x00,0x00]
+
+v_cvt_f16_fp8 v150, 0x1234
+// GFX1250: v_cvt_f16_fp8_e64 v150, 0x1234 ; encoding:
[0x96,0x00,0xf7,0xd5,0xff,0x00,0x00,0x00,0x34,0x12,0x00,0x00]
+
+v_cvt_f16_fp8 v1, v2 byte_sel:2
+// GFX1250: v_cvt_f16_fp8_e64 v1, v2 byte_sel:2 ; encoding:
[0x01,0x08,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v1, v2 byte_sel:1
+// GFX1250: v_cvt_f16_fp8_e64 v1, v2 byte_sel:1 ; encoding:
[0x01,0x10,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v1, v2 byte_sel:3
+// GFX1250: v_cvt_f16_fp8_e64 v1, v2 byte_sel:3 ; encoding:
[0x01,0x18,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v128, v2 op_sel:[0,1]
+// GFX1250: v_cvt_f16_fp8_e64 v128, v2 op_sel:[0,1] ; encoding:
[0x80,0x40,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v1, v2 op_sel:[0,1] byte_sel:2
+// GFX1250: v_cvt_f16_fp8_e64 v1, v2 op_sel:[0,1] byte_sel:2 ; encoding:
[0x01,0x48,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
v_cvt_pk_f16_bf8 v1, v150
// GFX1250: v_cvt_pk_f16_bf8 v1, v150 ; encoding:
[0x01,0x00,0xf6,0xd5,0x96,0x01,0x00,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
index 531d734a0683d..5546841e9154b 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1.s
@@ -79,6 +79,33 @@ v_cvt_f32_bf16_e64 v5, src_scc op_sel:[1]
v_cvt_f32_bf16_e64 v5, v128.h
// GFX1250: v_cvt_f32_bf16_e64 v5, v128.h op_sel:[1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0x80,0x01,0x00,0x00]
+v_cvt_f16_fp8 v150.l, v2
+// GFX1250: v_cvt_f16_fp8_e64 v150.l, v2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v150.l, s2
+// GFX1250: v_cvt_f16_fp8_e64 v150.l, s2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x00,0x00,0x00]
+
+v_cvt_f16_fp8 v150.l, 2
+// GFX1250: v_cvt_f16_fp8_e64 v150.l, 2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x82,0x00,0x00,0x00]
+
+v_cvt_f16_fp8 v150.l, 0x1234
+// GFX1250: v_cvt_f16_fp8_e64 v150.l, 0x1234 ; encoding:
[0x96,0x00,0xf7,0xd5,0xff,0x00,0x00,0x00,0x34,0x12,0x00,0x00]
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:2
+// GFX1250: v_cvt_f16_fp8_e64 v1.l, v2 byte_sel:2 ; encoding:
[0x01,0x08,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:1
+// GFX1250: v_cvt_f16_fp8_e64 v1.l, v2 byte_sel:1 ; encoding:
[0x01,0x10,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:3
+// GFX1250: v_cvt_f16_fp8_e64 v1.l, v2 byte_sel:3 ; encoding:
[0x01,0x18,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v128.h, v2
+// GFX1250: v_cvt_f16_fp8_e64 v128.h, v2 op_sel:[0,1] ; encoding:
[0x80,0x40,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+v_cvt_f16_fp8 v1.h, v2 byte_sel:2
+// GFX1250: v_cvt_f16_fp8_e64 v1.h, v2 op_sel:[0,1] byte_sel:2 ; encoding:
[0x01,0x48,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
v_cvt_pk_f16_bf8 v1, v150
// GFX1250: v_cvt_pk_f16_bf8 v1, v150 ; encoding:
[0x01,0x00,0xf6,0xd5,0x96,0x01,0x00,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
index 844b4259229ed..8f2bd6b9ddb77 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16-fake16.s
@@ -46,6 +46,26 @@ v_cvt_f32_bf16_e64_dpp v5, v1 row_share:0 row_mask:0xf
bank_mask:0xf
// GFX1250: v_cvt_f32_bf16_e64_dpp v5, v1 row_share:0 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0xf2,0xd5,0xfa,0x00,0x00,0x00,0x01,0x50,0x01,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_cvt_f16_fp8 v1, v2 byte_sel:2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:2 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x08,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 byte_sel:1 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:1 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x10,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 byte_sel:3 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:3 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x18,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v150, v2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x96,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 op_sel:[0,1] byte_sel:3 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 op_sel:[0,1] byte_sel:3
quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x58,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v128 quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf fi:1
// GFX1250: v_cvt_pk_f16_bf8_e64_dpp v1, v128 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf fi:1 ; encoding:
[0x01,0x00,0xf6,0xd5,0xfa,0x00,0x00,0x00,0x80,0xe4,0x04,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
index 32c2e54cf0e71..9f2cc29dcd0b5 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp16.s
@@ -50,6 +50,30 @@ v_cvt_f32_bf16_e64_dpp v5, v128.h quad_perm:[3,2,1,0]
// GFX1250: v_cvt_f32_bf16_e64_dpp v5, v128.h op_sel:[1,0] quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x08,0xf2,0xd5,0xfa,0x00,0x00,0x00,0x80,0x1b,0x00,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+v_cvt_f16_fp8 v1.l, v2 byte_sel:2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:2 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x08,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:1 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:1 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x10,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:3 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:3 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x18,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v150.l, v2 quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150.l, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x96,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.h, v2 byte_sel:3 quad_perm:[0,1,2,3]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.h, v2 op_sel:[0,1] byte_sel:3
quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x58,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v128.l, v2 quad_perm:[0,1,2,3]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v128.l, v2 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf ; encoding:
[0x80,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v128.l quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf fi:1
// GFX1250: v_cvt_pk_f16_bf8_e64_dpp v1, v128.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf fi:1 ; encoding:
[0x01,0x00,0xf6,0xd5,0xfa,0x00,0x00,0x00,0x80,0xe4,0x04,0xff]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
index 7ff69fad0fdf1..608e07b9f0f5c 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8-fake16.s
@@ -2,6 +2,34 @@
// RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16
-show-encoding < %s | FileCheck --check-prefix=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_cvt_f16_fp8 v150, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x96,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x01,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 byte_sel:2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x01,0x08,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 byte_sel:1 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:1 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x01,0x10,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v150, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x96,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v150, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0x96,0x00,0xf7,0xd5,0xea,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1, v2 op_sel:[0,1] byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1, v2 op_sel:[0,1] byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x58,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
v_cvt_pk_f16_bf8 v1, v128 dpp8:[7,6,5,4,3,2,1,0] fi:1
// GFX1250: v_cvt_pk_f16_bf8_e64_dpp v1, v128 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0x01,0x00,0xf6,0xd5,0xea,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
diff --git a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
index e307e407b8e8e..94e1d375630e5 100644
--- a/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
+++ b/llvm/test/MC/AMDGPU/gfx1250_asm_vop3_from_vop1_dpp8.s
@@ -2,12 +2,36 @@
// RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16
-show-encoding < %s | FileCheck --check-prefix=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_cvt_f32_bf16_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0]
-// GFX1250: v_cvt_f32_bf16_e64_dpp v5, v1 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0xf2,0xd5,0xe9,0x00,0x00,0x00,0x01,0x77,0x39,0x05]
+v_cvt_f16_fp8 v150.l, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150.l, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x96,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
-v_cvt_f32_bf16_e64_dpp v5, v128.h dpp8:[7,6,5,4,3,2,1,0]
-// GFX1250: v_cvt_f32_bf16_e64_dpp v5, v128.h op_sel:[1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0xe9,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
+v_cvt_f16_fp8 v1.l, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x01,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x01,0x08,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.l, v2 byte_sel:1 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:1 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x01,0x10,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v150.l, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150.l, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x96,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v150.l, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
+// GFX1250: v_cvt_f16_fp8_e64_dpp v150.l, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0x96,0x00,0xf7,0xd5,0xea,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v1.h, v2 byte_sel:3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v1.h, v2 op_sel:[0,1] byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x58,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
+
+v_cvt_f16_fp8 v128.l, v2 dpp8:[7,6,5,4,3,2,1,0]
+// GFX1250: v_cvt_f16_fp8_e64_dpp v128.l, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x80,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
// GFX12-ERR: :[[@LINE-2]]:1: error: instruction not supported on this GPU
v_cvt_pk_f16_bf8 v1, v128.l dpp8:[7,6,5,4,3,2,1,0] fi:1
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
index 47eebb9d44a95..9518a48120970 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1.txt
@@ -1,6 +1,6 @@
# 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 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250 %s
-# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-REAL16 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-FAKE16 %s
0xff,0xe4,0xfe,0x7e,0x00,0x80,0x00,0x00
# GFX1250: v_cvt_f32_bf16_e32 v127, 0x8000 ; encoding:
[0xff,0xe4,0xfe,0x7e,0x00,0x80,0x00,0x00]
@@ -50,6 +50,25 @@
0x81,0xe5,0x0a,0x7e
# GFX1250: v_cvt_f32_bf16_e32 v5, v1.h ; encoding:
[0x81,0xe5,0x0a,0x7e]
+0xff,0xee,0x02,0x7e,0x34,0x12,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e32 v1.l, 0x1234 ; encoding:
[0xff,0xee,0x02,0x7e,0x34,0x12,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e32 v1, 0x1234 ; encoding:
[0xff,0xee,0x02,0x7e,0x34,0x12,0x00,0x00]
+
+0x82,0xee,0x02,0x7e
+# GFX1250-REAL16: v_cvt_f16_fp8_e32 v1.l, 2 ; encoding:
[0x82,0xee,0x02,0x7e]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e32 v1, 2 ; encoding:
[0x82,0xee,0x02,0x7e]
+
+0x02,0xee,0x02,0x7e
+# GFX1250-REAL16: v_cvt_f16_fp8_e32 v1.l, s2 ; encoding:
[0x02,0xee,0x02,0x7e]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e32 v1, s2 ; encoding:
[0x02,0xee,0x02,0x7e]
+
+0x02,0xef,0x02,0x7e
+# GFX1250-REAL16: v_cvt_f16_fp8_e32 v1.l, v2 ; encoding:
[0x02,0xef,0x02,0x7e]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e32 v1, v2 ; encoding:
[0x02,0xef,0x02,0x7e]
+
+0x02,0xef,0x02,0x7f
+# GFX1250-REAL16: v_cvt_f16_fp8_e32 v1.h, v2 ; encoding:
[0x02,0xef,0x02,0x7f]
+
0xff,0xec,0x02,0x7e,0x64,0x00,0x00,0x00
# GFX1250: v_cvt_pk_f16_bf8 v1, 0x64 ; encoding:
[0xff,0xec,0x02,0x7e,0x64,0x00,0x00,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
index 25e982b7fd688..16bdfc0b3fdfd 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp16.txt
@@ -1,6 +1,6 @@
# 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 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250 %s
-# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-REAL16 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-FAKE16 %s
0xfa,0xe4,0xfe,0x7e,0x7f,0x6f,0x35,0x30
# GFX1250: v_cvt_f32_bf16_dpp v127, -|v127.l| row_xmask:15 row_mask:0x3
bank_mask:0x0 fi:1 ; encoding: [0xfa,0xe4,0xfe,0x7e,0x7f,0x6f,0x35,0x30]
@@ -47,6 +47,17 @@
0xfa,0xe4,0x0a,0x7e,0x81,0x1b,0x00,0xff
# GFX1250: v_cvt_f32_bf16_dpp v5, v1.h quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xe4,0x0a,0x7e,0x81,0x1b,0x00,0xff]
+0xfa,0xee,0x02,0x7e,0x02,0x39,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_dpp v1.l, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_dpp v1, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x00,0xff]
+
+0xfa,0xee,0x02,0x7e,0x02,0x39,0x04,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_dpp v1.l, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x04,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_dpp v1, v2 quad_perm:[1,2,3,0] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xee,0x02,0x7e,0x02,0x39,0x04,0xff]
+
+0xfa,0xee,0x02,0x7f,0x02,0xe4,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_dpp v1.h, v2 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf ; encoding: [0xfa,0xee,0x02,0x7f,0x02,0xe4,0x00,0xff]
+
0xfa,0xec,0x02,0x7e,0x02,0xe4,0x04,0xff
# GFX1250-REAL16: v_cvt_pk_f16_bf8_dpp v1, v2.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf fi:1 ; encoding:
[0xfa,0xec,0x02,0x7e,0x02,0xe4,0x04,0xff]
# GFX1250-FAKE16: v_cvt_pk_f16_bf8_dpp v1, v2 quad_perm:[0,1,2,3] row_mask:0xf
bank_mask:0xf fi:1 ; encoding: [0xfa,0xec,0x02,0x7e,0x02,0xe4,0x04,0xff]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
index bd524af907ee0..694417f233ae9 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop1_dpp8.txt
@@ -1,6 +1,6 @@
# 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 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250 %s
-# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=+real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-REAL16 %s
+# RUN: llvm-mc -triple=amdgcn -mcpu=gfx1250 -mattr=-real-true16 -disassemble
-show-encoding < %s | FileCheck -check-prefixes=GFX1250,GFX1250-FAKE16 %s
0xe9,0xe4,0xfe,0x7e,0x7f,0x00,0x00,0x00
# GFX1250: v_cvt_f32_bf16_dpp v127, v127.l dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xe9,0xe4,0xfe,0x7e,0x7f,0x00,0x00,0x00]
@@ -14,6 +14,17 @@
0xe9,0xe4,0x0a,0x7e,0x81,0x77,0x39,0x05
# GFX1250: v_cvt_f32_bf16_dpp v5, v1.h dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0xe9,0xe4,0x0a,0x7e,0x81,0x77,0x39,0x05]
+0xe9,0xee,0x02,0x7e,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_dpp v1.l, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0xe9,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0xe9,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+
+0xea,0xee,0x02,0x7e,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_dpp v1.l, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0xea,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0xea,0xee,0x02,0x7e,0x02,0x77,0x39,0x05]
+
+0xe9,0xee,0x02,0x7f,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_dpp v1.h, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0xe9,0xee,0x02,0x7f,0x02,0x77,0x39,0x05]
+
0xea,0xec,0x02,0x7e,0x02,0x77,0x39,0x05
# GFX1250-REAL16: v_cvt_pk_f16_bf8_dpp v1, v2.l dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0xea,0xec,0x02,0x7e,0x02,0x77,0x39,0x05]
# GFX1250-FAKE16: v_cvt_pk_f16_bf8_dpp v1, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0xea,0xec,0x02,0x7e,0x02,0x77,0x39,0x05]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
index 70abf4289ac11..fc544304bf023 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1.txt
@@ -48,6 +48,42 @@
# GFX1250-REAL16: v_cvt_f32_bf16_e64 v5, v255.h op_sel:[1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0xff,0x01,0x00,0x00]
# GFX1250-FAKE16: v_cvt_f32_bf16_e64 v5, v255 op_sel:[1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0xff,0x01,0x00,0x00]
+0x01,0x10,0xf7,0xd5,0x02,0x01,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v1.l, v2 byte_sel:1 ; encoding:
[0x01,0x10,0xf7,0xd5,0x02,0x01,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v1, v2 byte_sel:1 ; encoding:
[0x01,0x10,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+0x01,0x08,0xf7,0xd5,0x02,0x01,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v1.l, v2 byte_sel:2 ; encoding:
[0x01,0x08,0xf7,0xd5,0x02,0x01,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v1, v2 byte_sel:2 ; encoding:
[0x01,0x08,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+0x01,0x18,0xf7,0xd5,0x02,0x01,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v1.l, v2 byte_sel:3 ; encoding:
[0x01,0x18,0xf7,0xd5,0x02,0x01,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v1, v2 byte_sel:3 ; encoding:
[0x01,0x18,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+0x96,0x00,0xf7,0xd5,0xff,0x00,0x00,0x00,0x34,0x12,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v150.l, 0x1234 ; encoding:
[0x96,0x00,0xf7,0xd5,0xff,0x00,0x00,0x00,0x34,0x12,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v150, 0x1234 ; encoding:
[0x96,0x00,0xf7,0xd5,0xff,0x00,0x00,0x00,0x34,0x12,0x00,0x00]
+
+0x96,0x00,0xf7,0xd5,0x82,0x00,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v150.l, 2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x82,0x00,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v150, 2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x82,0x00,0x00,0x00]
+
+0x96,0x00,0xf7,0xd5,0x02,0x00,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v150.l, s2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x00,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v150, s2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x00,0x00,0x00]
+
+0x96,0x00,0xf7,0xd5,0x02,0x01,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v150.l, v2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x01,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v150, v2 ; encoding:
[0x96,0x00,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+0x80,0x40,0xf7,0xd5,0x02,0x01,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v128.h, v2 op_sel:[0,1] ; encoding:
[0x80,0x40,0xf7,0xd5,0x02,0x01,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v128, v2 op_sel:[0,1] ; encoding:
[0x80,0x40,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
+0x01,0x48,0xf7,0xd5,0x02,0x01,0x00,0x00
+# GFX1250-REAL16: v_cvt_f16_fp8_e64 v1.h, v2 op_sel:[0,1] byte_sel:2 ;
encoding: [0x01,0x48,0xf7,0xd5,0x02,0x01,0x00,0x00]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64 v1, v2 op_sel:[0,1] byte_sel:2 ; encoding:
[0x01,0x48,0xf7,0xd5,0x02,0x01,0x00,0x00]
+
0x01,0x08,0xf6,0xd5,0x02,0x00,0x00,0x00
# GFX1250: v_cvt_pk_f16_bf8 v1, s2 op_sel:[1,0] ; encoding:
[0x01,0x08,0xf6,0xd5,0x02,0x00,0x00,0x00]
diff --git
a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
index d53d532eef804..57135e8356bfd 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp16.txt
@@ -50,6 +50,30 @@
# GFX1250-REAL16: v_cvt_f32_bf16_e64_dpp v5, v128.h op_sel:[1,0]
quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x08,0xf2,0xd5,0xfa,0x00,0x00,0x00,0x80,0x1b,0x00,0xff]
# GFX1250-FAKE16: v_cvt_f32_bf16_e64_dpp v5, v128 op_sel:[1,0]
quad_perm:[3,2,1,0] row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x08,0xf2,0xd5,0xfa,0x00,0x00,0x00,0x80,0x1b,0x00,0xff]
+0x01,0x10,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:1
quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x10,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:1 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x10,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+
+0x01,0x08,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:2
quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x08,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:2 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x08,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+
+0x01,0x18,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:3
quad_perm:[1,2,3,0] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x18,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:3 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x18,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+
+0x96,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v150.l, v2 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x96,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v150, v2 quad_perm:[1,2,3,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x96,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0x39,0x00,0xff]
+
+0x01,0x58,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.h, v2 op_sel:[0,1] byte_sel:3
quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x58,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 op_sel:[0,1] byte_sel:3
quad_perm:[0,1,2,3] row_mask:0xf bank_mask:0xf ; encoding:
[0x01,0x58,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+
+0x80,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v128.l, v2 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x80,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v128, v2 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x80,0x00,0xf7,0xd5,0xfa,0x00,0x00,0x00,0x02,0xe4,0x00,0xff]
+
0x01,0x00,0xf6,0xd5,0xfa,0x00,0x00,0x00,0x80,0xe4,0x04,0xff
# GFX1250-REAL16: v_cvt_pk_f16_bf8_e64_dpp v1, v128.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf fi:1 ; encoding:
[0x01,0x00,0xf6,0xd5,0xfa,0x00,0x00,0x00,0x80,0xe4,0x04,0xff]
# GFX1250-FAKE16: v_cvt_pk_f16_bf8_e64_dpp v1, v128 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf fi:1 ; encoding:
[0x01,0x00,0xf6,0xd5,0xfa,0x00,0x00,0x00,0x80,0xe4,0x04,0xff]
diff --git
a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
index 8df21f3f5e4df..9cd2bd088f889 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx1250_dasm_vop3_from_vop1_dpp8.txt
@@ -10,6 +10,38 @@
# GFX1250-REAL16: v_cvt_f32_bf16_e64_dpp v5, v128.h op_sel:[1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0xe9,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
# GFX1250-FAKE16: v_cvt_f32_bf16_e64_dpp v5, v128 op_sel:[1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x08,0xf2,0xd5,0xe9,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
+0x01,0x10,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:1
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x10,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:1
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x10,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x01,0x08,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:2
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x08,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:2
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x08,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x01,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.l, v2 byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x96,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v150.l, v2 byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x96,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v150, v2 byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x96,0x18,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x96,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v150.l, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x96,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v150, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x96,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x01,0x58,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v1.h, v2 op_sel:[0,1] byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x58,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v1, v2 op_sel:[0,1] byte_sel:3
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x01,0x58,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x80,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v128.l, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x80,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v128, v2 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x80,0x00,0xf7,0xd5,0xe9,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
+0x96,0x00,0xf7,0xd5,0xea,0x00,0x00,0x00,0x02,0x77,0x39,0x05
+# GFX1250-REAL16: v_cvt_f16_fp8_e64_dpp v150.l, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1
; encoding: [0x96,0x00,0xf7,0xd5,0xea,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+# GFX1250-FAKE16: v_cvt_f16_fp8_e64_dpp v150, v2 dpp8:[7,6,5,4,3,2,1,0] fi:1 ;
encoding: [0x96,0x00,0xf7,0xd5,0xea,0x00,0x00,0x00,0x02,0x77,0x39,0x05]
+
0x01,0x00,0xf6,0xd5,0xea,0x00,0x00,0x00,0x80,0x77,0x39,0x05
# GFX1250-REAL16: v_cvt_pk_f16_bf8_e64_dpp v1, v128.l dpp8:[7,6,5,4,3,2,1,0]
fi:1 ; encoding: [0x01,0x00,0xf6,0xd5,0xea,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
# GFX1250-FAKE16: v_cvt_pk_f16_bf8_e64_dpp v1, v128 dpp8:[7,6,5,4,3,2,1,0]
fi:1 ; encoding: [0x01,0x00,0xf6,0xd5,0xea,0x00,0x00,0x00,0x80,0x77,0x39,0x05]
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
