Skip to content

AMDGPU: Verify f8f6f4 formats in assembler #117826

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Nov 27, 2024
Merged

AMDGPU: Verify f8f6f4 formats in assembler #117826

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
32 changes: 32 additions & 0 deletions llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4196,6 +4196,38 @@ bool AMDGPUAsmParser::validateMFMA(const MCInst &Inst,
if ((Desc.TSFlags & SIInstrFlags::IsMAI) == 0)
return true;

int BlgpIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::blgp);
if (BlgpIdx != -1) {
if (const MFMA_F8F6F4_Info *Info = AMDGPU::isMFMA_F8F6F4(Opc)) {
int CbszIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::cbsz);

unsigned CBSZ = Inst.getOperand(CbszIdx).getImm();
unsigned BLGP = Inst.getOperand(BlgpIdx).getImm();

// Validate the correct register size was used for the floating point
// format operands

bool Success = true;
if (Info->NumRegsSrcA != mfmaScaleF8F6F4FormatToNumRegs(CBSZ)) {
int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
Error(getRegLoc(mc2PseudoReg(Inst.getOperand(Src0Idx).getReg()),
Operands),
"wrong register tuple size for cbsz value " + Twine(CBSZ));
Success = false;
}

if (Info->NumRegsSrcB != mfmaScaleF8F6F4FormatToNumRegs(BLGP)) {
int Src1Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src1);
Error(getRegLoc(mc2PseudoReg(Inst.getOperand(Src1Idx).getReg()),
Operands),
"wrong register tuple size for blgp value " + Twine(BLGP));
Success = false;
}

return Success;
}
}

const int Src2Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src2);
if (Src2Idx == -1)
return true;
Expand Down
9 changes: 9 additions & 0 deletions llvm/lib/Target/AMDGPU/SIInstrInfo.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3230,6 +3230,15 @@ def getMFMA_F8F6F4_WithSize : GenericTable {
let PrimaryKeyName = "getMFMA_F8F6F4_InstWithNumRegs" ;
}

def isMFMA_F8F6F4Table : GenericTable {
let FilterClass = "MFMA_F8F6F4_WithSizeTable";
let CppTypeName = "MFMA_F8F6F4_Info";
// let Fields = [ "Opcode" ];
let Fields = [ "Opcode", "F8F8Opcode", "NumRegsSrcA", "NumRegsSrcB" ];
let PrimaryKey = [ "Opcode" ];
let PrimaryKeyName = "isMFMA_F8F6F4" ;
}

def FP8DstByteSelTable : GenericTable {
let FilterClass = "VOP3_Pseudo";
let CppTypeName = "FP8DstByteSelInfo";
Expand Down
3 changes: 2 additions & 1 deletion llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -417,6 +417,7 @@ struct FP8DstByteSelInfo {
#define GET_WMMAOpcode3AddrMappingTable_IMPL
#define GET_getMFMA_F8F6F4_WithSize_DECL
#define GET_getMFMA_F8F6F4_WithSize_IMPL
#define GET_isMFMA_F8F6F4Table_IMPL
#include "AMDGPUGenSearchableTables.inc"

int getMTBUFBaseOpcode(unsigned Opc) {
Expand Down Expand Up @@ -525,7 +526,7 @@ bool getMAIIsGFX940XDL(unsigned Opc) {
return Info ? Info->is_gfx940_xdl : false;
}

static uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal) {
uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal) {
switch (EncodingVal) {
case MFMAScaleFormats::FP6_E2M3:
case MFMAScaleFormats::FP6_E3M2:
Expand Down
17 changes: 11 additions & 6 deletions llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -96,13 +96,22 @@ struct MAIInstInfo {
bool is_gfx940_xdl;
};

struct MFMA_F8F6F4_Info {
unsigned Opcode;
unsigned F8F8Opcode;
uint8_t NumRegsSrcA;
uint8_t NumRegsSrcB;
};

#define GET_MIMGBaseOpcode_DECL
#define GET_MIMGDim_DECL
#define GET_MIMGEncoding_DECL
#define GET_MIMGLZMapping_DECL
#define GET_MIMGMIPMapping_DECL
#define GET_MIMGBiASMapping_DECL
#define GET_MAIInstInfoTable_DECL
#define GET_MAIInstInfoTable_DECL
#define GET_isMFMA_F8F6F4Table_DECL
#include "AMDGPUGenSearchableTables.inc"

namespace IsaInfo {
Expand Down Expand Up @@ -581,12 +590,8 @@ unsigned getVOPDEncodingFamily(const MCSubtargetInfo &ST);
LLVM_READONLY
CanBeVOPD getCanBeVOPD(unsigned Opc);

struct MFMA_F8F6F4_Info {
unsigned Opcode;
unsigned F8F8Opcode;
uint8_t NumRegsSrcA;
uint8_t NumRegsSrcB;
};
LLVM_READNONE
uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal);

LLVM_READONLY
const MFMA_F8F6F4_Info *getMFMA_F8F6F4_WithFormatArgs(unsigned CBSZ,
Expand Down
127 changes: 127 additions & 0 deletions llvm/test/MC/AMDGPU/mai-gfx950-err.s
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -29,3 +29,130 @@ v_mfma_ld_scale_b32 v0, v0 neg_hi:[0,1]

v_mfma_ld_scale_b32 v0, v0 neg_lo:[0,1] neg_hi:[0,1]
// CHECK: :[[@LINE-1]]:28: error: not a valid operand


v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] blgp:2
// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 2

v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2 blgp:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 2


v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] blgp:2
// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 2

v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2 blgp:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] blgp:2
// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 2

v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] cbsz:2
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2

v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] blgp:2
// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] v32, v33 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] v32, v33 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2



v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:1
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3


v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:1
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:1
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:2
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:1
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:2
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2

v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:3
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:7], v[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:3
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[0:3] v20, v21 cbsz:4
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:4
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:9], v[0:3] v20, v21 blgp:4
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4

v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:4
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4
Loading