[AMDGPU] Allow accvgpr_read/write decode with opsel

[rampitec]

Summary

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

[dp]

It is a pity we have to use such hacks but I do not see how to make it cleaner.

[thakis]

This seems to break tests on windows: http://45.33.8.238/win/33173/step_11.txt

PTAL, and revert for now if it takes a while to fix.

[rampitec]

In D96543#2560839, @thakis wrote:

This seems to break tests on windows: http://45.33.8.238/win/33173/step_11.txt

PTAL, and revert for now if it takes a while to fix.

Ouch! Fix is trivial, constant shall be 1ull and not 1ul. I will fix it now.

[rampitec]

In D96543#2560862, @rampitec wrote:

In D96543#2560839, @thakis wrote:

This seems to break tests on windows: http://45.33.8.238/win/33173/step_11.txt

PTAL, and revert for now if it takes a while to fix.

Ouch! Fix is trivial, constant shall be 1ull and not 1ul. I will fix it now.

Sorry, fixed: https://reviews.llvm.org/rGc96e214b9ca7739737cf257baf32ff819854027d

llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp

void SIMCCodeEmitter::encodeInstruction(const MCInst &MI, raw_ostream &OS,
                                        const MCSubtargetInfo &STI) const {
   verifyInstructionPredicates(MI,
                               computeAvailableFeatures(STI.getFeatureBits()));
 
   uint64_t Encoding = getBinaryCodeForInstr(MI, Fixups, STI);
   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));
   }
 
   // NSA encoding.
   if (AMDGPU::isGFX10Plus(STI) && Desc.TSFlags & SIInstrFlags::MIMG) {
     int vaddr0 = AMDGPU::getNamedOperandIdx(MI.getOpcode(),
                                             AMDGPU::OpName::vaddr0);

llvm/lib/Target/AMDGPU/VOP3PInstructions.td

multiclass VOP3P_Real_vi<bits<7> op> {
   }
 }
 
 multiclass VOP3P_Real_MAI<bits<7> op> {
   def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,
             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)
   }
 }
 
 multiclass VOP3P_Real_MFMA<bits<7> op> {
   def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME#"_e64"), SIEncodingFamily.VI>,
             VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME#"_e64").Pfl> {
     let AssemblerPredicate = HasMAIInsts;
     let DecoderNamespace = "GFX8";

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

 # 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,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
 
 # 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,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
 0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18
 
 # 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,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
 
 # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] cbsz:3 abid:2 blgp:7 ; encoding: [0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xe4]
 0x00,0x13,0xc0,0xd3,0x00,0x03,0x06,0xe4
 
 # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x14]
 0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x14