Skip to content

Commit d9c4e9f

Browse files
arsenmarsenm
authored
AMDGPU: Verify f8f6f4 formats in assembler (#117826)
Verify the register widths of the corresponding operands match the floating point format expected size.
1 parent a2c3e0c commit d9c4e9f

File tree

6 files changed

+224
-55
lines changed

6 files changed

+224
-55
lines changed

‎llvm/lib/Target/AMDGPU/AsmParser/AMDGPUAsmParser.cpp

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4196,6 +4196,38 @@ bool AMDGPUAsmParser::validateMFMA(const MCInst &Inst,
41964196
if ((Desc.TSFlags & SIInstrFlags::IsMAI) == 0)
41974197
return true;
41984198

4199+
int BlgpIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::blgp);
4200+
if (BlgpIdx != -1) {
4201+
if (const MFMA_F8F6F4_Info *Info = AMDGPU::isMFMA_F8F6F4(Opc)) {
4202+
int CbszIdx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::cbsz);
4203+
4204+
unsigned CBSZ = Inst.getOperand(CbszIdx).getImm();
4205+
unsigned BLGP = Inst.getOperand(BlgpIdx).getImm();
4206+
4207+
// Validate the correct register size was used for the floating point
4208+
// format operands
4209+
4210+
bool Success = true;
4211+
if (Info->NumRegsSrcA != mfmaScaleF8F6F4FormatToNumRegs(CBSZ)) {
4212+
int Src0Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src0);
4213+
Error(getRegLoc(mc2PseudoReg(Inst.getOperand(Src0Idx).getReg()),
4214+
Operands),
4215+
"wrong register tuple size for cbsz value " + Twine(CBSZ));
4216+
Success = false;
4217+
}
4218+
4219+
if (Info->NumRegsSrcB != mfmaScaleF8F6F4FormatToNumRegs(BLGP)) {
4220+
int Src1Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src1);
4221+
Error(getRegLoc(mc2PseudoReg(Inst.getOperand(Src1Idx).getReg()),
4222+
Operands),
4223+
"wrong register tuple size for blgp value " + Twine(BLGP));
4224+
Success = false;
4225+
}
4226+
4227+
return Success;
4228+
}
4229+
}
4230+
41994231
const int Src2Idx = AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::src2);
42004232
if (Src2Idx == -1)
42014233
return true;

‎llvm/lib/Target/AMDGPU/SIInstrInfo.td

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3230,6 +3230,15 @@ def getMFMA_F8F6F4_WithSize : GenericTable {
32303230
let PrimaryKeyName = "getMFMA_F8F6F4_InstWithNumRegs" ;
32313231
}
32323232

3233+
def isMFMA_F8F6F4Table : GenericTable {
3234+
let FilterClass = "MFMA_F8F6F4_WithSizeTable";
3235+
let CppTypeName = "MFMA_F8F6F4_Info";
3236+
// let Fields = [ "Opcode" ];
3237+
let Fields = [ "Opcode", "F8F8Opcode", "NumRegsSrcA", "NumRegsSrcB" ];
3238+
let PrimaryKey = [ "Opcode" ];
3239+
let PrimaryKeyName = "isMFMA_F8F6F4" ;
3240+
}
3241+
32333242
def FP8DstByteSelTable : GenericTable {
32343243
let FilterClass = "VOP3_Pseudo";
32353244
let CppTypeName = "FP8DstByteSelInfo";

‎llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -417,6 +417,7 @@ struct FP8DstByteSelInfo {
417417
#define GET_WMMAOpcode3AddrMappingTable_IMPL
418418
#define GET_getMFMA_F8F6F4_WithSize_DECL
419419
#define GET_getMFMA_F8F6F4_WithSize_IMPL
420+
#define GET_isMFMA_F8F6F4Table_IMPL
420421
#include "AMDGPUGenSearchableTables.inc"
421422

422423
int getMTBUFBaseOpcode(unsigned Opc) {
@@ -525,7 +526,7 @@ bool getMAIIsGFX940XDL(unsigned Opc) {
525526
return Info ? Info->is_gfx940_xdl : false;
526527
}
527528

528-
static uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal) {
529+
uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal) {
529530
switch (EncodingVal) {
530531
case MFMAScaleFormats::FP6_E2M3:
531532
case MFMAScaleFormats::FP6_E3M2:

‎llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

Lines changed: 11 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -96,13 +96,22 @@ struct MAIInstInfo {
9696
bool is_gfx940_xdl;
9797
};
9898

99+
struct MFMA_F8F6F4_Info {
100+
unsigned Opcode;
101+
unsigned F8F8Opcode;
102+
uint8_t NumRegsSrcA;
103+
uint8_t NumRegsSrcB;
104+
};
105+
99106
#define GET_MIMGBaseOpcode_DECL
100107
#define GET_MIMGDim_DECL
101108
#define GET_MIMGEncoding_DECL
102109
#define GET_MIMGLZMapping_DECL
103110
#define GET_MIMGMIPMapping_DECL
104111
#define GET_MIMGBiASMapping_DECL
105112
#define GET_MAIInstInfoTable_DECL
113+
#define GET_MAIInstInfoTable_DECL
114+
#define GET_isMFMA_F8F6F4Table_DECL
106115
#include "AMDGPUGenSearchableTables.inc"
107116

108117
namespace IsaInfo {
@@ -581,12 +590,8 @@ unsigned getVOPDEncodingFamily(const MCSubtargetInfo &ST);
581590
LLVM_READONLY
582591
CanBeVOPD getCanBeVOPD(unsigned Opc);
583592

584-
struct MFMA_F8F6F4_Info {
585-
unsigned Opcode;
586-
unsigned F8F8Opcode;
587-
uint8_t NumRegsSrcA;
588-
uint8_t NumRegsSrcB;
589-
};
593+
LLVM_READNONE
594+
uint8_t mfmaScaleF8F6F4FormatToNumRegs(unsigned EncodingVal);
590595

591596
LLVM_READONLY
592597
const MFMA_F8F6F4_Info *getMFMA_F8F6F4_WithFormatArgs(unsigned CBSZ,

‎llvm/test/MC/AMDGPU/mai-gfx950-err.s

Lines changed: 127 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,3 +29,130 @@ v_mfma_ld_scale_b32 v0, v0 neg_hi:[0,1]
2929

3030
v_mfma_ld_scale_b32 v0, v0 neg_lo:[0,1] neg_hi:[0,1]
3131
// CHECK: :[[@LINE-1]]:28: error: not a valid operand
32+
33+
34+
v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2
35+
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
36+
37+
v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] blgp:2
38+
// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 2
39+
40+
v_mfma_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] cbsz:2 blgp:2
41+
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
42+
// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 2
43+
44+
45+
v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2
46+
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
47+
48+
v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] blgp:2
49+
// CHECK: :[[@LINE-1]]:46: error: wrong register tuple size for blgp value 2
50+
51+
v_mfma_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] cbsz:2 blgp:2
52+
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
53+
// CHECK: :[[@LINE-2]]:46: error: wrong register tuple size for blgp value 2
54+
55+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:2
56+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2
57+
58+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 cbsz:2
59+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2
60+
61+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 blgp:2
62+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2
63+
64+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:19], a[0:3] v20, v21 blgp:2
65+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2
66+
67+
v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] cbsz:2
68+
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
69+
70+
v_mfma_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] blgp:2
71+
// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 2
72+
73+
v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] cbsz:2
74+
// CHECK: :[[@LINE-1]]:37: error: wrong register tuple size for cbsz value 2
75+
76+
v_mfma_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] blgp:2
77+
// CHECK: :[[@LINE-1]]:47: error: wrong register tuple size for blgp value 2
78+
79+
v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] v32, v33 cbsz:2
80+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2
81+
82+
v_mfma_scale_f32_32x32x64_f8f6f4 a[0:15], a[16:23], a[24:31], a[0:15] v32, v33 cbsz:2
83+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2
84+
85+
86+
87+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21
88+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0
89+
90+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:1
91+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1
92+
93+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:2
94+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2
95+
96+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3
97+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3
98+
99+
100+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21
101+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0
102+
103+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:1
104+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1
105+
106+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:2
107+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2
108+
109+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:15], v[0:3] v20, v21 blgp:3
110+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3
111+
112+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21
113+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 0
114+
115+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:1
116+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 1
117+
118+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:2
119+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 2
120+
121+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:7], a[12:19], a[0:3] v20, v21 cbsz:3
122+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3
123+
124+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21
125+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 0
126+
127+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:1
128+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 1
129+
130+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:2
131+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 2
132+
133+
v_mfma_scale_f32_16x16x128_f8f6f4 a[0:3], a[4:11], a[12:15], a[0:3] v20, v21 blgp:3
134+
// CHECK: :[[@LINE-1]]:52: error: wrong register tuple size for blgp value 3
135+
136+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:7], v[12:19], v[0:3] v20, v21 cbsz:3
137+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3
138+
139+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:3
140+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 3
141+
142+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:7], v[0:3] v20, v21 blgp:3
143+
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3
144+
145+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:3
146+
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 3
147+
148+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:9], v[12:19], v[0:3] v20, v21 cbsz:4
149+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4
150+
151+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[4:11], v[12:19], v[0:3] v20, v21 cbsz:4
152+
// CHECK: :[[@LINE-1]]:43: error: wrong register tuple size for cbsz value 4
153+
154+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:9], v[0:3] v20, v21 blgp:4
155+
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4
156+
157+
v_mfma_scale_f32_16x16x128_f8f6f4 v[0:3], v[12:19], v[4:11], v[0:3] v20, v21 blgp:4
158+
// CHECK: :[[@LINE-1]]:53: error: wrong register tuple size for blgp value 4

0 commit comments

Comments
 (0)