[PATCH] D122044: [AMDGPU] New gfx940 mfma instructions

Thu Mar 24 09:43:05 PDT 2022

foad added a comment.

In D122044#3405607 <https://reviews.llvm.org/D122044#3405607>, @rampitec wrote:

> In D122044#3404924 <https://reviews.llvm.org/D122044#3404924>, @foad wrote:
>
>>> 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.

================
Comment at: llvm/lib/Target/AMDGPU/VOP3PInstructions.td:539
+let Predicates = [isGFX940Plus] in {
+  defm V_MFMA_I32_32X32X16I8       : MAIInst<"v_mfma_i32_32x32x16i8",       "I32_I64_X32",    int_amdgcn_mfma_i32_32x32x16_i8>;
+  defm V_MFMA_I32_16X16X32I8       : MAIInst<"v_mfma_i32_16x16x32i8",       "I32_I64_X16",    int_amdgcn_mfma_i32_16x16x32_i8>;
----------------
I am still confused why the instruction name in `MAIInst<"v_mfma_i32_32x32x16i8"` has no underscore before `i8`...

================
Comment at: llvm/lib/Target/AMDGPU/VOP3PInstructions.td:745

+defm V_MFMA_I32_32X32X16I8       : VOP3P_Real_MFMA_gfx940 <0x56, "v_mfma_i32_32x32x16_i8">;
+defm V_MFMA_I32_16X16X32I8       : VOP3P_Real_MFMA_gfx940 <0x57, "v_mfma_i32_16x16x32_i8">;
----------------
... but here the same instruction name does have an underscore.

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D122044/new/

https://reviews.llvm.org/D122044