This is an archive of the discontinued LLVM Phabricator instance.

[AMDGPU] New gfx940 mfma instructions
ClosedPublic

Authored by rampitec on Mar 18 2022, 3:06 PM.

Diff Detail

Unit TestsFailed

Event Timeline

rampitec created this revision.Mar 18 2022, 3:06 PM
Herald added a project: Restricted Project. · View Herald TranscriptMar 18 2022, 3:06 PM
rampitec requested review of this revision.Mar 18 2022, 3:06 PM
Herald added a project: Restricted Project. · View Herald TranscriptMar 18 2022, 3:06 PM
Herald added a subscriber: wdng. · View Herald Transcript
foad added inline comments.Mar 21 2022, 8:02 AM
clang/include/clang/Basic/BuiltinsAMDGPU.def
308

Why do the new ones have _ before the i8/xf32 suffix? None of the old ones have it. What does xf32 mean?

rampitec marked an inline comment as done.Mar 21 2022, 10:20 AM
rampitec added inline comments.
clang/include/clang/Basic/BuiltinsAMDGPU.def
308

xf32 suffix indicates data in TF32 format: 8-bit exponent with FP16’s 10-bit mantissa. Uses 32-bits of storage for 19 bits of data.

MFMA instructions got new names in the gfx940 (see D121741). It now includes block size which was implicit before. That is because some of them operate on a 2x-16x blocks. Together adding an 'x' suffix and new size factor made the names ambiguous so the decision was made to separate fields with an underscore. The old names are preserved as aliases for compatibility with the existing programs. The same is true for the builtins as these are used in the existing programs. Going forward a new names will be used.

rampitec updated this revision to Diff 417103.Mar 21 2022, 2:43 PM
rampitec marked an inline comment as done.

Added 2 more tests.

rampitec updated this revision to Diff 417113.Mar 21 2022, 3:05 PM

One more clang test.

foad added a comment.Mar 24 2022, 3:51 AM

xf32 suffix indicates data in TF32 format

Why not tf32 suffix?

xf32 suffix indicates data in TF32 format

Why not tf32 suffix?

The names come from the HW spec, it is not something I can change.

foad added a comment.Mar 24 2022, 9:43 AM

xf32 suffix indicates data in TF32 format

Why not tf32 suffix?

The names come from the HW spec, it is not something I can change.

OK, if the intrinsic name matches the instruction name then that is obviously fine.

llvm/lib/Target/AMDGPU/VOP3PInstructions.td
539

I am still confused why the instruction name in MAIInst<"v_mfma_i32_32x32x16i8" has no underscore before i8...

745

... but here the same instruction name does have an underscore.

rampitec marked an inline comment as done.Mar 24 2022, 10:00 AM
rampitec added inline comments.
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
745

These are asm aliases. The same aliases with and w/o underscores supported by the SP3. In fact Justin and me decided we want these aliases to make it easier to adopt old kernels.

rampitec marked an inline comment as done.Mar 24 2022, 10:09 AM
rampitec added inline comments.
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
745

Look at the ai-gfx940.s lines 282 and 286 which tests both names are accepted.

foad added inline comments.Mar 24 2022, 10:15 AM
llvm/lib/Target/AMDGPU/VOP3PInstructions.td
745

OK. No further comments from me :)

kzhuravl accepted this revision.Mar 24 2022, 11:50 AM

LGTM, unless @foad has any additional questions

This revision is now accepted and ready to land.Mar 24 2022, 11:50 AM
This revision was landed with ongoing or failed builds.Mar 24 2022, 12:13 PM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptMar 24 2022, 12:13 PM
Herald added a subscriber: cfe-commits. · View Herald Transcript