[mlir][Transform] Introduce nvgpu transform extensions
authorNicolas Vasilache <nicolasvasilache@users.noreply.github.com>
Wed, 21 Jun 2023 12:01:15 +0000 (12:01 +0000)
committerNicolas Vasilache <nicolasvasilache@users.noreply.github.com>
Mon, 26 Jun 2023 16:21:28 +0000 (16:21 +0000)
commit40deed40ae77ba22f7c72693903752ab6bfeb4e7
treee608bb8f076a9a91c5847adedda44e5136df9598
parent7d2d693cac998391ac1d0df1e98d655e21c24f53
[mlir][Transform] Introduce nvgpu transform extensions

Mapping to NVGPU operations such as mma.sync with mixed precision and ldmatrix with transposes and
various data types involves complex matchings from low-level IR.
This is akin to raising complex patterns after unnecessarily having lost structural information.
To avoid such unnecessary complexity, introduce a direct mapping step from a matmul on memrefs
to distributed NVGPU vector abstractions.
In this context, mapping to specific mma.sync operations is trivial and consists in simply
translating the documentation into indexing expressions.

Correctness is demonstrated with an end-to-end integration test.

Differential Revision: https://reviews.llvm.org/D153420
mlir/include/mlir/Dialect/NVGPU/CMakeLists.txt
mlir/include/mlir/Dialect/NVGPU/TransformOps/CMakeLists.txt [new file with mode: 0644]
mlir/include/mlir/Dialect/NVGPU/TransformOps/NVGPUTransformOps.h [new file with mode: 0644]
mlir/include/mlir/Dialect/NVGPU/TransformOps/NVGPUTransformOps.td [new file with mode: 0644]
mlir/include/mlir/InitAllDialects.h
mlir/lib/Dialect/NVGPU/CMakeLists.txt
mlir/lib/Dialect/NVGPU/TransformOps/CMakeLists.txt [new file with mode: 0644]
mlir/lib/Dialect/NVGPU/TransformOps/NVGPUTransformOps.cpp [new file with mode: 0644]
mlir/test/Dialect/NVGPU/transform-matmul-to-nvvm.mlir [new file with mode: 0644]
mlir/test/Integration/GPU/CUDA/TensorCore/transform-mma-sync-matmul-f32.mlir [new file with mode: 0644]
utils/bazel/llvm-project-overlay/mlir/BUILD.bazel