https://github.com/broxigarchen updated
https://github.com/llvm/llvm-project/pull/119750
>From 68ba9f46992b683dcff947e929e5050691acf0f1 Mon Sep 17 00:00:00 2001
From: guochen2 <guoch...@amd.com>
Date: Thu, 12 Dec 2024 13:33:14 -0500
Subject: [PATCH 1/4] True16 for v_alignbyte_b32 in MC
---
clang/lib/CodeGen/CGBuiltin.cpp | 8 ++++
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 2 +-
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 8 +++-
llvm/test/MC/AMDGPU/gfx11_asm_vop3.s | 11 +++--
llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp16.s | 42 +++++++++++++------
llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp8.s | 17 ++++++--
llvm/test/MC/AMDGPU/gfx12_asm_vop3.s | 3 ++
llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp16.s | 3 ++
llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp8.s | 3 ++
.../Disassembler/AMDGPU/gfx11_dasm_vop3.txt | 16 ++++++-
.../AMDGPU/gfx11_dasm_vop3_dpp16.txt | 31 +++++++++++---
.../AMDGPU/gfx11_dasm_vop3_dpp8.txt | 16 ++++++-
.../Disassembler/AMDGPU/gfx12_dasm_vop3.txt | 16 ++++++-
.../AMDGPU/gfx12_dasm_vop3_dpp16.txt | 36 +++++++++++++---
.../AMDGPU/gfx12_dasm_vop3_dpp8.txt | 21 ++++++++--
15 files changed, 190 insertions(+), 43 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 7ec9d59bfed5cf1..92e8ea9d23e12b1 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19792,6 +19792,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;
llvm::SyncScope::ID SSID;
switch (BuiltinID) {
+ case AMDGPU::BI__builtin_amdgcn_alignbyte: {
+ llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
+ llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
+ llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
+ llvm::Function *F =
+ CGM.getIntrinsic(Intrinsic::amdgcn_alignbyte, Src2->getType());
+ return Builder.CreateCall(F, {Src0, Src1, Src2});
+ }
case AMDGPU::BI__builtin_amdgcn_div_scale:
case AMDGPU::BI__builtin_amdgcn_div_scalef: {
// Translate from the intrinsics's struct return to the builtin's out
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index f721d5267cd2a08..fbbf9f74cc9c633 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2353,7 +2353,7 @@ def int_amdgcn_writelane :
>;
def int_amdgcn_alignbyte : ClangBuiltin<"__builtin_amdgcn_alignbyte">,
- DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty,
llvm_anyint_ty],
[IntrNoMem, IntrSpeculatable]
>;
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index c06c932a5375e81..a0d904ba9a3353e 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -212,7 +212,11 @@ defm V_BFE_U32 : VOP3Inst <"v_bfe_u32",
VOP3_Profile<VOP_I32_I32_I32_I32>, AMDGP
defm V_BFE_I32 : VOP3Inst <"v_bfe_i32", VOP3_Profile<VOP_I32_I32_I32_I32>,
AMDGPUbfe_i32>;
defm V_BFI_B32 : VOP3Inst <"v_bfi_b32", VOP3_Profile<VOP_I32_I32_I32_I32>,
AMDGPUbfi>;
defm V_ALIGNBIT_B32 : VOP3Inst <"v_alignbit_b32",
VOP3_Profile<VOP_I32_I32_I32_I32>, fshr>;
-defm V_ALIGNBYTE_B32 : VOP3Inst <"v_alignbyte_b32",
VOP3_Profile<VOP_I32_I32_I32_I32>, int_amdgcn_alignbyte>;
+defm V_ALIGNBYTE_B32 : VOP3Inst_t16_with_profiles <"v_alignbyte_b32",
+
VOP3_Profile<VOP_I32_I32_I32_I32>,
+
VOP3_Profile_True16<VOP_I32_I32_I32_I16, VOP3_OPSEL>,
+
VOP3_Profile_Fake16<VOP_I32_I32_I32_I16, VOP3_OPSEL>,
+ int_amdgcn_alignbyte>;
// XXX - No FPException seems suspect but manual doesn't say it does
let mayRaiseFPException = 0 in {
@@ -1690,7 +1694,7 @@ defm V_FMA_F32 :
VOP3_Realtriple_gfx11_gfx12<0x213>;
defm V_FMA_F64 : VOP3_Real_Base_gfx11_gfx12<0x214>;
defm V_LERP_U8 : VOP3_Realtriple_gfx11_gfx12<0x215>;
defm V_ALIGNBIT_B32 : VOP3_Realtriple_gfx11_gfx12<0x216>;
-defm V_ALIGNBYTE_B32 : VOP3_Realtriple_gfx11_gfx12<0x217>;
+defm V_ALIGNBYTE_B32 : VOP3_Realtriple_t16_and_fake16_gfx11_gfx12<0x217,
"v_alignbyte_b32">;
defm V_MULLIT_F32 : VOP3_Realtriple_gfx11_gfx12<0x218>;
defm V_MIN3_F32 : VOP3_Realtriple_gfx11<0x219>;
defm V_MIN3_I32 : VOP3_Realtriple_gfx11_gfx12<0x21a>;
diff --git a/llvm/test/MC/AMDGPU/gfx11_asm_vop3.s
b/llvm/test/MC/AMDGPU/gfx11_asm_vop3.s
index 40e3fbda47787a6..c166cb58d5cb52d 100644
--- a/llvm/test/MC/AMDGPU/gfx11_asm_vop3.s
+++ b/llvm/test/MC/AMDGPU/gfx11_asm_vop3.s
@@ -461,11 +461,11 @@ v_alignbyte_b32 v5, s1, v255, s3
v_alignbyte_b32 v5, s105, s105, s105
// GFX11: v_alignbyte_b32 v5, s105, s105, s105 ; encoding:
[0x05,0x00,0x17,0xd6,0x69,0xd2,0xa4,0x01]
-v_alignbyte_b32 v5, vcc_lo, ttmp15, v3
-// GFX11: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+v_alignbyte_b32 v5, vcc_lo, ttmp15, v3.l
+// GFX11: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
-v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255
-// GFX11: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.l
+// GFX11: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
v_alignbyte_b32 v5, ttmp15, src_scc, ttmp15
// GFX11: v_alignbyte_b32 v5, ttmp15, src_scc, ttmp15 ; encoding:
[0x05,0x00,0x17,0xd6,0x7b,0xfa,0xed,0x01]
@@ -494,6 +494,9 @@ v_alignbyte_b32 v5, src_scc, vcc_lo, -1
v_alignbyte_b32 v255, 0xaf123456, vcc_hi, null
// GFX11: v_alignbyte_b32 v255, 0xaf123456, vcc_hi, null ; encoding:
[0xff,0x00,0x17,0xd6,0xff,0xd6,0xf0,0x01,0x56,0x34,0x12,0xaf]
+v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h
+// GFX11: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h op_sel:[0,0,1,0] ;
encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+
v_and_b16 v5.l, v1.l, v2.l
// GFX11: v_and_b16 v5.l, v1.l, v2.l ; encoding:
[0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp16.s
b/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp16.s
index 2bff644605ff604..30df47aa0252f1e 100644
--- a/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp16.s
@@ -363,22 +363,22 @@ v_alignbit_b32_e64_dpp v5, v1, v2, -1 row_xmask:0
row_mask:0x1 bank_mask:0x3 bou
v_alignbit_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15 row_mask:0x3
bank_mask:0x0 bound_ctrl:0 fi:1
// GFX11: v_alignbit_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 fi:1 ; encoding:
[0xff,0x00,0x16,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x05,0x30]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[3,2,1,0]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[0,1,2,3]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_mirror row_mask:0xf bank_mask:0xf
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_half_mirror row_mask:0xf
bank_mask:0xf
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_shl:1 row_mask:0xf bank_mask:0xf
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
-v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:15
+v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:15 row_mask:0xf bank_mask:0xf
// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:15 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xa6,0x01,0x01,0x0f,0x01,0xff]
v_alignbyte_b32_e64_dpp v5, v1, v2, vcc_hi row_shr:1
@@ -387,7 +387,7 @@ v_alignbyte_b32_e64_dpp v5, v1, v2, vcc_hi row_shr:1
v_alignbyte_b32_e64_dpp v5, v1, v2, vcc_lo row_shr:15
// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, vcc_lo row_shr:15 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xaa,0x01,0x01,0x1f,0x01,0xff]
-v_alignbyte_b32_e64_dpp v5, v1, v2, ttmp15 row_ror:1
+v_alignbyte_b32_e64_dpp v5, v1, v2, ttmp15 row_ror:1 row_mask:0xf bank_mask:0xf
// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, ttmp15 row_ror:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xee,0x01,0x01,0x21,0x01,0xff]
v_alignbyte_b32_e64_dpp v5, v1, v2, exec_hi row_ror:15
@@ -405,6 +405,24 @@ v_alignbyte_b32_e64_dpp v5, v1, v2, -1 row_xmask:0
row_mask:0x1 bank_mask:0x3 bo
v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15 row_mask:0x3
bank_mask:0x0 bound_ctrl:0 fi:1
// GFX11: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 fi:1 ; encoding:
[0xff,0x00,0x17,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x05,0x30]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_mirror
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, s3 row_half_mirror
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s3 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x00,0x01,0x41,0x01,0xff]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:1
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xa6,0x01,0x01,0x01,0x01,0xff]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, ttmp15 row_shl:15
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, ttmp15 row_shl:15 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xee,0x01,0x01,0x0f,0x01,0xff]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, m0 row_ror:1
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, m0 row_ror:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xf6,0x01,0x01,0x21,0x01,0xff]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h row_mirror
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+
v_and_b16_e64_dpp v5.l, v1.l, v2.l quad_perm:[3,2,1,0]
// GFX11: v_and_b16_e64_dpp v5.l, v1.l, v2.l quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff]
diff --git a/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp8.s
b/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp8.s
index 2f9b5efca9e1728..9af925ae73e961d 100644
--- a/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp8.s
+++ b/llvm/test/MC/AMDGPU/gfx11_asm_vop3_dpp8.s
@@ -187,11 +187,11 @@ v_alignbit_b32_e64_dpp v5, v1, v2, -1
dpp8:[7,6,5,4,3,2,1,0] fi:1
v_alignbit_b32_e64_dpp v255, v255, v255, src_scc dpp8:[0,0,0,0,0,0,0,0] fi:0
// GFX11: v_alignbit_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xff,0x00,0x16,0xd6,0xe9,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0]
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l dpp8:[7,6,5,4,3,2,1,0]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
-v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0]
-// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l dpp8:[7,6,5,4,3,2,1,0]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
v_alignbyte_b32_e64_dpp v5, v1, v2, s105 dpp8:[7,6,5,4,3,2,1,0]
// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xa6,0x01,0x01,0x77,0x39,0x05]
@@ -220,6 +220,15 @@ v_alignbyte_b32_e64_dpp v5, v1, v2, -1
dpp8:[7,6,5,4,3,2,1,0] fi:1
v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc dpp8:[0,0,0,0,0,0,0,0] fi:0
// GFX11: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xff,0x00,0x17,0xd6,0xe9,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
+v_alignbyte_b32_e64_dpp v5, v1, v2, s3 dpp8:[7,6,5,4,3,2,1,0]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x00,0x01,0x77,0x39,0x05]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, m0 dpp8:[7,6,5,4,3,2,1,0]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, m0 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xf6,0x01,0x01,0x77,0x39,0x05]
+
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h dpp8:[7,6,5,4,3,2,1,0]
+// GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+
v_and_b16_e64_dpp v5.l, v1.l, v2.l dpp8:[7,6,5,4,3,2,1,0]
// GFX11: v_and_b16_e64_dpp v5.l, v1.l, v2.l dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05]
diff --git a/llvm/test/MC/AMDGPU/gfx12_asm_vop3.s
b/llvm/test/MC/AMDGPU/gfx12_asm_vop3.s
index ff24cbbbd2bea0b..e199ae73ffc54a6 100644
--- a/llvm/test/MC/AMDGPU/gfx12_asm_vop3.s
+++ b/llvm/test/MC/AMDGPU/gfx12_asm_vop3.s
@@ -452,6 +452,9 @@ v_alignbyte_b32 v5, src_scc, vcc_lo, -1
v_alignbyte_b32 v255, 0xaf123456, vcc_hi, null
// GFX12: v_alignbyte_b32 v255, 0xaf123456, vcc_hi, null ; encoding:
[0xff,0x00,0x17,0xd6,0xff,0xd6,0xf0,0x01,0x56,0x34,0x12,0xaf]
+v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h
+// GFX12: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h op_sel:[0,0,1,0] ;
encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+
v_and_b16 v5.l, v1.l, v2.l
// GFX12: v_and_b16 v5.l, v1.l, v2.l ; encoding:
[0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00]
diff --git a/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp16.s
b/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp16.s
index 64b21385e0ed242..89ff2ced2b131c0 100644
--- a/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp16.s
+++ b/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp16.s
@@ -485,6 +485,9 @@ v_alignbyte_b32_e64_dpp v5, v1, v2, -1 row_xmask:0
row_mask:0x1 bank_mask:0x3 bo
v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15 row_mask:0x3
bank_mask:0x0 bound_ctrl:0 fi:1
// GFX12: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 fi:1 ; encoding:
[0xff,0x00,0x17,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x05,0x30]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h row_mirror
+// GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+
v_and_b16_e64_dpp v5.l, v1.l, v2.l quad_perm:[3,2,1,0]
// GFX12: v_and_b16_e64_dpp v5.l, v1.l, v2.l quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff]
diff --git a/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp8.s
b/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp8.s
index c910e2020ee123b..738ad59ca6bfe18 100644
--- a/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp8.s
+++ b/llvm/test/MC/AMDGPU/gfx12_asm_vop3_dpp8.s
@@ -288,6 +288,9 @@ v_alignbyte_b32_e64_dpp v5, v1, v2, -1
dpp8:[7,6,5,4,3,2,1,0] fi:1
v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc dpp8:[0,0,0,0,0,0,0,0] fi:0
// GFX12: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] ; encoding:
[0xff,0x00,0x17,0xd6,0xe9,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
+v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h dpp8:[7,6,5,4,3,2,1,0]
+// GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+
v_and_b16_e64_dpp v5.l, v1.l, v2.l dpp8:[7,6,5,4,3,2,1,0]
// GFX12: v_and_b16_e64_dpp v5.l, v1.l, v2.l dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3.txt
index 05174e312891997..380d8a0d622ddc9 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3.txt
@@ -508,10 +508,16 @@
# GFX11: v_alignbyte_b32 v5, s105, s105, s105 ; encoding:
[0x05,0x00,0x17,0xd6,0x69,0xd2,0xa4,0x01]
0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04
-# GFX11: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W32-REAL16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W32-FAKE16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W64-REAL16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W64-FAKE16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf
-# GFX11: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W32-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W32-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
0x05,0x00,0x17,0xd6,0x7b,0xfa,0xed,0x01
# GFX11: v_alignbyte_b32 v5, ttmp15, src_scc, ttmp15 ; encoding:
[0x05,0x00,0x17,0xd6,0x7b,0xfa,0xed,0x01]
@@ -540,6 +546,12 @@
0xff,0x00,0x17,0xd6,0xff,0xd6,0xf0,0x01,0x56,0x34,0x12,0xaf
# GFX11: v_alignbyte_b32 v255, 0xaf123456, vcc_hi, null ; encoding:
[0xff,0x00,0x17,0xd6,0xff,0xd6,0xf0,0x01,0x56,0x34,0x12,0xaf]
+0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf
+# W32-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h op_sel:[0,0,1,0]
; encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W32-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 op_sel:[0,0,1,0] ;
encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h op_sel:[0,0,1,0]
; encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 op_sel:[0,0,1,0] ;
encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+
0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00
# W32-REAL16: v_and_b16 v5.l, v1.l, v2.l ; encoding:
[0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00]
# W32-FAKE16: v_and_b16 v5, v1, v2 ; encoding:
[0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp16.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp16.txt
index c9ef3c714213d1a..e4f2ee27ac9e06c 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp16.txt
@@ -228,19 +228,34 @@
# GFX11: v_alignbit_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 bound_ctrl:1 fi:1 ; encoding:
[0xff,0x00,0x16,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x0d,0x30]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_shl:1
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_shl:1
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0xa6,0x01,0x01,0x0f,0x01,0xff
# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:15 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xa6,0x01,0x01,0x0f,0x01,0xff]
@@ -269,6 +284,12 @@
0xff,0x00,0x17,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x0d,0x30
# GFX11: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 bound_ctrl:1 fi:1 ; encoding:
[0xff,0x00,0x17,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x0d,0x30]
+0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+
0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff
# W32-REAL16: v_and_b16_e64_dpp v5.l, v1.l, v2.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff]
# W32-FAKE16: v_and_b16_e64_dpp v5, v1, v2 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp8.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp8.txt
index 1e74b5aec0cf31f..337e3197d66f89d 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp8.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx11_dasm_vop3_dpp8.txt
@@ -135,10 +135,16 @@
# GFX11: v_alignbit_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] fi:1 ; encoding:
[0xff,0x00,0x16,0xd6,0xea,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05
-# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
0x05,0x00,0x17,0xd6,0xe9,0x04,0xa6,0x01,0x01,0x77,0x39,0x05
# GFX11: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xa6,0x01,0x01,0x77,0x39,0x05]
@@ -167,6 +173,12 @@
0xff,0x00,0x17,0xd6,0xea,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00
# GFX11: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] fi:1 ; encoding:
[0xff,0x00,0x17,0xd6,0xea,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
+0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+
0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05
# W32-REAL16: v_and_b16_e64_dpp v5.l, v1.l, v2.l dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05]
# W32-FAKE16: v_and_b16_e64_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3.txt
index 4fa7dd54205713a..261d794af6a9bc8 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3.txt
@@ -472,10 +472,16 @@
# GFX12: v_alignbyte_b32 v5, s105, s105, s105 ; encoding:
[0x05,0x00,0x17,0xd6,0x69,0xd2,0xa4,0x01]
0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04
-# GFX12: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W32-REAL16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W32-FAKE16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W64-REAL16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
+# W64-FAKE16: v_alignbyte_b32 v5, vcc_lo, ttmp15, v3 ; encoding:
[0x05,0x00,0x17,0xd6,0x6a,0xf6,0x0c,0x04]
0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf
-# GFX12: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W32-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W32-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.l ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 ; encoding:
[0x05,0x00,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
0x05,0x00,0x17,0xd6,0x7b,0xfa,0xed,0x01
# GFX12: v_alignbyte_b32 v5, ttmp15, src_scc, ttmp15 ; encoding:
[0x05,0x00,0x17,0xd6,0x7b,0xfa,0xed,0x01]
@@ -504,6 +510,12 @@
0xff,0x00,0x17,0xd6,0xff,0xd6,0xf0,0x01,0x56,0x34,0x12,0xaf
# GFX12: v_alignbyte_b32 v255, 0xaf123456, vcc_hi, null ; encoding:
[0xff,0x00,0x17,0xd6,0xff,0xd6,0xf0,0x01,0x56,0x34,0x12,0xaf]
+0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf
+# W32-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h op_sel:[0,0,1,0]
; encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W32-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 op_sel:[0,0,1,0] ;
encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-REAL16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255.h op_sel:[0,0,1,0]
; encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+# W64-FAKE16: v_alignbyte_b32 v5, vcc_hi, 0xaf123456, v255 op_sel:[0,0,1,0] ;
encoding: [0x05,0x20,0x17,0xd6,0x6b,0xfe,0xfd,0x07,0x56,0x34,0x12,0xaf]
+
0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00
# W32-REAL16: v_and_b16 v5.l, v1.l, v2.l ; encoding:
[0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00]
# W32-FAKE16: v_and_b16 v5, v1, v2 ; encoding:
[0x05,0x00,0x62,0xd7,0x01,0x05,0x02,0x00]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp16.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp16.txt
index e44808607a2b3f2..e69daefeb4d2220 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp16.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp16.txt
@@ -240,22 +240,40 @@
# GFX12: v_alignbit_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 bound_ctrl:1 fi:1 ; encoding:
[0xff,0x00,0x16,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x0d,0x30]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x1b,0x00,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x06,0x0c,0x04,0x01,0x1b,0x00,0xff
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, s3, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x06,0x0c,0x04,0x01,0x1b,0x00,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x06,0x0c,0x04,0x01,0x1b,0x00,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x06,0x0c,0x04,0x01,0x1b,0x00,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x06,0x0c,0x04,0x01,0x1b,0x00,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3 quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x06,0x0c,0x04,0x01,0x1b,0x00,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 quad_perm:[0,1,2,3]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0xe4,0x00,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x40,0x01,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 row_half_mirror
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0x0e,0x04,0x01,0x41,0x01,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_shl:1
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l row_shl:1
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 row_shl:1 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x01,0x01,0xff]
0x05,0x00,0x17,0xd6,0xfa,0x04,0xa6,0x01,0x01,0x0f,0x01,0xff
# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 row_shl:15 row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x17,0xd6,0xfa,0x04,0xa6,0x01,0x01,0x0f,0x01,0xff]
@@ -284,6 +302,12 @@
0xff,0x00,0x17,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x0d,0x30
# GFX12: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc row_xmask:15
row_mask:0x3 bank_mask:0x0 bound_ctrl:1 fi:1 ; encoding:
[0xff,0x00,0x17,0xd6,0xfa,0xfe,0xf7,0x03,0xff,0x6f,0x0d,0x30]
+0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
row_mirror row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x20,0x17,0xd6,0xfa,0x04,0xfe,0x07,0x01,0x40,0x01,0xff]
+
0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff
# W32-REAL16: v_and_b16_e64_dpp v5.l, v1.l, v2.l quad_perm:[3,2,1,0]
row_mask:0xf bank_mask:0xf ; encoding:
[0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff]
# W32-FAKE16: v_and_b16_e64_dpp v5, v1, v2 quad_perm:[3,2,1,0] row_mask:0xf
bank_mask:0xf ; encoding:
[0x05,0x00,0x62,0xd7,0xfa,0x04,0x02,0x00,0x01,0x1b,0x00,0xff]
diff --git a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp8.txt
b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp8.txt
index 6fbdc85da43c4a9..4c39e09a1805bca 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp8.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/gfx12_dasm_vop3_dpp8.txt
@@ -147,13 +147,22 @@
# GFX12: v_alignbit_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] fi:1 ; encoding:
[0xff,0x00,0x16,0xd6,0xea,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3.l dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0x0e,0x04,0x01,0x77,0x39,0x05]
0x05,0x00,0x17,0xd6,0xe9,0x06,0x0c,0x04,0x01,0x77,0x39,0x05
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, s3, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x06,0x0c,0x04,0x01,0x77,0x39,0x05]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3.l dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x06,0x0c,0x04,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x06,0x0c,0x04,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3.l dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x06,0x0c,0x04,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, s3, v3 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x06,0x0c,0x04,0x01,0x77,0x39,0x05]
0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05
-# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.l
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 dpp8:[7,6,5,4,3,2,1,0]
; encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
0x05,0x00,0x17,0xd6,0xe9,0x04,0xa6,0x01,0x01,0x77,0x39,0x05
# GFX12: v_alignbyte_b32_e64_dpp v5, v1, v2, s105 dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x17,0xd6,0xe9,0x04,0xa6,0x01,0x01,0x77,0x39,0x05]
@@ -182,6 +191,12 @@
0xff,0x00,0x17,0xd6,0xea,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00
# GFX12: v_alignbyte_b32_e64_dpp v255, v255, v255, src_scc
dpp8:[0,0,0,0,0,0,0,0] fi:1 ; encoding:
[0xff,0x00,0x17,0xd6,0xea,0xfe,0xf7,0x03,0xff,0x00,0x00,0x00]
+0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05
+# W32-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W32-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-REAL16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255.h op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+# W64-FAKE16: v_alignbyte_b32_e64_dpp v5, v1, v2, v255 op_sel:[0,0,1,0]
dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x20,0x17,0xd6,0xe9,0x04,0xfe,0x07,0x01,0x77,0x39,0x05]
+
0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05
# W32-REAL16: v_and_b16_e64_dpp v5.l, v1.l, v2.l dpp8:[7,6,5,4,3,2,1,0] ;
encoding: [0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05]
# W32-FAKE16: v_and_b16_e64_dpp v5, v1, v2 dpp8:[7,6,5,4,3,2,1,0] ; encoding:
[0x05,0x00,0x62,0xd7,0xe9,0x04,0x02,0x00,0x01,0x77,0x39,0x05]
>From 6e51adf7a82fc3c483a9f3c9c1a8d3f5442c0e8e Mon Sep 17 00:00:00 2001
From: guochen2 <guoch...@amd.com>
Date: Fri, 13 Dec 2024 13:26:30 -0500
Subject: [PATCH 2/4] remove clang change
---
clang/lib/CodeGen/CGBuiltin.cpp | 16 ++++++++--------
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 2 +-
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 4 ++--
llvm/unittests/IR/IntrinsicsTest.cpp | 1 -
4 files changed, 11 insertions(+), 12 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 92e8ea9d23e12b1..2076f6be7e9c56d 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -19792,14 +19792,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
llvm::AtomicOrdering AO = llvm::AtomicOrdering::SequentiallyConsistent;
llvm::SyncScope::ID SSID;
switch (BuiltinID) {
- case AMDGPU::BI__builtin_amdgcn_alignbyte: {
- llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
- llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
- llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
- llvm::Function *F =
- CGM.getIntrinsic(Intrinsic::amdgcn_alignbyte, Src2->getType());
- return Builder.CreateCall(F, {Src0, Src1, Src2});
- }
case AMDGPU::BI__builtin_amdgcn_div_scale:
case AMDGPU::BI__builtin_amdgcn_div_scalef: {
// Translate from the intrinsics's struct return to the builtin's out
@@ -20426,6 +20418,14 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());
return Builder.CreateCall(F, { Src0, Src1, Src2 });
}
+ case AMDGPU::BI__builtin_amdgcn_alignbyte: {
+ llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
+ llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
+ llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
+ llvm::Function *F =
+ CGM.getIntrinsic(Intrinsic::amdgcn_alignbyte, Src2->getType());
+ return Builder.CreateCall(F, {Src0, Src1, Src2});
+ }
case AMDGPU::BI__builtin_amdgcn_fence: {
ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),
EmitScalarExpr(E->getArg(1)), AO, SSID);
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index ded5f6b5ac4fd3c..db816da45b8b35f 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -734,7 +734,7 @@ kernel void test_alignbit(global uint* out, uint src0, uint
src1, uint src2) {
}
// CHECK-LABEL: @test_alignbyte(
-// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte(i32 %src0, i32 %src1, i32
%src2)
+// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte.i32(i32 %src0, i32 %src1,
i32 %src2)
kernel void test_alignbyte(global uint* out, uint src0, uint src1, uint src2) {
*out = __builtin_amdgcn_alignbyte(src0, src1, src2);
}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index fbbf9f74cc9c633..b80151c72927020 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2352,8 +2352,8 @@ def int_amdgcn_writelane :
[IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]
>;
-def int_amdgcn_alignbyte : ClangBuiltin<"__builtin_amdgcn_alignbyte">,
- DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty,
llvm_anyint_ty],
+def int_amdgcn_alignbyte : DefaultAttrsIntrinsic<[llvm_i32_ty],
+ [llvm_i32_ty, llvm_i32_ty, llvm_anyint_ty],
[IntrNoMem, IntrSpeculatable]
>;
diff --git a/llvm/unittests/IR/IntrinsicsTest.cpp
b/llvm/unittests/IR/IntrinsicsTest.cpp
index ad9af2075e371d4..a5b0360198362a3 100644
--- a/llvm/unittests/IR/IntrinsicsTest.cpp
+++ b/llvm/unittests/IR/IntrinsicsTest.cpp
@@ -87,7 +87,6 @@ TEST(IntrinsicNameLookup, ClangBuiltinLookup) {
{"__builtin_adjust_trampoline", "", adjust_trampoline},
{"__builtin_trap", "", trap},
{"__builtin_arm_chkfeat", "aarch64", aarch64_chkfeat},
- {"__builtin_amdgcn_alignbyte", "amdgcn", amdgcn_alignbyte},
{"__builtin_amdgcn_workgroup_id_z", "amdgcn", amdgcn_workgroup_id_z},
{"__builtin_arm_cdp", "arm", arm_cdp},
{"__builtin_bpf_preserve_type_info", "bpf", bpf_preserve_type_info},
>From 051d98e218cba53cc4487d85e07362395cfe54c3 Mon Sep 17 00:00:00 2001
From: guochen2 <guoch...@amd.com>
Date: Wed, 29 Jan 2025 10:56:42 -0500
Subject: [PATCH 3/4] revert instrinsic change
---
clang/lib/CodeGen/CGBuiltin.cpp | 8 --------
clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 2 +-
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 4 ++--
llvm/unittests/IR/IntrinsicsTest.cpp | 1 +
4 files changed, 4 insertions(+), 11 deletions(-)
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 2076f6be7e9c56d..7ec9d59bfed5cf1 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -20418,14 +20418,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned
BuiltinID,
Function *F = CGM.getIntrinsic(Intrinsic::fshr, Src0->getType());
return Builder.CreateCall(F, { Src0, Src1, Src2 });
}
- case AMDGPU::BI__builtin_amdgcn_alignbyte: {
- llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
- llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
- llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
- llvm::Function *F =
- CGM.getIntrinsic(Intrinsic::amdgcn_alignbyte, Src2->getType());
- return Builder.CreateCall(F, {Src0, Src1, Src2});
- }
case AMDGPU::BI__builtin_amdgcn_fence: {
ProcessOrderScopeAMDGCN(EmitScalarExpr(E->getArg(0)),
EmitScalarExpr(E->getArg(1)), AO, SSID);
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index db816da45b8b35f..ded5f6b5ac4fd3c 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -734,7 +734,7 @@ kernel void test_alignbit(global uint* out, uint src0, uint
src1, uint src2) {
}
// CHECK-LABEL: @test_alignbyte(
-// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte.i32(i32 %src0, i32 %src1,
i32 %src2)
+// CHECK: tail call{{.*}} i32 @llvm.amdgcn.alignbyte(i32 %src0, i32 %src1, i32
%src2)
kernel void test_alignbyte(global uint* out, uint src0, uint src1, uint src2) {
*out = __builtin_amdgcn_alignbyte(src0, src1, src2);
}
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index b80151c72927020..f721d5267cd2a08 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2352,8 +2352,8 @@ def int_amdgcn_writelane :
[IntrNoMem, IntrConvergent, IntrWillReturn, IntrNoCallback, IntrNoFree]
>;
-def int_amdgcn_alignbyte : DefaultAttrsIntrinsic<[llvm_i32_ty],
- [llvm_i32_ty, llvm_i32_ty, llvm_anyint_ty],
+def int_amdgcn_alignbyte : ClangBuiltin<"__builtin_amdgcn_alignbyte">,
+ DefaultAttrsIntrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
[IntrNoMem, IntrSpeculatable]
>;
diff --git a/llvm/unittests/IR/IntrinsicsTest.cpp
b/llvm/unittests/IR/IntrinsicsTest.cpp
index a5b0360198362a3..ad9af2075e371d4 100644
--- a/llvm/unittests/IR/IntrinsicsTest.cpp
+++ b/llvm/unittests/IR/IntrinsicsTest.cpp
@@ -87,6 +87,7 @@ TEST(IntrinsicNameLookup, ClangBuiltinLookup) {
{"__builtin_adjust_trampoline", "", adjust_trampoline},
{"__builtin_trap", "", trap},
{"__builtin_arm_chkfeat", "aarch64", aarch64_chkfeat},
+ {"__builtin_amdgcn_alignbyte", "amdgcn", amdgcn_alignbyte},
{"__builtin_amdgcn_workgroup_id_z", "amdgcn", amdgcn_workgroup_id_z},
{"__builtin_arm_cdp", "arm", arm_cdp},
{"__builtin_bpf_preserve_type_info", "bpf", bpf_preserve_type_info},
>From 3681a8c3ad4e6dc335705265faf6c8d698c9b7d9 Mon Sep 17 00:00:00 2001
From: guochen2 <guoch...@amd.com>
Date: Wed, 29 Jan 2025 12:46:08 -0500
Subject: [PATCH 4/4] revert intrinsic change, update pattern
---
llvm/lib/Target/AMDGPU/SIInstructions.td | 20 +++++++++
llvm/lib/Target/AMDGPU/VOP3Instructions.td | 12 +++---
.../CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll | 42 ++++++++++++++++++-
3 files changed, 67 insertions(+), 7 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td
b/llvm/lib/Target/AMDGPU/SIInstructions.td
index aa81d9b7e22a7e8..41614e023fae99f 100644
--- a/llvm/lib/Target/AMDGPU/SIInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SIInstructions.td
@@ -3802,6 +3802,26 @@ def : FPMinCanonMaxPat<V_MINIMUMMAXIMUM_F16_fake16_e64,
f16, DivergentBinFrag<fm
def : FPMinCanonMaxPat<V_MAXIMUMMINIMUM_F16_fake16_e64, f16,
DivergentBinFrag<fminimum>, fmaximum_oneuse>;
}
+let True16Predicate = UseFakeTrue16Insts in
+def : GCNPat <
+(i32 (int_amdgcn_alignbyte (i32 (VOP3OpSelMods i32:$src0,
i32:$src0_modifiers)),
+ (i32 (VOP3OpSelMods i32:$src1,
i32:$src1_modifiers)),
+ (i32 (VOP3OpSelMods i32:$src2,
i32:$src2_modifiers)))),
+(V_ALIGNBYTE_B32_fake16_e64 i32:$src0_modifiers, VSrc_b32:$src0,
+ i32:$src1_modifiers, VSrc_b32:$src1,
+ i32:$src2_modifiers, VGPR_32:$src2)
+>;
+
+let True16Predicate = UseRealTrue16Insts in
+def : GCNPat <
+(i32 (int_amdgcn_alignbyte (i32 (VOP3OpSelMods i32:$src0,
i32:$src0_modifiers)),
+ (i32 (VOP3OpSelMods i32:$src1,
i32:$src1_modifiers)),
+ (i32 (VOP3OpSelMods i32:$src2,
i32:$src2_modifiers)))),
+(V_ALIGNBYTE_B32_t16_e64 i32:$src0_modifiers, VSrc_b32:$src0,
+ i32:$src1_modifiers, VSrc_b32:$src1,
+ i32:$src2_modifiers, (i16 (EXTRACT_SUBREG
VGPR_32:$src2, lo16)))
+>;
+
// Convert a floating-point power of 2 to the integer exponent.
def FPPow2ToExponentXForm : SDNodeXForm<fpimm, [{
const auto &APF = N->getValueAPF();
diff --git a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
index a0d904ba9a3353e..ab4baae69140a05 100644
--- a/llvm/lib/Target/AMDGPU/VOP3Instructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3Instructions.td
@@ -212,11 +212,13 @@ defm V_BFE_U32 : VOP3Inst <"v_bfe_u32",
VOP3_Profile<VOP_I32_I32_I32_I32>, AMDGP
defm V_BFE_I32 : VOP3Inst <"v_bfe_i32", VOP3_Profile<VOP_I32_I32_I32_I32>,
AMDGPUbfe_i32>;
defm V_BFI_B32 : VOP3Inst <"v_bfi_b32", VOP3_Profile<VOP_I32_I32_I32_I32>,
AMDGPUbfi>;
defm V_ALIGNBIT_B32 : VOP3Inst <"v_alignbit_b32",
VOP3_Profile<VOP_I32_I32_I32_I32>, fshr>;
-defm V_ALIGNBYTE_B32 : VOP3Inst_t16_with_profiles <"v_alignbyte_b32",
-
VOP3_Profile<VOP_I32_I32_I32_I32>,
-
VOP3_Profile_True16<VOP_I32_I32_I32_I16, VOP3_OPSEL>,
-
VOP3_Profile_Fake16<VOP_I32_I32_I32_I16, VOP3_OPSEL>,
- int_amdgcn_alignbyte>;
+
+let True16Predicate = NotHasTrue16BitInsts in
+defm V_ALIGNBYTE_B32 : VOP3Inst <"v_alignbyte_b32",
VOP3_Profile<VOP_I32_I32_I32_I32>, int_amdgcn_alignbyte>;
+let True16Predicate = UseRealTrue16Insts in
+defm V_ALIGNBYTE_B32_t16 : VOP3Inst <"v_alignbyte_b32_t16",
VOP3_Profile_True16<VOP_I32_I32_I32_I16, VOP3_OPSEL>>;
+let True16Predicate = UseFakeTrue16Insts in
+defm V_ALIGNBYTE_B32_fake16 : VOP3Inst <"v_alignbyte_b32_fake16",
VOP3_Profile_Fake16<VOP_I32_I32_I32_I16, VOP3_OPSEL>>;
// XXX - No FPException seems suspect but manual doesn't say it does
let mayRaiseFPException = 0 in {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
index 8b16fef915a79d6..15e99be5bdfe96a 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.alignbyte.ll
@@ -1,10 +1,48 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
UTC_ARGS: --version 5
; RUN: llc -mtriple=amdgcn -verify-machineinstrs < %s | FileCheck
-check-prefix=GCN %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -mattr=+real-true16
-verify-machineinstrs < %s | FileCheck -check-prefixes=GFX11-TRUE16 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -mattr=-real-true16
-verify-machineinstrs < %s | FileCheck -check-prefixes=GFX11-FAKE16 %s
declare i32 @llvm.amdgcn.alignbyte(i32, i32, i32) #0
-; GCN-LABEL: {{^}}v_alignbyte_b32:
-; GCN: v_alignbyte_b32 {{[vs][0-9]+}}, {{[vs][0-9]+}}, {{[vs][0-9]+}}
define amdgpu_kernel void @v_alignbyte_b32(ptr addrspace(1) %out, i32 %src1,
i32 %src2, i32 %src3) #1 {
+; GCN-LABEL: v_alignbyte_b32:
+; GCN: ; %bb.0:
+; GCN-NEXT: s_load_dwordx4 s[0:3], s[4:5], 0xb
+; GCN-NEXT: s_load_dwordx2 s[4:5], s[4:5], 0x9
+; GCN-NEXT: s_mov_b32 s7, 0xf000
+; GCN-NEXT: s_mov_b32 s6, -1
+; GCN-NEXT: s_waitcnt lgkmcnt(0)
+; GCN-NEXT: v_mov_b32_e32 v0, s1
+; GCN-NEXT: v_mov_b32_e32 v1, s2
+; GCN-NEXT: v_alignbyte_b32 v0, s0, v0, v1
+; GCN-NEXT: buffer_store_dword v0, off, s[4:7], 0
+; GCN-NEXT: s_endpgm
+;
+; GFX11-TRUE16-LABEL: v_alignbyte_b32:
+; GFX11-TRUE16: ; %bb.0:
+; GFX11-TRUE16-NEXT: s_clause 0x1
+; GFX11-TRUE16-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c
+; GFX11-TRUE16-NEXT: s_load_b64 s[4:5], s[4:5], 0x24
+; GFX11-TRUE16-NEXT: v_mov_b32_e32 v1, 0
+; GFX11-TRUE16-NEXT: s_waitcnt lgkmcnt(0)
+; GFX11-TRUE16-NEXT: v_mov_b16_e32 v0.l, s2
+; GFX11-TRUE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX11-TRUE16-NEXT: v_alignbyte_b32 v0, s0, s1, v0.l
+; GFX11-TRUE16-NEXT: global_store_b32 v1, v0, s[4:5]
+; GFX11-TRUE16-NEXT: s_endpgm
+;
+; GFX11-FAKE16-LABEL: v_alignbyte_b32:
+; GFX11-FAKE16: ; %bb.0:
+; GFX11-FAKE16-NEXT: s_clause 0x1
+; GFX11-FAKE16-NEXT: s_load_b128 s[0:3], s[4:5], 0x2c
+; GFX11-FAKE16-NEXT: s_load_b64 s[4:5], s[4:5], 0x24
+; GFX11-FAKE16-NEXT: s_waitcnt lgkmcnt(0)
+; GFX11-FAKE16-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_mov_b32 v0, s2
+; GFX11-FAKE16-NEXT: s_delay_alu instid0(VALU_DEP_1)
+; GFX11-FAKE16-NEXT: v_alignbyte_b32 v0, s0, s1, v0
+; GFX11-FAKE16-NEXT: global_store_b32 v1, v0, s[4:5]
+; GFX11-FAKE16-NEXT: s_endpgm
%val = call i32 @llvm.amdgcn.alignbyte(i32 %src1, i32 %src2, i32 %src3) #0
store i32 %val, ptr addrspace(1) %out
ret void
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
