tree e08b00e2e149e437475c3f207f76b05a9b280edc
parent 6c0bfc0226e15f75121783dc754a78a1e100cfaf
author Manish Gupta <manigupta@google.com> 1674458785 -0800
committer GitHub <noreply@github.com> 1674458785 -0800
gpgsig -----BEGIN PGP SIGNATURE-----
 
 wsBcBAABCAAQBQJjzjahCRBK7hj4Ov3rIwAAfiYIAJ8996Z9od+ry2uUS9Ba7PnD
 /WZGcBNxTGMj/ZadjNJV0/5j9Ul5ItodbG3SSn1DTafXE4qppApkyaf6qnSJHCjo
 Wx7rXz4Pb9OJkPD5MGQP1MNB8TCu+acj1+capTSDGgiWhJr/VZqn199bR9McKreD
 fAJYnlpUFTZ++5F5qpHWsJdsHS4ERZOMuCbyzpx6V2bW6EQwmXGPmdoawJHa0qJC
 P/Giq2Br+0yDLCN9+ilRMSpgssu7dRwWZXQhAyY32m2uvz1Bc8FgNuU5abxcfX1r
 O8WGSl/n5yObZwD1zlUOpLOMvmhKAHIJAKFFv5+q9sOZPGCvQjoQIiHVEm6s71s=
 =YQYA
 -----END PGP SIGNATURE-----
 

Adds Native Tensor Core (F16) Support [mma.sync.16816.f16.f16 and ldmatrix] (#11817)

Adds native Tensor Core (F16) support for NVIDIA GPUs. Specifically,
adds support to target
`mma.sync.aligned.m16n8k16.row.col.f16.f16.f16.f16` and `ldmatrix`
instructions on NVIDIA A100 GPUs. For now, these instructions are not
enabled by default and requires using
(`--iree-codegen-llvmgpu-use-mma-sync=true`).

IREE runs row-row matmul and the native Tensor Core support requires
`ldmatrix` for operandA and `ldmatrix.trans` for operandB.
`ldmatrix.trans` transposes the data while loading to use Tensor Core
`mma*row.col*` instruction.

There are LLVM/MLIR upstream pathes and previous IREE PRs that enable
targeting native Tensor Core F16 support.

A few LLVM/MLIR upstream patches:
- [[mlir][NVGPU] Adding Support for cp_async_zfill via Inline
Asm](https://reviews.llvm.org/D132269)
- [[mlir][NVGPU] Handle Native mma.sync sizes and ldmatrix(x4) for
matrixB](https://reviews.llvm.org/D135749)
- [[mlir][NVGPU] Fix affine maps computing indices for LdMatrixOp
srcMemref](https://reviews.llvm.org/D138978)
- [[mlir][NVGPU] Fix NVGPU Bazel
Layering](https://reviews.llvm.org/D141982)


Previous IREE PRs:
- [Support GEMM Pipelining *without* Epilogue
Peeling](https://github.com/iree-org/iree/pull/10388)
- [Add support for GEMM e2e Test For CUDA backend on F16
input](https://github.com/iree-org/iree/pull/10842)