[PATCH] D96543: [AMDGPU] Allow accvgpr_read/write decode with opsel

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rGc0d7a8bc6241: [AMDGPU] Allow accvgpr_read/write decode with opsel (authored by rampitec).

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D96543/new/

https://reviews.llvm.org/D96543

Files:
  llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp
  llvm/lib/Target/AMDGPU/VOP3PInstructions.td
  llvm/test/MC/Disassembler/AMDGPU/mai.txt


Index: llvm/test/MC/Disassembler/AMDGPU/mai.txt
===================================================================
--- llvm/test/MC/Disassembler/AMDGPU/mai.txt
+++ llvm/test/MC/Disassembler/AMDGPU/mai.txt
@@ -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
 
Index: llvm/lib/Target/AMDGPU/VOP3PInstructions.td
===================================================================
--- llvm/lib/Target/AMDGPU/VOP3PInstructions.td
+++ llvm/lib/Target/AMDGPU/VOP3PInstructions.td
@@ -438,9 +438,9 @@
             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)
   }
 }
 
Index: llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp
+++ llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp
@@ -284,6 +284,20 @@
   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));
   }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D96543.323380.patch
Type: text/x-patch
Size: 3176 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20210212/04ff8667/attachment-0001.bin>