[PATCH] D122044: [AMDGPU] New gfx940 mfma instructions
Jay Foad via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
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
More information about the llvm-commits
mailing list