Add support for TensorCore F32 emulation through TF32

[ThomasRaoux]

Cutlass added support for float32 emulation using TF32 tensorcore operations. In MLIR we have representations for mma.sync for TF32. We should differentiate mma.sync for float32 and tf32 and have a lowering pattern from mma.sycn f32 to a code sequence of mma.sync tf32. This would go in nvgpu dialect transformations in MLIR and can then be used for IREE.

Possible break down:

1. Add an optional attribute in mma.sync op to specify that the op uses TF32 precision
2. Create a pattern in NVGPU dialect transforms to lower mma.sync f32 without the tf32 attribute to a code sequence of mma.sync with the tf32 attribute.

[manishucsd]

self tag @manishucsd

[manishucsd]

OptionalAttr: Present OR Not Present. Follow bypassL1 which uses the same datatype

Present : Allowed to use TF32 lowering given that the data type for the F32 (update the verifier) Not present : TF32 lowering is not allowed

Precision information comes from the users:

• Per operations

  • This is not possible as we do not have a native TF32 type in LLVM/MLIR

• Global model level. There are three choices for F32 input:

  1. No Tensor cores
     • This decision happens much earlier and the codegen shouldn't even reach nvgpu.mmasync for it
  2. TF32 Tensor cores
  3. 3xTF32 Tensor cores

Choice and 2. and 3. are enabled using an enum passed to the pattern rewriter in populateMmaSyncF32ToT32Patterns

[manishucsd]

In progress here: https://reviews.llvm.org/D130294

[manishucsd]

The next steps here are to use the added OptionalAttr tf32Enabled and enum MmaSyncF32Lowering to enable support for TF32x3 a.k.a. F32 emulation through TensorCores.

[allieculp]

@manishucsd Is this still open? Still P1 work?