Step 4/n (n=4). Rename Conversion folder to Codegen and clean up pass names. (#6267)
The Conversion folder is renamed to Codegen (Conversion was a legacy
artefact, and Codegen is a more accurate description of the
folder). Also
CPU, SPIRV and LLVMGPU folders/passes are renamed to drop the
"LinalgTo" prefix.
Flags are renamed to use
"iree-codegen-*" for common passes
"iree-llvmcpu-*" for CPU backend passes
"iree-spirv-*" for SPIRV backend passes
"iree-llvmgpu-*" for LLVMGPU backend passes.
Filed renamed to be consistent
File name reflects name of the pass implemented in the file.
Nit: declerations in .h and .td file made alphabetical.
diff --git a/integrations/tensorflow/iree_tf_compiler/MHLO/BUILD b/integrations/tensorflow/iree_tf_compiler/MHLO/BUILD
index 6f00831..1b07b9c 100644
--- a/integrations/tensorflow/iree_tf_compiler/MHLO/BUILD
+++ b/integrations/tensorflow/iree_tf_compiler/MHLO/BUILD
@@ -24,7 +24,7 @@
"IREE_COMPILER_TENSORFLOW_ENABLED",
],
deps = [
- "@iree//iree/compiler/Conversion:PassHeaders",
+ "@iree//iree/compiler/Codegen:PassHeaders",
"@iree//iree/compiler/Dialect/Flow/IR",
"@iree//iree/compiler/Dialect/Flow/Transforms",
"@iree//iree/compiler/Dialect/IREE/IR",
diff --git a/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp b/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp
index 52b0a95..5248fa7 100644
--- a/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp
+++ b/integrations/tensorflow/iree_tf_compiler/MHLO/Passes.cpp
@@ -6,7 +6,7 @@
#include "iree_tf_compiler/MHLO/Passes.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/Flow/Transforms/Passes.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "iree/compiler/InputConversion/MHLO/Passes.h"
diff --git a/iree/compiler/Conversion/BUILD b/iree/compiler/Codegen/BUILD
similarity index 83%
rename from iree/compiler/Conversion/BUILD
rename to iree/compiler/Codegen/BUILD
index 0dc9fb3..7034811 100644
--- a/iree/compiler/Conversion/BUILD
+++ b/iree/compiler/Codegen/BUILD
@@ -44,16 +44,16 @@
)
cc_library(
- name = "Conversion",
+ name = "Codegen",
srcs = [
"Passes.cpp",
],
deps = [
":PassHeaders",
":PassesIncGen",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/LinalgToLLVM",
- "//iree/compiler/Conversion/LinalgToLLVMGPU",
- "//iree/compiler/Conversion/LinalgToSPIRV",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/LLVMCPU",
+ "//iree/compiler/Codegen/LLVMGPU",
+ "//iree/compiler/Codegen/SPIRV",
],
)
diff --git a/iree/compiler/Conversion/CMakeLists.txt b/iree/compiler/Codegen/CMakeLists.txt
similarity index 78%
rename from iree/compiler/Conversion/CMakeLists.txt
rename to iree/compiler/Codegen/CMakeLists.txt
index 55e5bb1..1bca592 100644
--- a/iree/compiler/Conversion/CMakeLists.txt
+++ b/iree/compiler/Codegen/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/BUILD #
+# iree/compiler/Codegen/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -36,21 +36,21 @@
iree_cc_library(
NAME
- Conversion
+ Codegen
SRCS
"Passes.cpp"
DEPS
::PassHeaders
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::LinalgToLLVM
- iree::compiler::Conversion::LinalgToLLVMGPU
- iree::compiler::Conversion::LinalgToSPIRV
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::LLVMCPU
+ iree::compiler::Codegen::LLVMGPU
+ iree::compiler::Codegen::SPIRV
PUBLIC
)
### BAZEL_TO_CMAKE_PRESERVES_ALL_CONTENT_BELOW_THIS_LINE ###
# TODO: For some reason, these dependencies are not being added automatically.
add_dependencies(
- iree_compiler_Conversion_PassHeaders
- iree_compiler_Conversion_PassesIncGen
+ iree_compiler_Codegen_PassHeaders
+ iree_compiler_Codegen_PassesIncGen
)
diff --git a/iree/compiler/Conversion/Common/BUILD b/iree/compiler/Codegen/Common/BUILD
similarity index 91%
rename from iree/compiler/Conversion/Common/BUILD
rename to iree/compiler/Codegen/Common/BUILD
index 0d81579..afd1b47 100644
--- a/iree/compiler/Conversion/Common/BUILD
+++ b/iree/compiler/Codegen/Common/BUILD
@@ -43,10 +43,10 @@
"VectorizeConv.cpp",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common:FoldTensorExtractOpIncGen",
- "//iree/compiler/Conversion/Transforms",
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common:FoldTensorExtractOpIncGen",
+ "//iree/compiler/Codegen/Transforms",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/IR",
"//iree/compiler/Dialect/IREE/IR",
diff --git a/iree/compiler/Conversion/Common/CMakeLists.txt b/iree/compiler/Codegen/Common/CMakeLists.txt
similarity index 86%
rename from iree/compiler/Conversion/Common/CMakeLists.txt
rename to iree/compiler/Codegen/Common/CMakeLists.txt
index 8dc80bb..ae91a28 100644
--- a/iree/compiler/Conversion/Common/CMakeLists.txt
+++ b/iree/compiler/Codegen/Common/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/Common/BUILD #
+# iree/compiler/Codegen/Common/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -51,10 +51,10 @@
MLIRTensor
MLIRTransforms
MLIRVector
- iree::compiler::Conversion::Common::FoldTensorExtractOpIncGen
- iree::compiler::Conversion::PassHeaders
- iree::compiler::Conversion::Transforms
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Common::FoldTensorExtractOpIncGen
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::Transforms
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::IREE::IR
diff --git a/iree/compiler/Conversion/Common/CleanupBufferAllocViewPass.cpp b/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
similarity index 97%
rename from iree/compiler/Conversion/Common/CleanupBufferAllocViewPass.cpp
rename to iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
index 6dc13a6..fffe268 100644
--- a/iree/compiler/Conversion/Common/CleanupBufferAllocViewPass.cpp
+++ b/iree/compiler/Codegen/Common/CleanupBufferAllocViewPass.cpp
@@ -12,8 +12,8 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
diff --git a/iree/compiler/Conversion/Common/DemoteF32ToF16.cpp b/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
similarity index 98%
rename from iree/compiler/Conversion/Common/DemoteF32ToF16.cpp
rename to iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
index e62fb71..d5487c6 100644
--- a/iree/compiler/Conversion/Common/DemoteF32ToF16.cpp
+++ b/iree/compiler/Codegen/Common/DemoteF32ToF16.cpp
@@ -7,8 +7,8 @@
#include <memory>
#include <utility>
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/IREE/IR/IREETypes.h"
#include "llvm/ADT/APFloat.h"
diff --git a/iree/compiler/Conversion/Common/FlattenMemRefSubspanPass.cpp b/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
similarity index 99%
rename from iree/compiler/Conversion/Common/FlattenMemRefSubspanPass.cpp
rename to iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
index ca4d3b8..a120d8a 100644
--- a/iree/compiler/Conversion/Common/FlattenMemRefSubspanPass.cpp
+++ b/iree/compiler/Codegen/Common/FlattenMemRefSubspanPass.cpp
@@ -33,8 +33,8 @@
#include <memory>
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
diff --git a/iree/compiler/Conversion/Common/FoldTensorExtractOp.td b/iree/compiler/Codegen/Common/FoldTensorExtractOp.td
similarity index 73%
rename from iree/compiler/Conversion/Common/FoldTensorExtractOp.td
rename to iree/compiler/Codegen/Common/FoldTensorExtractOp.td
index bf75e7b..84d8be1 100644
--- a/iree/compiler/Conversion/Common/FoldTensorExtractOp.td
+++ b/iree/compiler/Codegen/Common/FoldTensorExtractOp.td
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_CONVERSION_COMMON_FOLDTENSOREXTRACTOP
-#define IREE_COMPILER_CONVERSION_COMMON_FOLDTENSOREXTRACTOP
+#ifndef IREE_COMPILER_CODEGEN_COMMON_FOLDTENSOREXTRACTOP
+#define IREE_COMPILER_CODEGEN_COMMON_FOLDTENSOREXTRACTOP
include "mlir/Dialect/MemRef/IR/MemRefOps.td"
include "mlir/Dialect/Tensor/IR/TensorOps.td"
@@ -15,4 +15,4 @@
def : Pat<(Tensor_ExtractOp (TensorLoadOp $value), $indices),
(LoadOp $value, $indices)>;
-#endif // IREE_COMPILER_CONVERSION_COMMON_FOLDTENSOREXTRACTOP
+#endif // IREE_COMPILER_CODEGEN_COMMON_FOLDTENSOREXTRACTOP
diff --git a/iree/compiler/Conversion/Common/FoldTensorExtractOpPass.cpp b/iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp
similarity index 93%
rename from iree/compiler/Conversion/Common/FoldTensorExtractOpPass.cpp
rename to iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp
index 14e212d..ac4b3e7 100644
--- a/iree/compiler/Conversion/Common/FoldTensorExtractOpPass.cpp
+++ b/iree/compiler/Codegen/Common/FoldTensorExtractOpPass.cpp
@@ -3,8 +3,8 @@
// Licensed 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 "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/StandardOps/IR/Ops.h"
#include "mlir/Dialect/Tensor/IR/Tensor.h"
@@ -16,7 +16,7 @@
namespace iree_compiler {
namespace {
-#include "iree/compiler/Conversion/Common/FoldTensorExtractOp.cpp.inc"
+#include "iree/compiler/Codegen/Common/FoldTensorExtractOp.cpp.inc"
}
namespace {
diff --git a/iree/compiler/Conversion/Common/ForOpCanonicalizationPass.cpp b/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
similarity index 98%
rename from iree/compiler/Conversion/Common/ForOpCanonicalizationPass.cpp
rename to iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
index aa93b50..4bbb3e2 100644
--- a/iree/compiler/Conversion/Common/ForOpCanonicalizationPass.cpp
+++ b/iree/compiler/Codegen/Common/ForOpCanonicalizationPass.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Dialect/SCF/SCF.h"
#include "mlir/Dialect/Vector/VectorOps.h"
#include "mlir/IR/BlockAndValueMapping.h"
diff --git a/iree/compiler/Conversion/Common/LinalgBufferizePass.cpp b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
similarity index 99%
rename from iree/compiler/Conversion/Common/LinalgBufferizePass.cpp
rename to iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
index 7edcdfd..172621c 100644
--- a/iree/compiler/Conversion/Common/LinalgBufferizePass.cpp
+++ b/iree/compiler/Codegen/Common/LinalgBufferizePass.cpp
@@ -37,10 +37,10 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/Flow/IR/FlowTypes.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
diff --git a/iree/compiler/Conversion/Common/OptimizeVectorTransferPass.cpp b/iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp
similarity index 97%
rename from iree/compiler/Conversion/Common/OptimizeVectorTransferPass.cpp
rename to iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp
index da04a20..a392321 100644
--- a/iree/compiler/Conversion/Common/OptimizeVectorTransferPass.cpp
+++ b/iree/compiler/Codegen/Common/OptimizeVectorTransferPass.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/StandardOps/IR/Ops.h"
#include "mlir/Dialect/Vector/VectorOps.h"
diff --git a/iree/compiler/Conversion/Common/SetNumWorkgroupsPass.cpp b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
similarity index 94%
rename from iree/compiler/Conversion/Common/SetNumWorkgroupsPass.cpp
rename to iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
index 67af962..7f3a8b3 100644
--- a/iree/compiler/Conversion/Common/SetNumWorkgroupsPass.cpp
+++ b/iree/compiler/Codegen/Common/SetNumWorkgroupsPass.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
diff --git a/iree/compiler/Conversion/Common/VectorizeConv.cpp b/iree/compiler/Codegen/Common/VectorizeConv.cpp
similarity index 99%
rename from iree/compiler/Conversion/Common/VectorizeConv.cpp
rename to iree/compiler/Codegen/Common/VectorizeConv.cpp
index 8f02df5..181b44e 100644
--- a/iree/compiler/Conversion/Common/VectorizeConv.cpp
+++ b/iree/compiler/Codegen/Common/VectorizeConv.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/Support/Debug.h"
diff --git a/iree/compiler/Conversion/Common/test/BUILD b/iree/compiler/Codegen/Common/test/BUILD
similarity index 97%
rename from iree/compiler/Conversion/Common/test/BUILD
rename to iree/compiler/Codegen/Common/test/BUILD
index 402c0e5..0a2a705 100644
--- a/iree/compiler/Conversion/Common/test/BUILD
+++ b/iree/compiler/Codegen/Common/test/BUILD
@@ -21,6 +21,7 @@
[
"affinemin_canonicalization.mlir",
"canonicalize_interface_load_store.mlir",
+ "dead_alloc.mlir",
"f32Tof16.mlir",
"flatten_memref_subspan.mlir",
"fold_tensor_extract_op.mlir",
diff --git a/iree/compiler/Conversion/Common/test/CMakeLists.txt b/iree/compiler/Codegen/Common/test/CMakeLists.txt
similarity index 93%
rename from iree/compiler/Conversion/Common/test/CMakeLists.txt
rename to iree/compiler/Codegen/Common/test/CMakeLists.txt
index bc5e14d..394633b 100644
--- a/iree/compiler/Conversion/Common/test/CMakeLists.txt
+++ b/iree/compiler/Codegen/Common/test/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/Common/test/BUILD #
+# iree/compiler/Codegen/Common/test/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -16,6 +16,7 @@
SRCS
"affinemin_canonicalization.mlir"
"canonicalize_interface_load_store.mlir"
+ "dead_alloc.mlir"
"f32Tof16.mlir"
"flatten_memref_subspan.mlir"
"fold_tensor_extract_op.mlir"
diff --git a/iree/compiler/Conversion/Common/test/affinemin_canonicalization.mlir b/iree/compiler/Codegen/Common/test/affinemin_canonicalization.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/affinemin_canonicalization.mlir
rename to iree/compiler/Codegen/Common/test/affinemin_canonicalization.mlir
diff --git a/iree/compiler/Conversion/Common/test/canonicalize_interface_load_store.mlir b/iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/canonicalize_interface_load_store.mlir
rename to iree/compiler/Codegen/Common/test/canonicalize_interface_load_store.mlir
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/dead_alloc.mlir b/iree/compiler/Codegen/Common/test/dead_alloc.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/dead_alloc.mlir
rename to iree/compiler/Codegen/Common/test/dead_alloc.mlir
diff --git a/iree/compiler/Conversion/Common/test/f32Tof16.mlir b/iree/compiler/Codegen/Common/test/f32Tof16.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/f32Tof16.mlir
rename to iree/compiler/Codegen/Common/test/f32Tof16.mlir
diff --git a/iree/compiler/Conversion/Common/test/flatten_memref_subspan.mlir b/iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/flatten_memref_subspan.mlir
rename to iree/compiler/Codegen/Common/test/flatten_memref_subspan.mlir
diff --git a/iree/compiler/Conversion/Common/test/fold_tensor_extract_op.mlir b/iree/compiler/Codegen/Common/test/fold_tensor_extract_op.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/fold_tensor_extract_op.mlir
rename to iree/compiler/Codegen/Common/test/fold_tensor_extract_op.mlir
diff --git a/iree/compiler/Conversion/Common/test/forop_canonicalization.mlir b/iree/compiler/Codegen/Common/test/forop_canonicalization.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/forop_canonicalization.mlir
rename to iree/compiler/Codegen/Common/test/forop_canonicalization.mlir
diff --git a/iree/compiler/Conversion/Common/test/linalg_bufferize.mlir b/iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/linalg_bufferize.mlir
rename to iree/compiler/Codegen/Common/test/linalg_bufferize.mlir
diff --git a/iree/compiler/Conversion/Common/test/remove_dead_allocs.mlir b/iree/compiler/Codegen/Common/test/remove_dead_allocs.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/remove_dead_allocs.mlir
rename to iree/compiler/Codegen/Common/test/remove_dead_allocs.mlir
diff --git a/iree/compiler/Conversion/Common/test/transpose_canonicalization.mlir b/iree/compiler/Codegen/Common/test/transpose_canonicalization.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/transpose_canonicalization.mlir
rename to iree/compiler/Codegen/Common/test/transpose_canonicalization.mlir
diff --git a/iree/compiler/Conversion/Common/test/vectorize_linalg_conv.mlir b/iree/compiler/Codegen/Common/test/vectorize_linalg_conv.mlir
similarity index 100%
rename from iree/compiler/Conversion/Common/test/vectorize_linalg_conv.mlir
rename to iree/compiler/Codegen/Common/test/vectorize_linalg_conv.mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVM/BUILD b/iree/compiler/Codegen/LLVMCPU/BUILD
similarity index 80%
rename from iree/compiler/Conversion/LinalgToLLVM/BUILD
rename to iree/compiler/Codegen/LLVMCPU/BUILD
index f18729e..746369e 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/BUILD
+++ b/iree/compiler/Codegen/LLVMCPU/BUILD
@@ -11,28 +11,27 @@
)
cc_library(
- name = "LinalgToLLVM",
+ name = "LLVMCPU",
srcs = [
"ConvertToLLVM.cpp",
"KernelDispatch.cpp",
- "LinalgToLLVMWorkgroupsVectorizationPass.cpp",
- "LinalgVectorizePass.cpp",
- "LowerExecutableTargetPass.cpp",
- "PadLinalgWorkgroupTiles.cpp",
+ "LLVMCPULowerExecutableTarget.cpp",
+ "LLVMCPUPadWorkgroupTiles.cpp",
+ "LLVMCPUPlanConvLoopOrder.cpp",
+ "LLVMCPUTilePadAndVectorize.cpp",
+ "LLVMCPUUnfuseFMAOps.cpp",
+ "LLVMCPUVectorization.cpp",
"Passes.cpp",
- "PlanConvLoopOrder.cpp",
- "TilePadAndVectorizeWorkgroups.cpp",
- "UnfuseFMAOps.cpp",
"VectorContractToAArch64InlineAsmOp.cpp",
],
hdrs = [
"KernelDispatch.h",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/Transforms",
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Transforms",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/IR",
"//iree/compiler/Dialect/HAL/IR:HALDialect",
diff --git a/iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt b/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
similarity index 77%
rename from iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt
rename to iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
index 38d3b40..78ced25 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/CMakeLists.txt
+++ b/iree/compiler/Codegen/LLVMCPU/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/LinalgToLLVM/BUILD #
+# iree/compiler/Codegen/LLVMCPU/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -12,20 +12,19 @@
iree_cc_library(
NAME
- LinalgToLLVM
+ LLVMCPU
HDRS
"KernelDispatch.h"
SRCS
"ConvertToLLVM.cpp"
"KernelDispatch.cpp"
- "LinalgToLLVMWorkgroupsVectorizationPass.cpp"
- "LinalgVectorizePass.cpp"
- "LowerExecutableTargetPass.cpp"
- "PadLinalgWorkgroupTiles.cpp"
+ "LLVMCPULowerExecutableTarget.cpp"
+ "LLVMCPUPadWorkgroupTiles.cpp"
+ "LLVMCPUPlanConvLoopOrder.cpp"
+ "LLVMCPUTilePadAndVectorize.cpp"
+ "LLVMCPUUnfuseFMAOps.cpp"
+ "LLVMCPUVectorization.cpp"
"Passes.cpp"
- "PlanConvLoopOrder.cpp"
- "TilePadAndVectorizeWorkgroups.cpp"
- "UnfuseFMAOps.cpp"
"VectorContractToAArch64InlineAsmOp.cpp"
DEPS
LLVMSupport
@@ -53,10 +52,10 @@
MLIRVector
MLIRVectorToLLVM
MLIRVectorToSCF
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::PassHeaders
- iree::compiler::Conversion::Transforms
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::Transforms
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::HAL::IR::HALDialect
diff --git a/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
similarity index 99%
rename from iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp
rename to iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
index 75eae97..23ededa 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/ConvertToLLVM.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/ConvertToLLVM.cpp
@@ -4,9 +4,9 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/IREE/IR/IREEDialect.h"
diff --git a/iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.cpp b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
similarity index 97%
rename from iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.cpp
rename to iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
index 435d684..1b77438 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.cpp
@@ -4,11 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h"
+#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "llvm/ADT/TypeSwitch.h"
diff --git a/iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.h
similarity index 79%
rename from iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h
rename to iree/compiler/Codegen/LLVMCPU/KernelDispatch.h
index 306d666..3275030 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h
+++ b/iree/compiler/Codegen/LLVMCPU/KernelDispatch.h
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_CONVERSION_LINALGTOLLVM_KERNELDISPATCH_H_
-#define IREE_COMPILER_CONVERSION_LINALGTOLLVM_KERNELDISPATCH_H_
+#ifndef IREE_COMPILER_CODEGEN_LLVMCPU_KERNELDISPATCH_H_
+#define IREE_COMPILER_CODEGEN_LLVMCPU_KERNELDISPATCH_H_
#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "mlir/IR/BuiltinOps.h"
@@ -28,4 +28,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_LINALGTOLLVM_KERNELDISPATCH_H_
+#endif // IREE_COMPILER_CODEGEN_LLVMCPU_KERNELDISPATCH_H_
diff --git a/iree/compiler/Conversion/LinalgToLLVM/LowerExecutableTargetPass.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
similarity index 89%
rename from iree/compiler/Conversion/LinalgToLLVM/LowerExecutableTargetPass.cpp
rename to iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
index 75a0017..ace9266 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/LowerExecutableTargetPass.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPULowerExecutableTarget.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
@@ -24,17 +24,19 @@
/// - first lower to scalar/native-vector code
/// - then convert to LLVM dialect.
/// In due course this could be used to generate code for all backends.
-class LowerExecutableTargetPass
- : public LowerExecutableTargetBase<LowerExecutableTargetPass> {
+class LLVMCPULowerExecutableTargetPass
+ : public LLVMCPULowerExecutableTargetBase<
+ LLVMCPULowerExecutableTargetPass> {
public:
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<IREE::HAL::HALDialect, linalg::LinalgDialect,
LLVM::LLVMDialect, vector::VectorDialect>();
}
- LowerExecutableTargetPass(bool vectorize = true)
+ LLVMCPULowerExecutableTargetPass(bool vectorize = true)
: lowerToVectors(vectorize) {}
- LowerExecutableTargetPass(const LowerExecutableTargetPass &pass) {}
+ LLVMCPULowerExecutableTargetPass(
+ const LLVMCPULowerExecutableTargetPass &pass) {}
void runOnOperation() override;
@@ -86,7 +88,7 @@
return input;
}
-void LowerExecutableTargetPass::runOnOperation() {
+void LLVMCPULowerExecutableTargetPass::runOnOperation() {
IREE::HAL::ExecutableTargetOp targetOp = getOperation();
ModuleOp moduleOp = targetOp.getInnerModule();
@@ -160,8 +162,8 @@
}
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLowerExecutableTargetPass(bool lowerToVectors) {
- return std::make_unique<LowerExecutableTargetPass>(lowerToVectors);
+createLLVMCPULowerExecutableTargetPass(bool lowerToVectors) {
+ return std::make_unique<LLVMCPULowerExecutableTargetPass>(lowerToVectors);
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVM/PadLinalgWorkgroupTiles.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUPadWorkgroupTiles.cpp
similarity index 94%
rename from iree/compiler/Conversion/LinalgToLLVM/PadLinalgWorkgroupTiles.cpp
rename to iree/compiler/Codegen/LLVMCPU/LLVMCPUPadWorkgroupTiles.cpp
index 1db1446..c5fad11 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/PadLinalgWorkgroupTiles.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUPadWorkgroupTiles.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "llvm/Support/Debug.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
@@ -16,7 +16,7 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-#define DEBUG_TYPE "iree-codegen-llvm-pad-linalg-workgroup-tiles"
+#define DEBUG_TYPE "iree-llvmcpu-pad-workgroup-tiles"
namespace mlir {
namespace iree_compiler {
@@ -224,8 +224,8 @@
}
};
-struct LLVMPadLinalgWorkgroupTilesPass
- : LLVMPadLinalgWorkgroupTilesBase<LLVMPadLinalgWorkgroupTilesPass> {
+struct LLVMCPUPadWorkgroupTilesPass
+ : LLVMCPUPadWorkgroupTilesBase<LLVMCPUPadWorkgroupTilesPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<linalg::LinalgDialect>();
}
@@ -238,8 +238,8 @@
};
} // namespace
-std::unique_ptr<OperationPass<FuncOp>> createLLVMPadLinalgWorkgroupTilesPass() {
- return std::make_unique<LLVMPadLinalgWorkgroupTilesPass>();
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUPadWorkgroupTilesPass() {
+ return std::make_unique<LLVMCPUPadWorkgroupTilesPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVM/PlanConvLoopOrder.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUPlanConvLoopOrder.cpp
similarity index 79%
rename from iree/compiler/Conversion/LinalgToLLVM/PlanConvLoopOrder.cpp
rename to iree/compiler/Codegen/LLVMCPU/LLVMCPUPlanConvLoopOrder.cpp
index c391d05..72a5ec5 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/PlanConvLoopOrder.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUPlanConvLoopOrder.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
#include "mlir/Pass/Pass.h"
@@ -16,8 +16,8 @@
namespace {
-struct LinalgToLLVMPlanConvLoopOrderPass
- : LinalgToLLVMPlanConvLoopOrderBase<LinalgToLLVMPlanConvLoopOrderPass> {
+struct LLVMCPUPlanConvLoopOrderPass
+ : LLVMCPUPlanConvLoopOrderBase<LLVMCPUPlanConvLoopOrderPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<linalg::LinalgDialect>();
}
@@ -26,7 +26,7 @@
} // namespace
-void LinalgToLLVMPlanConvLoopOrderPass::runOnOperation() {
+void LLVMCPUPlanConvLoopOrderPass::runOnOperation() {
auto funcOp = getOperation();
auto context = funcOp.getContext();
@@ -56,9 +56,8 @@
(void)applyPatternsAndFoldGreedily(funcOp, std::move(patterns));
}
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMPlanConvLoopOrderPass() {
- return std::make_unique<LinalgToLLVMPlanConvLoopOrderPass>();
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUPlanConvLoopOrderPass() {
+ return std::make_unique<LLVMCPUPlanConvLoopOrderPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVM/TilePadAndVectorizeWorkgroups.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp
similarity index 90%
rename from iree/compiler/Conversion/LinalgToLLVM/TilePadAndVectorizeWorkgroups.cpp
rename to iree/compiler/Codegen/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp
index f16e217..cd8fdfe 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/TilePadAndVectorizeWorkgroups.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUTilePadAndVectorize.cpp
@@ -4,11 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "mlir/Conversion/VectorToSCF/VectorToSCF.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
#include "mlir/Dialect/Linalg/Transforms/Hoisting.h"
@@ -49,9 +49,8 @@
} // namespace
namespace {
-struct LinalgToLLVMTilePadAndVectorizeWorkgroupsPass
- : public LinalgToLLVMTilePadAndVectorizeWorkgroupsBase<
- LinalgToLLVMTilePadAndVectorizeWorkgroupsPass> {
+struct LLVMCPUTilePadAndVectorizePass
+ : public LLVMCPUTilePadAndVectorizeBase<LLVMCPUTilePadAndVectorizePass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<linalg::LinalgDialect, memref::MemRefDialect,
vector::VectorDialect>();
@@ -60,7 +59,7 @@
};
} // namespace
-void LinalgToLLVMTilePadAndVectorizeWorkgroupsPass::runOnOperation() {
+void LLVMCPUTilePadAndVectorizePass::runOnOperation() {
MLIRContext *context = &getContext();
auto funcOp = getOperation();
@@ -201,9 +200,8 @@
}
}
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMTilePadAndVectorizeWorkgroupsPass() {
- return std::make_unique<LinalgToLLVMTilePadAndVectorizeWorkgroupsPass>();
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUTilePadAndVectorizePass() {
+ return std::make_unique<LLVMCPUTilePadAndVectorizePass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVM/UnfuseFMAOps.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUUnfuseFMAOps.cpp
similarity index 82%
rename from iree/compiler/Conversion/LinalgToLLVM/UnfuseFMAOps.cpp
rename to iree/compiler/Codegen/LLVMCPU/LLVMCPUUnfuseFMAOps.cpp
index f7e6fee..d4dbee5 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/UnfuseFMAOps.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUUnfuseFMAOps.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/Pass/Pass.h"
@@ -36,8 +36,8 @@
} // namespace
namespace {
-struct LinalgToLLVMUnfuseFMAOpsPass
- : LinalgToLLVMUnfuseFMAOpsBase<LinalgToLLVMUnfuseFMAOpsPass> {
+struct LLVMCPUUnfuseFMAOpsPass
+ : LLVMCPUUnfuseFMAOpsBase<LLVMCPUUnfuseFMAOpsPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<LLVM::LLVMDialect>();
}
@@ -50,7 +50,7 @@
patterns.insert<UnfusedFMAOpsPassConversion>(context);
}
-void LinalgToLLVMUnfuseFMAOpsPass::runOnOperation() {
+void LLVMCPUUnfuseFMAOpsPass::runOnOperation() {
auto funcOp = getOperation();
auto context = funcOp.getContext();
OwningRewritePatternList patterns(&getContext());
@@ -58,8 +58,8 @@
(void)applyPatternsAndFoldGreedily(funcOp, std::move(patterns));
}
-std::unique_ptr<OperationPass<FuncOp>> createLinalgToLLVMUnfuseFMAOpsPass() {
- return std::make_unique<LinalgToLLVMUnfuseFMAOpsPass>();
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUUnfuseFMAOpsPass() {
+ return std::make_unique<LLVMCPUUnfuseFMAOpsPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVM/LinalgToLLVMWorkgroupsVectorizationPass.cpp b/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
similarity index 91%
rename from iree/compiler/Conversion/LinalgToLLVM/LinalgToLLVMWorkgroupsVectorizationPass.cpp
rename to iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
index d30d8d5..7969658 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/LinalgToLLVMWorkgroupsVectorizationPass.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/LLVMCPUVectorization.cpp
@@ -4,11 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVM/KernelDispatch.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/LLVMCPU/KernelDispatch.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "mlir/Conversion/VectorToSCF/VectorToSCF.h"
#include "mlir/Dialect/Linalg/IR/LinalgInterfaces.h"
#include "mlir/Dialect/Linalg/Transforms/CodegenStrategy.h"
@@ -54,13 +54,10 @@
} // namespace
namespace {
-struct LinalgToLLVMWorkgroupsVectorizationPass
- : public LinalgToLLVMWorkgroupsVectorizationBase<
- LinalgToLLVMWorkgroupsVectorizationPass> {
- LinalgToLLVMWorkgroupsVectorizationPass(bool vectorize = true)
- : lowerToVectors(vectorize) {}
- LinalgToLLVMWorkgroupsVectorizationPass(
- const LinalgToLLVMWorkgroupsVectorizationPass &pass) {
+struct LLVMCPUVectorizationPass
+ : public LLVMCPUVectorizationBase<LLVMCPUVectorizationPass> {
+ LLVMCPUVectorizationPass(bool vectorize = true) : lowerToVectors(vectorize) {}
+ LLVMCPUVectorizationPass(const LLVMCPUVectorizationPass &pass) {
lowerToVectors = pass.lowerToVectors;
}
void getDependentDialects(DialectRegistry ®istry) const override {
@@ -146,7 +143,7 @@
} // namespace
-void LinalgToLLVMWorkgroupsVectorizationPass::runOnOperation() {
+void LLVMCPUVectorizationPass::runOnOperation() {
auto funcOp = getOperation();
MLIRContext *context = &getContext();
// Promotes workgroups subviews to a full-tile allocated on the stack.
@@ -296,10 +293,9 @@
}
}
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMWorkgroupsVectorizationPass(bool lowerToVectors) {
- return std::make_unique<LinalgToLLVMWorkgroupsVectorizationPass>(
- lowerToVectors);
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUVectorizationPass(
+ bool lowerToVectors) {
+ return std::make_unique<LLVMCPUVectorizationPass>(lowerToVectors);
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVM/Passes.cpp b/iree/compiler/Codegen/LLVMCPU/Passes.cpp
similarity index 85%
rename from iree/compiler/Conversion/LinalgToLLVM/Passes.cpp
rename to iree/compiler/Codegen/LLVMCPU/Passes.cpp
index d081bdf..c4bf585 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/Passes.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/Passes.cpp
@@ -4,9 +4,9 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
-#include "iree/compiler/Conversion/PassDetail.h"
+#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
#include "mlir/Dialect/Linalg/Passes.h"
@@ -40,8 +40,7 @@
if (clUseTensorPadTileAndVectorize) {
// Tile and vectorize linalg ops on tensors.
- passManager.addNestedPass<FuncOp>(
- createLinalgToLLVMTilePadAndVectorizeWorkgroupsPass());
+ passManager.addNestedPass<FuncOp>(createLLVMCPUTilePadAndVectorizePass());
passManager.addNestedPass<FuncOp>(createCSEPass());
passManager.addNestedPass<FuncOp>(createCanonicalizerPass());
}
@@ -54,26 +53,26 @@
if (!clUseTensorPadTileAndVectorize) {
// Tile and vectorize linalg ops on buffers.
passManager.addNestedPass<FuncOp>(
- createLinalgToLLVMWorkgroupsVectorizationPass(lowerToVectors));
+ createLLVMCPUVectorizationPass(lowerToVectors));
passManager.addNestedPass<FuncOp>(createCSEPass());
passManager.addNestedPass<FuncOp>(createCanonicalizerPass());
}
passManager.addNestedPass<FuncOp>(createForOpCanonicalizationPass());
- passManager.addNestedPass<FuncOp>(createLinalgToLLVMPlanConvLoopOrderPass());
+ passManager.addNestedPass<FuncOp>(createLLVMCPUPlanConvLoopOrderPass());
}
void addCPUDefaultPassPipeline(OpPassManager &passManager) {
passManager.addPass(createCanonicalizerPass());
// Use stack allocation on CPU side.
addLinalgBufferizePasses(passManager, cpuAllocationFunction);
- passManager.addNestedPass<FuncOp>(createLinalgToLLVMPlanConvLoopOrderPass());
+ passManager.addNestedPass<FuncOp>(createLLVMCPUPlanConvLoopOrderPass());
}
static void addLowerToLLVMPasses(
OpPassManager &passManager,
- const LLVMTransformPassPipelineOptions &options) {
+ const LLVMCPUCodegenPassPipelineOptions &options) {
// Linalg -> SCF
passManager.addNestedPass<FuncOp>(createConvertLinalgToLoopsPass());
passManager.addNestedPass<FuncOp>(createCanonicalizerPass());
@@ -96,9 +95,9 @@
passManager.addPass(createCSEPass());
}
-void buildLLVMTransformPassPipeline(
+void buildLLVMCPUCodegenPassPipeline(
OpPassManager &passManager,
- const LLVMTransformPassPipelineOptions &options) {
+ const LLVMCPUCodegenPassPipelineOptions &options) {
OpPassManager &nestedModulePM = passManager.nest<ModuleOp>();
addLowerToLLVMPasses(nestedModulePM, options);
}
diff --git a/iree/compiler/Conversion/LinalgToLLVM/VectorContractToAArch64InlineAsmOp.cpp b/iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp
similarity index 98%
rename from iree/compiler/Conversion/LinalgToLLVM/VectorContractToAArch64InlineAsmOp.cpp
rename to iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp
index bdf8d90..9dd3d36 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/VectorContractToAArch64InlineAsmOp.cpp
+++ b/iree/compiler/Codegen/LLVMCPU/VectorContractToAArch64InlineAsmOp.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/BUILD b/iree/compiler/Codegen/LLVMCPU/test/BUILD
similarity index 88%
rename from iree/compiler/Conversion/LinalgToLLVM/test/BUILD
rename to iree/compiler/Codegen/LLVMCPU/test/BUILD
index dd7b11c..680f1c1 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/BUILD
+++ b/iree/compiler/Codegen/LLVMCPU/test/BUILD
@@ -22,12 +22,11 @@
"hal_interface_bindings.mlir",
"hal_interface_constants.mlir",
"hal_interface_workgroup_info.mlir",
- "linalg_vectorize.mlir",
"materialize_launch_configuration.mlir",
"matmul_vectorization.mlir",
- "pad_linalg_workgroup_tiles.mlir",
+ "pad_workgroup_tiles.mlir",
"plan_conv_loop_order.mlir",
- "tile_pad_and_vectorize_workgroups.mlir",
+ "tile_pad_and_vectorize.mlir",
"unfused_fma.mlir",
"vector_contract_to_aarch64_asm.mlir",
],
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/CMakeLists.txt b/iree/compiler/Codegen/LLVMCPU/test/CMakeLists.txt
similarity index 86%
rename from iree/compiler/Conversion/LinalgToLLVM/test/CMakeLists.txt
rename to iree/compiler/Codegen/LLVMCPU/test/CMakeLists.txt
index a1d3f90..860d1a3 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/CMakeLists.txt
+++ b/iree/compiler/Codegen/LLVMCPU/test/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/LinalgToLLVM/test/BUILD #
+# iree/compiler/Codegen/LLVMCPU/test/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -17,12 +17,11 @@
"hal_interface_bindings.mlir"
"hal_interface_constants.mlir"
"hal_interface_workgroup_info.mlir"
- "linalg_vectorize.mlir"
"materialize_launch_configuration.mlir"
"matmul_vectorization.mlir"
- "pad_linalg_workgroup_tiles.mlir"
+ "pad_workgroup_tiles.mlir"
"plan_conv_loop_order.mlir"
- "tile_pad_and_vectorize_workgroups.mlir"
+ "tile_pad_and_vectorize.mlir"
"unfused_fma.mlir"
"vector_contract_to_aarch64_asm.mlir"
DATA
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_bindings.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
similarity index 96%
rename from iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_bindings.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
index a6ceebf..0c34259 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_bindings.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_bindings.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -allow-unregistered-dialect -iree-codegen-convert-to-llvm -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -allow-unregistered-dialect -iree-convert-to-llvm -split-input-file %s | IreeFileCheck %s
// CHECK-LABEL: llvm.func internal @binding_ptrs
func @binding_ptrs() {
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_constants.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir
similarity index 85%
rename from iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_constants.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir
index 8afc632..b5b3265 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_constants.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_constants.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -allow-unregistered-dialect -iree-codegen-convert-to-llvm -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -allow-unregistered-dialect -iree-convert-to-llvm -split-input-file %s | IreeFileCheck %s
// CHECK-LABEL: llvm.func internal @constant_values
func @constant_values() {
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_workgroup_info.mlir b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir
similarity index 93%
rename from iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_workgroup_info.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir
index 8e80761..631100e 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/hal_interface_workgroup_info.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/hal_interface_workgroup_info.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -allow-unregistered-dialect -iree-codegen-convert-to-llvm -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -allow-unregistered-dialect -iree-convert-to-llvm -split-input-file %s | IreeFileCheck %s
// CHECK-LABEL: llvm.func internal @workgroup_id
func @workgroup_id() {
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToLLVM/test/materialize_launch_configuration.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
index 0247bdd..b903ba5 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/materialize_launch_configuration.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/materialize_launch_configuration.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-lower-executable-target-pass{test-lowering-configuration=true}))" -cse -canonicalize -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-llvmcpu-lower-executable-target{test-lowering-configuration=true}))" -cse -canonicalize -split-input-file %s | IreeFileCheck %s
hal.executable @matmul_tensors attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/matmul_vectorization.mlir b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
similarity index 96%
rename from iree/compiler/Conversion/LinalgToLLVM/test/matmul_vectorization.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
index 090c695..ef5089d 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/matmul_vectorization.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/matmul_vectorization.mlir
@@ -1,5 +1,5 @@
-// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-lower-executable-target-pass{use-lowering-pipeline='func(iree-codegen-linalg-to-llvm-workgroups-vectorization-pass)'}))" -split-input-file %s | IreeFileCheck %s
-// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-lower-executable-target-pass{use-lowering-pipeline='func(iree-codegen-linalg-to-llvm-workgroups-vectorization-pass{promote-workgroup-to-full-tiles}),cse'}))" -split-input-file %s | IreeFileCheck %s -check-prefix=CHECK-PROMOTED
+// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-llvmcpu-lower-executable-target{use-lowering-pipeline='func(iree-llvmcpu-vectorization)'}))" -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-llvmcpu-lower-executable-target{use-lowering-pipeline='func(iree-llvmcpu-vectorization{promote-workgroup-to-full-tiles}),cse'}))" -split-input-file %s | IreeFileCheck %s -check-prefix=CHECK-PROMOTED
#config = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
hal.executable @dynamic_matmul attributes {sym_visibility = "private"} {
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/pad_linalg_workgroup_tiles.mlir b/iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToLLVM/test/pad_linalg_workgroup_tiles.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir
index d960b5e..9b2726e 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/pad_linalg_workgroup_tiles.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/pad_workgroup_tiles.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt %s -cse -iree-codegen-llvm-pad-linalg-workgroup-tiles -split-input-file | IreeFileCheck %s
+// RUN: iree-opt %s -cse -iree-llvmcpu-pad-workgroup-tiles -split-input-file | IreeFileCheck %s
#config0 = {tileSizes = [[64, 64]]}
#config1 = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/plan_conv_loop_order.mlir b/iree/compiler/Codegen/LLVMCPU/test/plan_conv_loop_order.mlir
similarity index 89%
rename from iree/compiler/Conversion/LinalgToLLVM/test/plan_conv_loop_order.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/plan_conv_loop_order.mlir
index 0cc8ace..8a99e9d 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/plan_conv_loop_order.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/plan_conv_loop_order.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-codegen-linalg-to-llvm-plan-conv-loop-order %s | IreeFileCheck %s
+// RUN: iree-opt -iree-llvmcpu-plan-conv-loop-order %s | IreeFileCheck %s
func @conv(%filter: memref<3x3x3x32xf32>, %input: memref<1x225x225x3xf32>, %output: memref<1x112x112x32xf32>) {
linalg.conv(%filter, %input, %output) {dilations = [1, 1], strides = [2, 2]} : memref<3x3x3x32xf32>, memref<1x225x225x3xf32>, memref<1x112x112x32xf32>
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/tile_pad_and_vectorize_workgroups.mlir b/iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir
similarity index 96%
rename from iree/compiler/Conversion/LinalgToLLVM/test/tile_pad_and_vectorize_workgroups.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir
index 69389d7..4a91e6e 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/tile_pad_and_vectorize_workgroups.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/tile_pad_and_vectorize.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt %s -cse -iree-codegen-linalg-to-llvm-tile-pad-and-vectorize-workgroups -cse -canonicalize -split-input-file | IreeFileCheck %s
+// RUN: iree-opt %s -cse -iree-llvmcpu-tile-pad-and-vectorize -cse -canonicalize -split-input-file | IreeFileCheck %s
#config0 = {tileSizes = [[64, 64]]}
#config1 = {nativeVectorSize = [4, 4, 4], tileSizes = [[64, 64], [32, 32, 32], [4, 4, 4]]}
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/unfused_fma.mlir b/iree/compiler/Codegen/LLVMCPU/test/unfused_fma.mlir
similarity index 80%
rename from iree/compiler/Conversion/LinalgToLLVM/test/unfused_fma.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/unfused_fma.mlir
index e2302ad..4e2a2e6 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/unfused_fma.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/unfused_fma.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-codegen-linalg-to-llvm-unfuse-fma-pass %s | IreeFileCheck %s
+// RUN: iree-opt -iree-llvmcpu-unfuse-fma-pass %s | IreeFileCheck %s
func @fma_unfused(%a : f32, %b: f32, %c: f32) -> f32 {
%0 = "llvm.intr.fma"(%a, %b, %c) : (f32, f32, f32) -> f32
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/vector_contract_to_aarch64_asm.mlir b/iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToLLVM/test/vector_contract_to_aarch64_asm.mlir
rename to iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir
index 6e0a960..185bac5 100644
--- a/iree/compiler/Conversion/LinalgToLLVM/test/vector_contract_to_aarch64_asm.mlir
+++ b/iree/compiler/Codegen/LLVMCPU/test/vector_contract_to_aarch64_asm.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-codegen-vector-to-aarch64-inline-asm %s | IreeFileCheck %s
+// RUN: iree-opt -iree-llvmcpu-vector-to-aarch64-inline-asm %s | IreeFileCheck %s
func @vector_matmul_to_aarch64_asm_vec_dot(%lhs: memref<4x4xi8>, %rhs: memref<4x4xi8>, %dst: memref<4x4xi32>) {
%c0 = constant 0 : index
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/BUILD b/iree/compiler/Codegen/LLVMGPU/BUILD
similarity index 83%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/BUILD
rename to iree/compiler/Codegen/LLVMGPU/BUILD
index 2675446..b348c5f 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/BUILD
+++ b/iree/compiler/Codegen/LLVMGPU/BUILD
@@ -11,27 +11,27 @@
)
cc_library(
- name = "LinalgToLLVMGPU",
+ name = "LLVMGPU",
srcs = [
"ConvertToLLVM.cpp",
"ConvertToNVVM.cpp",
"ConvertToROCDL.cpp",
"KernelConfig.cpp",
- "LowerExecutableTargetPass.cpp",
+ "LLVMGPULowerExecutableTarget.cpp",
+ "LLVMGPURemoveTrivialLoops.cpp",
+ "LLVMGPUTileAndDistribute.cpp",
+ "LLVMGPUVectorization.cpp",
"Passes.cpp",
- "RemoveTrivialLoops.cpp",
- "TileAndDistribute.cpp",
- "VectorizationPass.cpp",
],
hdrs = [
"ConvertToLLVM.h",
"KernelConfig.h",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/Transforms",
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Transforms",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/HAL/IR",
"//iree/compiler/Dialect/IREE/IR",
"//iree/compiler/Dialect/Shape/Transforms",
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/CMakeLists.txt b/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
similarity index 80%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/CMakeLists.txt
rename to iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
index 34b391f..d02522f 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/CMakeLists.txt
+++ b/iree/compiler/Codegen/LLVMGPU/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/LinalgToLLVMGPU/BUILD #
+# iree/compiler/Codegen/LLVMGPU/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -12,7 +12,7 @@
iree_cc_library(
NAME
- LinalgToLLVMGPU
+ LLVMGPU
HDRS
"ConvertToLLVM.h"
"KernelConfig.h"
@@ -21,11 +21,11 @@
"ConvertToNVVM.cpp"
"ConvertToROCDL.cpp"
"KernelConfig.cpp"
- "LowerExecutableTargetPass.cpp"
+ "LLVMGPULowerExecutableTarget.cpp"
+ "LLVMGPURemoveTrivialLoops.cpp"
+ "LLVMGPUTileAndDistribute.cpp"
+ "LLVMGPUVectorization.cpp"
"Passes.cpp"
- "RemoveTrivialLoops.cpp"
- "TileAndDistribute.cpp"
- "VectorizationPass.cpp"
DEPS
MLIRAffine
MLIRAffineToStandard
@@ -48,10 +48,10 @@
MLIRVector
MLIRVectorToLLVM
MLIRVectorToSCF
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::PassHeaders
- iree::compiler::Conversion::Transforms
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::Transforms
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::IREE::IR
iree::compiler::Dialect::Shape::Transforms
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
similarity index 95%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.cpp
rename to iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
index e140534..56c62b4 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.cpp
@@ -4,11 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.h"
+#include "iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
#include "mlir/Dialect/GPU/Passes.h"
@@ -61,9 +61,8 @@
};
/// Pass to test scalarization pattern.
-class TestLinalgToLLVMGPUScalarizeMathOpPass
- : public TestLinalgToLLVMGPUScalarizeMathOpBase<
- TestLinalgToLLVMGPUScalarizeMathOpPass> {
+class TestLLVMGPUScalarizeMathOpPass
+ : public TestLLVMGPUScalarizeMathOpBase<TestLLVMGPUScalarizeMathOpPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<vector::VectorDialect>();
}
@@ -289,9 +288,8 @@
patterns.getContext());
}
-std::unique_ptr<OperationPass<FuncOp>>
-createTestLinalgToLLVMGPUScalarizeMathOpPass() {
- return std::make_unique<TestLinalgToLLVMGPUScalarizeMathOpPass>();
+std::unique_ptr<OperationPass<FuncOp>> createTestLLVMGPUScalarizeMathOpPass() {
+ return std::make_unique<TestLLVMGPUScalarizeMathOpPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.h b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.h
similarity index 78%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.h
rename to iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.h
index 1ea4b81..be61344 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.h
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.h
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_CONVERSION_LINALGTOLLVMGPU_COMMON_H_
-#define IREE_COMPILER_CONVERSION_LINALGTOLLVMGPU_COMMON_H_
+#ifndef IREE_COMPILER_CODEGEN_LLVMGPU_COMMON_H_
+#define IREE_COMPILER_CODEGEN_LLVMGPU_COMMON_H_
#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
@@ -21,4 +21,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_LINALGTOLLVMGPU_COMMON_H_
+#endif // IREE_COMPILER_CODEGEN_LLVMGPU_COMMON_H_
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToNVVM.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
similarity index 86%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToNVVM.cpp
rename to iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
index 1685599..081f376 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToNVVM.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToNVVM.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
@@ -29,9 +29,7 @@
///
/// This pass only handles device code and is not meant to be run on GPU host
/// code.
-struct LinalgToLLVMGPUConvertToNVVMPass
- : public LinalgToLLVMGPUConvertToNVVMBase<
- LinalgToLLVMGPUConvertToNVVMPass> {
+struct ConvertToNVVMPass : public ConvertToNVVMBase<ConvertToNVVMPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<LLVM::LLVMDialect, NVVM::NVVMDialect>();
}
@@ -85,9 +83,8 @@
} // anonymous namespace
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToLLVMGPUConvertToNVVMPass() {
- return std::make_unique<LinalgToLLVMGPUConvertToNVVMPass>();
+std::unique_ptr<OperationPass<ModuleOp>> createConvertToNVVMPass() {
+ return std::make_unique<ConvertToNVVMPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToROCDL.cpp b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
similarity index 86%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToROCDL.cpp
rename to iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
index c753520..a3e727d 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToROCDL.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/ConvertToROCDL.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVMGPU/ConvertToLLVM.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/LLVMGPU/ConvertToLLVM.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "mlir/Conversion/GPUToROCDL/GPUToROCDLPass.h"
#include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
@@ -29,9 +29,7 @@
///
/// This pass only handles device code and is not meant to be run on GPU host
/// code.
-struct LinalgToLLVMGPUConvertToROCDLPass
- : public LinalgToLLVMGPUConvertToROCDLBase<
- LinalgToLLVMGPUConvertToROCDLPass> {
+struct ConvertToROCDLPass : public ConvertToROCDLBase<ConvertToROCDLPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<LLVM::LLVMDialect, ROCDL::ROCDLDialect>();
}
@@ -85,9 +83,8 @@
} // anonymous namespace
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToLLVMGPUConvertToROCDLPass() {
- return std::make_unique<LinalgToLLVMGPUConvertToROCDLPass>();
+std::unique_ptr<OperationPass<ModuleOp>> createConvertToROCDLPass() {
+ return std::make_unique<ConvertToROCDLPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.cpp b/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
similarity index 98%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.cpp
rename to iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
index 816d487..b3276ea 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp
@@ -4,9 +4,9 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.h"
+#include "iree/compiler/Codegen/LLVMGPU/KernelConfig.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/IR/Types.h"
#include "mlir/IR/Value.h"
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.h b/iree/compiler/Codegen/LLVMGPU/KernelConfig.h
similarity index 70%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.h
rename to iree/compiler/Codegen/LLVMGPU/KernelConfig.h
index f10af2f..2717b90 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.h
+++ b/iree/compiler/Codegen/LLVMGPU/KernelConfig.h
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_CONVERSION_LINALGTOLLVMGPU_KERNELCONFIG_H_
-#define IREE_COMPILER_CONVERSION_LINALGTOLLVMGPU_KERNELCONFIG_H_
+#ifndef IREE_COMPILER_CODEGEN_LLVMGPU_KERNELCONFIG_H_
+#define IREE_COMPILER_CODEGEN_LLVMGPU_KERNELCONFIG_H_
#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "mlir/IR/BuiltinOps.h"
@@ -17,4 +17,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_LINALGTOLLVMGPU_KERNELCONFIG_H_
+#endif // IREE_COMPILER_CODEGEN_LLVMGPU_KERNELCONFIG_H_
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/LowerExecutableTargetPass.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
similarity index 83%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/LowerExecutableTargetPass.cpp
rename to iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
index dfa4c3e..b3e8e7c 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/LowerExecutableTargetPass.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPULowerExecutableTarget.cpp
@@ -4,10 +4,10 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/LLVMGPU/KernelConfig.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/GPU/GPUDialect.h"
@@ -25,24 +25,24 @@
/// - then convert to NVVM/ROCDL dialect.
/// This should be merged with the equivalent pass in LinalgToLLVM. Fo
/// simplicity it is currently a separate pass.
-class LinalgToLLVMGPULowerExecutableTargetPass
- : public LinalgToLLVMGPULowerExecutableTargetBase<
- LinalgToLLVMGPULowerExecutableTargetPass> {
+class LLVMGPULowerExecutableTargetPass
+ : public LLVMGPULowerExecutableTargetBase<
+ LLVMGPULowerExecutableTargetPass> {
public:
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<IREE::HAL::HALDialect, linalg::LinalgDialect,
vector::VectorDialect, gpu::GPUDialect>();
}
- LinalgToLLVMGPULowerExecutableTargetPass() = default;
- LinalgToLLVMGPULowerExecutableTargetPass(
- const LinalgToLLVMGPULowerExecutableTargetPass &pass) = default;
+ LLVMGPULowerExecutableTargetPass() = default;
+ LLVMGPULowerExecutableTargetPass(
+ const LLVMGPULowerExecutableTargetPass &pass) = default;
void runOnOperation() override;
};
} // namespace
-void LinalgToLLVMGPULowerExecutableTargetPass::runOnOperation() {
+void LLVMGPULowerExecutableTargetPass::runOnOperation() {
IREE::HAL::ExecutableTargetOp targetOp = getOperation();
ModuleOp moduleOp = targetOp.getInnerModule();
@@ -112,8 +112,8 @@
}
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToLLVMGPULowerExecutableTargetPass() {
- return std::make_unique<LinalgToLLVMGPULowerExecutableTargetPass>();
+createLLVMGPULowerExecutableTargetPass() {
+ return std::make_unique<LLVMGPULowerExecutableTargetPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/RemoveTrivialLoops.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp
similarity index 82%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/RemoveTrivialLoops.cpp
rename to iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp
index 5d49a1c..2d8873b 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/RemoveTrivialLoops.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPURemoveTrivialLoops.cpp
@@ -4,9 +4,9 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "mlir/Dialect/GPU/Passes.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/MLIRContext.h"
@@ -35,9 +35,9 @@
namespace {
-class LinalgToLLVMGPURemoveSingleIterationLoopPass
- : public LinalgToLLVMGPURemoveSingleIterationLoopBase<
- LinalgToLLVMGPURemoveSingleIterationLoopPass> {
+class LLVMGPURemoveSingleIterationLoopPass
+ : public LLVMGPURemoveSingleIterationLoopBase<
+ LLVMGPURemoveSingleIterationLoopPass> {
void runOnOperation() override {
FuncOp funcOp = getOperation();
std::array<int32_t, 3> workgroupSize;
@@ -61,8 +61,8 @@
} // namespace
std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMGPURemoveSingleIterationLoopPass() {
- return std::make_unique<LinalgToLLVMGPURemoveSingleIterationLoopPass>();
+createLLVMGPURemoveSingleIterationLoopPass() {
+ return std::make_unique<LLVMGPURemoveSingleIterationLoopPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/TileAndDistribute.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
similarity index 96%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/TileAndDistribute.cpp
rename to iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
index 8180644..e9ccb36 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/TileAndDistribute.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUTileAndDistribute.cpp
@@ -4,12 +4,12 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/LinalgToLLVMGPU/KernelConfig.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/LLVMGPU/KernelConfig.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/LoweringConfig.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
@@ -318,9 +318,8 @@
ArrayRef<int64_t> tileSize;
};
-struct LinalgToLLVMGPUTileAndDistributePass
- : public LinalgToLLVMGPUTileAndDistributeBase<
- LinalgToLLVMGPUTileAndDistributePass> {
+struct LLVMGPUTileAndDistributePass
+ : public LLVMGPUTileAndDistributeBase<LLVMGPUTileAndDistributePass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<AffineDialect, gpu::GPUDialect>();
}
@@ -432,8 +431,8 @@
} // namespace
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToLLVMGPUTileAndDistributeToThreads() {
- return std::make_unique<LinalgToLLVMGPUTileAndDistributePass>();
+createLLVMGPUTileAndDistributeToThreads() {
+ return std::make_unique<LLVMGPUTileAndDistributePass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/VectorizationPass.cpp b/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp
similarity index 90%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/VectorizationPass.cpp
rename to iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp
index b902940..541f1f8 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/VectorizationPass.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/LLVMGPUVectorization.cpp
@@ -4,11 +4,11 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "mlir/Conversion/VectorToSCF/VectorToSCF.h"
#include "mlir/Dialect/Linalg/Transforms/Hoisting.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
@@ -74,9 +74,8 @@
}
namespace {
-struct LinalgToLLVMGPUVectorizationPass
- : public LinalgToLLVMGPUVectorizationBase<
- LinalgToLLVMGPUVectorizationPass> {
+struct LLVMGPUVectorizationPass
+ : public LLVMGPUVectorizationBase<LLVMGPUVectorizationPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<vector::VectorDialect>();
}
@@ -148,9 +147,8 @@
};
} // namespace
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMGPUVectorizationPass() {
- return std::make_unique<LinalgToLLVMGPUVectorizationPass>();
+std::unique_ptr<OperationPass<FuncOp>> createLLVMGPUVectorizationPass() {
+ return std::make_unique<LLVMGPUVectorizationPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/Passes.cpp b/iree/compiler/Codegen/LLVMGPU/Passes.cpp
similarity index 86%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/Passes.cpp
rename to iree/compiler/Codegen/LLVMGPU/Passes.cpp
index 3d53d1f..acaa233 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/Passes.cpp
+++ b/iree/compiler/Codegen/LLVMGPU/Passes.cpp
@@ -4,9 +4,9 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
-#include "iree/compiler/Conversion/PassDetail.h"
+#include "iree/compiler/Codegen/PassDetail.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
#include "mlir/Conversion/SCFToStandard/SCFToStandard.h"
@@ -38,16 +38,15 @@
pm.addNestedPass<ModuleOp>(createCSEPass());
// Distribute linalg onto threads within the workgroup.
- pm.addPass(createLinalgToLLVMGPUTileAndDistributeToThreads());
+ pm.addPass(createLLVMGPUTileAndDistributeToThreads());
pm.addNestedPass<ModuleOp>(createCanonicalizerPass());
pm.addNestedPass<ModuleOp>(createCSEPass());
pm.nest<ModuleOp>().addNestedPass<FuncOp>(
- createLinalgToLLVMGPURemoveSingleIterationLoopPass());
+ createLLVMGPURemoveSingleIterationLoopPass());
// Linalg -> vector
- pm.nest<ModuleOp>().addNestedPass<FuncOp>(
- createLinalgToLLVMGPUVectorizationPass());
+ pm.nest<ModuleOp>().addNestedPass<FuncOp>(createLLVMGPUVectorizationPass());
pm.nest<ModuleOp>().addNestedPass<FuncOp>(createCanonicalizerPass());
pm.nest<ModuleOp>().addNestedPass<FuncOp>(createCSEPass());
}
@@ -63,12 +62,12 @@
pm.addNestedPass<ModuleOp>(createCSEPass());
// Distribute linalg onto threads within the workgroup.
- pm.addPass(createLinalgToLLVMGPUTileAndDistributeToThreads());
+ pm.addPass(createLLVMGPUTileAndDistributeToThreads());
pm.addNestedPass<ModuleOp>(createCanonicalizerPass());
pm.addNestedPass<ModuleOp>(createCSEPass());
pm.nest<ModuleOp>().addNestedPass<FuncOp>(
- createLinalgToLLVMGPURemoveSingleIterationLoopPass());
+ createLLVMGPURemoveSingleIterationLoopPass());
}
static void addLowerToLLVMGPUPasses(OpPassManager &pm, bool useROCM) {
@@ -97,15 +96,15 @@
pm.addNestedPass<ModuleOp>(createStripDebugInfoPass());
if (useROCM) {
// convert to ROCDL.
- pm.addNestedPass<ModuleOp>(createLinalgToLLVMGPUConvertToROCDLPass());
+ pm.addNestedPass<ModuleOp>(createConvertToROCDLPass());
} else {
// convert to NVVM.
- pm.addNestedPass<ModuleOp>(createLinalgToLLVMGPUConvertToNVVMPass());
+ pm.addNestedPass<ModuleOp>(createConvertToNVVMPass());
}
}
void buildLLVMGPUTransformPassPipeline(OpPassManager &pm, bool useROCM) {
- pm.addPass(createLinalgToLLVMGPULowerExecutableTargetPass());
+ pm.addPass(createLLVMGPULowerExecutableTargetPass());
//===--------------------------------------------------------------------===//
// Convert Linalg ops to LLVM+NVVM/ROCDL ops.
//
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/BUILD b/iree/compiler/Codegen/LLVMGPU/test/BUILD
similarity index 100%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/BUILD
rename to iree/compiler/Codegen/LLVMGPU/test/BUILD
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/CMakeLists.txt b/iree/compiler/Codegen/LLVMGPU/test/CMakeLists.txt
similarity index 94%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/CMakeLists.txt
rename to iree/compiler/Codegen/LLVMGPU/test/CMakeLists.txt
index 362b91f..b18ba76 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/test/CMakeLists.txt
+++ b/iree/compiler/Codegen/LLVMGPU/test/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/LinalgToLLVMGPU/test/BUILD #
+# iree/compiler/Codegen/LLVMGPU/test/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/convert_to_nvvm.mlir b/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
similarity index 96%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/convert_to_nvvm.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
index e133f39..851db76 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/test/convert_to_nvvm.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/convert_to_nvvm.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-codegen-convert-to-nvvm %s | IreeFileCheck %s
+// RUN: iree-opt -iree-convert-to-nvvm %s | IreeFileCheck %s
// Test that that standard and GPU ops are converted to LLVM and NVVM.
func @abs_ex_dispatch_0() {
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/convert_to_rocdl.mlir b/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
similarity index 94%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/convert_to_rocdl.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
index 9a120a4..016403f 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/test/convert_to_rocdl.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/convert_to_rocdl.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-codegen-convert-to-rocdl %s | IreeFileCheck %s
+// RUN: iree-opt -iree-convert-to-rocdl %s | IreeFileCheck %s
// Test that that standard and GPU ops are converted to LLVM and NVVM.
func @abs_ex_dispatch_0() {
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/distribute_to_thread.mlir b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/distribute_to_thread.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
index 9cd22e5..d009cb8 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/test/distribute_to_thread.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/distribute_to_thread.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-codegen-llvmgpu-tile-and-distribute))" %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-llvmgpu-tile-and-distribute))" %s | IreeFileCheck %s
hal.executable @add_dispatch_0 attributes {sym_visibility = "private"} {
hal.executable.target @cuda, filter="cuda" {
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/nvvm_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/nvvm_pipeline_test.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/nvvm_pipeline_test.mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/remove_loops.mlir b/iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/remove_loops.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/remove_loops.mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/rocdl_pipeline_test.mlir b/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/rocdl_pipeline_test.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/scalarize.mlir b/iree/compiler/Codegen/LLVMGPU/test/scalarize.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/scalarize.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/scalarize.mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVMGPU/test/vectorization.mlir b/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToLLVMGPU/test/vectorization.mlir
rename to iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir
index 2029618..f0574b6 100644
--- a/iree/compiler/Conversion/LinalgToLLVMGPU/test/vectorization.mlir
+++ b/iree/compiler/Codegen/LLVMGPU/test/vectorization.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -iree-codegen-llvmgpu-vectorization %s | IreeFileCheck %s
+// RUN: iree-opt -iree-llvmgpu-vectorization %s | IreeFileCheck %s
func @add_dispatch_0() attributes {cuda_workgroup_size = dense<[32, 1, 1]> : vector<3xi64>} {
%c128 = constant 128 : index
diff --git a/iree/compiler/Conversion/PassDetail.h b/iree/compiler/Codegen/PassDetail.h
similarity index 91%
rename from iree/compiler/Conversion/PassDetail.h
rename to iree/compiler/Codegen/PassDetail.h
index b272b97..4ac9280 100644
--- a/iree/compiler/Conversion/PassDetail.h
+++ b/iree/compiler/Codegen/PassDetail.h
@@ -14,7 +14,7 @@
namespace iree_compiler {
#define GEN_PASS_CLASSES
-#include "iree/compiler/Conversion/Passes.h.inc"
+#include "iree/compiler/Codegen/Passes.h.inc"
} // namespace iree_compiler
} // namespace mlir
diff --git a/iree/compiler/Conversion/Passes.cpp b/iree/compiler/Codegen/Passes.cpp
similarity index 79%
rename from iree/compiler/Conversion/Passes.cpp
rename to iree/compiler/Codegen/Passes.cpp
index 8d4f0fd..07f7d19 100644
--- a/iree/compiler/Conversion/Passes.cpp
+++ b/iree/compiler/Codegen/Passes.cpp
@@ -4,27 +4,27 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
namespace mlir {
namespace iree_compiler {
namespace {
#define GEN_PASS_REGISTRATION
-#include "iree/compiler/Conversion/Passes.h.inc"
+#include "iree/compiler/Codegen/Passes.h.inc"
} // namespace
-void registerConversionPasses() {
+void registerCodegenPasses() {
// Generated.
registerPasses();
- static PassPipelineRegistration<LLVMTransformPassPipelineOptions>
+ static PassPipelineRegistration<LLVMCPUCodegenPassPipelineOptions>
linalgLLVMVPipeline(
"iree-codegen-linalg-to-llvm-pipeline",
"Runs the progressive lowering pipeline from Linalg to LLVM",
[](OpPassManager &passManager,
- const LLVMTransformPassPipelineOptions &options) {
- buildLLVMTransformPassPipeline(passManager, options);
+ const LLVMCPUCodegenPassPipelineOptions &options) {
+ buildLLVMCPUCodegenPassPipeline(passManager, options);
});
static PassPipelineRegistration<> LinalgNVVMPipeline(
@@ -54,8 +54,8 @@
"Runs the progressive lowering pipeline from XLA HLO to Linalg to "
"SPIR-V",
[](OpPassManager &passManager) {
- buildSPIRVTransformPassPipeline(
- passManager, SPIRVCodegenOptions::getFromCLOptions());
+ buildSPIRVCodegenPassPipeline(passManager,
+ SPIRVCodegenOptions::getFromCLOptions());
});
}
diff --git a/iree/compiler/Conversion/Passes.h b/iree/compiler/Codegen/Passes.h
similarity index 80%
rename from iree/compiler/Conversion/Passes.h
rename to iree/compiler/Codegen/Passes.h
index 5ec2d4a..b433622 100644
--- a/iree/compiler/Conversion/Passes.h
+++ b/iree/compiler/Codegen/Passes.h
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_CONVERSION_PASSES_H_
-#define IREE_COMPILER_CONVERSION_PASSES_H_
+#ifndef IREE_COMPILER_CODEGEN_PASSES_H_
+#define IREE_COMPILER_CODEGEN_PASSES_H_
#include <memory>
@@ -19,16 +19,12 @@
namespace iree_compiler {
// Registers all conversion passes in this directory.
-void registerConversionPasses();
+void registerCodegenPasses();
//------------------------------------------------------------------------------
// Misc/common conversions
//------------------------------------------------------------------------------
-/// Create a pass to convert a model using f32 type to the equivalent one
-/// using f16.
-std::unique_ptr<OperationPass<ModuleOp>> createDemoteF32ToF16Pass();
-
/// Alias for callback function that allocates workgroup level memory. The
/// callback expects the static shape and element-type of the result memref
/// type. Also expects values for the dynamic dimension of the allocated memref,
@@ -44,6 +40,28 @@
OpPassManager &passManager,
WorkgroupMemoryAllocationFn allocationFn = nullptr);
+/// Pass to perform canonicalizations/cleanups related to HAL interface/buffer
+/// allocations and view operations.
+std::unique_ptr<OperationPass<FuncOp>> createCleanupBufferAllocViewPass();
+
+/// Create a pass to convert a model using f32 type to the equivalent one
+/// using f16.
+std::unique_ptr<OperationPass<ModuleOp>> createDemoteF32ToF16Pass();
+
+/// Flattens n-D MemRef subspan ops to 1-D MemRef and folds the byte offsets on
+/// subspan ops to the consumer load/store ops, in preparation for lowering to
+/// backends that require linearized access.
+std::unique_ptr<OperationPass<ModuleOp>> createFlattenMemRefSubspanPass();
+
+/// After running the upstream TensorConstantBufferize pass, remove tensor_loads
+/// introduced for use only in tensor_extract. These can be folded to use a load
+/// of the created memref object that holds the constant values.
+std::unique_ptr<OperationPass<>> createFoldTensorExtractOpPass();
+
+/// An ad-hoc pass to canonicalize selected loop carried dependencies on
+/// scf.for.
+std::unique_ptr<OperationPass<FuncOp>> createForOpCanonicalizationPass();
+
/// Pass to perform linalg on tensor bufferization. The function passed into the
/// pass through the `allocationFn` argument is invoked whenever a new buffer is
/// to be created. The callback will be passed the Values for the dynamic
@@ -55,73 +73,79 @@
std::unique_ptr<OperationPass<FuncOp>> createLinalgBufferizePass(
WorkgroupMemoryAllocationFn allocationFn = nullptr);
+/// Creates a pass to vectorize a very specific form of linalg.conv ops.
+std::unique_ptr<OperationPass<FuncOp>> createLinalgToVectorVectorizeConvPass();
+
/// Pass to optimize vector transfer_read and transfer_write.
std::unique_ptr<OperationPass<FuncOp>> createOptimizeVectorTransferPass();
-/// An ad-hoc pass to canonicalize selected loop carried dependencies on
-/// scf.for.
-std::unique_ptr<OperationPass<FuncOp>> createForOpCanonicalizationPass();
-
-/// Pass to perform canonicalizations/cleanups related to HAL interface/buffer
-/// allocations and view operations.
-std::unique_ptr<OperationPass<FuncOp>> createCleanupBufferAllocViewPass();
-
-/// Flattens n-D MemRef subspan ops to 1-D MemRef and folds the byte offsets on
-/// subspan ops to the consumer load/store ops, in preparation for lowering to
-/// backends that require linearized access.
-std::unique_ptr<OperationPass<ModuleOp>> createFlattenMemRefSubspanPass();
-
/// Sets the number of workgroups to use for each entry point in the dispatch
/// region.
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
createSetNumWorkgroupsPass(ArrayRef<int64_t> workgroupSize = {});
-/// After running the upstream TensorConstantBufferize pass, remove tensor_loads
-/// introduced for use only in tensor_extract. These can be folded to use a load
-/// of the created memref object that holds the constant values.
-std::unique_ptr<OperationPass<>> createFoldTensorExtractOpPass();
+//----------------------------------------------------------------------------//
+// Common codegen patterns.
+//----------------------------------------------------------------------------//
+
+/// Populates `patterns` with a very specific pattern that vectorizes a
+/// linalg.conv op for a single thread. The linalg.conv should compute on
+/// static-sized subviews. To match, output shape must be 1x1xWoxCo, where Co
+/// Co is a multiple of 4, and filter shape must be 1x1x4xCo.
+void populateLinalgToVectorVectorizeConvPatterns(
+ MLIRContext *context, OwningRewritePatternList &patterns);
//------------------------------------------------------------------------------
-// LinalgToLinalg
+// LLVMCPU
//------------------------------------------------------------------------------
-//------------------------------------------------------------------------------
-// LinalgToLLVM
-//------------------------------------------------------------------------------
-
-/// Converts linalg.conv into linalg.generic with a CPU-friendly iteration
-/// order.
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMPlanConvLoopOrderPass();
-
-/// Pad linalg ops workgroup tiles into the next integer multiple of the target
-/// vector size.
-std::unique_ptr<OperationPass<FuncOp>> createLLVMPadLinalgWorkgroupTilesPass();
-
-/// Multi-level tiling, padding and vectorization of linalg ops on tensors.
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMTilePadAndVectorizeWorkgroupsPass();
-
-/// Vectorizes linalg ops executed in the same hal.interface.workgroup.
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMWorkgroupsVectorizationPass(bool lowerToVectors = true);
-
-/// Replaces llvm.intr.fma with its unfused mul and add ops.
-std::unique_ptr<OperationPass<FuncOp>> createLinalgToLLVMUnfuseFMAOpsPass();
-
-void populateUnfusedFMAOpsPassPatterns(MLIRContext *context,
- OwningRewritePatternList &patterns);
-
/// Performs the final conversion to LLVM dialect.
std::unique_ptr<OperationPass<ModuleOp>> createConvertToLLVMPass(
std::string targetTriple = "", std::string targetDataLayout = "",
bool unfuseFMAOps = false);
-/// Pass to convert Linalg ops into vector operations.
-std::unique_ptr<OperationPass<FuncOp>> createLinalgVectorizePass();
+/// Pass to lower the module an hal.executable.target operation to external
+/// dialect. Currently this pass lowers to LLVM dialect, but could be
+/// generalized to lower to any "final" dialect like SPIR-V/NVVM, etc.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
+createLLVMCPULowerExecutableTargetPass(bool lowerToVectors = true);
+
+/// Pad linalg ops workgroup tiles into the next integer multiple of the target
+/// vector size.
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUPadWorkgroupTilesPass();
+
+/// Converts linalg.conv into linalg.generic with a CPU-friendly iteration
+/// order.
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUPlanConvLoopOrderPass();
+
+/// Multi-level tiling, padding and vectorization of linalg ops on tensors.
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUTilePadAndVectorizePass();
+
+/// Vectorizes linalg ops executed in the same hal.interface.workgroup.
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUVectorizationPass(
+ bool lowerToVectors = true);
+
+/// Replaces llvm.intr.fma with its unfused mul and add ops.
+std::unique_ptr<OperationPass<FuncOp>> createLLVMCPUUnfuseFMAOpsPass();
+
+/// A pass that converts vector dialect operations to inline assembly
+std::unique_ptr<OperationPass<FuncOp>>
+createVectorToAArch64InlineAssemblyPass();
+
+//------------------------------------------------------------------------------
+// LLVMCPU Codegen specific patterns.
+//------------------------------------------------------------------------------
+
+/// Populates `patterns` to convert vector.contract op to a sequence
+/// of AArch64 inline assembly operations.
+void populateVectorContractToAArch64InlineAsm(
+ OwningRewritePatternList &patterns, MLIRContext *context);
+
+void populateUnfusedFMAOpsPassPatterns(MLIRContext *context,
+ OwningRewritePatternList &patterns);
//----------------------------------------------------------------------------//
-// LinalgToLLVM Pass Pipelines for CPU Lowering.
+// LLVMCPU backend Pass Pipelines.
//----------------------------------------------------------------------------//
/// Populates the passes to lower to scalars operations for linalg based
@@ -134,27 +158,20 @@
void addCPUVectorizationPassPipeline(OpPassManager &passManager,
bool lowerToVectors = true);
-/// Pass to lower the module an hal.executable.target operation to external
-/// dialect. Currently this pass lowers to LLVM dialect, but could be
-/// generalized to lower to any "final" dialect like SPIR-V/NVVM, etc.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLowerExecutableTargetPass(bool lowerToVectors = true);
-
//----------------------------------------------------------------------------//
-// LinalgToLLVM Pass Pipelines for lowering to LLVM dialect.
+// LLVMCPU Pass Pipelines for lowering to LLVM dialect.
//----------------------------------------------------------------------------//
/// Options for LLVM pipeline.
-struct LLVMTransformPassPipelineOptions
- : public PassPipelineOptions<LLVMTransformPassPipelineOptions> {
- Option<std::string> targetTriple{
- *this, "target-triple", llvm::cl::desc("Code generation target triple."),
- llvm::cl::init("")};
+struct LLVMCPUCodegenPassPipelineOptions
+ : public PassPipelineOptions<LLVMCPUCodegenPassPipelineOptions> {
Option<std::string> targetDataLayout{
*this, "target-data-layout",
llvm::cl::desc("Code generation target data layout."),
llvm::cl::init("")};
-
+ Option<std::string> targetTriple{
+ *this, "target-triple", llvm::cl::desc("Code generation target triple."),
+ llvm::cl::init("")};
Option<bool> unfuseFMAOps{
*this, "unfuse-fma-ops",
llvm::cl::desc("Enable rewriting llvm.fma to its unfused version."),
@@ -164,36 +181,14 @@
/// Populates passes needed to lower a XLA HLO op to LLVM dialect via the
/// structured ops path. The pass manager `pm` in here should operate on the
/// module within the IREE::HAL::ExecutableOp.
-void buildLLVMTransformPassPipeline(
+void buildLLVMCPUCodegenPassPipeline(
OpPassManager &passManager,
- const LLVMTransformPassPipelineOptions &options);
+ const LLVMCPUCodegenPassPipelineOptions &options);
//------------------------------------------------------------------------------
-// LinalgToGPU
+// LLVMGPU
//------------------------------------------------------------------------------
-/// Performs the final conversion to NNVM+LLVM dialect.
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToLLVMGPUConvertToNVVMPass();
-
-/// Performs the final conversion to ROCDL+LLVM dialect.
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToLLVMGPUConvertToROCDLPass();
-
-/// Convert Linalg ops to Vector.
-std::unique_ptr<OperationPass<FuncOp>> createLinalgToLLVMGPUVectorizationPass();
-
-/// Perform tiling and distribution to threads.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToLLVMGPUTileAndDistributeToThreads();
-
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToLLVMGPURemoveSingleIterationLoopPass();
-
-/// Create pass calling the dynamic pipeline for LLVMGPU.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToLLVMGPULowerExecutableTargetPass();
-
/// Lowering calling vectorization patterns.
void addGPUVectorizationPassPipeline(OpPassManager &passManager);
@@ -206,8 +201,28 @@
/// module within the IREE::HAL::ExecutableOp.
void buildLLVMGPUTransformPassPipeline(OpPassManager &pm, bool useROCM);
+/// Performs the final conversion to NNVM+LLVM dialect.
+std::unique_ptr<OperationPass<ModuleOp>> createConvertToNVVMPass();
+
+/// Performs the final conversion to ROCDL+LLVM dialect.
+std::unique_ptr<OperationPass<ModuleOp>> createConvertToROCDLPass();
+
+/// Perform tiling and distribution to threads.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
+createLLVMGPUTileAndDistributeToThreads();
+
+std::unique_ptr<OperationPass<FuncOp>>
+createLLVMGPURemoveSingleIterationLoopPass();
+
+/// Create pass calling the dynamic pipeline for LLVMGPU.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
+createLLVMGPULowerExecutableTargetPass();
+
+/// Convert Linalg ops to Vector.
+std::unique_ptr<OperationPass<FuncOp>> createLLVMGPUVectorizationPass();
+
//------------------------------------------------------------------------------
-// LinalgToSPIRV Passes
+// SPIRV Passes
//------------------------------------------------------------------------------
// Options that can be used to configure SPIR-V code generation.
@@ -221,57 +236,47 @@
static SPIRVCodegenOptions getFromCLOptions();
};
-/// Pass to tile and vectorize Linalg operations on buffers in a single
-/// workgroup.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVTileAndVectorizeOneWorkgroupPass(
- const SPIRVCodegenOptions &options);
-
-/// Pass to add the synchronizations and attributes needed to lower from PLoops
-/// to GPU dialect.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVConvertToGPUPass();
-
/// Pass to perform the final conversion to SPIR-V dialect.
/// This pass converts remaining interface ops into SPIR-V global variables,
/// GPU processor ID ops into SPIR-V global variables, loop/standard ops into
/// corresponding SPIR-V ops.
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToSPIRVConvertToSPIRVPass();
-
-/// Pass to convert vector operations to GPU level operations. Instructions of
-/// vector size equal to subgroup size are distributed across the subgroup.
-std::unique_ptr<OperationPass<FuncOp>> createLinalgToSPIRVConvertVectorToGPU();
-
-/// Pass to convert vector read/write/arithmetic operations to the corresponding
-/// cooperative matrix ops when possible.
-std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToSPIRVVectorToCooperativeMatrixPass();
-
-/// Converts memref of scalar to memref of vector of efficent size. This will
-/// allow to convert memory accesses to vector load/store in SPIR-V without
-/// having pointer bitcast.
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToSPIRVVectorizeMemRefLoadStore();
-
-/// Creates a pass to fold processor ID uses where possible.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVFoldProcessorIDUsesPass();
-
-/// Pass that materializes new hal.executable.entry_point ops for
-/// spv.EntryPoints that are added by other passes.
-/// To be removed along with SplitDispatchFunctionPass.
-std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createMaterializeEntryPointsPass();
+std::unique_ptr<OperationPass<ModuleOp>> createConvertToSPIRVPass();
/// Creates a pass to concretize hal.interface.workgroup.* ops with concrete
/// tiling and distribution scheme.
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVConcretizeTileAmongWorkgroupsPass(
- const SPIRVCodegenOptions &options);
+createSPIRVConcretizeWorkgroupTilesPass(const SPIRVCodegenOptions &options);
+
+/// Pass to add the synchronizations and attributes needed to lower from PLoops
+/// to GPU dialect.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
+createSPIRVConvertToGPUPass();
+
+/// Creates a pass to fold processor ID uses where possible.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
+createSPIRVFoldProcessorIDUsesPass();
+
+/// Pass to tile and vectorize Linalg operations on buffers in a single
+/// workgroup.
+std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
+createSPIRVTileAndVectorizePass(const SPIRVCodegenOptions &options);
+
+/// Pass to convert vector read/write/arithmetic operations to the corresponding
+/// cooperative matrix ops when possible.
+std::unique_ptr<OperationPass<FuncOp>>
+createSPIRVVectorToCooperativeMatrixPass();
+
+/// Pass to convert vector operations to GPU level operations. Instructions of
+/// vector size equal to subgroup size are distributed across the subgroup.
+std::unique_ptr<OperationPass<FuncOp>> createSPIRVVectorToGPUPass();
+
+/// Converts memref of scalar to memref of vector of efficent size. This will
+/// allow to convert memory accesses to vector load/store in SPIR-V without
+/// having pointer bitcast.
+std::unique_ptr<OperationPass<ModuleOp>> createSPIRVVectorizeLoadStore();
//----------------------------------------------------------------------------//
-// LinalgToSPIRV Pipelines
+// SPIRV Codegen Pass Pipelines.
//----------------------------------------------------------------------------//
/// Populates passes need to lower from Linalf to SPIR-V.
@@ -285,11 +290,11 @@
/// testing purposes only. The pass pipeline will set an appropriate workgroup
/// size.
/// TODO: Are both of these needed and does this one still work on HLO?
-void buildSPIRVTransformPassPipeline(OpPassManager &pm,
- const SPIRVCodegenOptions &options);
+void buildSPIRVCodegenPassPipeline(OpPassManager &pm,
+ const SPIRVCodegenOptions &options);
//----------------------------------------------------------------------------//
-// LinalgToSPIRV Patterns
+// SPIRV Codegen specific patterns.
//----------------------------------------------------------------------------//
/// Populates patterns to tile and distribute linalg.copy operations.
@@ -302,40 +307,12 @@
OwningRewritePatternList &patterns);
//------------------------------------------------------------------------------
-// LinalgToVector
-//------------------------------------------------------------------------------
-
-/// Creates a pass to vectorize a very specific form of linalg.conv ops.
-std::unique_ptr<OperationPass<FuncOp>> createLinalgToVectorVectorizeConvPass();
-
-/// Populates `patterns` with a very specific pattern that vectorizes a
-/// linalg.conv op for a single thread. The linalg.conv should compute on
-/// static-sized subviews. To match, output shape must be 1x1xWoxCo, where Co
-/// Co is a multiple of 4, and filter shape must be 1x1x4xCo.
-void populateLinalgToVectorVectorizeConvPatterns(
- MLIRContext *context, OwningRewritePatternList &patterns);
-
-//------------------------------------------------------------------------------
-// VectorToLLVM
-//------------------------------------------------------------------------------
-
-/// A pass that converts vector dialect operations to inline assembly
-std::unique_ptr<OperationPass<FuncOp>>
-createVectorToAArch64InlineAssemblyPass();
-
-/// Populates `patterns` to convert vector.contract op to a sequence
-/// of AArch64 inline assembly operations.
-void populateVectorContractToAArch64InlineAsm(
- OwningRewritePatternList &patterns, MLIRContext *context);
-
-//------------------------------------------------------------------------------
// Test passes
//------------------------------------------------------------------------------
-std::unique_ptr<OperationPass<FuncOp>>
-createTestLinalgToLLVMGPUScalarizeMathOpPass();
+std::unique_ptr<OperationPass<FuncOp>> createTestLLVMGPUScalarizeMathOpPass();
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_PASSES_H_
+#endif // IREE_COMPILER_CODEGEN_PASSES_H_
diff --git a/iree/compiler/Codegen/Passes.td b/iree/compiler/Codegen/Passes.td
new file mode 100644
index 0000000..3a53b3b
--- /dev/null
+++ b/iree/compiler/Codegen/Passes.td
@@ -0,0 +1,262 @@
+// Copyright 2021 The IREE Authors
+//
+// Licensed 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
+
+#ifndef IREE_CODEGEN_PASSES
+#define IREE_CODEGEN_PASSES
+
+include "mlir/Pass/PassBase.td"
+
+//------------------------------------------------------------------------------
+// Common/misc passes
+//------------------------------------------------------------------------------
+
+def CleanupBufferAllocView :
+ Pass<"iree-codegen-cleanup-buffer-alloc-view", "FuncOp"> {
+ let summary =
+ "Performs cleanups over HAL interface/buffer allocation/view operations";
+ let constructor = "mlir::iree_compiler::createCleanupBufferAllocViewPass()";
+}
+
+def DemoteF32ToF16 :
+ Pass<"iree-convert-f32-to-f16", "ModuleOp"> {
+ let summary = "Convert f32 operations and values into equivalent f16 ones.";
+ let constructor = "mlir::iree_compiler::createDemoteF32ToF16Pass()";
+}
+
+def FlattenMemRefSubspan :
+ Pass<"iree-codegen-flatten-memref-subspan", "ModuleOp"> {
+ let summary =
+ "Flatten n-D MemRef subspan ops to 1-D ones and fold byte offsets";
+ let constructor = "mlir::iree_compiler::createFlattenMemRefSubspanPass()";
+}
+
+def FoldTensorExtractOp :
+ Pass<"iree-codegen-fold-tensor-extract-op", ""> {
+ let summary = "Fold `tensor.extract` operations prior to lowering to LLVM";
+ let constructor = "mlir::iree_compiler::createFoldTensorExtractOpPass()";
+}
+
+def ForOpCanonicalization :
+ Pass<"iree-codegen-canonicalize-scf-for", "FuncOp"> {
+ let summary =
+ "Adhoc canonicalization of selected loop-carried values/dependencies for scf.for ops";
+ let constructor = "mlir::iree_compiler::createForOpCanonicalizationPass()";
+}
+
+def LinalgBufferize :
+ Pass<"iree-codegen-linalg-bufferize", "FuncOp"> {
+ let summary = "Convert from to Linalg ops on tensors to buffers";
+ let constructor = "mlir::iree_compiler::createLinalgBufferizePass(nullptr)";
+}
+
+def OptimizeVectorTransfer :
+ Pass<"iree-codegen-optimize-vector-transfer", "FuncOp"> {
+ let summary =
+ "Run optimization transformations on vector transfer operations";
+ let constructor = "mlir::iree_compiler::createOptimizeVectorTransferPass()";
+}
+
+// TODO: Consider removing or moving to HAL/Transforms in order to avoid
+// polluting common pass declarations with HAL specific ops.
+def SetNumWorkgroups :
+ Pass<"iree-set-num-workgroups",
+ "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary =
+ "Set the number of workgroups for entry point functions";
+ let constructor = "mlir::iree_compiler::createSetNumWorkgroupsPass()";
+}
+
+// TODO: Rename argument to be fully qualified.
+def LinalgToVectorVectorizeConv :
+ Pass<"iree-codegen-vectorize-linalg-conv", "FuncOp"> {
+ let summary = "Vectorize a very specific form of linalg.conv";
+ let constructor =
+ "mlir::iree_compiler::createLinalgToVectorVectorizeConvPass();";
+}
+
+//------------------------------------------------------------------------------
+// LLVMCPU
+//------------------------------------------------------------------------------
+
+def ConvertToLLVM :
+ Pass<"iree-convert-to-llvm", "ModuleOp"> {
+ let summary =
+ "Perform final conversion from Linalg/HAL/Shape/Vector/Standard to LLVMIR dialect";
+ let constructor = "mlir::iree_compiler::createConvertToLLVMPass()";
+}
+
+def LLVMCPULowerExecutableTarget :
+ Pass<"iree-llvmcpu-lower-executable-target",
+ "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary =
+ "Lower executable target using an IREE::HAL::DispatchLoweringPassPipeline";
+ let constructor =
+ "mlir::iree_compiler::createLLVMCPULowerExecutableTargetPass()";
+}
+
+def LLVMCPUPadWorkgroupTiles :
+ Pass<"iree-llvmcpu-pad-workgroup-tiles", "FuncOp"> {
+ let summary =
+ "Pad workgroup tiles to an integer multiple of tiling parameters.";
+ let constructor = "mlir::iree_compiler::createLLVMCPUPadWorkgroupTilesPass()";
+}
+
+def LLVMCPUPlanConvLoopOrder :
+ Pass<"iree-llvmcpu-plan-conv-loop-order", "FuncOp"> {
+ let summary =
+ "Convert linalg.conv to linalg.generic with a CPU-friendly iterator order";
+ let constructor = "mlir::iree_compiler::createLLVMCPUPlanConvLoopOrderPass()";
+}
+
+def LLVMCPUTilePadAndVectorize :
+ Pass<"iree-llvmcpu-tile-pad-and-vectorize", "FuncOp"> {
+ let summary = "Tile and pad workgroups tiles";
+ let constructor =
+ "mlir::iree_compiler::createLLVMCPUTilePadAndVectorizePass()";
+}
+
+def LLVMCPUUnfuseFMAOps :
+ Pass<"iree-llvmcpu-unfuse-fma-pass", "FuncOp"> {
+ let summary = "Convert llvm.fma into unfused mulf and addf ops";
+ let constructor = "mlir::iree_compiler::createLLVMCPUUnfuseFMAOpsPass()";
+}
+
+def LLVMCPUVectorization :
+ Pass<"iree-llvmcpu-vectorization", "FuncOp"> {
+ let summary = "Tile and vectorize for CPU backends";
+ let constructor = "mlir::iree_compiler::createLLVMCPUVectorizationPass()";
+}
+
+def VectorToAArch64InlineAsm :
+ Pass<"iree-llvmcpu-vector-to-aarch64-inline-asm", "FuncOp"> {
+ let summary = "Convert vector operations to aarch64 inline asm LLVMIR dialect";
+ let constructor = "mlir::iree_compiler::createVectorToAArch64InlineAssemblyPass()";
+}
+
+//------------------------------------------------------------------------------
+// LLVMGPU
+//------------------------------------------------------------------------------
+
+// TODO: Bring the argument in line with the names used elsewhere.
+def ConvertToROCDL :
+ Pass<"iree-convert-to-rocdl", "ModuleOp"> {
+ let summary = "Perform final conversion from builtin/GPU/HAL/standard dialect to LLVM "
+ "and ROCDL dialects";
+ let constructor = "mlir::iree_compiler::createConvertToROCDLPass()";
+}
+
+// TODO: Bring the argument in line with the names used elsewhere.
+def ConvertToNVVM :
+ Pass<"iree-convert-to-nvvm", "ModuleOp"> {
+ let summary = "Perform final conversion from builtin/GPU/HAL/standard dialect to LLVM "
+ "and NVVM dialects";
+ let constructor = "mlir::iree_compiler::createConvertToNVVMPass()";
+}
+
+// TODO: Bring the argument in line with the names used elsewhere.
+def LLVMGPULowerExecutableTarget :
+ Pass<"iree-llvmgpu-lower-executable-target-pass", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary = "Perform lowering of executable target using one of the IREE::HAL::DispatchLoweringPassPipeline";
+ let constructor = "mlir::iree_compiler::createLLVMGPULowerExecutableTargetPass()";
+}
+
+def LLVMGPURemoveSingleIterationLoop :
+ Pass<"iree-llvmgpu-remove-single-iteration-loop", "FuncOp"> {
+ let summary = "Remove distributed loop with single iteration.";
+ let constructor = "mlir::iree_compiler::createLLVMGPURemoveSingleIterationLoopPass()";
+}
+
+def LLVMGPUTileAndDistribute :
+ Pass<"iree-llvmgpu-tile-and-distribute", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary = "Pass to tile and distribute linalg ops within a workgroup.";
+ let constructor = "mlir::iree_compiler::createLLVMGPUTileAndDistributeToThreads()";
+}
+
+def LLVMGPUVectorization :
+ Pass<"iree-llvmgpu-vectorization", "FuncOp"> {
+ let summary = "Pass to convert linalg into Vector.";
+ let constructor = "mlir::iree_compiler::createLLVMGPUVectorizationPass()";
+}
+
+//------------------------------------------------------------------------------
+// SPIRV
+//------------------------------------------------------------------------------
+
+// TODO: Rename argument to be fully qualified.
+def ConvertToSPIRV :
+ Pass<"iree-convert-to-spirv", "ModuleOp"> {
+ let summary = "Perform final conversion to SPIR-V dialect";
+ let constructor = "mlir::iree_compiler::createConvertToSPIRVPass()";
+}
+
+def SPIRVConcretizeWorkgroupTiles :
+ Pass<"iree-spirv-concretize-workgroup-tiles",
+ "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary = "Replace hal.interface.workgroup.* ops with constant values";
+ let constructor =
+ "mlir::iree_compiler::createSPIRVConcretizeWorkgroupTilesPass(mlir::iree_compiler::SPIRVCodegenOptions::getFromCLOptions())";
+}
+
+// TODO: Rename argument to be fully qualified.
+def SPIRVConvertToGPU :
+ Pass<"iree-spirv-convert-to-gpu",
+ "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary = "Map tiled linalg and loop ops to GPU";
+ let constructor = "mlir::iree_compiler::createSPIRVConvertToGPUPass()";
+}
+
+// TODO: Rename argument to be fully qualified.
+// TODO: Does not appear used?
+def SPIRVFoldProcessorIDUses :
+ Pass<"iree-spirv-fold-gpu-procid-uses",
+ "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary = "Fold GPU processor ID uses where possible";
+ let constructor = "mlir::iree_compiler::createSPIRVFoldProcessorIDUsesPass()";
+}
+
+// TODO: Rename argument to be fully qualified.
+def SPIRVTileAndVectorize :
+ Pass<"iree-spirv-tile-and-vectorize",
+ "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
+ let summary =
+ "Tile and vectorize Linalg operations on buffers in one workgroup";
+ let constructor =
+ "mlir::iree_compiler::createSPIRVTileAndVectorizePass(mlir::iree_compiler::SPIRVCodegenOptions::getFromCLOptions())";
+}
+
+// TODO: Rename argument to be fully qualified.
+def SPIRVVectorizeLoadStore :
+ Pass<"iree-spirv-vectorize-load-store", "ModuleOp"> {
+ let summary = "Vectorize load/store of memrefs for better memory access";
+ let constructor = "mlir::iree_compiler::createSPIRVVectorizeLoadStore()";
+}
+
+// TODO: Rename argument to be fully qualified.
+def SPIRVVectorToCooperativeMatrix :
+ Pass<"iree-spirv-vector-to-cooperative-matrix", "FuncOp"> {
+ let summary = "Generate cooperative matrix ops when possible";
+ let constructor =
+ "mlir::iree_compiler::createSPIRVVectorToCooperativeMatrixPass()";
+}
+
+// TODO: Rename argument to be fully qualified.
+def SPIRVVectorToGPU :
+ Pass<"iree-spirv-vector-to-gpu", "FuncOp"> {
+ let summary = "Convert vector dialect to gpu subgroup level GPU instructions";
+ let constructor = "mlir::iree_compiler::createSPIRVVectorToGPUPass()";
+}
+
+//------------------------------------------------------------------------------
+// Test passes
+//------------------------------------------------------------------------------
+
+def TestLLVMGPUScalarizeMathOp :
+ Pass<"iree-test-llvmgpu-scalarize-math-op", "FuncOp"> {
+ let summary = "Test pass for scalarization patterns.";
+ let constructor = "mlir::iree_compiler::createTestLLVMGPUScalarizeMathOpPass()";
+}
+
+#endif // IREE_DIALECT_FLOW_PASSES
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/BUILD b/iree/compiler/Codegen/SPIRV/BUILD
similarity index 81%
rename from iree/compiler/Conversion/LinalgToSPIRV/BUILD
rename to iree/compiler/Codegen/SPIRV/BUILD
index 77199a6..e77a60c 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/BUILD
+++ b/iree/compiler/Codegen/SPIRV/BUILD
@@ -11,21 +11,21 @@
)
cc_library(
- name = "LinalgToSPIRV",
+ name = "SPIRV",
srcs = [
"CodeGenOptionUtils.cpp",
- "ConcretizeTileAmongWorkgroupsPass.cpp",
- "ConvertToGPUPass.cpp",
"ConvertToSPIRVPass.cpp",
- "FoldGPUProcessorIDUses.cpp",
"KernelDispatchUtils.cpp",
"LaunchConfig.cpp",
"Passes.cpp",
- "TileAndVectorizeInOneWorkgroupPass.cpp",
+ "SPIRVConcretizeWorkgroupTiles.cpp",
+ "SPIRVConvertToGPU.cpp",
+ "SPIRVFoldGPUProcessorIDUses.cpp",
+ "SPIRVTileAndVectorize.cpp",
+ "SPIRVVectorToCooperativeMatrix.cpp",
+ "SPIRVVectorToGPU.cpp",
+ "SPIRVVectorizeLoadStore.cpp",
"Utils.cpp",
- "VectorToCooperativeMatrixPass.cpp",
- "VectorToGPUPass.cpp",
- "VectorizeMemrefLoadStorePass.cpp",
],
hdrs = [
"KernelDispatchUtils.h",
@@ -34,10 +34,10 @@
"Utils.h",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/Transforms",
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/Transforms",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/IR",
"//iree/compiler/Dialect/HAL/IR:HALDialect",
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt b/iree/compiler/Codegen/SPIRV/CMakeLists.txt
similarity index 78%
rename from iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt
rename to iree/compiler/Codegen/SPIRV/CMakeLists.txt
index b6a39c7..fffa6fd 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/CMakeLists.txt
+++ b/iree/compiler/Codegen/SPIRV/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/LinalgToSPIRV/BUILD #
+# iree/compiler/Codegen/SPIRV/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -12,7 +12,7 @@
iree_cc_library(
NAME
- LinalgToSPIRV
+ SPIRV
HDRS
"KernelDispatchUtils.h"
"LaunchConfig.h"
@@ -20,18 +20,18 @@
"Utils.h"
SRCS
"CodeGenOptionUtils.cpp"
- "ConcretizeTileAmongWorkgroupsPass.cpp"
- "ConvertToGPUPass.cpp"
"ConvertToSPIRVPass.cpp"
- "FoldGPUProcessorIDUses.cpp"
"KernelDispatchUtils.cpp"
"LaunchConfig.cpp"
"Passes.cpp"
- "TileAndVectorizeInOneWorkgroupPass.cpp"
+ "SPIRVConcretizeWorkgroupTiles.cpp"
+ "SPIRVConvertToGPU.cpp"
+ "SPIRVFoldGPUProcessorIDUses.cpp"
+ "SPIRVTileAndVectorize.cpp"
+ "SPIRVVectorToCooperativeMatrix.cpp"
+ "SPIRVVectorToGPU.cpp"
+ "SPIRVVectorizeLoadStore.cpp"
"Utils.cpp"
- "VectorToCooperativeMatrixPass.cpp"
- "VectorToGPUPass.cpp"
- "VectorizeMemrefLoadStorePass.cpp"
DEPS
LLVMSupport
MLIRAffine
@@ -60,10 +60,10 @@
MLIRTransforms
MLIRVector
MLIRVectorToSPIRV
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::PassHeaders
- iree::compiler::Conversion::Transforms
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::Transforms
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::HAL::IR::HALDialect
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/CodeGenOptionUtils.cpp b/iree/compiler/Codegen/SPIRV/CodeGenOptionUtils.cpp
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/CodeGenOptionUtils.cpp
rename to iree/compiler/Codegen/SPIRV/CodeGenOptionUtils.cpp
index 53472b0..9d16e61 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/CodeGenOptionUtils.cpp
+++ b/iree/compiler/Codegen/SPIRV/CodeGenOptionUtils.cpp
@@ -4,7 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "llvm/Support/CommandLine.h"
namespace mlir {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
similarity index 95%
rename from iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp
rename to iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
index 264f5b0..b1affc9 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToSPIRVPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/ConvertToSPIRVPass.cpp
@@ -13,9 +13,9 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "llvm/ADT/DenseMapInfo.h"
@@ -222,16 +222,14 @@
/// This pass converts remaining interface ops into SPIR-V global variables,
/// GPU processor ID ops into SPIR-V global variables, loop/standard ops into
/// corresponding SPIR-V ops.
-struct LinalgToSPIRVConvertToSPIRVPass
- : public LinalgToSPIRVConvertToSPIRVBase<LinalgToSPIRVConvertToSPIRVPass> {
+struct ConvertToSPIRVPass : public ConvertToSPIRVBase<ConvertToSPIRVPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<spirv::SPIRVDialect>();
}
void runOnOperation() override;
- LinalgToSPIRVConvertToSPIRVPass() {}
- LinalgToSPIRVConvertToSPIRVPass(const LinalgToSPIRVConvertToSPIRVPass &pass) {
- }
+ ConvertToSPIRVPass() {}
+ ConvertToSPIRVPass(const ConvertToSPIRVPass &pass) {}
};
} // namespace
@@ -263,7 +261,7 @@
return success();
}
-void LinalgToSPIRVConvertToSPIRVPass::runOnOperation() {
+void ConvertToSPIRVPass::runOnOperation() {
MLIRContext *context = &getContext();
ModuleOp moduleOp = getOperation();
@@ -354,9 +352,8 @@
// Pass entry point and registration
//===----------------------------------------------------------------------===//
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToSPIRVConvertToSPIRVPass() {
- return std::make_unique<LinalgToSPIRVConvertToSPIRVPass>();
+std::unique_ptr<OperationPass<ModuleOp>> createConvertToSPIRVPass() {
+ return std::make_unique<ConvertToSPIRVPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.cpp b/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.cpp
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.cpp
rename to iree/compiler/Codegen/SPIRV/KernelDispatchUtils.cpp
index 0600b69..aa10372 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.cpp
+++ b/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.cpp
@@ -13,12 +13,12 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h"
+#include "iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/Utils.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/SPIRV/LaunchConfig.h"
+#include "iree/compiler/Codegen/SPIRV/Utils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeOps.h"
#include "llvm/Support/Debug.h"
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h b/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h
similarity index 82%
rename from iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h
rename to iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h
index 9e35d6a..a93bb3b 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h
+++ b/iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h
@@ -12,13 +12,13 @@
//
//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_LINALGTOSPIRV_KERNELDISPATCHUTILS_H_
-#define IREE_COMPILER_CONVERSION_LINALGTOSPIRV_KERNELDISPATCHUTILS_H_
+#ifndef IREE_COMPILER_CODEGEN_SPIRV_KERNELDISPATCHUTILS_H_
+#define IREE_COMPILER_CODEGEN_SPIRV_KERNELDISPATCHUTILS_H_
#include <array>
-#include "iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/SPIRV/LaunchConfig.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StringMap.h"
#include "llvm/Support/FormatVariadic.h"
@@ -46,4 +46,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_LINALGTOSPIRV_DISPATCHUTILS_H_
+#endif // IREE_COMPILER_CODEGEN_SPIRV_DISPATCHUTILS_H_
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.cpp b/iree/compiler/Codegen/SPIRV/LaunchConfig.cpp
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.cpp
rename to iree/compiler/Codegen/SPIRV/LaunchConfig.cpp
index 605db44..ce31d1e 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.cpp
+++ b/iree/compiler/Codegen/SPIRV/LaunchConfig.cpp
@@ -13,9 +13,9 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h"
+#include "iree/compiler/Codegen/SPIRV/LaunchConfig.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "llvm/Support/FormatVariadic.h"
#include "mlir/Dialect/Linalg/Analysis/DependenceAnalysis.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h b/iree/compiler/Codegen/SPIRV/LaunchConfig.h
similarity index 96%
rename from iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h
rename to iree/compiler/Codegen/SPIRV/LaunchConfig.h
index 26319d5..99f30fc 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h
+++ b/iree/compiler/Codegen/SPIRV/LaunchConfig.h
@@ -12,8 +12,8 @@
// desired code. This allows sharing codegen infra between different backends.
//
//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_COMMON_LAUNCHCONFIG_H_
-#define IREE_COMPILER_CONVERSION_COMMON_LAUNCHCONFIG_H_
+#ifndef IREE_COMPILER_CODEGEN_COMMON_LAUNCHCONFIG_H_
+#define IREE_COMPILER_CODEGEN_COMMON_LAUNCHCONFIG_H_
#include <array>
#include "llvm/ADT/SmallVector.h"
@@ -141,4 +141,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_COMMON_LAUNCHCONFIG_H_
+#endif // IREE_COMPILER_CODEGEN_COMMON_LAUNCHCONFIG_H_
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h b/iree/compiler/Codegen/SPIRV/MemorySpace.h
similarity index 84%
rename from iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h
rename to iree/compiler/Codegen/SPIRV/MemorySpace.h
index 13a1854..4183e59 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h
+++ b/iree/compiler/Codegen/SPIRV/MemorySpace.h
@@ -13,8 +13,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_LINALGTOSPIRV_MEMORYSPACE_H_
-#define IREE_COMPILER_CONVERSION_LINALGTOSPIRV_MEMORYSPACE_H_
+#ifndef IREE_COMPILER_CODEGEN_SPIRV_MEMORYSPACE_H_
+#define IREE_COMPILER_CODEGEN_SPIRV_MEMORYSPACE_H_
#include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h"
@@ -29,4 +29,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_LINALGTOSPIRV_MEMORYSPACE_H_
+#endif // IREE_COMPILER_CODEGEN_SPIRV_MEMORYSPACE_H_
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp b/iree/compiler/Codegen/SPIRV/Passes.cpp
similarity index 89%
rename from iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp
rename to iree/compiler/Codegen/SPIRV/Passes.cpp
index 0bda998..e65bab6 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/Passes.cpp
+++ b/iree/compiler/Codegen/SPIRV/Passes.cpp
@@ -10,11 +10,11 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/SPIRV/MemorySpace.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "llvm/Support/CommandLine.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
@@ -76,9 +76,9 @@
// flow.dispatch.workgroups performed abstract tiling and distribution. Make
// them concrete now since we know the target and settings now.
- pm.addPass(createLinalgToSPIRVConcretizeTileAmongWorkgroupsPass(options));
+ pm.addPass(createSPIRVConcretizeWorkgroupTilesPass(options));
- pm.addPass(createLinalgToSPIRVTileAndVectorizeOneWorkgroupPass(options));
+ pm.addPass(createSPIRVTileAndVectorizePass(options));
pm.nest<ModuleOp>().addPass(createCanonicalizerPass());
//===--------------------------------------------------------------------===//
@@ -89,9 +89,8 @@
// workgroups.
// - Linalg ops are converted to loop.for ops and mapped to workitems.
//===--------------------------------------------------------------------===//
- pm.addPass(createLinalgToSPIRVConvertToGPUPass());
- pm.nest<ModuleOp>().addNestedPass<FuncOp>(
- createLinalgToSPIRVConvertVectorToGPU());
+ pm.addPass(createSPIRVConvertToGPUPass());
+ pm.nest<ModuleOp>().addNestedPass<FuncOp>(createSPIRVVectorToGPUPass());
pm.nest<ModuleOp>().addPass(createLowerAffinePass());
pm.nest<ModuleOp>().addPass(createCanonicalizerPass());
pm.nest<ModuleOp>().addPass(createCSEPass());
@@ -107,9 +106,9 @@
pm.nest<ModuleOp>().addPass(memref::createFoldSubViewOpsPass());
pm.nest<ModuleOp>().addPass(createCanonicalizerPass());
pm.nest<ModuleOp>().addPass(createCSEPass());
- pm.nest<ModuleOp>().addPass(createLinalgToSPIRVVectorizeMemRefLoadStore());
+ pm.nest<ModuleOp>().addPass(createSPIRVVectorizeLoadStore());
pm.nest<ModuleOp>().addNestedPass<FuncOp>(
- createLinalgToSPIRVVectorToCooperativeMatrixPass());
+ createSPIRVVectorToCooperativeMatrixPass());
pm.nest<ModuleOp>().addNestedPass<FuncOp>(createForOpCanonicalizationPass());
pm.nest<ModuleOp>().addPass(createCanonicalizerPass());
pm.nest<ModuleOp>().addPass(createCSEPass());
@@ -126,7 +125,7 @@
// - All ops are converted to SPIR-V counterparts.
// - spv.module ops are formed to hold all SPIR-V ops.
//===--------------------------------------------------------------------===//
- pm.nest<ModuleOp>().addPass(createLinalgToSPIRVConvertToSPIRVPass());
+ pm.nest<ModuleOp>().addPass(createConvertToSPIRVPass());
//===--------------------------------------------------------------------===//
// SPIR-V dialect level conversions.
@@ -142,8 +141,8 @@
spirvModulePM.addPass(spirv::createUpdateVersionCapabilityExtensionPass());
}
-void buildSPIRVTransformPassPipeline(OpPassManager &pm,
- const SPIRVCodegenOptions &options) {
+void buildSPIRVCodegenPassPipeline(OpPassManager &pm,
+ const SPIRVCodegenOptions &options) {
//===--------------------------------------------------------------------===//
// Inline the impl dispatch function into the wrapper dispatch function.
//
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/ConcretizeTileAmongWorkgroupsPass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVConcretizeWorkgroupTiles.cpp
similarity index 92%
rename from iree/compiler/Conversion/LinalgToSPIRV/ConcretizeTileAmongWorkgroupsPass.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVConcretizeWorkgroupTiles.cpp
index 461fade..4104270 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/ConcretizeTileAmongWorkgroupsPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVConcretizeWorkgroupTiles.cpp
@@ -4,7 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//===- ConcretizeTileAmongWorkgroupsPass.cpp ------------------------------===//
+//===- SPIRVConcretizeWorkgroupTiles.cpp ----------------------------------===//
//
// This pass concretizes hal.interface.workgroup ops by replacing them with
// constant values from the chosen tiling and distribution scheme.
@@ -32,13 +32,13 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/LaunchConfig.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/Utils.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h"
+#include "iree/compiler/Codegen/SPIRV/LaunchConfig.h"
+#include "iree/compiler/Codegen/SPIRV/Utils.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "llvm/ADT/STLExtras.h"
@@ -61,7 +61,7 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-#define DEBUG_TYPE "iree-spirv-concretize-tile-among-workgroups"
+#define DEBUG_TYPE "iree-spirv-concretize-workgroups-tile"
namespace mlir {
namespace iree_compiler {
@@ -279,15 +279,14 @@
/// Concretizes hal.interface.workgroup.* ops with constants from the chosen
/// tiling sheme when possible and perform loop canonicalization afterwards.
-class LinalgToSPIRVConcretizeTileAmongWorkgroupsPass
- : public LinalgToSPIRVConcretizeTileAmongWorkgroupsBase<
- LinalgToSPIRVConcretizeTileAmongWorkgroupsPass> {
+class SPIRVConcretizeWorkgroupTilesPass
+ : public SPIRVConcretizeWorkgroupTilesBase<
+ SPIRVConcretizeWorkgroupTilesPass> {
public:
- LinalgToSPIRVConcretizeTileAmongWorkgroupsPass(
- const SPIRVCodegenOptions &options)
+ SPIRVConcretizeWorkgroupTilesPass(const SPIRVCodegenOptions &options)
: options(options) {}
- LinalgToSPIRVConcretizeTileAmongWorkgroupsPass(
- const LinalgToSPIRVConcretizeTileAmongWorkgroupsPass &that)
+ SPIRVConcretizeWorkgroupTilesPass(
+ const SPIRVConcretizeWorkgroupTilesPass &that)
: options(that.options) {
inlineTripOneLoops = that.inlineTripOneLoops;
}
@@ -404,10 +403,8 @@
} // namespace
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVConcretizeTileAmongWorkgroupsPass(
- const SPIRVCodegenOptions &options) {
- return std::make_unique<LinalgToSPIRVConcretizeTileAmongWorkgroupsPass>(
- options);
+createSPIRVConcretizeWorkgroupTilesPass(const SPIRVCodegenOptions &options) {
+ return std::make_unique<SPIRVConcretizeWorkgroupTilesPass>(options);
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp
similarity index 96%
rename from iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp
index 360a539..15298b2 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/ConvertToGPUPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVConvertToGPU.cpp
@@ -4,7 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//===- ConvertToGPUPass.cpp -----------------------------------------------===//
+//===- SPIRVConvertToGPUPass.cpp ------------------------------------------===//
//
// Partition computation within dispatch function to workgroups/workitems.
//
@@ -13,14 +13,14 @@
#include <array>
#include <numeric>
-#include "iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/Utils.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h"
+#include "iree/compiler/Codegen/SPIRV/MemorySpace.h"
+#include "iree/compiler/Codegen/SPIRV/Utils.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeDialect.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
@@ -441,8 +441,8 @@
namespace {
/// Pass to convert from tiled and fused linalg ops into gpu.func.
-struct LinalgToSPIRVConvertToGPUPass
- : public LinalgToSPIRVConvertToGPUBase<LinalgToSPIRVConvertToGPUPass> {
+struct SPIRVConvertToGPUPass
+ : public SPIRVConvertToGPUBase<SPIRVConvertToGPUPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<AffineDialect, gpu::GPUDialect, memref::MemRefDialect,
scf::SCFDialect, ShapeDialect>();
@@ -644,7 +644,7 @@
patterns.insert<TileAndDistributeCopyOp>(context);
}
-void LinalgToSPIRVConvertToGPUPass::runOnOperation() {
+void SPIRVConvertToGPUPass::runOnOperation() {
MLIRContext *context = &getContext();
ConversionTarget target(*context);
// After this pass Linalg and scf.parallel ops should be gone.
@@ -678,8 +678,8 @@
}
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVConvertToGPUPass() {
- return std::make_unique<LinalgToSPIRVConvertToGPUPass>();
+createSPIRVConvertToGPUPass() {
+ return std::make_unique<SPIRVConvertToGPUPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/FoldGPUProcessorIDUses.cpp b/iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp
similarity index 94%
rename from iree/compiler/Conversion/LinalgToSPIRV/FoldGPUProcessorIDUses.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp
index 2c1dcd0..61f68ef 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/FoldGPUProcessorIDUses.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVFoldGPUProcessorIDUses.cpp
@@ -4,14 +4,14 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//===- FoldGPUProcessorIDUses.cpp -----------------------------------------===//
+//===- SPIRVFoldGPUProcessorIDUses.cpp ------------------------------------===//
//
// This file implements patterns and passes for folding GPU processor ID uses.
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/Support/Debug.h"
@@ -29,7 +29,7 @@
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-#define DEBUG_TYPE "iree-fold-gpu-procid-uses"
+#define DEBUG_TYPE "iree-spirv-fold-gpu-procid-uses"
namespace mlir {
namespace iree_compiler {
@@ -255,12 +255,10 @@
};
/// Tests processor ID use folding patterns.
-struct LinalgToSPIRVFoldProcessorIDUsesPass
- : public LinalgToSPIRVFoldProcessorIDUsesBase<
- LinalgToSPIRVFoldProcessorIDUsesPass> {
- LinalgToSPIRVFoldProcessorIDUsesPass() = default;
- LinalgToSPIRVFoldProcessorIDUsesPass(
- const LinalgToSPIRVFoldProcessorIDUsesPass &pass) {}
+struct SPIRVFoldProcessorIDUsesPass
+ : public SPIRVFoldProcessorIDUsesBase<SPIRVFoldProcessorIDUsesPass> {
+ SPIRVFoldProcessorIDUsesPass() = default;
+ SPIRVFoldProcessorIDUsesPass(const SPIRVFoldProcessorIDUsesPass &pass) {}
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<AffineDialect, gpu::GPUDialect>();
@@ -284,8 +282,8 @@
}
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVFoldProcessorIDUsesPass() {
- return std::make_unique<LinalgToSPIRVFoldProcessorIDUsesPass>();
+createSPIRVFoldProcessorIDUsesPass() {
+ return std::make_unique<SPIRVFoldProcessorIDUsesPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/TileAndVectorizeInOneWorkgroupPass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorize.cpp
similarity index 95%
rename from iree/compiler/Conversion/LinalgToSPIRV/TileAndVectorizeInOneWorkgroupPass.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorize.cpp
index 76d607b..58aa32c 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/TileAndVectorizeInOneWorkgroupPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVTileAndVectorize.cpp
@@ -4,21 +4,21 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//===- TileAndVectorizeInOneWorkgroup.cpp ---------------------------------===//
+//===- SPIRVTileAndVectorize.cpp ------------------------------------------===//
//
// This pass tiles and vectorizes Linalg ops on buffers within in a single
// workgroup.
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/LinalgToSPIRV/KernelDispatchUtils.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/Utils.h"
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/SPIRV/KernelDispatchUtils.h"
+#include "iree/compiler/Codegen/SPIRV/MemorySpace.h"
+#include "iree/compiler/Codegen/SPIRV/Utils.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALDialect.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/Shape/IR/ShapeDialect.h"
@@ -42,7 +42,7 @@
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
#include "mlir/Transforms/LoopUtils.h"
-#define DEBUG_TYPE "iree-spirv-tile-and-vectorize-in-one-workgroup"
+#define DEBUG_TYPE "iree-spirv-tile-and-vectorize"
namespace mlir {
namespace iree_compiler {
@@ -76,15 +76,12 @@
namespace {
/// Function pass that implements tiling and fusion in Linalg on buffers.
-class LinalgToSPIRVTileAndVectorizeOneWorkgroupPass
- : public LinalgToSPIRVTileAndVectorizeOneWorkgroupBase<
- LinalgToSPIRVTileAndVectorizeOneWorkgroupPass> {
+class SPIRVTileAndVectorizePass
+ : public SPIRVTileAndVectorizeBase<SPIRVTileAndVectorizePass> {
public:
- LinalgToSPIRVTileAndVectorizeOneWorkgroupPass(
- const SPIRVCodegenOptions &passOptions)
+ SPIRVTileAndVectorizePass(const SPIRVCodegenOptions &passOptions)
: options(passOptions) {}
- LinalgToSPIRVTileAndVectorizeOneWorkgroupPass(
- const LinalgToSPIRVTileAndVectorizeOneWorkgroupPass &pass)
+ SPIRVTileAndVectorizePass(const SPIRVTileAndVectorizePass &pass)
: options(pass.options) {}
void getDependentDialects(DialectRegistry ®istry) const override {
@@ -538,7 +535,7 @@
// Main pass implementation
//====---------------------------------------------------------------------===//
-void LinalgToSPIRVTileAndVectorizeOneWorkgroupPass::runOnOperation() {
+void SPIRVTileAndVectorizePass::runOnOperation() {
MLIRContext *context = &getContext();
IREE::HAL::ExecutableTargetOp targetOp = getOperation();
ModuleOp module = targetOp.getInnerModule();
@@ -751,10 +748,8 @@
//===----------------------------------------------------------------------===//
std::unique_ptr<OperationPass<IREE::HAL::ExecutableTargetOp>>
-createLinalgToSPIRVTileAndVectorizeOneWorkgroupPass(
- const SPIRVCodegenOptions &options) {
- return std::make_unique<LinalgToSPIRVTileAndVectorizeOneWorkgroupPass>(
- options);
+createSPIRVTileAndVectorizePass(const SPIRVCodegenOptions &options) {
+ return std::make_unique<SPIRVTileAndVectorizePass>(options);
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/VectorToCooperativeMatrixPass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp
similarity index 96%
rename from iree/compiler/Conversion/LinalgToSPIRV/VectorToCooperativeMatrixPass.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp
index 67e1d79..5974b8c 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/VectorToCooperativeMatrixPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorToCooperativeMatrix.cpp
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "mlir/Analysis/SliceAnalysis.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SCF/SCF.h"
@@ -250,9 +250,9 @@
const CooperativeMatrixAnalysis &cooperativeMatrixAnalysis;
};
-struct LinalgToSPIRVVectorToCooperativeMatrixPass final
- : public LinalgToSPIRVVectorToCooperativeMatrixBase<
- LinalgToSPIRVVectorToCooperativeMatrixPass> {
+struct SPIRVVectorToCooperativeMatrixPass final
+ : public SPIRVVectorToCooperativeMatrixBase<
+ SPIRVVectorToCooperativeMatrixPass> {
void runOnOperation() override {
FuncOp funcOp = getOperation();
auto targetAttr = spirv::lookupTargetEnv(funcOp);
@@ -301,8 +301,8 @@
} // namespace
std::unique_ptr<OperationPass<FuncOp>>
-createLinalgToSPIRVVectorToCooperativeMatrixPass() {
- return std::make_unique<LinalgToSPIRVVectorToCooperativeMatrixPass>();
+createSPIRVVectorToCooperativeMatrixPass() {
+ return std::make_unique<SPIRVVectorToCooperativeMatrixPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/VectorToGPUPass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorToGPU.cpp
similarity index 83%
rename from iree/compiler/Conversion/LinalgToSPIRV/VectorToGPUPass.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVVectorToGPU.cpp
index 32b4e75..51abf8d 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/VectorToGPUPass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorToGPU.cpp
@@ -4,7 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//===---- VectorToGPUPass.cpp - Pass for the final SPIR-V conversion ------===//
+//===---- SPIRVVectorToGPUPass.cpp - Pass for the final SPIR-V conversion -===//
//
// This file implement a pass to convert vector dialect operations to GPU
// operations distributed across a subgroup.
@@ -12,11 +12,11 @@
//===----------------------------------------------------------------------===//
#include <memory>
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "llvm/ADT/STLExtras.h"
#include "llvm/Support/FormatVariadic.h"
#include "mlir/Dialect/GPU/GPUDialect.h"
@@ -42,9 +42,8 @@
namespace iree_compiler {
namespace {
-struct LinalgToSPIRVConvertVectorToGPUPass
- : public LinalgToSPIRVConvertVectorToGPUBase<
- LinalgToSPIRVConvertVectorToGPUPass> {
+struct SPIRVVectorToGPUPass
+ : public SPIRVVectorToGPUBase<SPIRVVectorToGPUPass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<AffineDialect, gpu::GPUDialect, memref::MemRefDialect,
scf::SCFDialect, vector::VectorDialect>();
@@ -73,8 +72,8 @@
}
};
-void LinalgToSPIRVConvertVectorToGPUPass::tileAndVectorizeLinalgCopy(
- FuncOp funcOp, MLIRContext *context) {
+void SPIRVVectorToGPUPass::tileAndVectorizeLinalgCopy(FuncOp funcOp,
+ MLIRContext *context) {
// 1. Tile linalg and distribute it on invocations.
std::unique_ptr<ConversionTarget> target =
std::make_unique<ConversionTarget>(*context);
@@ -107,7 +106,7 @@
(void)applyPatternsAndFoldGreedily(funcOp, std::move(vectorizationPatterns));
}
-void LinalgToSPIRVConvertVectorToGPUPass::runOnOperation() {
+void SPIRVVectorToGPUPass::runOnOperation() {
MLIRContext *context = &getContext();
FuncOp funcOp = getOperation();
tileAndVectorizeLinalgCopy(funcOp, context);
@@ -117,8 +116,8 @@
//===----------------------------------------------------------------------===//
// Pass entry point and registration
//===----------------------------------------------------------------------===//
-std::unique_ptr<OperationPass<FuncOp>> createLinalgToSPIRVConvertVectorToGPU() {
- return std::make_unique<LinalgToSPIRVConvertVectorToGPUPass>();
+std::unique_ptr<OperationPass<FuncOp>> createSPIRVVectorToGPUPass() {
+ return std::make_unique<SPIRVVectorToGPUPass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/VectorizeMemrefLoadStorePass.cpp b/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
similarity index 96%
rename from iree/compiler/Conversion/LinalgToSPIRV/VectorizeMemrefLoadStorePass.cpp
rename to iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
index a2d2ccb..b821de9 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/VectorizeMemrefLoadStorePass.cpp
+++ b/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp
@@ -12,9 +12,9 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/PassDetail.h"
+#include "iree/compiler/Codegen/Passes.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/IREE/IR/IREEOps.h"
#include "llvm/ADT/TypeSwitch.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
@@ -408,12 +408,12 @@
}
};
-class LinalgToSPIRVVectorizeMemRefLoadStorePass final
- : public LinalgToSPIRVVectorizeMemRefLoadStoreBase<
- LinalgToSPIRVVectorizeMemRefLoadStorePass> {
+class SPIRVVectorizeLoadStorePass final
+ : public SPIRVVectorizeLoadStoreBase<SPIRVVectorizeLoadStorePass> {
void getDependentDialects(DialectRegistry ®istry) const override {
registry.insert<memref::MemRefDialect>();
}
+
void runOnOperation() override;
private:
@@ -449,7 +449,7 @@
return success();
}
-void LinalgToSPIRVVectorizeMemRefLoadStorePass::runOnOperation() {
+void SPIRVVectorizeLoadStorePass::runOnOperation() {
// Uses the signature conversion methodology of the dialect conversion
// framework to implement the conversion.
ModuleOp module = getOperation();
@@ -495,9 +495,8 @@
}
}
-std::unique_ptr<OperationPass<ModuleOp>>
-createLinalgToSPIRVVectorizeMemRefLoadStore() {
- return std::make_unique<LinalgToSPIRVVectorizeMemRefLoadStorePass>();
+std::unique_ptr<OperationPass<ModuleOp>> createSPIRVVectorizeLoadStore() {
+ return std::make_unique<SPIRVVectorizeLoadStorePass>();
}
} // namespace iree_compiler
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/Utils.cpp b/iree/compiler/Codegen/SPIRV/Utils.cpp
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/Utils.cpp
rename to iree/compiler/Codegen/SPIRV/Utils.cpp
index c62a000..2a47649 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/Utils.cpp
+++ b/iree/compiler/Codegen/SPIRV/Utils.cpp
@@ -10,10 +10,10 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/LinalgToSPIRV/Utils.h"
+#include "iree/compiler/Codegen/SPIRV/Utils.h"
-#include "iree/compiler/Conversion/LinalgToSPIRV/MemorySpace.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/SPIRV/MemorySpace.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "mlir/Dialect/GPU/GPUDialect.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
#include "mlir/Dialect/Linalg/Utils/Utils.h"
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/Utils.h b/iree/compiler/Codegen/SPIRV/Utils.h
similarity index 92%
rename from iree/compiler/Conversion/LinalgToSPIRV/Utils.h
rename to iree/compiler/Codegen/SPIRV/Utils.h
index d9805c7..f0d8635 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/Utils.h
+++ b/iree/compiler/Codegen/SPIRV/Utils.h
@@ -9,8 +9,8 @@
// Utility functions used while lowering from Linalg to SPIRV.
//
//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_LINALGTOSPIRV_UTILS_H_
-#define IREE_COMPILER_CONVERSION_LINALGTOSPIRV_UTILS_H_
+#ifndef IREE_COMPILER_CODEGEN_SPIRV_UTILS_H_
+#define IREE_COMPILER_CODEGEN_SPIRV_UTILS_H_
#include "mlir/Dialect/Linalg/Utils/Utils.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
@@ -56,4 +56,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_LINALGTOSPIRV_UTILS_H_
+#endif // IREE_COMPILER_CODEGEN_SPIRV_UTILS_H_
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/BUILD b/iree/compiler/Codegen/SPIRV/test/BUILD
similarity index 84%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/BUILD
rename to iree/compiler/Codegen/SPIRV/test/BUILD
index a675f42..55f3118 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/BUILD
+++ b/iree/compiler/Codegen/SPIRV/test/BUILD
@@ -19,26 +19,25 @@
name = "lit",
srcs = enforce_glob(
[
- "concretize_tile_among_workgroups.mlir",
- "concretize_tile_among_workgroups_dynamic.mlir",
+ "concretize_workgroup_tiles.mlir",
+ "concretize_workgroup_tiles_dynamic.mlir",
"convert_to_gpu.mlir",
"convert_to_spirv.mlir",
- "dead_alloc.mlir",
"fold_gpu_procid_uses.mlir",
"materialize_launch_configuration.mlir",
"materialize_launch_configuration2.mlir",
"pipeline_matmul_cooperative_matrix.mlir",
"pipeline_matmul_vectorization.mlir",
"promote_workgroup_memory.mlir",
+ "tile_and_vectorize.mlir",
"tile_and_vectorize_batch_matmul.mlir",
"tile_and_vectorize_conv.mlir",
- "tile_and_vectorize_in_one_workgroup.mlir",
"tile_and_vectorize_matmul.mlir",
"vector_to_cooperative_matrix.mlir",
"vector_to_gpu.mlir",
"vectorize_elementwise_ops.mlir",
"vectorize_matmul.mlir",
- "vectorize_memref_load_store.mlir",
+ "vectorize_load_store.mlir",
],
include = ["*.mlir"],
),
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/CMakeLists.txt b/iree/compiler/Codegen/SPIRV/test/CMakeLists.txt
similarity index 83%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/CMakeLists.txt
rename to iree/compiler/Codegen/SPIRV/test/CMakeLists.txt
index a85b005..3022f00 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/CMakeLists.txt
+++ b/iree/compiler/Codegen/SPIRV/test/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/LinalgToSPIRV/test/BUILD #
+# iree/compiler/Codegen/SPIRV/test/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -14,26 +14,25 @@
NAME
lit
SRCS
- "concretize_tile_among_workgroups.mlir"
- "concretize_tile_among_workgroups_dynamic.mlir"
+ "concretize_workgroup_tiles.mlir"
+ "concretize_workgroup_tiles_dynamic.mlir"
"convert_to_gpu.mlir"
"convert_to_spirv.mlir"
- "dead_alloc.mlir"
"fold_gpu_procid_uses.mlir"
"materialize_launch_configuration.mlir"
"materialize_launch_configuration2.mlir"
"pipeline_matmul_cooperative_matrix.mlir"
"pipeline_matmul_vectorization.mlir"
"promote_workgroup_memory.mlir"
+ "tile_and_vectorize.mlir"
"tile_and_vectorize_batch_matmul.mlir"
"tile_and_vectorize_conv.mlir"
- "tile_and_vectorize_in_one_workgroup.mlir"
"tile_and_vectorize_matmul.mlir"
"vector_to_cooperative_matrix.mlir"
"vector_to_gpu.mlir"
"vectorize_elementwise_ops.mlir"
+ "vectorize_load_store.mlir"
"vectorize_matmul.mlir"
- "vectorize_memref_load_store.mlir"
DATA
iree::tools::IreeFileCheck
iree::tools::iree-opt
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/concretize_tile_among_workgroups.mlir b/iree/compiler/Codegen/SPIRV/test/concretize_workgroup_tiles.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/concretize_tile_among_workgroups.mlir
rename to iree/compiler/Codegen/SPIRV/test/concretize_workgroup_tiles.mlir
index f973749..eea0181 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/concretize_tile_among_workgroups.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/concretize_workgroup_tiles.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-spirv-workgroup-tile-size=0,4,4,16 -iree-spirv-workgroup-size=4,4,1 -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-tile-among-workgroups))" -canonicalize -cse %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-spirv-workgroup-tile-size=0,4,4,16 -iree-spirv-workgroup-size=4,4,1 -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-workgroup-tiles))" -canonicalize -cse %s | IreeFileCheck %s
hal.executable @conv2d_static_shape attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/concretize_tile_among_workgroups_dynamic.mlir b/iree/compiler/Codegen/SPIRV/test/concretize_workgroup_tiles_dynamic.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/concretize_tile_among_workgroups_dynamic.mlir
rename to iree/compiler/Codegen/SPIRV/test/concretize_workgroup_tiles_dynamic.mlir
index 30068e5..4826f23 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/concretize_tile_among_workgroups_dynamic.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/concretize_workgroup_tiles_dynamic.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-spirv-workgroup-tile-size=4,16 -iree-spirv-workgroup-size=4,4,1 -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-tile-among-workgroups))" -canonicalize -cse %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-spirv-workgroup-tile-size=4,16 -iree-spirv-workgroup-size=4,4,1 -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-workgroup-tiles))" -canonicalize -cse %s | IreeFileCheck %s
hal.executable @matmul_dynamic_shape attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir b/iree/compiler/Codegen/SPIRV/test/convert_to_gpu.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
rename to iree/compiler/Codegen/SPIRV/test/convert_to_gpu.mlir
index 05ef416..0c47061 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_gpu.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/convert_to_gpu.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.target(iree-codegen-convert-to-gpu))' -canonicalize -cse %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline='hal.executable(hal.executable.target(iree-spirv-convert-to-gpu))' -canonicalize -cse %s | IreeFileCheck %s
#map0 = affine_map<(d0, d1, d2, d3) -> (d0, d1, d2, d3)>
hal.executable @parallel_4D attributes {sym_visibility = "private"} {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
rename to iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
index cb94378..e8d1b3b 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/convert_to_spirv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/convert_to_spirv.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-codegen-convert-to-spirv %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-convert-to-spirv %s | IreeFileCheck %s
module attributes {spv.target_env = #spv.target_env<#spv.vce<v1.3, [Shader], []>, {}>} {
// CHECK-LABEL: spv.module
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/fold_gpu_procid_uses.mlir b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/fold_gpu_procid_uses.mlir
rename to iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
index b20254e..d294e78 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/fold_gpu_procid_uses.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/fold_gpu_procid_uses.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-codegen-fold-gpu-procid-uses))" %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-fold-gpu-procid-uses))" %s | IreeFileCheck %s
hal.executable @fold_block_id attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/materialize_launch_configuration.mlir b/iree/compiler/Codegen/SPIRV/test/materialize_launch_configuration.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/materialize_launch_configuration.mlir
rename to iree/compiler/Codegen/SPIRV/test/materialize_launch_configuration.mlir
index b4188b1..f363d37 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/materialize_launch_configuration.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/materialize_launch_configuration.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-tile-among-workgroups))" -canonicalize -cse -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-workgroup-tiles))" -canonicalize -cse -split-input-file %s | IreeFileCheck %s
hal.executable @matmul_tensors attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/materialize_launch_configuration2.mlir b/iree/compiler/Codegen/SPIRV/test/materialize_launch_configuration2.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/materialize_launch_configuration2.mlir
rename to iree/compiler/Codegen/SPIRV/test/materialize_launch_configuration2.mlir
index bed2164..759cae6 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/materialize_launch_configuration2.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/materialize_launch_configuration2.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-codegen-convert-to-gpu))" -canonicalize -cse -split-input-file %s | IreeFileCheck %s
+// RUN: iree-opt -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-convert-to-gpu))" -canonicalize -cse -split-input-file %s | IreeFileCheck %s
hal.executable @add attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/pipeline_matmul_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/pipeline_matmul_cooperative_matrix.mlir
rename to iree/compiler/Codegen/SPIRV/test/pipeline_matmul_cooperative_matrix.mlir
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/pipeline_matmul_vectorization.mlir b/iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/pipeline_matmul_vectorization.mlir
rename to iree/compiler/Codegen/SPIRV/test/pipeline_matmul_vectorization.mlir
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/promote_workgroup_memory.mlir b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/promote_workgroup_memory.mlir
rename to iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
index 49cc6ba..0937ec3 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/promote_workgroup_memory.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/promote_workgroup_memory.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize-in-one-workgroup,canonicalize,cse))" -iree-spirv-use-workgroup-memory %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize,canonicalize,cse))" -iree-spirv-use-workgroup-memory %s | IreeFileCheck %s
hal.executable @matmul_promote_workgroup_memory attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_in_one_workgroup.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_in_one_workgroup.mlir
rename to iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
index 06f3ee6..fd2101c 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_in_one_workgroup.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize-in-one-workgroup,canonicalize,cse))" %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize,canonicalize,cse))" %s | IreeFileCheck %s
#map0 = affine_map<()[s0] -> (s0 * 8)>
#map1 = affine_map<()[s0, s1] -> (8, s1 - s0 * 8)>
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_batch_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_batch_matmul.mlir
rename to iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
index 161853a..b538796 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_batch_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_batch_matmul.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-tile-among-workgroups,iree-spirv-tile-and-vectorize-in-one-workgroup))" -canonicalize -cse -iree-spirv-workgroup-tile-size=1,8,64,4 -iree-spirv-invocation-tile-size=1,8,4,4 -iree-spirv-workgroup-size=16,1,1 %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-workgroup-tiles,iree-spirv-tile-and-vectorize))" -canonicalize -cse -iree-spirv-workgroup-tile-size=1,8,64,4 -iree-spirv-invocation-tile-size=1,8,4,4 -iree-spirv-workgroup-size=16,1,1 %s | IreeFileCheck %s
hal.executable @batch_matmul_static_shape attributes {sym_visibility = "private"} {
hal.interface @io attributes {sym_visibility = "private"} {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_conv.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_conv.mlir
rename to iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
index d864f6c..a4d3751 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_conv.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_conv.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-tile-among-workgroups,iree-spirv-tile-and-vectorize-in-one-workgroup))" -canonicalize -cse %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-workgroup-tiles,iree-spirv-tile-and-vectorize))" -canonicalize -cse %s | IreeFileCheck %s
hal.executable @conv_static_shape_f32 attributes {sym_visibility = "private"} {
hal.interface @io {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
similarity index 96%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_matmul.mlir
rename to iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
index f915a81..0cecd5d 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/tile_and_vectorize_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/tile_and_vectorize_matmul.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-tile-among-workgroups,iree-spirv-tile-and-vectorize-in-one-workgroup))" -canonicalize -cse -iree-spirv-workgroup-tile-size=8,64,4 -iree-spirv-invocation-tile-size=8,4,4 -iree-spirv-workgroup-size=16,1,1 %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-concretize-workgroup-tiles,iree-spirv-tile-and-vectorize))" -canonicalize -cse -iree-spirv-workgroup-tile-size=8,64,4 -iree-spirv-invocation-tile-size=8,4,4 -iree-spirv-workgroup-size=16,1,1 %s | IreeFileCheck %s
hal.executable @matmul_static_shape_f16 attributes {sym_visibility = "private"} {
hal.interface @io attributes {sym_visibility = "private"} {
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/vector_to_cooperative_matrix.mlir b/iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
similarity index 100%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/vector_to_cooperative_matrix.mlir
rename to iree/compiler/Codegen/SPIRV/test/vector_to_cooperative_matrix.mlir
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/vector_to_gpu.mlir b/iree/compiler/Codegen/SPIRV/test/vector_to_gpu.mlir
similarity index 96%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/vector_to_gpu.mlir
rename to iree/compiler/Codegen/SPIRV/test/vector_to_gpu.mlir
index c92aa33..e808be4 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/vector_to_gpu.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vector_to_gpu.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-codegen-vector-to-gpu %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-spirv-vector-to-gpu %s | IreeFileCheck %s
#map0 = affine_map<(d0, d1)[s0] -> (d0 * 4096 + s0 + d1)>
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_elementwise_ops.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
similarity index 97%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_elementwise_ops.mlir
rename to iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
index 82bddfb..b583f79 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_elementwise_ops.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_elementwise_ops.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize-in-one-workgroup,canonicalize,cse))" %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize,canonicalize,cse))" %s | IreeFileCheck %s
// CHECK-LABEL: func @elementwise_static_shape
// CHECK: vector.transfer_read %{{.+}}[%c0], {{.+}} memref<4xf32, #{{.+}}>, vector<4xf32>
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_memref_load_store.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_memref_load_store.mlir
rename to iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
index 5fdaba3..7a3af96 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_memref_load_store.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_load_store.mlir
@@ -1,4 +1,4 @@
-// RUN: iree-opt -split-input-file -iree-spirv-vectorize-memref-load-store -canonicalize %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -iree-spirv-vectorize-load-store -canonicalize %s | IreeFileCheck %s
// CHECK-LABEL: func @copy
// CHECK-SAME: (%[[ARG0:.+]]: memref<4096x1024xvector<4xf32>>
diff --git a/iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_matmul.mlir b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
similarity index 98%
rename from iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_matmul.mlir
rename to iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
index 3ffa05b..bffa63a 100644
--- a/iree/compiler/Conversion/LinalgToSPIRV/test/vectorize_matmul.mlir
+++ b/iree/compiler/Codegen/SPIRV/test/vectorize_matmul.mlir
@@ -1,5 +1,5 @@
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize-in-one-workgroup,canonicalize,cse))" %s | IreeFileCheck %s
-// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize-in-one-workgroup,canonicalize,cse))" -iree-spirv-use-workgroup-memory %s | IreeFileCheck %s -check-prefix=PROMOTE
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize,canonicalize,cse))" %s | IreeFileCheck %s
+// RUN: iree-opt -split-input-file -pass-pipeline="hal.executable(hal.executable.target(iree-spirv-tile-and-vectorize,canonicalize,cse))" -iree-spirv-use-workgroup-memory %s | IreeFileCheck %s -check-prefix=PROMOTE
hal.executable @matmul_static_shape attributes {sym_visibility = "private"} {
hal.interface @io attributes {sym_visibility = "private"} {
diff --git a/iree/compiler/Conversion/Transforms/AffineMinCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
similarity index 98%
rename from iree/compiler/Conversion/Transforms/AffineMinCanonicalization.cpp
rename to iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
index 214a579..e2ac924 100644
--- a/iree/compiler/Conversion/Transforms/AffineMinCanonicalization.cpp
+++ b/iree/compiler/Codegen/Transforms/AffineMinCanonicalization.cpp
@@ -10,7 +10,7 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Affine/Utils.h"
#include "mlir/Dialect/StandardOps/IR/Ops.h"
diff --git a/iree/compiler/Conversion/Transforms/AffineMinDistributedSCFCanonicalization.cpp b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
similarity index 98%
rename from iree/compiler/Conversion/Transforms/AffineMinDistributedSCFCanonicalization.cpp
rename to iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
index 6bfa6f2..32c8500 100644
--- a/iree/compiler/Conversion/Transforms/AffineMinDistributedSCFCanonicalization.cpp
+++ b/iree/compiler/Codegen/Transforms/AffineMinDistributedSCFCanonicalization.cpp
@@ -10,7 +10,7 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "mlir/Dialect/Affine/Utils.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/PatternMatch.h"
diff --git a/iree/compiler/Conversion/Transforms/BUILD b/iree/compiler/Codegen/Transforms/BUILD
similarity index 96%
rename from iree/compiler/Conversion/Transforms/BUILD
rename to iree/compiler/Codegen/Transforms/BUILD
index 19a5694..97b1d67 100644
--- a/iree/compiler/Conversion/Transforms/BUILD
+++ b/iree/compiler/Codegen/Transforms/BUILD
@@ -22,7 +22,7 @@
"Transforms.h",
],
deps = [
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/HAL/IR",
"@llvm-project//llvm:Support",
"@llvm-project//mlir:Affine",
diff --git a/iree/compiler/Conversion/Transforms/CMakeLists.txt b/iree/compiler/Codegen/Transforms/CMakeLists.txt
similarity index 91%
rename from iree/compiler/Conversion/Transforms/CMakeLists.txt
rename to iree/compiler/Codegen/Transforms/CMakeLists.txt
index f366a76..0bc1a05 100644
--- a/iree/compiler/Conversion/Transforms/CMakeLists.txt
+++ b/iree/compiler/Codegen/Transforms/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/Transforms/BUILD #
+# iree/compiler/Codegen/Transforms/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
@@ -32,7 +32,7 @@
MLIRSupport
MLIRTransforms
MLIRVector
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::IR
PUBLIC
)
diff --git a/iree/compiler/Conversion/Transforms/RemoveSingleIterationLoop.cpp b/iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp
similarity index 98%
rename from iree/compiler/Conversion/Transforms/RemoveSingleIterationLoop.cpp
rename to iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp
index 0041d90..ef6398c 100644
--- a/iree/compiler/Conversion/Transforms/RemoveSingleIterationLoop.cpp
+++ b/iree/compiler/Codegen/Transforms/RemoveSingleIterationLoop.cpp
@@ -11,7 +11,7 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
#include "mlir/Dialect/Affine/Utils.h"
#include "mlir/IR/BuiltinOps.h"
#include "mlir/IR/PatternMatch.h"
diff --git a/iree/compiler/Conversion/Transforms/Transforms.cpp b/iree/compiler/Codegen/Transforms/Transforms.cpp
similarity index 97%
rename from iree/compiler/Conversion/Transforms/Transforms.cpp
rename to iree/compiler/Codegen/Transforms/Transforms.cpp
index f4bd358..5c01ec1 100644
--- a/iree/compiler/Conversion/Transforms/Transforms.cpp
+++ b/iree/compiler/Codegen/Transforms/Transforms.cpp
@@ -10,10 +10,10 @@
//
//===----------------------------------------------------------------------===//
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
+#include "iree/compiler/Codegen/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "mlir/Dialect/Linalg/IR/LinalgInterfaces.h"
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
diff --git a/iree/compiler/Conversion/Transforms/Transforms.h b/iree/compiler/Codegen/Transforms/Transforms.h
similarity index 95%
rename from iree/compiler/Conversion/Transforms/Transforms.h
rename to iree/compiler/Codegen/Transforms/Transforms.h
index 69fb111..025b19b 100644
--- a/iree/compiler/Conversion/Transforms/Transforms.h
+++ b/iree/compiler/Codegen/Transforms/Transforms.h
@@ -9,8 +9,8 @@
// Defines transformations that are common to backends
//
//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_COMMON_TRANSFORMS_H_
-#define IREE_COMPILER_CONVERSION_COMMON_TRANSFORMS_H_
+#ifndef IREE_COMPILER_CODEGEN_COMMON_TRANSFORMS_H_
+#define IREE_COMPILER_CODEGEN_COMMON_TRANSFORMS_H_
#include "mlir/Dialect/Linalg/IR/LinalgOps.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
@@ -94,4 +94,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_COMMON_TRANSFORMS_H_
+#endif // IREE_COMPILER_CODEGEN_COMMON_TRANSFORMS_H_
diff --git a/iree/compiler/Conversion/Utils/BUILD b/iree/compiler/Codegen/Utils/BUILD
similarity index 100%
rename from iree/compiler/Conversion/Utils/BUILD
rename to iree/compiler/Codegen/Utils/BUILD
diff --git a/iree/compiler/Conversion/Utils/CMakeLists.txt b/iree/compiler/Codegen/Utils/CMakeLists.txt
similarity index 93%
rename from iree/compiler/Conversion/Utils/CMakeLists.txt
rename to iree/compiler/Codegen/Utils/CMakeLists.txt
index 5a19817..73966ad 100644
--- a/iree/compiler/Conversion/Utils/CMakeLists.txt
+++ b/iree/compiler/Codegen/Utils/CMakeLists.txt
@@ -1,6 +1,6 @@
################################################################################
# Autogenerated by build_tools/bazel_to_cmake/bazel_to_cmake.py from #
-# iree/compiler/Conversion/Utils/BUILD #
+# iree/compiler/Codegen/Utils/BUILD #
# #
# Use iree_cmake_extra_content from iree/build_defs.oss.bzl to add arbitrary #
# CMake-only content. #
diff --git a/iree/compiler/Conversion/Utils/MarkerUtils.cpp b/iree/compiler/Codegen/Utils/MarkerUtils.cpp
similarity index 97%
rename from iree/compiler/Conversion/Utils/MarkerUtils.cpp
rename to iree/compiler/Codegen/Utils/MarkerUtils.cpp
index 454d6a1..4f3ed32 100644
--- a/iree/compiler/Conversion/Utils/MarkerUtils.cpp
+++ b/iree/compiler/Codegen/Utils/MarkerUtils.cpp
@@ -4,7 +4,7 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
#include "mlir/IR/Attributes.h"
diff --git a/iree/compiler/Conversion/Utils/MarkerUtils.h b/iree/compiler/Codegen/Utils/MarkerUtils.h
similarity index 92%
rename from iree/compiler/Conversion/Utils/MarkerUtils.h
rename to iree/compiler/Codegen/Utils/MarkerUtils.h
index f4d0b65..05d73db 100644
--- a/iree/compiler/Conversion/Utils/MarkerUtils.h
+++ b/iree/compiler/Codegen/Utils/MarkerUtils.h
@@ -11,8 +11,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef IREE_COMPILER_CONVERSION_CODEGENUTILS_MARKERUTILS_H_
-#define IREE_COMPILER_CONVERSION_CODEGENUTILS_MARKERUTILS_H_
+#ifndef IREE_COMPILER_CODEGEN_CODEGENUTILS_MARKERUTILS_H_
+#define IREE_COMPILER_CODEGEN_CODEGENUTILS_MARKERUTILS_H_
#include "llvm/ADT/ArrayRef.h"
#include "mlir/IR/Operation.h"
@@ -66,4 +66,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_CODEGENUTILS_MARKERUTILS_H_
+#endif // IREE_COMPILER_CODEGEN_CODEGENUTILS_MARKERUTILS_H_
diff --git a/iree/compiler/Conversion/Utils/Utils.cpp b/iree/compiler/Codegen/Utils/Utils.cpp
similarity index 97%
rename from iree/compiler/Conversion/Utils/Utils.cpp
rename to iree/compiler/Codegen/Utils/Utils.cpp
index e1a1440..28af457 100644
--- a/iree/compiler/Conversion/Utils/Utils.cpp
+++ b/iree/compiler/Codegen/Utils/Utils.cpp
@@ -4,9 +4,9 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#include "iree/compiler/Conversion/Utils/Utils.h"
+#include "iree/compiler/Codegen/Utils/Utils.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
+#include "iree/compiler/Codegen/Utils/MarkerUtils.h"
#include "iree/compiler/Dialect/Flow/IR/FlowOps.h"
#include "llvm/ADT/TypeSwitch.h"
#include "mlir/Dialect/Linalg/Utils/Utils.h"
diff --git a/iree/compiler/Conversion/Utils/Utils.h b/iree/compiler/Codegen/Utils/Utils.h
similarity index 94%
rename from iree/compiler/Conversion/Utils/Utils.h
rename to iree/compiler/Codegen/Utils/Utils.h
index 0d8046a..e579f0b 100644
--- a/iree/compiler/Conversion/Utils/Utils.h
+++ b/iree/compiler/Codegen/Utils/Utils.h
@@ -4,8 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-#ifndef IREE_COMPILER_CONVERSION_UTILS_UTILS_H_
-#define IREE_COMPILER_CONVERSION_UTILS_UTILS_H_
+#ifndef IREE_COMPILER_CODEGEN_UTILS_UTILS_H_
+#define IREE_COMPILER_CODEGEN_UTILS_UTILS_H_
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "llvm/ADT/StringMap.h"
@@ -69,4 +69,4 @@
} // namespace iree_compiler
} // namespace mlir
-#endif // IREE_COMPILER_CONVERSION_UTILS_UTILS_H_
+#endif // IREE_COMPILER_CODEGEN_UTILS_UTILS_H_
diff --git a/iree/compiler/Conversion/LinalgToLLVM/LinalgVectorizePass.cpp b/iree/compiler/Conversion/LinalgToLLVM/LinalgVectorizePass.cpp
deleted file mode 100644
index c1bf301..0000000
--- a/iree/compiler/Conversion/LinalgToLLVM/LinalgVectorizePass.cpp
+++ /dev/null
@@ -1,122 +0,0 @@
-// Copyright 2021 The IREE Authors
-//
-// Licensed 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 "iree/compiler/Conversion/PassDetail.h"
-#include "iree/compiler/Conversion/Passes.h"
-#include "iree/compiler/Conversion/Transforms/Transforms.h"
-#include "iree/compiler/Conversion/Utils/MarkerUtils.h"
-#include "llvm/ADT/SmallVector.h"
-#include "mlir/Dialect/Linalg/Transforms/Hoisting.h"
-#include "mlir/Dialect/Linalg/Transforms/Transforms.h"
-#include "mlir/Dialect/Vector/VectorOps.h"
-#include "mlir/Dialect/Vector/VectorTransforms.h"
-#include "mlir/IR/PatternMatch.h"
-#include "mlir/Pass/Pass.h"
-#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
-
-#define DEBUG_TYPE "iree-linalg-vectorization"
-
-static llvm::cl::list<int64_t> clVectorUnrollShape(
- "iree-codegen-llvm-vector-unroll-shape",
- llvm::cl::desc("Comma-separated list of vector sizes for vector unrolling"),
- llvm::cl::CommaSeparated);
-
-namespace mlir {
-namespace iree_compiler {
-
-static Optional<SmallVector<int64_t, 4>> getShape(mlir::Operation *op) {
- auto unrollOp = dyn_cast<VectorUnrollOpInterface>(op);
- if (!unrollOp) return None;
- SmallVector<int64_t, 4> shape(clVectorUnrollShape.begin(),
- clVectorUnrollShape.end());
- shape.resize(unrollOp.getShapeForUnroll()->size(),
- std::numeric_limits<int64_t>::max());
- return shape;
-}
-
-namespace {
-struct LinalgVectorizationPass
- : public LinalgVectorizationBase<LinalgVectorizationPass> {
- void getDependentDialects(DialectRegistry ®istry) const override {
- registry.insert<vector::VectorDialect, AffineDialect>();
- }
- void runOnOperation() override;
-};
-} // namespace
-
-void LinalgVectorizationPass::runOnOperation() {
- FuncOp funcOp = getOperation();
- MLIRContext *context = &getContext();
- // Apply vectorization patterns.
- {
- OwningRewritePatternList vectorizationPatterns(&getContext());
- linalg::insertVectorizationPatterns<linalg::GenericOp,
- linalg::ContractionOpInterface>(
- vectorizationPatterns, linalg::LinalgVectorizationOptions(),
- linalg::LinalgTransformationFilter(ArrayRef<Identifier>(
- Identifier::get(getWorkgroupMarker(), context))));
- (void)applyPatternsAndFoldGreedily(funcOp,
- std::move(vectorizationPatterns));
-
- LLVM_DEBUG({
- llvm::dbgs() << "--- After Vectorization ---\n";
- funcOp.print(llvm::dbgs(), OpPrintingFlags().useLocalScope());
- llvm::dbgs() << "\n\n";
- });
- }
-
- // TODO: This should be a folding of Add into Contract in core but while they
- // live in different dialects, it is not possible without unnatural
- // dependencies.
- funcOp.walk([&](Operation *op) {
- if (auto contract = canonicalizeContractionAdd(op))
- op->replaceAllUsesWith(contract);
- });
-
- // Apply unrolling patterns.
- {
- OwningRewritePatternList vectorUnrollPatterns(&getContext());
- vectorUnrollPatterns.insert<vector::UnrollVectorPattern>(
- context, vector::UnrollVectorOptions().setNativeShapeFn(getShape));
- (void)applyPatternsAndFoldGreedily(funcOp, std::move(vectorUnrollPatterns));
-
- OwningRewritePatternList canonicalizationPatterns1(&getContext());
- vector::populateVectorToVectorCanonicalizationPatterns(
- canonicalizationPatterns1);
- vector::populateVectorToVectorTransformationPatterns(
- canonicalizationPatterns1);
- (void)applyPatternsAndFoldGreedily(funcOp,
- std::move(canonicalizationPatterns1));
-
- OwningRewritePatternList canonicalizationPatterns2(&getContext());
- vector::populateVectorSlicesLoweringPatterns(canonicalizationPatterns2);
- (void)applyPatternsAndFoldGreedily(funcOp,
- std::move(canonicalizationPatterns2));
-
- LLVM_DEBUG({
- llvm::dbgs() << "--- After Vector Unroll ---\n";
- funcOp.print(llvm::dbgs(), OpPrintingFlags().useLocalScope());
- llvm::dbgs() << "\n\n";
- });
- }
- // Apply hoisting patterns.
- {
- linalg::hoistRedundantVectorTransfersOnTensor(funcOp);
-
- LLVM_DEBUG({
- llvm::dbgs() << "--- After Hoisting ---\n";
- funcOp.print(llvm::dbgs(), OpPrintingFlags().useLocalScope());
- llvm::dbgs() << "\n\n";
- });
- }
-}
-
-std::unique_ptr<OperationPass<FuncOp>> createLinalgVectorizePass() {
- return std::make_unique<LinalgVectorizationPass>();
-}
-
-} // namespace iree_compiler
-} // namespace mlir
diff --git a/iree/compiler/Conversion/LinalgToLLVM/test/linalg_vectorize.mlir b/iree/compiler/Conversion/LinalgToLLVM/test/linalg_vectorize.mlir
deleted file mode 100644
index 57c8c97..0000000
--- a/iree/compiler/Conversion/LinalgToLLVM/test/linalg_vectorize.mlir
+++ /dev/null
@@ -1,46 +0,0 @@
-// RUN: iree-opt %s -iree-codegen-linalg-vectorization-pass -iree-codegen-llvm-vector-unroll-shape=1,1,1 -split-input-file | IreeFileCheck %s
-
-// CHECK-LABEL: func @tensor_dispatch_0
-// CHECK-DAG: %[[C0:.*]] = constant 0 : index
-// CHECK-DAG: %[[C1:.*]] = constant 1 : index
-// CHECK-DAG: %[[C2:.*]] = constant 2 : index
-// CHECK: %[[I0:.*]] = flow.dispatch.tensor.load {{.*}} : !flow.dispatch.tensor<readonly:2x3xf32> -> tensor<2x3xf32>
-// CHECK: %[[I1:.*]] = flow.dispatch.tensor.load {{.*}} : !flow.dispatch.tensor<readonly:3x4xf32> -> tensor<3x1xf32>
-// CHECK: %[[I2:.*]] = flow.dispatch.tensor.load {{.*}} : !flow.dispatch.tensor<readonly:2x4xf32> -> tensor<2x1xf32>
-// CHECK: %[[V0:.*]] = vector.transfer_read %[[I0]][%[[C0]], %[[C0]]], {{.*}} : tensor<2x3xf32>, vector<1x1xf32>
-// CHECK: %[[V1:.*]] = vector.transfer_read %[[I0]][%[[C0]], %[[C1]]], {{.*}} : tensor<2x3xf32>, vector<1x1xf32>
-// CHECK: %[[V2:.*]] = vector.transfer_read %[[I0]][%[[C0]], %[[C2]]], {{.*}} : tensor<2x3xf32>, vector<1x1xf32>
-// CHECK: %[[V3:.*]] = vector.transfer_read %[[I0]][%[[C1]], %[[C0]]], {{.*}} : tensor<2x3xf32>, vector<1x1xf32>
-// CHECK: %[[V4:.*]] = vector.transfer_read %[[I0]][%[[C1]], %[[C1]]], {{.*}} : tensor<2x3xf32>, vector<1x1xf32>
-// CHECK: %[[V5:.*]] = vector.transfer_read %[[I0]][%[[C1]], %[[C2]]], {{.*}} : tensor<2x3xf32>, vector<1x1xf32>
-// CHECK: %[[V6:.*]] = vector.transfer_read %[[I1]][%[[C0]], %[[C0]]], {{.*}} : tensor<3x1xf32>, vector<1x1xf32>
-// CHECK: %[[V7:.*]] = vector.transfer_read %[[I1]][%[[C1]], %[[C0]]], {{.*}} : tensor<3x1xf32>, vector<1x1xf32>
-// CHECK: %[[V8:.*]] = vector.transfer_read %[[I1]][%[[C2]], %[[C0]]], {{.*}} : tensor<3x1xf32>, vector<1x1xf32>
-// CHECK: %[[V9:.*]] = vector.transfer_read %[[I2]][%[[C0]], %[[C0]]], {{.*}} : tensor<2x1xf32>, vector<1x1xf32>
-// CHECK: %[[VA:.*]] = vector.transfer_read %[[I2]][%[[C1]], %[[C0]]], {{.*}} : tensor<2x1xf32>, vector<1x1xf32>
-// CHECK: %[[D0:.*]] = vector.contract {{.*}} %[[V0]], %[[V6]], %[[V9]] : vector<1x1xf32>, vector<1x1xf32> into vector<1x1xf32>
-// CHECK: %[[D1:.*]] = vector.contract {{.*}} %[[V1]], %[[V7]], %[[D0]] : vector<1x1xf32>, vector<1x1xf32> into vector<1x1xf32>
-// CHECK: %[[D2:.*]] = vector.contract {{.*}} %[[V2]], %[[V8]], %[[D1]] : vector<1x1xf32>, vector<1x1xf32> into vector<1x1xf32>
-// CHECK: %[[D3:.*]] = vector.contract {{.*}} %[[V3]], %[[V6]], %[[VA]] : vector<1x1xf32>, vector<1x1xf32> into vector<1x1xf32>
-// CHECK: %[[D4:.*]] = vector.contract {{.*}} %[[V4]], %[[V7]], %[[D3]] : vector<1x1xf32>, vector<1x1xf32> into vector<1x1xf32>
-// CHECK: %[[D5:.*]] = vector.contract {{.*}} %[[V5]], %[[V8]], %[[D4]] : vector<1x1xf32>, vector<1x1xf32> into vector<1x1xf32>
-// CHECK: %[[W0:.*]] = vector.transfer_write %[[D2]], %[[I2]][%[[C0]], %[[C0]]] {in_bounds = [true, true]} : vector<1x1xf32>, tensor<2x1xf32>
-// CHECK: %[[W1:.*]] = vector.transfer_write %[[D5]], %[[W0]][%[[C1]], %[[C0]]] {in_bounds = [true, true]} : vector<1x1xf32>, tensor<2x1xf32>
-// CHECK: flow.dispatch.tensor.store %[[W1]]
-
-func @tensor_dispatch_0() {
- %c0 = constant 0 : index
- %c3 = constant 3 : index
- %c1 = constant 1 : index
- %c2 = constant 1 : index
- %0 = hal.interface.binding.subspan @io::@arg0[%c0] : !flow.dispatch.tensor<readonly:2x3xf32>
- %1 = hal.interface.binding.subspan @io::@arg1[%c0] : !flow.dispatch.tensor<readonly:3x4xf32>
- %2 = hal.interface.binding.subspan @io::@arg2[%c0] : !flow.dispatch.tensor<readonly:2x4xf32>
- %3 = hal.interface.binding.subspan @io::@ret0[%c0] : !flow.dispatch.tensor<writeonly:2x4xf32>
- %4 = flow.dispatch.tensor.load %0, offsets = [%c0, %c0], sizes = [%c2, %c3], strides = [%c1, %c1] : !flow.dispatch.tensor<readonly:2x3xf32> -> tensor<2x3xf32>
- %5 = flow.dispatch.tensor.load %1, offsets = [%c0, %c0], sizes = [%c3, %c1], strides = [%c1, %c1] : !flow.dispatch.tensor<readonly:3x4xf32> -> tensor<3x1xf32>
- %6 = flow.dispatch.tensor.load %2, offsets = [%c0, %c0], sizes = [%c2, %c1], strides = [%c1, %c1] : !flow.dispatch.tensor<readonly:2x4xf32> -> tensor<2x1xf32>
- %7 = linalg.matmul {__internal_linalg_transform__ = "workgroup"} ins(%4, %5 : tensor<2x3xf32>, tensor<3x1xf32>) outs(%6 : tensor<2x1xf32>) -> tensor<2x1xf32>
- flow.dispatch.tensor.store %7, %3, offsets = [%c0, %c0], sizes = [%c2, %c1], strides = [%c1, %c1] : tensor<2x1xf32> -> !flow.dispatch.tensor<writeonly:2x4xf32>
- return
-}
diff --git a/iree/compiler/Conversion/Passes.td b/iree/compiler/Conversion/Passes.td
deleted file mode 100644
index ed849d3..0000000
--- a/iree/compiler/Conversion/Passes.td
+++ /dev/null
@@ -1,258 +0,0 @@
-// Copyright 2021 The IREE Authors
-//
-// Licensed 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
-
-#ifndef IREE_CONVERSION_PASSES
-#define IREE_CONVERSION_PASSES
-
-include "mlir/Pass/PassBase.td"
-
-//------------------------------------------------------------------------------
-// Common/misc passes
-//------------------------------------------------------------------------------
-
-def CleanupBufferAllocView :
- Pass<"iree-codegen-cleanup-buffer-alloc-view", "FuncOp"> {
- let summary = "Performs cleanups over HAL interface/buffer allocation/view operations";
- let constructor = "mlir::iree_compiler::createCleanupBufferAllocViewPass()";
-}
-
-def DemoteF32ToF16 :
- Pass<"iree-convert-f32-to-f16", "ModuleOp"> {
- let summary = "Convert f32 operations and values into equivalent f16 ones.";
- let constructor = "mlir::iree_compiler::createDemoteF32ToF16Pass()";
-}
-
-def FlattenMemRefSubspan :
- Pass<"iree-codegen-flatten-memref-subspan", "ModuleOp"> {
- let summary = "Flatten n-D MemRef subspan ops to 1-D ones and fold byte offsets on subspan ops to the consumer load/store ops";
- let constructor = "mlir::iree_compiler::createFlattenMemRefSubspanPass()";
-}
-
-def FoldTensorExtractOp :
- Pass<"iree-codegen-fold-tensor-extract-op", ""> {
- let summary = "Fold `tensor.extract` operations prior to lowering to LLVM";
- let constructor = "mlir::iree_compiler::createFoldTensorExtractOpPass()";
-}
-
-def ForOpCanonicalization :
- Pass<"iree-codegen-canonicalize-scf-for", "FuncOp"> {
- let summary = "An ad-hoc pass to canonicalize selected loop-carried values and dependencies around scf.for";
- let constructor = "mlir::iree_compiler::createForOpCanonicalizationPass()";
-}
-
-def LinalgBufferize :
- Pass<"iree-codegen-linalg-bufferize", "FuncOp"> {
- let summary = "Convert from to Linalg ops on tensors to buffers";
- let constructor = "mlir::iree_compiler::createLinalgBufferizePass(nullptr)";
-}
-
-def OptimizeVectorTransfer :
- Pass<"iree-codegen-optimize-vector-transfer", "FuncOp"> {
- let summary = "Run optimization transformations on vector transfer operations";
- let constructor = "mlir::iree_compiler::createOptimizeVectorTransferPass()";
-}
-
-// TODO: Consider removing or moving to HAL/Transforms in order to avoid
-// polluting common pass declarations with HAL specific ops.
-def SetNumWorkgroups :
- Pass<"iree-set-num-workgroups", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Set the number of workgroups to use for every entry point function in the dispatch region";
- let constructor = "mlir::iree_compiler::createSetNumWorkgroupsPass()";
-}
-
-//------------------------------------------------------------------------------
-// LinalgToLinalg
-//------------------------------------------------------------------------------
-
-//------------------------------------------------------------------------------
-// LinalgToLLVM
-//------------------------------------------------------------------------------
-
-def LinalgToLLVMWorkgroupsVectorization :
- Pass<"iree-codegen-linalg-to-llvm-workgroups-vectorization-pass", "FuncOp"> {
- let summary = "Tile and vectorize llvm workgroups";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMWorkgroupsVectorizationPass()";
-}
-
-def ConvertToLLVM :
- Pass<"iree-codegen-convert-to-llvm", "ModuleOp"> {
- let summary = "Perform final conversion from Linalg/HAL/Shape/Vector/Standard to LLVMIR dialect";
- let constructor = "mlir::iree_compiler::createConvertToLLVMPass()";
-}
-
-def LinalgVectorization :
- Pass<"iree-codegen-linalg-vectorization-pass", "FuncOp"> {
- let summary = "Vectorize linalg ops";
- let constructor = "mlir::iree_compiler::createLinalgVectorizePass()";
-}
-
-def LowerExecutableTarget :
- Pass<"iree-lower-executable-target-pass", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Perform lowering of executable target using one of the IREE::HAL::DispatchLoweringPassPipeline";
- let constructor = "mlir::iree_compiler::createLowerExecutableTargetPass()";
-}
-
-def LLVMPadLinalgWorkgroupTiles :
- Pass<"iree-codegen-llvm-pad-linalg-workgroup-tiles", "FuncOp"> {
- let summary = "Padding linalg workgroup tiles into an integer multiple of tiling parameters.";
- let constructor = "mlir::iree_compiler::createLLVMPadLinalgWorkgroupTilesPass()";
-}
-
-def LinalgToLLVMPlanConvLoopOrder :
- Pass<"iree-codegen-linalg-to-llvm-plan-conv-loop-order", "FuncOp"> {
- let summary = "Convert linalg.conv to linalg.generic with a CPU-friendly iterator order";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMPlanConvLoopOrderPass()";
-}
-
-def LinalgToLLVMTilePadAndVectorizeWorkgroups :
- Pass<"iree-codegen-linalg-to-llvm-tile-pad-and-vectorize-workgroups", "FuncOp"> {
- let summary = "Tile and pad workgroups tiles";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMTilePadAndVectorizeWorkgroupsPass()";
-}
-
-def LinalgToLLVMUnfuseFMAOps :
- Pass<"iree-codegen-linalg-to-llvm-unfuse-fma-pass", "FuncOp"> {
- let summary = "Convert llvm.fma into unfused mulf and addf ops";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMUnfuseFMAOpsPass()";
-}
-
-//------------------------------------------------------------------------------
-// LinalgToLLVMGPU
-//------------------------------------------------------------------------------
-
-def LinalgToLLVMGPUVectorization :
- Pass<"iree-codegen-llvmgpu-vectorization", "FuncOp"> {
- let summary = "Pass to convert linalg into Vector.";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMGPUVectorizationPass()";
-}
-
-def LinalgToLLVMGPUTileAndDistribute :
- Pass<"iree-codegen-llvmgpu-tile-and-distribute", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Pass to tile and distribute linalg ops within a workgroup.";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMGPUTileAndDistributeToThreads()";
-}
-
-def LinalgToLLVMGPURemoveSingleIterationLoop :
- Pass<"iree-llvmgpu-remove-single-iteration-loop", "FuncOp"> {
- let summary = "Remove distributed loop with single iteration.";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMGPURemoveSingleIterationLoopPass()";
-}
-
-// TODO: Bring the argument in line with the names used elsewhere.
-def LinalgToLLVMGPULowerExecutableTarget :
- Pass<"iree-lower-executable-target-gpu-pass", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Perform lowering of executable target using one of the IREE::HAL::DispatchLoweringPassPipeline";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMGPULowerExecutableTargetPass()";
-}
-
-// TODO: Bring the argument in line with the names used elsewhere.
-def LinalgToLLVMGPUConvertToROCDL :
- Pass<"iree-codegen-convert-to-rocdl", "ModuleOp"> {
- let summary = "Perform final conversion from builtin/GPU/HAL/standard dialect to LLVM "
- "and ROCDL dialects";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMGPUConvertToROCDLPass()";
-}
-
-// TODO: Bring the argument in line with the names used elsewhere.
-def LinalgToLLVMGPUConvertToNVVM :
- Pass<"iree-codegen-convert-to-nvvm", "ModuleOp"> {
- let summary = "Perform final conversion from builtin/GPU/HAL/standard dialect to LLVM "
- "and NVVM dialects";
- let constructor = "mlir::iree_compiler::createLinalgToLLVMGPUConvertToNVVMPass()";
-}
-
-//------------------------------------------------------------------------------
-// LinalgToSPIRV
-//------------------------------------------------------------------------------
-
-def LinalgToSPIRVConcretizeTileAmongWorkgroups :
- Pass<"iree-spirv-concretize-tile-among-workgroups", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Replace hal.interface.workgroup.* ops with constant values from chosen tiling and distribution scheme";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVConcretizeTileAmongWorkgroupsPass(mlir::iree_compiler::SPIRVCodegenOptions::getFromCLOptions())";
-}
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToSPIRVConvertToGPU :
- Pass<"iree-codegen-convert-to-gpu", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Map tiled linalg and loop ops to GPU";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVConvertToGPUPass()";
-}
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToSPIRVConvertToSPIRV :
- Pass<"iree-codegen-convert-to-spirv", "ModuleOp"> {
- let summary = "Perform final conversion from builtin/GPU/HAL/standard dialect to SPIR-V dialect";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVConvertToSPIRVPass()";
-}
-
-// TODO: Rename argument to be fully qualified.
-// TODO: Does not appear used?
-def LinalgToSPIRVFoldProcessorIDUses :
- Pass<"iree-codegen-fold-gpu-procid-uses", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Fold GPU processor ID uses where possible";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVFoldProcessorIDUsesPass()";
-}
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToSPIRVTileAndVectorizeOneWorkgroup :
- Pass<"iree-spirv-tile-and-vectorize-in-one-workgroup", "mlir::iree_compiler::IREE::HAL::ExecutableTargetOp"> {
- let summary = "Tile and vectorize Linalg operations on buffers in one workgroup";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVTileAndVectorizeOneWorkgroupPass(mlir::iree_compiler::SPIRVCodegenOptions::getFromCLOptions())";
-}
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToSPIRVVectorizeMemRefLoadStore :
- Pass<"iree-spirv-vectorize-memref-load-store", "ModuleOp"> {
- let summary = "Vectorize interface memrefs and their load/store for better memory access";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVVectorizeMemRefLoadStore()";
-}
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToSPIRVVectorToCooperativeMatrix :
- Pass<"iree-spirv-vector-to-cooperative-matrix", "FuncOp"> {
- let summary = "Convert vector load/store/arithmetic ops to cooperative matrix ops when possible";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVVectorToCooperativeMatrixPass()";
-}
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToSPIRVConvertVectorToGPU :
- Pass<"iree-codegen-vector-to-gpu", "FuncOp"> {
- let summary = "Convert vector dialect to gpu subgroup level GPU instructions";
- let constructor = "mlir::iree_compiler::createLinalgToSPIRVConvertVectorToGPU()";
-}
-
-//------------------------------------------------------------------------------
-// LinalgToVector Passes
-//------------------------------------------------------------------------------
-
-// TODO: Rename argument to be fully qualified.
-def LinalgToVectorVectorizeConv :
- Pass<"iree-codegen-vectorize-linalg-conv", "FuncOp"> {
- let summary = "Vectorize a very specific form of linalg.conv";
- let constructor = "mlir::iree_compiler::createLinalgToVectorVectorizeConvPass();";
-}
-
-//------------------------------------------------------------------------------
-// Target specific VectorToLLVM Passes
-//------------------------------------------------------------------------------
-
-def VectorToAArch64InlineAsm :
- Pass<"iree-codegen-vector-to-aarch64-inline-asm", "FuncOp"> {
- let summary = "Convert vector operations to aarch64 inline asm LLVMIR dialect";
- let constructor = "mlir::iree_compiler::createVectorToAArch64InlineAssemblyPass()";
-}
-
-//------------------------------------------------------------------------------
-// Test passes
-//------------------------------------------------------------------------------
-
-def TestLinalgToLLVMGPUScalarizeMathOp :
- Pass<"iree-test-llvmgpu-scalarize-math-op", "FuncOp"> {
- let summary = "Test pass for scalarization patterns.";
- let constructor = "mlir::iree_compiler::createTestLinalgToLLVMGPUScalarizeMathOpPass()";
-}
-
-#endif // IREE_DIALECT_FLOW_PASSES
diff --git a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
index ff98366..5aaad87 100644
--- a/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
+++ b/iree/compiler/Dialect/Flow/Transforms/DispatchLinalgOnTensors.cpp
@@ -122,7 +122,7 @@
/// Returns the number of consecutive outer loops that are "parallel". This is a
/// copy of the function from
-/// iree/compiler/Conversion/CodegenUtils/FunctionUtils.h that is duplicated
+/// iree/compiler/Codegen/CodegenUtils/FunctionUtils.h that is duplicated
/// here to avoid adding an build dependency.
static size_t getNumOuterParallelLoops(linalg::LinalgOp op) {
return op.iterator_types()
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/BUILD b/iree/compiler/Dialect/HAL/Target/CUDA/BUILD
index 5dfe3fd..282aea8 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/BUILD
@@ -40,8 +40,8 @@
deps = [
":cuda_libdevice",
"//iree/base/internal:flatcc",
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/LinalgToLLVMGPU",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/LLVMGPU",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Utils",
"//iree/schemas:cuda_executable_def_c_fbs",
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/CUDA/CMakeLists.txt
index bf839cd..7ebad2b 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/CMakeLists.txt
@@ -1,3 +1,11 @@
+# Copyright 2021 The IREE Authors
+#
+# Licensed 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
+
+# Doesn't use bazel_to_cmake because of various special logic throughout.
+
if(NOT "${IREE_TARGET_BACKEND_CUDA}")
return()
endif()
@@ -44,7 +52,7 @@
MLIRSupport
MLIRTargetLLVMIRExport
iree::base::internal::flatcc
- iree::compiler::Conversion::LinalgToLLVMGPU
+ iree::compiler::Codegen::LLVMGPU
iree::compiler::Dialect::HAL::Target
iree::compiler::Utils
iree::schemas::cuda_executable_def_c_fbs
diff --git a/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp b/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
index f3eb89b..0ed02fa 100644
--- a/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.cpp
@@ -6,7 +6,7 @@
#include "iree/compiler/Dialect/HAL/Target/CUDA/CUDATarget.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/CUDA/libdevice.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
index 022f8e0..01c1c13 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/BUILD
@@ -36,10 +36,10 @@
":LinkerTool",
":StaticLibraryGenerator",
"//iree/base/internal:flatcc",
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/LinalgToLLVM",
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/LLVMCPU",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Utils",
"//iree/schemas:dylib_executable_def_c_fbs",
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
index 049b5c8..f12301b 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/CMakeLists.txt
@@ -45,10 +45,10 @@
MLIRLLVMToLLVMIRTranslation
MLIRTargetLLVMIRExport
iree::base::internal::flatcc
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::LinalgToLLVM
- iree::compiler::Conversion::PassHeaders
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::LLVMCPU
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::Target
iree::compiler::Utils
iree::schemas::dylib_executable_def_c_fbs
diff --git a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
index 9ab52f0..650875a 100644
--- a/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/LLVM/LLVMAOTTarget.cpp
@@ -8,7 +8,7 @@
#include <cstdlib>
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/LLVM/LLVMIRPasses.h"
#include "iree/compiler/Dialect/HAL/Target/LLVM/LibraryBuilder.h"
#include "iree/compiler/Dialect/HAL/Target/LLVM/LinkerTool.h"
@@ -83,11 +83,9 @@
<< options_.targetTriple << "'";
return;
}
-
- passManager.addPass(createLowerExecutableTargetPass());
-
+ passManager.addPass(createLLVMCPULowerExecutableTargetPass());
// Set target specific options.
- LLVMTransformPassPipelineOptions codeGenOptions;
+ LLVMCPUCodegenPassPipelineOptions codeGenOptions;
codeGenOptions.targetTriple = options_.targetTriple;
codeGenOptions.targetDataLayout =
targetMachine->createDataLayout().getStringRepresentation();
@@ -98,7 +96,7 @@
codeGenOptions.unfuseFMAOps = true;
}
- buildLLVMTransformPassPipeline(passManager, codeGenOptions);
+ buildLLVMCPUCodegenPassPipeline(passManager, codeGenOptions);
}
LogicalResult linkExecutables(mlir::ModuleOp moduleOp) override {
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD
index c97eb0c..759dcb7 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/BUILD
@@ -26,8 +26,8 @@
hdrs = ["MetalSPIRVTarget.h"],
deps = [
":SPIRVToMSL",
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Dialect/HAL/Target/SPIRVCommon",
"//iree/compiler/Utils",
diff --git a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt
index 53135d4..d131f39 100644
--- a/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/MetalSPIRV/CMakeLists.txt
@@ -29,8 +29,8 @@
MLIRSPIRV
MLIRSPIRVSerialization
MLIRVector
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::PassHeaders
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::PassHeaders
iree::compiler::Dialect::HAL::Target
iree::compiler::Dialect::HAL::Target::SPIRVCommon
iree::compiler::Utils
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/BUILD b/iree/compiler/Dialect/HAL/Target/ROCM/BUILD
index 9e25df4..244ef51 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/BUILD
@@ -31,8 +31,8 @@
],
deps = [
"//iree/base/internal:flatcc",
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/LinalgToLLVMGPU",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/LLVMGPU",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Utils",
"//iree/schemas:rocm_executable_def_c_fbs",
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt
index 6bb6fd6..27fd906 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/CMakeLists.txt
@@ -38,8 +38,8 @@
MLIRSupport
MLIRTargetLLVMIRExport
iree::base::internal::flatcc
- iree::compiler::Conversion::LinalgToLLVMGPU
- iree::compiler::Conversion::PassHeaders
+ iree::compiler::Codegen::LLVMGPU
+ iree::compiler::Codegen::PassHeaders
iree::compiler::Dialect::HAL::Target
iree::compiler::Utils
iree::schemas::rocm_executable_def_c_fbs
diff --git a/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp b/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
index 9473149..95062fd 100644
--- a/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.cpp
@@ -6,7 +6,7 @@
#include "iree/compiler/Dialect/HAL/Target/ROCM/ROCMTarget.h"
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
#include "iree/schemas/rocm_executable_def_builder.h"
diff --git a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/BUILD b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/BUILD
index 897e51c..713bfbf 100644
--- a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/BUILD
@@ -30,10 +30,10 @@
"SPIRVTarget.h",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/LinalgToSPIRV",
- "//iree/compiler/Conversion/Utils",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/SPIRV",
+ "//iree/compiler/Codegen/Utils",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Dialect/IREE/IR",
diff --git a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/CMakeLists.txt
index 58a633c..f4640ed 100644
--- a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/CMakeLists.txt
@@ -29,10 +29,10 @@
MLIRSPIRVConversion
MLIRSPIRVSerialization
MLIRSupport
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::LinalgToSPIRV
- iree::compiler::Conversion::PassHeaders
- iree::compiler::Conversion::Utils
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::SPIRV
+ iree::compiler::Codegen::Utils
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::Target
iree::compiler::Dialect::IREE::IR
diff --git a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.cpp b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.cpp
index 3c03e96..ad39544 100644
--- a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.cpp
+++ b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.cpp
@@ -40,7 +40,7 @@
void SPIRVTargetBackend::buildTranslationPassPipeline(
OpPassManager &passManager) {
- buildSPIRVTransformPassPipeline(passManager, spvCodeGenOptions_);
+ buildSPIRVCodegenPassPipeline(passManager, spvCodeGenOptions_);
}
} // namespace HAL
diff --git a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.h b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.h
index 7fe3655..a7896ae 100644
--- a/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.h
+++ b/iree/compiler/Dialect/HAL/Target/SPIRVCommon/SPIRVTarget.h
@@ -9,7 +9,7 @@
#include <string>
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/TargetBackend.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVAttributes.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/BUILD b/iree/compiler/Dialect/HAL/Target/VMVX/BUILD
index 6f83d09..fdf19ce 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/BUILD
@@ -29,7 +29,7 @@
"VMVXTarget.h",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
+ "//iree/compiler/Codegen:PassHeaders",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Dialect/Modules/VMVX/IR:VMVXDialect",
diff --git a/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt
index ccc74da..f21e2c0 100644
--- a/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/VMVX/CMakeLists.txt
@@ -26,7 +26,7 @@
MLIRIR
MLIRPass
MLIRSupport
- iree::compiler::Conversion::PassHeaders
+ iree::compiler::Codegen::PassHeaders
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::Target
iree::compiler::Dialect::Modules::VMVX::IR::VMVXDialect
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
index fff2be4..e503870 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/BUILD
@@ -30,9 +30,9 @@
],
deps = [
"//iree/base/internal:flatcc",
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/LinalgToSPIRV",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/SPIRV",
"//iree/compiler/Dialect/Flow/IR",
"//iree/compiler/Dialect/HAL/Target",
"//iree/compiler/Dialect/HAL/Target/SPIRVCommon",
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt
index a9c2ac9..17e41eb 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/CMakeLists.txt
@@ -33,9 +33,9 @@
MLIRSupport
MLIRVector
iree::base::internal::flatcc
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::LinalgToSPIRV
- iree::compiler::Conversion::PassHeaders
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::PassHeaders
+ iree::compiler::Codegen::SPIRV
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::Target
iree::compiler::Dialect::HAL::Target::SPIRVCommon
diff --git a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h
index 719788d..677a549 100644
--- a/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h
+++ b/iree/compiler/Dialect/HAL/Target/VulkanSPIRV/VulkanSPIRVTarget.h
@@ -10,7 +10,7 @@
#include <functional>
#include <string>
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
namespace mlir {
namespace iree_compiler {
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD b/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD
index e76f5b7..2acdc18 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD
+++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/BUILD
@@ -20,9 +20,9 @@
"Passes.h",
],
deps = [
- "//iree/compiler/Conversion:PassHeaders",
- "//iree/compiler/Conversion/Common",
- "//iree/compiler/Conversion/LinalgToLLVM",
+ "//iree/compiler/Codegen:PassHeaders",
+ "//iree/compiler/Codegen/Common",
+ "//iree/compiler/Codegen/LLVMCPU",
"//iree/compiler/Dialect/HAL/IR:HALDialect",
"//iree/compiler/Dialect/HAL/Transforms",
"//iree/compiler/Dialect/IREE/IR",
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt b/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt
index 69243fa..d8af381 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt
+++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/CMakeLists.txt
@@ -41,9 +41,9 @@
MLIRTransforms
MLIRVector
MLIRVectorToSCF
- iree::compiler::Conversion::Common
- iree::compiler::Conversion::LinalgToLLVM
- iree::compiler::Conversion::PassHeaders
+ iree::compiler::Codegen::Common
+ iree::compiler::Codegen::LLVMCPU
+ iree::compiler::Codegen::PassHeaders
iree::compiler::Dialect::HAL::IR::HALDialect
iree::compiler::Dialect::HAL::Transforms
iree::compiler::Dialect::IREE::IR
diff --git a/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp b/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp
index 707d868..962ee99 100644
--- a/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp
+++ b/iree/compiler/Dialect/Modules/VMVX/Transforms/Passes.cpp
@@ -8,7 +8,7 @@
#include <memory>
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Transforms/Passes.h"
#include "iree/compiler/Dialect/Shape/Transforms/Passes.h"
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
@@ -30,7 +30,7 @@
// For now lower using the default CPU pass-pipeline which doesn't
// vectorize. When VMVX can lower vector operations, this can be relaxed.
passManager.addPass(
- createLowerExecutableTargetPass(/*lowerToVectors=*/false));
+ createLLVMCPULowerExecutableTargetPass(/*lowerToVectors=*/false));
OpPassManager &nestedModulePM = passManager.nest<ModuleOp>();
diff --git a/iree/tools/BUILD b/iree/tools/BUILD
index 5ee6e9b..f3a17bc 100644
--- a/iree/tools/BUILD
+++ b/iree/tools/BUILD
@@ -149,7 +149,7 @@
":init_iree_passes_and_dialects",
":init_mlir_passes_and_dialects",
":init_xla_dialects",
- "//iree/compiler/Conversion",
+ "//iree/compiler/Codegen",
"//iree/compiler/Dialect/HAL/Conversion:Passes",
],
)
@@ -319,7 +319,7 @@
":init_targets",
":init_translations",
":init_xla_dialects",
- "//iree/compiler/Conversion",
+ "//iree/compiler/Codegen",
"//iree/compiler/Dialect/VM/Target:init_targets",
"//iree/compiler/Dialect/VM/Target/Bytecode",
"//iree/compiler/Translation:IREEVM",
diff --git a/iree/tools/CMakeLists.txt b/iree/tools/CMakeLists.txt
index 34f6cee..b4f992d 100644
--- a/iree/tools/CMakeLists.txt
+++ b/iree/tools/CMakeLists.txt
@@ -225,7 +225,7 @@
::init_iree_passes_and_dialects
::init_mlir_passes_and_dialects
::init_xla_dialects
- iree::compiler::Conversion::Conversion
+ iree::compiler::Codegen::Codegen
iree::compiler::Dialect::HAL::Conversion::Passes
INCLUDES
"${IREE_EMITC_INCLUDES}"
@@ -311,7 +311,7 @@
MLIRSupport
MLIRTargetLLVMIRExport
MLIRTranslation
- iree::compiler::Conversion::Conversion
+ iree::compiler::Codegen::Codegen
iree::compiler::Dialect::VM::Target::Bytecode
iree::compiler::Dialect::VM::Target::init_targets
iree::compiler::Translation::IREEVM
diff --git a/iree/tools/init_passes.h b/iree/tools/init_passes.h
index 1f536ba..e675dfd 100644
--- a/iree/tools/init_passes.h
+++ b/iree/tools/init_passes.h
@@ -11,7 +11,7 @@
#include <cstdlib>
-#include "iree/compiler/Conversion/Passes.h"
+#include "iree/compiler/Codegen/Passes.h"
#include "iree/compiler/Dialect/HAL/Conversion/Passes.h"
#include "iree/tools/init_iree_passes.h"
#include "iree/tools/init_mlir_passes.h"
@@ -22,7 +22,7 @@
// Registers IREE core passes and other important passes to the global registry.
inline void registerAllPasses() {
registerAllIreePasses();
- registerConversionPasses();
+ registerCodegenPasses();
registerMlirPasses();
registerHALConversionPasses();
}
