caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[AMDGPU] Allow accvgpr_read/write decode with opsel

Browse Source 
These two instructions are VOP3P and have op_sel_hi bits,
however do not use op_sel_hi. That is recommended to set
unused op_sel_hi bits to 1. However, we cannot decode
both representations with 1 and 0 if bits are set to
default value 1. If bits are set to be ignored with '?'
initializer then encoding defaults them to 0.

The patch is a hack to force ignored '?' bits to 1 on
encoding for these instructions.

There is still canonicalization happens on disasm print
if incoming values are non-default, so that disasm output
does not match binary input, but this is pre-existing
problem for all instructions with '?' bits.

Fixes: SWDEV-272540

Differential Revision: https://reviews.llvm.org/D96543
This commit is contained in:
Stanislav Mekhanoshin
2021-02-11 13:31:35 -08:00
parent 6577cef9b0
commit c0d7a8bc62
3 changed files with 35 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp
View File

@@ -284,6 +284,20 @@ void SIMCCodeEmitter::encodeInstruction(const MCInst &MI, raw_ostream &OS,
const MCInstrDesc &Desc = MCII.get(MI.getOpcode());
unsigned bytes = Desc.getSize();
switch (MI.getOpcode()) {
case AMDGPU::V_ACCVGPR_READ_B32_vi:
case AMDGPU::V_ACCVGPR_WRITE_B32_vi:
// Set unused op_sel_hi bits to 1.
// FIXME: This shall be done for all VOP3P but not MAI instructions with
// unused op_sel_hi bits if corresponding operands do not exist.
// accvgpr_read/write are different, however. These are VOP3P, MAI, have
// src0, but do not use op_sel.
Encoding |= (1ul << 14) | (1ul << 59) | (1ul << 60);
break;
default:
break;
}
for (unsigned i = 0; i < bytes; i++) {
OS.write((uint8_t) ((Encoding >> (8 * i)) & 0xff));
}

6
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
View File

@@ -438,9 +438,9 @@ multiclass VOP3P_Real_MAI<bits<7> op> {
VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").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
let Inst{14} = ?; // op_sel_hi(2)
let Inst{59} = ?; // op_sel_hi(0)
let Inst{60} = ?; // op_sel_hi(1)
}
}

18
llvm/test/MC/Disassembler/AMDGPU/mai.txt
View File

@@ -3,6 +3,15 @@
# GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18
# Check the alternative encoding with unused op_sel_hi bits set to zero
# and not to default 1 is accepted.
#
# FIXME: Encoding is canonicalized when printing. It is valid but encoding
# bits are not the same as in input.
# GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x00
# GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]
0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18
@@ -18,6 +27,15 @@
# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18
# Check the alternative encoding with unused op_sel_hi bits set to zero
# and not to default 1 is accepted.
#
# FIXME: Encoding is canonicalized when printing. It is valid but encoding
# bits are not the same as in input.
# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00
# 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