This is an archive of the discontinued LLVM Phabricator instance.

Differential D91202

[AMDGPU] Set default op_sel_hi on accvgpr read/write
ClosedPublic

Authored by rampitec on Nov 10 2020, 12:38 PM.

Details

Reviewers
kzhuravl
Commits
rG544ef42e40aa: [AMDGPU] Set default op_sel_hi on accvgpr read/write
Summary

These are opsel opcodes with op_sel actually being ignored.
As a such op_sel_hi needs to be set to default 1 even though
these bits are ignored. This is compatibility change.

Diff Detail

Event Timeline

rampitec created this revision.Nov 10 2020, 12:38 PM
Herald added a project: Restricted Project. · View Herald TranscriptNov 10 2020, 12:38 PM
Herald added subscribers: kerbowa, hiraditya, t-tye and 6 others. · View Herald Transcript
rampitec requested review of this revision.Nov 10 2020, 12:38 PM
Herald added a subscriber: wdng. · View Herald TranscriptNov 10 2020, 12:38 PM
kzhuravl accepted this revision.Nov 10 2020, 12:49 PM

LGTM

This revision is now accepted and ready to land.Nov 10 2020, 12:49 PM
Closed by commit rG544ef42e40aa: [AMDGPU] Set default op_sel_hi on accvgpr read/write (authored by rampitec). · Explain WhyNov 10 2020, 1:08 PM
This revision was automatically updated to reflect the committed changes.
rampitec added a commit: rG544ef42e40aa: [AMDGPU] Set default op_sel_hi on accvgpr read/write.

Revision Contents

PathSize
llvm/
lib/
Target/
AMDGPU/
51 lines
test/
MC/
AMDGPU/
4 lines
16 lines
Disassembler/
AMDGPU/
24 lines

Diff 304313

llvm/lib/Target/AMDGPU/VOP3PInstructions.td

Show First 20 Lines • Show All 425 Lines • ▼ Show 20 Linesmulticlass VOP3P_Real_vi<bits<7> op> {
} }
} }
multiclass VOP3P_Real_MAI<bits<7> op> { multiclass VOP3P_Real_MAI<bits<7> op> {
def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>, def _vi : VOP3P_Real<!cast<VOP3_Pseudo>(NAME), SIEncodingFamily.VI>,
VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME).Pfl> { VOP3Pe_MAI <op, !cast<VOP3_Pseudo>(NAME).Pfl> {
let AssemblerPredicate = HasMAIInsts; let AssemblerPredicate = HasMAIInsts;
let DecoderNamespace = "GFX8"; 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
}
}
multiclass VOP3P_Real_MFMA<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 <0x00>; defm V_PK_MAD_I16 : VOP3P_Real_vi <0x00>;
defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>; defm V_PK_MUL_LO_U16 : VOP3P_Real_vi <0x01>;
defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>; defm V_PK_ADD_I16 : VOP3P_Real_vi <0x02>;
defm V_PK_SUB_I16 : VOP3P_Real_vi <0x03>; defm V_PK_SUB_I16 : VOP3P_Real_vi <0x03>;
defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x04>; defm V_PK_LSHLREV_B16 : VOP3P_Real_vi <0x04>;
▲ Show 20 Lines • Show All 48 Lines • ▼ Show 20 Lines
defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>; defm V_DOT8_I32_I4 : VOP3P_Real_vi <0x2a>;
} // End SubtargetPredicate = HasDot1Insts } // End SubtargetPredicate = HasDot1Insts
let SubtargetPredicate = HasMAIInsts in { let SubtargetPredicate = HasMAIInsts in {
defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>; defm V_ACCVGPR_READ_B32 : VOP3P_Real_MAI <0x58>;
defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>; defm V_ACCVGPR_WRITE_B32 : VOP3P_Real_MAI <0x59>;
defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MAI <0x40>; defm V_MFMA_F32_32X32X1F32 : VOP3P_Real_MFMA <0x40>;
defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MAI <0x41>; defm V_MFMA_F32_16X16X1F32 : VOP3P_Real_MFMA <0x41>;
defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MAI <0x42>; defm V_MFMA_F32_4X4X1F32 : VOP3P_Real_MFMA <0x42>;
defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MAI <0x44>; defm V_MFMA_F32_32X32X2F32 : VOP3P_Real_MFMA <0x44>;
defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MAI <0x45>; defm V_MFMA_F32_16X16X4F32 : VOP3P_Real_MFMA <0x45>;
defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MAI <0x48>; defm V_MFMA_F32_32X32X4F16 : VOP3P_Real_MFMA <0x48>;
defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MAI <0x49>; defm V_MFMA_F32_16X16X4F16 : VOP3P_Real_MFMA <0x49>;
defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MAI <0x4a>; defm V_MFMA_F32_4X4X4F16 : VOP3P_Real_MFMA <0x4a>;
defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MAI <0x4c>; defm V_MFMA_F32_32X32X8F16 : VOP3P_Real_MFMA <0x4c>;
defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MAI <0x4d>; defm V_MFMA_F32_16X16X16F16 : VOP3P_Real_MFMA <0x4d>;
defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MAI <0x50>; defm V_MFMA_I32_32X32X4I8 : VOP3P_Real_MFMA <0x50>;
defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MAI <0x51>; defm V_MFMA_I32_16X16X4I8 : VOP3P_Real_MFMA <0x51>;
defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MAI <0x52>; defm V_MFMA_I32_4X4X4I8 : VOP3P_Real_MFMA <0x52>;
defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MAI <0x54>; defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MFMA <0x55>;
defm V_MFMA_I32_16X16X16I8 : VOP3P_Real_MAI <0x55>; defm V_MFMA_I32_32X32X8I8 : VOP3P_Real_MFMA <0x54>;
defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MAI <0x68>; defm V_MFMA_F32_32X32X2BF16 : VOP3P_Real_MFMA <0x68>;
defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MAI <0x69>; defm V_MFMA_F32_16X16X2BF16 : VOP3P_Real_MFMA <0x69>;
defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MAI <0x6b>; defm V_MFMA_F32_4X4X2BF16 : VOP3P_Real_MFMA <0x6b>;
defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MAI <0x6c>; defm V_MFMA_F32_32X32X4BF16 : VOP3P_Real_MFMA <0x6c>;
defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MAI <0x6d>; defm V_MFMA_F32_16X16X8BF16 : VOP3P_Real_MFMA <0x6d>;
} // End SubtargetPredicate = HasMAIInsts } // End SubtargetPredicate = HasMAIInsts
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
// GFX10. // GFX10.
//===----------------------------------------------------------------------===// //===----------------------------------------------------------------------===//
let AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10" in { let AssemblerPredicate = isGFX10Plus, DecoderNamespace = "GFX10" in {
▲ Show 20 Lines • Show All 45 Lines • Show Last 20 Lines

llvm/test/MC/AMDGPU/accvgpr-altnames.s

// RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s // RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s
v_accvgpr_read_b32 v2, acc0 v_accvgpr_read_b32 v2, acc0
// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] // GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
v_accvgpr_write_b32 acc2, -2.0 v_accvgpr_write_b32 acc2, -2.0
// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00] // GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
v_mfma_f32_32x32x1f32 acc[0:31], acc0, acc1, acc[1:32] v_mfma_f32_32x32x1f32 acc[0:31], acc0, acc1, acc[1:32]
// GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x1c] // GFX908: v_mfma_f32_32x32x1f32 a[0:31], a0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x1c]

llvm/test/MC/AMDGPU/mai.s

// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding %s | FileCheck -check-prefix=GFX908 %s
// RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefix=NOGFX908 --implicit-check-not=error: %s // RUN: not llvm-mc -arch=amdgcn -mcpu=gfx908 %s 2>&1 | FileCheck -check-prefix=NOGFX908 --implicit-check-not=error: %s
v_accvgpr_read_b32 v2, a0 v_accvgpr_read_b32 v2, a0
// GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] // GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
v_accvgpr_read_b32 v2, a1 v_accvgpr_read_b32 v2, a1
// GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08] // GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]
v_accvgpr_read_b32 v2, a255 v_accvgpr_read_b32 v2, a255
// GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08] // GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18]
v_accvgpr_read v2, a10 v_accvgpr_read v2, a10
// GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x00,0xd8,0xd3,0x0a,0x01,0x00,0x08] // GFX908: v_accvgpr_read_b32 v2, a10 ; encoding: [0x02,0x40,0xd8,0xd3,0x0a,0x01,0x00,0x18]
v_accvgpr_write_b32 a2, -2.0 v_accvgpr_write_b32 a2, -2.0
// GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00] // GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
v_accvgpr_write_b32 a2, -2 v_accvgpr_write_b32 a2, -2
// GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00] // GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18]
v_accvgpr_write_b32 a2, v1 v_accvgpr_write_b32 a2, v1
// GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00] // GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
v_accvgpr_write a2, v255 v_accvgpr_write a2, v255
// GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x00,0xd9,0xd3,0xff,0x01,0x00,0x00] // GFX908: v_accvgpr_write_b32 a2, v255 ; encoding: [0x02,0x40,0xd9,0xd3,0xff,0x01,0x00,0x18]
v_accvgpr_write a2, 100 v_accvgpr_write a2, 100
// NOGFX908: error: invalid operand for instruction // NOGFX908: error: invalid operand for instruction
v_accvgpr_write a2, execz v_accvgpr_write a2, execz
// NOGFX908: error: source operand must be either a VGPR or an inline constant // NOGFX908: error: source operand must be either a VGPR or an inline constant
v_accvgpr_write a2, vccz v_accvgpr_write a2, vccz
▲ Show 20 Lines • Show All 490 Lines • Show Last 20 Lines

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

# RUN: llvm-mc -arch=amdgcn -mcpu=gfx908 -show-encoding -disassemble %s | FileCheck -check-prefix=GFX908 %s # 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,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08] # GFX908: v_accvgpr_read_b32 v2, a0 ; encoding: [0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18]
0x02,0x00,0xd8,0xd3,0x00,0x01,0x00,0x08 0x02,0x40,0xd8,0xd3,0x00,0x01,0x00,0x18
# GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08] # GFX908: v_accvgpr_read_b32 v2, a1 ; encoding: [0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18]
0x02,0x00,0xd8,0xd3,0x01,0x01,0x00,0x08 0x02,0x40,0xd8,0xd3,0x01,0x01,0x00,0x18
# GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08] # GFX908: v_accvgpr_read_b32 v2, a255 ; encoding: [0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18]
0x02,0x00,0xd8,0xd3,0xff,0x01,0x00,0x08 0x02,0x40,0xd8,0xd3,0xff,0x01,0x00,0x18
# GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00] # GFX908: v_accvgpr_write_b32 a2, -2.0 ; encoding: [0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18]
0x02,0x00,0xd9,0xd3,0xf5,0x00,0x00,0x00 0x02,0x40,0xd9,0xd3,0xf5,0x00,0x00,0x18
# GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00] # GFX908: v_accvgpr_write_b32 a2, -2 ; encoding: [0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18]
0x02,0x00,0xd9,0xd3,0xc2,0x00,0x00,0x00 0x02,0x40,0xd9,0xd3,0xc2,0x00,0x00,0x18
# GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00] # GFX908: v_accvgpr_write_b32 a2, v1 ; encoding: [0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18]
0x02,0x00,0xd9,0xd3,0x01,0x01,0x00,0x00 0x02,0x40,0xd9,0xd3,0x01,0x01,0x00,0x18
# GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, v1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x04] # 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 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] # 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 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] # GFX908: v_mfma_f32_32x32x1f32 a[0:31], v0, a1, a[1:32] ; encoding: [0x00,0x00,0xc0,0xd3,0x00,0x03,0x06,0x14]
▲ Show 20 Lines • Show All 952 Lines • Show Last 20 Lines