[llvm] 544ef42 - [AMDGPU] Set default op_sel_hi on accvgpr read/write
Stanislav Mekhanoshin via llvm-commits
llvm-commits at lists.llvm.org
Tue Nov 10 13:07:51 PST 2020
Author: Stanislav Mekhanoshin
Date: 2020-11-10T13:07:29-08:00
New Revision: 544ef42e40aa87d076595c275edabf13951aadf5
URL: https://github.com/llvm/llvm-project/commit/544ef42e40aa87d076595c275edabf13951aadf5
DIFF: https://github.com/llvm/llvm-project/commit/544ef42e40aa87d076595c275edabf13951aadf5.diff
LOG: [AMDGPU] Set default op_sel_hi on accvgpr read/write
These are opsel opcodes with op_sel actually being ignored.
As a such op_sel_hi needs to be set to default 1 even though
these bits are ignored. This is compatibility change.
Differential Revision: https://reviews.llvm.org/D91202
Added:
Modified:
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
llvm/test/MC/AMDGPU/accvgpr-altnames.s
llvm/test/MC/AMDGPU/mai.s
llvm/test/MC/Disassembler/AMDGPU/mai.txt
Removed:
################################################################################
diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
index 159731707eab..2a9992087ca9 100644
--- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -427,6 +427,17 @@ multiclass VOP3P_Real_vi<bits<7> op> {
}
multiclass VOP3P_Real_MAI<bits<7> op> {
+ def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>,
+ VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME).Pfl> {
+ let AssemblerPredicate = HasMAIInsts;
+ let DecoderNamespace = "GFX8";
+ let Inst{14} = 1; // op_sel_hi(2) default value
+ let Inst{59} = 1; // op_sel_hi(0) default value
+ let Inst{60} = 1; // op_sel_hi(1) default value
+ }
+}
+
+multiclass VOP3P_Real_MFMA<bits<7> op> {
def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>,
VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME).Pfl> {
let AssemblerPredicate = HasMAIInsts;
@@ -495,26 +506,26 @@ let SubtargetPredicate = HasMAIInsts in {
defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>;
defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>;
-defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MAI <0x40>;
-defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MAI <0x41>;
-defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MAI <0x42>;
-defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MAI <0x44>;
-defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MAI <0x45>;
-defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MAI <0x48>;
-defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MAI <0x49>;
-defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MAI <0x4a>;
-defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MAI <0x4c>;
-defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MAI <0x4d>;
-defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MAI <0x50>;
-defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MAI <0x51>;
-defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MAI <0x52>;
-defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MAI <0x54>;
-defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MAI <0x55>;
-defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MAI <0x68>;
-defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MAI <0x69>;
-defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MAI <0x6b>;
-defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MAI <0x6c>;
-defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MAI <0x6d>;
+defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40>;
+defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41>;
+defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42>;
+defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44>;
+defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45>;
+defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48>;
+defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49>;
+defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a>;
+defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c>;
+defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d>;
+defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50>;
+defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51>;
+defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52>;
+defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA <0x55>;
+defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA <0x54>;
+defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA <0x68>;
+defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA <0x69>;
+defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA <0x6b>;
+defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA <0x6c>;
+defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA <0x6d>;
} // End SubtargetPredicate = HasMAIInsts
diff --git a/llvm/test/MC/AMDGPU/accvgpr-altnames.s b/llvm/test/MC/AMDGPU/accvgpr-altnames.s
index fda18aad8906..cc52162480a5 100644
--- a/llvm/test/MC/AMDGPU/accvgpr-altnames.s
+++ b/llvm/test/MC/AMDGPU/accvgpr-altnames.s
@@ -1,10 +1,10 @@
// RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s
v_accvgpr_read_b32 v2, acc0
-// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08]
+// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
v_accvgpr_write_b32 acc2, -2.0
-// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00]
+// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
v_mfma_f32_32x32x1f32 acc[0:31], acc0, acc1, acc[1:32]
// GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x1c]
diff --git a/llvm/test/MC/AMDGPU/mai.s b/llvm/test/MC/AMDGPU/mai.s
index c02139a616fc..5629df067590 100644
--- a/llvm/test/MC/AMDGPU/mai.s
+++ b/llvm/test/MC/AMDGPU/mai.s
@@ -2,28 +2,28 @@
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefix=NOGFX908 --implicit-check-not=error: %s
v_accvgpr_read_b32 v2, a0
-// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08]
+// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
v_accvgpr_read_b32 v2, a1
-// GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08]
+// GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]
v_accvgpr_read_b32 v2, a255
-// GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08]
+// GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18]
v_accvgpr_read v2, a10
-// GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x00,0xd8,0xd3,0x0a,0x01,0x00,0x08]
+// GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x40,0xd8,0xd3,0x0a,0x01,0x00,0x18]
v_accvgpr_write_b32 a2, -2.0
-// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00]
+// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
v_accvgpr_write_b32 a2, -2
-// GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00]
+// GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18]
v_accvgpr_write_b32 a2, v1
-// GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00]
+// GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
v_accvgpr_write a2, v255
-// GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x00,0xd9,0xd3,0xff,0x01,0x00,0x00]
+// GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x40,0xd9,0xd3,0xff,0x01,0x00,0x18]
v_accvgpr_write a2, 100
// NOGFX908: error: invalid operand for instruction
diff --git a/llvm/test/MC/Disassembler/AMDGPU/mai.txt b/llvm/test/MC/Disassembler/AMDGPU/mai.txt
index ef2acaf2c13d..f279bde02d34 100644
--- a/llvm/test/MC/Disassembler/AMDGPU/mai.txt
+++ b/llvm/test/MC/Disassembler/AMDGPU/mai.txt
@@ -1,22 +1,22 @@
# RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding -disassemble %s | FileCheck -check-prefix=GFX908 %s
-# GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08]
-0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08
+# GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
+0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18
-# GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08]
-0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08
+# GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]
+0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18
-# GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08]
-0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08
+# GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18]
+0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18
-# GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00]
-0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00
+# GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
+0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18
-# GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00]
-0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00
+# GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18]
+0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18
-# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00]
-0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00
+# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
+0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18
# GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04]
0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04
More information about the llvm-commits
mailing list