[mlir][MemRef] Add patterns to extract address computations
authorQuentin Colombet <quentin.colombet@gmail.com>
Thu, 23 Mar 2023 14:41:14 +0000 (15:41 +0100)
committerQuentin Colombet <quentin.colombet@gmail.com>
Tue, 28 Mar 2023 11:52:29 +0000 (13:52 +0200)
commit54cda2ec976a89fcf5157d78479a576b09922df7
treeac350e2c68074b3853a67aedbcc9a65fb15a57e8
parent86ce609d3f420e55fc2c282560acf8f71deaddbf
[mlir][MemRef] Add patterns to extract address computations

This patch adds patterns to rewrite memory accesses such that the resulting
accesses are only using a base pointer.
E.g.,
```mlir
memref.load %base[%off0, ...]
```

Will be rewritten in:
```mlir
%new_base = memref.subview %base[%off0,...][1,...][1,...]
memref.load %new_base[%c0,...]
```

The idea behind these patterns is to offer a way to more gradually lower
address computations.

These patterns are the exact opposite of FoldMemRefAliasOps.
I've implemented the support of only five operations in this patch:
- memref.load
- memref.store
- nvgpu.ldmatrix
- vector.transfer_read
- vector.transfer_write

Going forward we may want to provide an interface for these rewritings (and
the ones in FoldMemRefAliasOps.)
One step at a time!

Differential Revision: https://reviews.llvm.org/D146724
mlir/include/mlir/Dialect/MemRef/TransformOps/MemRefTransformOps.td
mlir/include/mlir/Dialect/MemRef/Transforms/Transforms.h [new file with mode: 0644]
mlir/lib/Dialect/MemRef/TransformOps/CMakeLists.txt
mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp
mlir/lib/Dialect/MemRef/Transforms/CMakeLists.txt
mlir/lib/Dialect/MemRef/Transforms/ExtractAddressComputations.cpp [new file with mode: 0644]
mlir/test/Dialect/MemRef/extract-address-computations.mlir [new file with mode: 0644]
utils/bazel/llvm-project-overlay/mlir/BUILD.bazel