diff --git a/llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp b/llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp --- a/llvm/lib/Target/AMDGPU/MCTargetDesc/SIMCCodeEmitter.cpp +++ b/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)); } diff --git a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td --- a/llvm/lib/Target/AMDGPU/VOP3PInstructions.td +++ b/llvm/lib/Target/AMDGPU/VOP3PInstructions.td @@ -438,9 +438,9 @@ VOP3Pe_MAI (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) } } diff --git a/llvm/test/MC/Disassembler/AMDGPU/mai.txt b/llvm/test/MC/Disassembler/AMDGPU/mai.txt --- a/llvm/test/MC/Disassembler/AMDGPU/mai.txt +++ b/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