[AMDGPU] Reorganize VOP3P encoding

[rampitec]

Summary

This changes width of encoding and opcode fields to match the
documentation.

[kzhuravl]

LGTM

llvm/lib/Target/AMDGPU/VOP3PInstructions.td

 def V_MFMA_F32_32X32X4BF16 : VOP3Inst<"v_mfma_f32_32x32x4bf16", VOPProfileMAI_F32_V2I16_X16, int_amdgcn_mfma_f32_32x32x4bf16>;
 } // End isConvergent = 1, mayRaiseFPException = 0, ReadsModeReg = 1
 
 } // End SubtargetPredicate = HasMAIInsts
 
 def : MnemonicAlias<"v_accvgpr_read",  "v_accvgpr_read_b32">;
 def : MnemonicAlias<"v_accvgpr_write", "v_accvgpr_write_b32">;
 
-multiclass VOP3P_Real_vi<bits<10> op> {
+multiclass VOP3P_Real_vi<bits<7> op> {
   def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>,
             VOP3Pe <op, !cast<VOP3_Pseudo>(NAME).Pfl> {
     let AssemblerPredicate = HasVOP3PInsts;
     let DecoderNamespace = "GFX8";
   }
 }
 
-multiclass VOP3P_Real_MAI<bits<10> 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";
   }
 }
 
-defm V_PK_MAD_I16 : VOP3P_Real_vi <0x380>;
-defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x381>;
-defm V_PK_ADD_I16 : VOP3P_Real_vi <0x382>;
-defm V_PK_SUB_I16 : VOP3P_Real_vi <0x383>;
-defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x384>;
-defm V_PK_LSHRREV_B16 : VOP3P_Real_vi <0x385>;
-defm V_PK_ASHRREV_I16 : VOP3P_Real_vi <0x386>;
-defm V_PK_MAX_I16 : VOP3P_Real_vi <0x387>;
-defm V_PK_MIN_I16 : VOP3P_Real_vi <0x388>;
-defm V_PK_MAD_U16 : VOP3P_Real_vi <0x389>;
-
-defm V_PK_ADD_U16 : VOP3P_Real_vi <0x38a>;
-defm V_PK_SUB_U16 : VOP3P_Real_vi <0x38b>;
-defm V_PK_MAX_U16 : VOP3P_Real_vi <0x38c>;
-defm V_PK_MIN_U16 : VOP3P_Real_vi <0x38d>;
-defm V_PK_FMA_F16 : VOP3P_Real_vi <0x38e>;
-defm V_PK_ADD_F16 : VOP3P_Real_vi <0x38f>;
-defm V_PK_MUL_F16 : VOP3P_Real_vi <0x390>;
-defm V_PK_MIN_F16 : VOP3P_Real_vi <0x391>;
-defm V_PK_MAX_F16 : VOP3P_Real_vi <0x392>;
+defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>;
+defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>;
+defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>;
+defm V_PK_SUB_I16 : VOP3P_Real_vi <0x03>;
+defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x04>;
+defm V_PK_LSHRREV_B16 : VOP3P_Real_vi <0x05>;
+defm V_PK_ASHRREV_I16 : VOP3P_Real_vi <0x06>;
+defm V_PK_MAX_I16 : VOP3P_Real_vi <0x07>;
+defm V_PK_MIN_I16 : VOP3P_Real_vi <0x08>;
+defm V_PK_MAD_U16 : VOP3P_Real_vi <0x09>;
+
+defm V_PK_ADD_U16 : VOP3P_Real_vi <0x0a>;
+defm V_PK_SUB_U16 : VOP3P_Real_vi <0x0b>;
+defm V_PK_MAX_U16 : VOP3P_Real_vi <0x0c>;
+defm V_PK_MIN_U16 : VOP3P_Real_vi <0x0d>;
+defm V_PK_FMA_F16 : VOP3P_Real_vi <0x0e>;
+defm V_PK_ADD_F16 : VOP3P_Real_vi <0x0f>;
+defm V_PK_MUL_F16 : VOP3P_Real_vi <0x10>;
+defm V_PK_MIN_F16 : VOP3P_Real_vi <0x11>;
+defm V_PK_MAX_F16 : VOP3P_Real_vi <0x12>;
 
 
 let SubtargetPredicate = HasMadMixInsts in {
-defm V_MAD_MIX_F32 : VOP3P_Real_vi <0x3a0>;
-defm V_MAD_MIXLO_F16 : VOP3P_Real_vi <0x3a1>;
-defm V_MAD_MIXHI_F16 : VOP3P_Real_vi <0x3a2>;
+defm V_MAD_MIX_F32 : VOP3P_Real_vi <0x20>;
+defm V_MAD_MIXLO_F16 : VOP3P_Real_vi <0x21>;
+defm V_MAD_MIXHI_F16 : VOP3P_Real_vi <0x22>;
 }
 
 let SubtargetPredicate = HasFmaMixInsts in {
 let DecoderNamespace = "GFX9_DL" in {
 // The mad_mix instructions were renamed and their behaviors changed,
 // but the opcode stayed the same so we need to put these in a
 // different DecoderNamespace to avoid the ambiguity.
-defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x3a0>;
-defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x3a1>;
-defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x3a2>;
+defm V_FMA_MIX_F32 : VOP3P_Real_vi <0x20>;
+defm V_FMA_MIXLO_F16 : VOP3P_Real_vi <0x21>;
+defm V_FMA_MIXHI_F16 : VOP3P_Real_vi <0x22>;
 }
 }
 
 
 let SubtargetPredicate = HasDot2Insts in {
 
-defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x3a3>;
-defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x3a6>;
-defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x3a7>;
-defm V_DOT4_U32_U8  : VOP3P_Real_vi <0x3a9>;
-defm V_DOT8_U32_U4  : VOP3P_Real_vi <0x3ab>;
+defm V_DOT2_F32_F16 : VOP3P_Real_vi <0x23>;
+defm V_DOT2_I32_I16 : VOP3P_Real_vi <0x26>;
+defm V_DOT2_U32_U16 : VOP3P_Real_vi <0x27>;
+defm V_DOT4_U32_U8  : VOP3P_Real_vi <0x29>;
+defm V_DOT8_U32_U4  : VOP3P_Real_vi <0x2b>;
 
 } // End SubtargetPredicate = HasDot2Insts
 
 let SubtargetPredicate = HasDot1Insts in {
 
-defm V_DOT4_I32_I8  : VOP3P_Real_vi <0x3a8>;
-defm V_DOT8_I32_I4  : VOP3P_Real_vi <0x3aa>;
+defm V_DOT4_I32_I8  : VOP3P_Real_vi <0x28>;
+defm V_DOT8_I32_I4  : VOP3P_Real_vi <0x2a>;
 
 } // End SubtargetPredicate = HasDot1Insts
 
 let SubtargetPredicate = HasMAIInsts in {
 
-defm V_ACCVGPR_READ_B32  : VOP3P_Real_MAI <0x3d8>;
-defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x3d9>;
-defm V_MFMA_F32_32X32X1F32  : VOP3P_Real_MAI <0x3c0>;
-defm V_MFMA_F32_16X16X1F32  : VOP3P_Real_MAI <0x3c1>;
-defm V_MFMA_F32_4X4X1F32    : VOP3P_Real_MAI <0x3c2>;
-defm V_MFMA_F32_32X32X2F32  : VOP3P_Real_MAI <0x3c4>;
-defm V_MFMA_F32_16X16X4F32  : VOP3P_Real_MAI <0x3c5>;
-defm V_MFMA_F32_32X32X4F16  : VOP3P_Real_MAI <0x3c8>;
-defm V_MFMA_F32_16X16X4F16  : VOP3P_Real_MAI <0x3c9>;
-defm V_MFMA_F32_4X4X4F16    : VOP3P_Real_MAI <0x3ca>;
-defm V_MFMA_F32_32X32X8F16  : VOP3P_Real_MAI <0x3cc>;
-defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MAI <0x3cd>;
-defm V_MFMA_I32_32X32X4I8   : VOP3P_Real_MAI <0x3d0>;
-defm V_MFMA_I32_16X16X4I8   : VOP3P_Real_MAI <0x3d1>;
-defm V_MFMA_I32_4X4X4I8     : VOP3P_Real_MAI <0x3d2>;
-defm V_MFMA_I32_32X32X8I8   : VOP3P_Real_MAI <0x3d4>;
-defm V_MFMA_I32_16X16X16I8  : VOP3P_Real_MAI <0x3d5>;
-defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MAI <0x3e8>;
-defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MAI <0x3e9>;
-defm V_MFMA_F32_4X4X2BF16   : VOP3P_Real_MAI <0x3eb>;
-defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MAI <0x3ec>;
-defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MAI <0x3ed>;
+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>;
 
 } // End SubtargetPredicate = HasMAIInsts
 
 //===----------------------------------------------------------------------===//
 // GFX10.
 //===----------------------------------------------------------------------===//
 
 let AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10" in {
-  multiclass VOP3P_Real_gfx10<bits<10> op> {
+  multiclass VOP3P_Real_gfx10<bits<7> op> {
     def _gfx10 : VOP3P_Real<!cast<VOP3P_Pseudo>(NAME), SIEncodingFamily.GFX10>,
                  VOP3Pe_gfx10 <op, !cast<VOP3P_Pseudo>(NAME).Pfl>;
   }
 } // End AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10"
 
-defm V_PK_MAD_I16     : VOP3P_Real_gfx10<0x000>;
-defm V_PK_MUL_LO_U16  : VOP3P_Real_gfx10<0x001>;
-defm V_PK_ADD_I16     : VOP3P_Real_gfx10<0x002>;
-defm V_PK_SUB_I16     : VOP3P_Real_gfx10<0x003>;
-defm V_PK_LSHLREV_B16 : VOP3P_Real_gfx10<0x004>;
-defm V_PK_LSHRREV_B16 : VOP3P_Real_gfx10<0x005>;
-defm V_PK_ASHRREV_I16 : VOP3P_Real_gfx10<0x006>;
-defm V_PK_MAX_I16     : VOP3P_Real_gfx10<0x007>;
-defm V_PK_MIN_I16     : VOP3P_Real_gfx10<0x008>;
-defm V_PK_MAD_U16     : VOP3P_Real_gfx10<0x009>;
-defm V_PK_ADD_U16     : VOP3P_Real_gfx10<0x00a>;
-defm V_PK_SUB_U16     : VOP3P_Real_gfx10<0x00b>;
-defm V_PK_MAX_U16     : VOP3P_Real_gfx10<0x00c>;
-defm V_PK_MIN_U16     : VOP3P_Real_gfx10<0x00d>;
-defm V_PK_FMA_F16     : VOP3P_Real_gfx10<0x00e>;
-defm V_PK_ADD_F16     : VOP3P_Real_gfx10<0x00f>;
-defm V_PK_MUL_F16     : VOP3P_Real_gfx10<0x010>;
-defm V_PK_MIN_F16     : VOP3P_Real_gfx10<0x011>;
-defm V_PK_MAX_F16     : VOP3P_Real_gfx10<0x012>;
-defm V_FMA_MIX_F32    : VOP3P_Real_gfx10<0x020>;
-defm V_FMA_MIXLO_F16  : VOP3P_Real_gfx10<0x021>;
-defm V_FMA_MIXHI_F16  : VOP3P_Real_gfx10<0x022>;
+defm V_PK_MAD_I16     : VOP3P_Real_gfx10<0x00>;
+defm V_PK_MUL_LO_U16  : VOP3P_Real_gfx10<0x01>;
+defm V_PK_ADD_I16     : VOP3P_Real_gfx10<0x02>;
+defm V_PK_SUB_I16     : VOP3P_Real_gfx10<0x03>;
+defm V_PK_LSHLREV_B16 : VOP3P_Real_gfx10<0x04>;
+defm V_PK_LSHRREV_B16 : VOP3P_Real_gfx10<0x05>;
+defm V_PK_ASHRREV_I16 : VOP3P_Real_gfx10<0x06>;
+defm V_PK_MAX_I16     : VOP3P_Real_gfx10<0x07>;
+defm V_PK_MIN_I16     : VOP3P_Real_gfx10<0x08>;
+defm V_PK_MAD_U16     : VOP3P_Real_gfx10<0x09>;
+defm V_PK_ADD_U16     : VOP3P_Real_gfx10<0x0a>;
+defm V_PK_SUB_U16     : VOP3P_Real_gfx10<0x0b>;
+defm V_PK_MAX_U16     : VOP3P_Real_gfx10<0x0c>;
+defm V_PK_MIN_U16     : VOP3P_Real_gfx10<0x0d>;
+defm V_PK_FMA_F16     : VOP3P_Real_gfx10<0x0e>;
+defm V_PK_ADD_F16     : VOP3P_Real_gfx10<0x0f>;
+defm V_PK_MUL_F16     : VOP3P_Real_gfx10<0x10>;
+defm V_PK_MIN_F16     : VOP3P_Real_gfx10<0x11>;
+defm V_PK_MAX_F16     : VOP3P_Real_gfx10<0x12>;
+defm V_FMA_MIX_F32    : VOP3P_Real_gfx10<0x20>;
+defm V_FMA_MIXLO_F16  : VOP3P_Real_gfx10<0x21>;
+defm V_FMA_MIXHI_F16  : VOP3P_Real_gfx10<0x22>;
 
 let SubtargetPredicate = HasDot2Insts in {
 
-defm V_DOT2_F32_F16 : VOP3P_Real_gfx10 <0x013>;
-defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x014>;
-defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x015>;
-defm V_DOT4_U32_U8  : VOP3P_Real_gfx10 <0x017>;
-defm V_DOT8_U32_U4  : VOP3P_Real_gfx10 <0x019>;
+defm V_DOT2_F32_F16 : VOP3P_Real_gfx10 <0x13>;
+defm V_DOT2_I32_I16 : VOP3P_Real_gfx10 <0x14>;
+defm V_DOT2_U32_U16 : VOP3P_Real_gfx10 <0x15>;
+defm V_DOT4_U32_U8  : VOP3P_Real_gfx10 <0x17>;
+defm V_DOT8_U32_U4  : VOP3P_Real_gfx10 <0x19>;
 
 } // End SubtargetPredicate = HasDot2Insts
 
 let SubtargetPredicate = HasDot1Insts in {
 
-defm V_DOT4_I32_I8  : VOP3P_Real_gfx10 <0x016>;
-defm V_DOT8_I32_I4  : VOP3P_Real_gfx10 <0x018>;
+defm V_DOT4_I32_I8  : VOP3P_Real_gfx10 <0x16>;
+defm V_DOT8_I32_I4  : VOP3P_Real_gfx10 <0x18>;
 
 } // End SubtargetPredicate = HasDot1Insts

llvm/lib/Target/AMDGPU/VOPInstructions.td

class VOP3be <VOPProfile P> : Enc64 {
   let Inst{49-41} = !if(P.HasSrc1, src1, 0);
   let Inst{58-50} = !if(P.HasSrc2, src2, 0);
   let Inst{60-59} = !if(P.HasOMod, omod, 0);
   let Inst{61}    = !if(P.HasSrc0Mods, src0_modifiers{0}, 0);
   let Inst{62}    = !if(P.HasSrc1Mods, src1_modifiers{0}, 0);
   let Inst{63}    = !if(P.HasSrc2Mods, src2_modifiers{0}, 0);
 }
 
-class VOP3Pe <bits<10> op, VOPProfile P> : Enc64 {
+class VOP3Pe <bits<7> op, VOPProfile P> : Enc64 {
   bits<8> vdst;
   // neg, neg_hi, op_sel put in srcN_modifiers
   bits<4> src0_modifiers;
   bits<9> src0;
   bits<4> src1_modifiers;
   bits<9> src1;
   bits<4> src2_modifiers;
   bits<9> src2;
   bits<1> clamp;
 
   let Inst{7-0} = vdst;
   let Inst{8} = !if(P.HasSrc0Mods, src0_modifiers{1}, 0); // neg_hi src0
   let Inst{9} = !if(P.HasSrc1Mods, src1_modifiers{1}, 0); // neg_hi src1
   let Inst{10} = !if(P.HasSrc2Mods, src2_modifiers{1}, 0); // neg_hi src2
 
   let Inst{11} = !if(!and(P.HasSrc0, P.HasOpSel), src0_modifiers{2}, 0); // op_sel(0)
   let Inst{12} = !if(!and(P.HasSrc1, P.HasOpSel), src1_modifiers{2}, 0); // op_sel(1)
   let Inst{13} = !if(!and(P.HasSrc2, P.HasOpSel), src2_modifiers{2}, 0); // op_sel(2)
 
   let Inst{14} = !if(!and(P.HasSrc2, P.HasOpSel), src2_modifiers{3}, 0); // op_sel_hi(2)
 
   let Inst{15} = !if(P.HasClamp, clamp{0}, 0);
 
-  let Inst{25-16} = op;
-  let Inst{31-26} = 0x34; //encoding
+  let Inst{22-16} = op;
+  let Inst{31-23} = 0x1a7; //encoding
   let Inst{40-32} = !if(P.HasSrc0, src0, 0);
   let Inst{49-41} = !if(P.HasSrc1, src1, 0);
   let Inst{58-50} = !if(P.HasSrc2, src2, 0);
   let Inst{59}    = !if(!and(P.HasSrc0, P.HasOpSel), src0_modifiers{3}, 0); // op_sel_hi(0)
   let Inst{60}    = !if(!and(P.HasSrc1, P.HasOpSel), src1_modifiers{3}, 0); // op_sel_hi(1)
   let Inst{61}    = !if(P.HasSrc0Mods, src0_modifiers{0}, 0); // neg (lo)
   let Inst{62}    = !if(P.HasSrc1Mods, src1_modifiers{0}, 0); // neg (lo)
   let Inst{63}    = !if(P.HasSrc2Mods, src2_modifiers{0}, 0); // neg (lo)
 }
 
-class VOP3Pe_MAI <bits<10> op, VOPProfile P> : Enc64 {
+class VOP3Pe_MAI <bits<7> op, VOPProfile P> : Enc64 {
   bits<8> vdst;
   bits<10> src0;
   bits<10> src1;
   bits<9> src2;
   bits<3> blgp;
   bits<3> cbsz;
   bits<4> abid;
   bits<1> clamp;
 
   let Inst{7-0} = vdst;
 
   let Inst{10-8}  = !if(P.HasSrc1, cbsz, 0);
   let Inst{14-11} = !if(P.HasSrc1, abid, 0);
 
   let Inst{15} = !if(P.HasClamp, clamp{0}, 0);
 
-  let Inst{25-16} = op;
-  let Inst{31-26} = 0x34; //encoding
+  let Inst{22-16} = op;
+  let Inst{31-23} = 0x1a7; //encoding
   let Inst{40-32} = !if(P.HasSrc0, src0{8-0}, 0);
   let Inst{49-41} = !if(P.HasSrc1, src1{8-0}, 0);
   let Inst{58-50} = !if(P.HasSrc2, src2, 0);
 
   let Inst{59}    = !if(P.HasSrc0, src0{9}, 0); // acc(0)
   let Inst{60}    = !if(P.HasSrc1, src1{9}, 0); // acc(1)
 
   let Inst{63-61} = !if(P.HasSrc1, blgp, 0);
 }
 
 
-class VOP3Pe_gfx10 <bits<10> op, VOPProfile P> : VOP3Pe<op, P> {
-  let Inst{31-26} = 0x33; //encoding
+class VOP3Pe_gfx10 <bits<7> op, VOPProfile P> : VOP3Pe<op, P> {
+  let Inst{31-23} = 0x198; //encoding
 }
 
 class VOP3be_gfx6_gfx7<bits<9> op, VOPProfile p> : VOP3be<p> {
   let Inst{25-17} = op;
 }
 
 class VOP3be_gfx10<bits<10> op, VOPProfile p> : VOP3be<p> {
   bits<1> clamp;