Commit
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
[mlir][NVGPU][NFC] Clean up code structure
* Move passes to `Transforms` directory. * Add `Utils.h` (will be utilized in a subsequent change). Differential Revision: https://reviews.llvm.org/D155427
- Loading branch information
1 parent
92542f2
commit a4f4d82
Showing
11 changed files
with
95 additions
and
59 deletions.
There are no files selected for viewing
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -1,10 +1,3 @@ | ||
| add_subdirectory(IR) | ||
| add_subdirectory(TransformOps) | ||
|
|
||
| set(LLVM_TARGET_DEFINITIONS Passes.td) | ||
| mlir_tablegen(Passes.h.inc -gen-pass-decls -name NVGPU) | ||
| mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix NVGPU) | ||
| mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix NVGPU) | ||
| add_public_tablegen_target(MLIRNVGPUPassIncGen) | ||
|
|
||
| add_mlir_doc(Passes NVGPUPasses ./ -gen-pass-doc) | ||
| add_subdirectory(Transforms) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,7 @@ | ||
| set(LLVM_TARGET_DEFINITIONS Passes.td) | ||
| mlir_tablegen(Passes.h.inc -gen-pass-decls -name NVGPU) | ||
| mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix NVGPU) | ||
| mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix NVGPU) | ||
| add_public_tablegen_target(MLIRNVGPUPassIncGen) | ||
|
|
||
| add_mlir_doc(Passes NVGPUPasses ./ -gen-pass-doc) |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
File renamed without changes.
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,21 @@ | ||
| //===- Utils.h - Transform utilities -----------------------------*- C++-*-===// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include "mlir/IR/Operation.h" | ||
|
|
||
| namespace mlir { | ||
| namespace nvgpu { | ||
|
|
||
| /// Get the indices that the given load/store operation is operating on. | ||
| Operation::operand_range getIndices(Operation *op); | ||
|
|
||
| /// Set the indices that the given load/store operation is operating on. | ||
| void setIndices(Operation *op, ArrayRef<Value> indices); | ||
|
|
||
| } // namespace nvgpu | ||
| } // namespace mlir |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,48 @@ | ||
| //===- Utils.cpp - Transform utilities ------------------------------------===// | ||
| // | ||
| // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. | ||
| // See https://llvm.org/LICENSE.txt for license information. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| #include "mlir/Dialect/NVGPU/Transforms/Utils.h" | ||
|
|
||
| #include "mlir/Dialect/MemRef/IR/MemRef.h" | ||
| #include "mlir/Dialect/NVGPU/IR/NVGPUDialect.h" | ||
| #include "mlir/Dialect/Vector/IR/VectorOps.h" | ||
|
|
||
| using namespace mlir; | ||
| using namespace mlir::nvgpu; | ||
|
|
||
| Operation::operand_range nvgpu::getIndices(Operation *op) { | ||
| if (auto ldmatrixOp = dyn_cast<LdMatrixOp>(op)) | ||
| return ldmatrixOp.getIndices(); | ||
| if (auto copyOp = dyn_cast<DeviceAsyncCopyOp>(op)) | ||
| return copyOp.getDstIndices(); | ||
| if (auto loadOp = dyn_cast<memref::LoadOp>(op)) | ||
| return loadOp.getIndices(); | ||
| if (auto storeOp = dyn_cast<memref::StoreOp>(op)) | ||
| return storeOp.getIndices(); | ||
| if (auto vectorReadOp = dyn_cast<vector::LoadOp>(op)) | ||
| return vectorReadOp.getIndices(); | ||
| if (auto vectorStoreOp = dyn_cast<vector::StoreOp>(op)) | ||
| return vectorStoreOp.getIndices(); | ||
| llvm_unreachable("unsupported op type"); | ||
| } | ||
|
|
||
| void nvgpu::setIndices(Operation *op, ArrayRef<Value> indices) { | ||
| if (auto ldmatrixOp = dyn_cast<LdMatrixOp>(op)) | ||
| return ldmatrixOp.getIndicesMutable().assign(indices); | ||
| if (auto copyOp = dyn_cast<DeviceAsyncCopyOp>(op)) | ||
| return copyOp.getDstIndicesMutable().assign(indices); | ||
| if (auto loadOp = dyn_cast<memref::LoadOp>(op)) | ||
| return loadOp.getIndicesMutable().assign(indices); | ||
| if (auto storeOp = dyn_cast<memref::StoreOp>(op)) | ||
| return storeOp.getIndicesMutable().assign(indices); | ||
| if (auto vectorReadOp = dyn_cast<vector::LoadOp>(op)) | ||
| return vectorReadOp.getIndicesMutable().assign(indices); | ||
| if (auto vectorStoreOp = dyn_cast<vector::StoreOp>(op)) | ||
| return vectorStoreOp.getIndicesMutable().assign(indices); | ||
| llvm_unreachable("unsupported op type"); | ||
| } |
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters