[Mlir-commits] [mlir] d7ef488 - [mlir][gpu] Move GPU headers into IR/ and Transforms/
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Thu Jun 9 15:49:08 PDT 2022
Author: Mogball
Date: 2022-06-09T22:49:03Z
New Revision: d7ef488bb6914500050dfccf95c793148457e54b
URL: https://github.com/llvm/llvm-project/commit/d7ef488bb6914500050dfccf95c793148457e54b
DIFF: https://github.com/llvm/llvm-project/commit/d7ef488bb6914500050dfccf95c793148457e54b.diff
LOG: [mlir][gpu] Move GPU headers into IR/ and Transforms/
Depends on D127350
Reviewed By: rriddle
Differential Revision: https://reviews.llvm.org/D127352
Added:
mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
mlir/include/mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td
mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt
mlir/include/mlir/Dialect/GPU/Transforms/MemoryPromotion.h
mlir/include/mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h
mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
mlir/include/mlir/Dialect/GPU/Transforms/Utils.h
Modified:
mlir/include/mlir-c/Dialect/GPU.h
mlir/include/mlir/Dialect/GPU/CMakeLists.txt
mlir/include/mlir/InitAllDialects.h
mlir/include/mlir/InitAllPasses.h
mlir/lib/CAPI/Dialect/GPU.cpp
mlir/lib/CAPI/Dialect/GPUPasses.cpp
mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h
mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp
mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp
mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp
mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp
mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
mlir/lib/Dialect/GPU/Transforms/MemoryPromotion.cpp
mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
mlir/lib/Dialect/GPU/Transforms/PassDetail.h
mlir/lib/Dialect/GPU/Transforms/SerializeToBlob.cpp
mlir/lib/Dialect/GPU/Transforms/SerializeToCubin.cpp
mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToCubin.cpp
mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToHsaco.cpp
mlir/test/lib/Dialect/GPU/TestGpuMemoryPromotion.cpp
mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
mlir/test/lib/Dialect/Linalg/TestLinalgCodegenStrategy.cpp
mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp
mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp
mlir/test/lib/Dialect/Vector/TestVectorTransforms.cpp
mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
Removed:
mlir/include/mlir/Dialect/GPU/GPUBase.td
mlir/include/mlir/Dialect/GPU/GPUDialect.h
mlir/include/mlir/Dialect/GPU/GPUOps.td
mlir/include/mlir/Dialect/GPU/MemoryPromotion.h
mlir/include/mlir/Dialect/GPU/ParallelLoopMapper.h
mlir/include/mlir/Dialect/GPU/ParallelLoopMapperAttr.td
mlir/include/mlir/Dialect/GPU/Passes.h
mlir/include/mlir/Dialect/GPU/Passes.td
mlir/include/mlir/Dialect/GPU/Utils.h
################################################################################
diff --git a/mlir/include/mlir-c/Dialect/GPU.h b/mlir/include/mlir-c/Dialect/GPU.h
index e4797a7ee4e8..cf4e899c5535 100644
--- a/mlir/include/mlir-c/Dialect/GPU.h
+++ b/mlir/include/mlir-c/Dialect/GPU.h
@@ -23,6 +23,6 @@ MLIR_DECLARE_CAPI_DIALECT_REGISTRATION(GPU, gpu);
}
#endif
-#include "mlir/Dialect/GPU/Passes.capi.h.inc"
+#include "mlir/Dialect/GPU/Transforms/Passes.capi.h.inc"
#endif // MLIR_C_DIALECT_GPU_H
diff --git a/mlir/include/mlir/Dialect/GPU/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/CMakeLists.txt
index 59559ba4c1af..9f57627c321f 100644
--- a/mlir/include/mlir/Dialect/GPU/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/GPU/CMakeLists.txt
@@ -1,35 +1,2 @@
-add_mlir_dialect(GPUOps gpu)
-add_mlir_doc(GPUOps GPUOps Dialects/ -gen-op-doc)
-
-set(LLVM_TARGET_DEFINITIONS GPUBase.td)
-mlir_tablegen(GPUOpInterfaces.h.inc -gen-op-interface-decls)
-mlir_tablegen(GPUOpInterfaces.cpp.inc -gen-op-interface-defs)
-add_public_tablegen_target(MLIRGPUOpInterfacesIncGen)
-
-set(LLVM_TARGET_DEFINITIONS ParallelLoopMapperAttr.td)
-mlir_tablegen(ParallelLoopMapperAttr.h.inc -gen-struct-attr-decls)
-mlir_tablegen(ParallelLoopMapperAttr.cpp.inc -gen-struct-attr-defs)
-add_public_tablegen_target(MLIRParallelLoopMapperAttrGen)
-
-set(LLVM_TARGET_DEFINITIONS ParallelLoopMapperAttr.td)
-mlir_tablegen(ParallelLoopMapperEnums.h.inc -gen-enum-decls)
-mlir_tablegen(ParallelLoopMapperEnums.cpp.inc -gen-enum-defs)
-add_public_tablegen_target(MLIRParallelLoopMapperEnumsGen)
-
-set(LLVM_TARGET_DEFINITIONS Passes.td)
-mlir_tablegen(Passes.h.inc -gen-pass-decls -name GPU)
-mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix GPU)
-mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix GPU)
-add_public_tablegen_target(MLIRGPUPassIncGen)
-
-set(LLVM_TARGET_DEFINITIONS GPUOps.td)
-mlir_tablegen(GPUOpsEnums.h.inc -gen-enum-decls)
-mlir_tablegen(GPUOpsEnums.cpp.inc -gen-enum-defs)
-add_public_tablegen_target(MLIRGPUOpsEnumsGen)
-
-set(LLVM_TARGET_DEFINITIONS GPUOps.td)
-mlir_tablegen(GPUOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gpu)
-mlir_tablegen(GPUOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gpu)
-add_public_tablegen_target(MLIRGPUOpsAttributesIncGen)
-
-add_mlir_doc(Passes GPUPasses ./ -gen-pass-doc)
+add_subdirectory(IR)
+add_subdirectory(Transforms)
diff --git a/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
new file mode 100644
index 000000000000..d3ddd4bba354
--- /dev/null
+++ b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
@@ -0,0 +1,27 @@
+add_mlir_dialect(GPUOps gpu)
+add_mlir_doc(GPUOps GPUOps Dialects/ -gen-op-doc)
+
+set(LLVM_TARGET_DEFINITIONS GPUBase.td)
+mlir_tablegen(GPUOpInterfaces.h.inc -gen-op-interface-decls)
+mlir_tablegen(GPUOpInterfaces.cpp.inc -gen-op-interface-defs)
+add_public_tablegen_target(MLIRGPUOpInterfacesIncGen)
+
+set(LLVM_TARGET_DEFINITIONS ParallelLoopMapperAttr.td)
+mlir_tablegen(ParallelLoopMapperAttr.h.inc -gen-struct-attr-decls)
+mlir_tablegen(ParallelLoopMapperAttr.cpp.inc -gen-struct-attr-defs)
+add_public_tablegen_target(MLIRParallelLoopMapperAttrGen)
+
+set(LLVM_TARGET_DEFINITIONS ParallelLoopMapperAttr.td)
+mlir_tablegen(ParallelLoopMapperEnums.h.inc -gen-enum-decls)
+mlir_tablegen(ParallelLoopMapperEnums.cpp.inc -gen-enum-defs)
+add_public_tablegen_target(MLIRParallelLoopMapperEnumsGen)
+
+set(LLVM_TARGET_DEFINITIONS GPUOps.td)
+mlir_tablegen(GPUOpsEnums.h.inc -gen-enum-decls)
+mlir_tablegen(GPUOpsEnums.cpp.inc -gen-enum-defs)
+add_public_tablegen_target(MLIRGPUOpsEnumsGen)
+
+set(LLVM_TARGET_DEFINITIONS GPUOps.td)
+mlir_tablegen(GPUOpsAttributes.h.inc -gen-attrdef-decls -attrdefs-dialect=gpu)
+mlir_tablegen(GPUOpsAttributes.cpp.inc -gen-attrdef-defs -attrdefs-dialect=gpu)
+add_public_tablegen_target(MLIRGPUOpsAttributesIncGen)
diff --git a/mlir/include/mlir/Dialect/GPU/GPUBase.td b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
similarity index 100%
rename from mlir/include/mlir/Dialect/GPU/GPUBase.td
rename to mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
diff --git a/mlir/include/mlir/Dialect/GPU/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
similarity index 94%
rename from mlir/include/mlir/Dialect/GPU/GPUDialect.h
rename to mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
index d8dc67a21e2c..4637b5f2d3be 100644
--- a/mlir/include/mlir/Dialect/GPU/GPUDialect.h
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h
@@ -11,8 +11,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_GPU_GPUDIALECT_H
-#define MLIR_DIALECT_GPU_GPUDIALECT_H
+#ifndef MLIR_DIALECT_GPU_IR_GPUDIALECT_H
+#define MLIR_DIALECT_GPU_IR_GPUDIALECT_H
#include "mlir/Dialect/DLTI/Traits.h"
#include "mlir/IR/Builders.h"
@@ -164,16 +164,16 @@ void addAsyncDependency(Operation *op, Value token);
} // namespace gpu
} // namespace mlir
-#include "mlir/Dialect/GPU/GPUOpsEnums.h.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsEnums.h.inc"
-#include "mlir/Dialect/GPU/GPUOpsDialect.h.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsDialect.h.inc"
-#include "mlir/Dialect/GPU/GPUOpInterfaces.h.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.h.inc"
#define GET_ATTRDEF_CLASSES
-#include "mlir/Dialect/GPU/GPUOpsAttributes.h.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.h.inc"
#define GET_OP_CLASSES
-#include "mlir/Dialect/GPU/GPUOps.h.inc"
+#include "mlir/Dialect/GPU/IR/GPUOps.h.inc"
-#endif // MLIR_DIALECT_GPU_GPUDIALECT_H
+#endif // MLIR_DIALECT_GPU_IR_GPUDIALECT_H
diff --git a/mlir/include/mlir/Dialect/GPU/GPUOps.td b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
similarity index 99%
rename from mlir/include/mlir/Dialect/GPU/GPUOps.td
rename to mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
index e1e818fe33ce..064e646446cb 100644
--- a/mlir/include/mlir/Dialect/GPU/GPUOps.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
@@ -14,8 +14,8 @@
#define GPU_OPS
include "mlir/Dialect/DLTI/DLTIBase.td"
-include "mlir/Dialect/GPU/GPUBase.td"
-include "mlir/Dialect/GPU/ParallelLoopMapperAttr.td"
+include "mlir/Dialect/GPU/IR/GPUBase.td"
+include "mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td"
include "mlir/IR/EnumAttr.td"
include "mlir/IR/FunctionInterfaces.td"
include "mlir/IR/SymbolInterfaces.td"
diff --git a/mlir/include/mlir/Dialect/GPU/ParallelLoopMapperAttr.td b/mlir/include/mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td
similarity index 98%
rename from mlir/include/mlir/Dialect/GPU/ParallelLoopMapperAttr.td
rename to mlir/include/mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td
index 9f16365ca9b8..0a15095a7f14 100644
--- a/mlir/include/mlir/Dialect/GPU/ParallelLoopMapperAttr.td
+++ b/mlir/include/mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td
@@ -14,7 +14,7 @@
#ifndef PARALLEL_LOOP_MAPPER_ATTR
#define PARALLEL_LOOP_MAPPER_ATTR
-include "mlir/Dialect/GPU/GPUBase.td"
+include "mlir/Dialect/GPU/IR/GPUBase.td"
include "mlir/IR/EnumAttr.td"
def BlockX : I64EnumAttrCase<"BlockX", 0, "block_x">;
diff --git a/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt
new file mode 100644
index 000000000000..60daed4ee97e
--- /dev/null
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt
@@ -0,0 +1,7 @@
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name GPU)
+mlir_tablegen(Passes.capi.h.inc -gen-pass-capi-header --prefix GPU)
+mlir_tablegen(Passes.capi.cpp.inc -gen-pass-capi-impl --prefix GPU)
+add_public_tablegen_target(MLIRGPUPassIncGen)
+
+add_mlir_doc(Passes GPUPasses ./ -gen-pass-doc)
diff --git a/mlir/include/mlir/Dialect/GPU/MemoryPromotion.h b/mlir/include/mlir/Dialect/GPU/Transforms/MemoryPromotion.h
similarity index 84%
rename from mlir/include/mlir/Dialect/GPU/MemoryPromotion.h
rename to mlir/include/mlir/Dialect/GPU/Transforms/MemoryPromotion.h
index b26011480f7d..bf26286058eb 100644
--- a/mlir/include/mlir/Dialect/GPU/MemoryPromotion.h
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/MemoryPromotion.h
@@ -11,8 +11,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_GPU_MEMORYPROMOTION_H
-#define MLIR_DIALECT_GPU_MEMORYPROMOTION_H
+#ifndef MLIR_DIALECT_GPU_TRANSFORMS_MEMORYPROMOTION_H
+#define MLIR_DIALECT_GPU_TRANSFORMS_MEMORYPROMOTION_H
namespace mlir {
@@ -26,4 +26,4 @@ void promoteToWorkgroupMemory(gpu::GPUFuncOp op, unsigned arg);
} // namespace mlir
-#endif // MLIR_DIALECT_GPU_MEMORYPROMOTION_H
+#endif // MLIR_DIALECT_GPU_TRANSFORMS_MEMORYPROMOTION_H
diff --git a/mlir/include/mlir/Dialect/GPU/ParallelLoopMapper.h b/mlir/include/mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h
similarity index 86%
rename from mlir/include/mlir/Dialect/GPU/ParallelLoopMapper.h
rename to mlir/include/mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h
index 40798ea587ca..74b78f96649e 100644
--- a/mlir/include/mlir/Dialect/GPU/ParallelLoopMapper.h
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h
@@ -11,10 +11,10 @@
//
//===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_GPU_PARALLELLOOPMAPPER_H
-#define MLIR_DIALECT_GPU_PARALLELLOOPMAPPER_H
+#ifndef MLIR_DIALECT_GPU_TRANSFORMS_PARALLELLOOPMAPPER_H
+#define MLIR_DIALECT_GPU_TRANSFORMS_PARALLELLOOPMAPPER_H
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Support/LLVM.h"
#include "llvm/ADT/StringRef.h"
@@ -46,4 +46,4 @@ LogicalResult setMappingAttr(scf::ParallelOp ploopOp,
ArrayRef<ParallelLoopDimMappingAttr> mapping);
} // namespace gpu
} // namespace mlir
-#endif // MLIR_DIALECT_GPU_PARALLELLOOPMAPPER_H
+#endif // MLIR_DIALECT_GPU_TRANSFORMS_PARALLELLOOPMAPPER_H
diff --git a/mlir/include/mlir/Dialect/GPU/Passes.h b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
similarity index 95%
rename from mlir/include/mlir/Dialect/GPU/Passes.h
rename to mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
index 53f3f84efbaa..0ef83bc7a6ba 100644
--- a/mlir/include/mlir/Dialect/GPU/Passes.h
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.h
@@ -10,10 +10,10 @@
//
//===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_GPU_PASSES_H_
-#define MLIR_DIALECT_GPU_PASSES_H_
+#ifndef MLIR_DIALECT_GPU_TRANSFORMS_PASSES_H_
+#define MLIR_DIALECT_GPU_TRANSFORMS_PASSES_H_
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Pass/Pass.h"
namespace llvm {
@@ -126,8 +126,8 @@ std::unique_ptr<Pass> createGpuSerializeToHsacoPass(StringRef triple,
/// Generate the code for registering passes.
#define GEN_PASS_REGISTRATION
-#include "mlir/Dialect/GPU/Passes.h.inc"
+#include "mlir/Dialect/GPU/Transforms/Passes.h.inc"
} // namespace mlir
-#endif // MLIR_DIALECT_GPU_PASSES_H_
+#endif // MLIR_DIALECT_GPU_TRANSFORMS_PASSES_H_
diff --git a/mlir/include/mlir/Dialect/GPU/Passes.td b/mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
similarity index 100%
rename from mlir/include/mlir/Dialect/GPU/Passes.td
rename to mlir/include/mlir/Dialect/GPU/Transforms/Passes.td
diff --git a/mlir/include/mlir/Dialect/GPU/Utils.h b/mlir/include/mlir/Dialect/GPU/Transforms/Utils.h
similarity index 92%
rename from mlir/include/mlir/Dialect/GPU/Utils.h
rename to mlir/include/mlir/Dialect/GPU/Transforms/Utils.h
index aa867a58382a..e72d2bc41c52 100644
--- a/mlir/include/mlir/Dialect/GPU/Utils.h
+++ b/mlir/include/mlir/Dialect/GPU/Transforms/Utils.h
@@ -10,8 +10,8 @@
//
//===----------------------------------------------------------------------===//
-#ifndef MLIR_DIALECT_GPU_UTILS_H_
-#define MLIR_DIALECT_GPU_UTILS_H_
+#ifndef MLIR_DIALECT_GPU_TRANSFORMS_UTILS_H_
+#define MLIR_DIALECT_GPU_TRANSFORMS_UTILS_H_
#include "mlir/Support/LLVM.h"
@@ -44,4 +44,4 @@ LogicalResult sinkOperationsIntoLaunchOp(
llvm::function_ref<bool(Operation *)> isSinkingBeneficiary);
} // namespace mlir
-#endif // MLIR_DIALECT_GPU_UTILS_H_
+#endif // MLIR_DIALECT_GPU_TRANSFORMS_UTILS_H_
diff --git a/mlir/include/mlir/InitAllDialects.h b/mlir/include/mlir/InitAllDialects.h
index 0b569c424065..58b91d5a4b3f 100644
--- a/mlir/include/mlir/InitAllDialects.h
+++ b/mlir/include/mlir/InitAllDialects.h
@@ -30,7 +30,7 @@
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/EmitC/IR/EmitC.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
diff --git a/mlir/include/mlir/InitAllPasses.h b/mlir/include/mlir/InitAllPasses.h
index d6aacb4f6da8..bf8d96011f3f 100644
--- a/mlir/include/mlir/InitAllPasses.h
+++ b/mlir/include/mlir/InitAllPasses.h
@@ -20,7 +20,7 @@
#include "mlir/Dialect/Async/Passes.h"
#include "mlir/Dialect/Bufferization/Transforms/Passes.h"
#include "mlir/Dialect/Func/Transforms/Passes.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/Transforms/Passes.h"
#include "mlir/Dialect/Linalg/Passes.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
diff --git a/mlir/lib/CAPI/Dialect/GPU.cpp b/mlir/lib/CAPI/Dialect/GPU.cpp
index 0de2cfa330ef..cd58f0e249a9 100644
--- a/mlir/lib/CAPI/Dialect/GPU.cpp
+++ b/mlir/lib/CAPI/Dialect/GPU.cpp
@@ -8,6 +8,6 @@
#include "mlir-c/Dialect/GPU.h"
#include "mlir/CAPI/Registration.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
MLIR_DEFINE_CAPI_DIALECT_REGISTRATION(GPU, gpu, mlir::gpu::GPUDialect)
diff --git a/mlir/lib/CAPI/Dialect/GPUPasses.cpp b/mlir/lib/CAPI/Dialect/GPUPasses.cpp
index 4ec167f8854f..5128c63ecb51 100644
--- a/mlir/lib/CAPI/Dialect/GPUPasses.cpp
+++ b/mlir/lib/CAPI/Dialect/GPUPasses.cpp
@@ -7,11 +7,11 @@
//===----------------------------------------------------------------------===//
#include "mlir/CAPI/Pass.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Pass/Pass.h"
// Must include the declarations as they carry important visibility attributes.
-#include "mlir/Dialect/GPU/Passes.capi.h.inc"
+#include "mlir/Dialect/GPU/Transforms/Passes.capi.h.inc"
using namespace mlir;
@@ -19,7 +19,7 @@ using namespace mlir;
extern "C" {
#endif
-#include "mlir/Dialect/GPU/Passes.capi.cpp.inc"
+#include "mlir/Dialect/GPU/Transforms/Passes.capi.cpp.inc"
#ifdef __cplusplus
}
diff --git a/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h b/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h
index 0800a00ec0d8..58d8ceae2404 100644
--- a/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h
+++ b/mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h
@@ -9,7 +9,7 @@
#define MLIR_CONVERSION_GPUCOMMON_GPUOPSLOWERING_H_
#include "mlir/Conversion/LLVMCommon/Pattern.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
namespace mlir {
diff --git a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
index 49484717abdc..26fed8c1691e 100644
--- a/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
+++ b/mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
@@ -26,8 +26,8 @@
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
#include "mlir/Dialect/Async/IR/Async.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/IR/Attributes.h"
#include "mlir/IR/Builders.h"
diff --git a/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h b/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
index 97006b199b3f..938c3040e3bc 100644
--- a/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
+++ b/mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
@@ -9,7 +9,7 @@
#define MLIR_CONVERSION_GPUCOMMON_INDEXINTRINSICSOPLOWERING_H_
#include "mlir/Conversion/LLVMCommon/Pattern.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "llvm/ADT/StringSwitch.h"
diff --git a/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h b/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
index a08aafca9181..2cf16934acdc 100644
--- a/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
+++ b/mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
@@ -9,7 +9,7 @@
#define MLIR_CONVERSION_GPUCOMMON_OPTOFUNCCALLLOWERING_H_
#include "mlir/Conversion/LLVMCommon/Pattern.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/IR/Builders.h"
diff --git a/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td b/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
index dd5889bd1fab..f513bb1a0a82 100644
--- a/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
+++ b/mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
@@ -14,7 +14,7 @@
#define MLIR_CONVERSION_GPUTONVVM_TD
include "mlir/IR/PatternBase.td"
-include "mlir/Dialect/GPU/GPUOps.td"
+include "mlir/Dialect/GPU/IR/GPUOps.td"
include "mlir/Dialect/LLVMIR/NVVMOps.td"
def : Pat<(GPU_BarrierOp), (NVVM_Barrier0Op)>;
diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
index 56b3733c5e45..e6ae7c8126c0 100644
--- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
@@ -23,8 +23,8 @@
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/ControlFlow/IR/ControlFlow.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Dialect/Math/IR/Math.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
diff --git a/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp b/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
index 1a9254e9db89..912d350e6bd5 100644
--- a/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
@@ -13,7 +13,7 @@
#include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
#include "mlir/Conversion/LLVMCommon/Pattern.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/IR/TypeUtilities.h"
diff --git a/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td b/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
index 92da56046e6d..8d2f30a9a168 100644
--- a/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
+++ b/mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
@@ -14,7 +14,7 @@
#define MLIR_CONVERSION_GPUTOROCDL_TD
include "mlir/IR/PatternBase.td"
-include "mlir/Dialect/GPU/GPUOps.td"
+include "mlir/Dialect/GPU/IR/GPUOps.td"
include "mlir/Dialect/LLVMIR/ROCDLOps.td"
def : Pat<(GPU_BarrierOp), (ROCDL_BarrierOp)>;
diff --git a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
index 9aec6afb90cc..6b48b4ee4230 100644
--- a/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
+++ b/mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
@@ -25,8 +25,8 @@
#include "mlir/Conversion/VectorToLLVM/ConvertVectorToLLVM.h"
#include "mlir/Conversion/VectorToROCDL/VectorToROCDL.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
#include "mlir/Dialect/Math/IR/Math.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
index 514330a5a925..dacee9972e3d 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
@@ -11,7 +11,7 @@
//===----------------------------------------------------------------------===//
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVEnums.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
index c28e89898125..5eaa099278b2 100644
--- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
+++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
@@ -18,7 +18,7 @@
#include "mlir/Conversion/FuncToSPIRV/FuncToSPIRV.h"
#include "mlir/Conversion/GPUToSPIRV/GPUToSPIRV.h"
#include "mlir/Conversion/MemRefToSPIRV/MemRefToSPIRV.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/Dialect/SPIRV/Transforms/SPIRVConversion.h"
diff --git a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
index a0c1eb027e22..361dae64ceaf 100644
--- a/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
+++ b/mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
@@ -16,7 +16,7 @@
#include "../PassDetail.h"
#include "mlir/Conversion/GPUToVulkan/ConvertGPUToVulkanPass.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/IR/Attributes.h"
diff --git a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
index 7ee7dce361f3..0c9de3254dd1 100644
--- a/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
+++ b/mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
@@ -10,7 +10,7 @@
#include "../PassDetail.h"
#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
#include "mlir/Conversion/LLVMCommon/Pattern.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/NVVMDialect.h"
#include "mlir/Dialect/NVGPU/NVGPUDialect.h"
diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
index 901810ec4958..6e74a90bc03d 100644
--- a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
+++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
@@ -17,8 +17,8 @@
#include "mlir/Conversion/AffineToStandard/AffineToStandard.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/ParallelLoopMapper.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SCF/SCF.h"
#include "mlir/IR/AffineExpr.h"
diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp
index c0e983b455b0..84a3af5229bf 100644
--- a/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp
+++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp
@@ -12,7 +12,7 @@
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/Complex/IR/Complex.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SCF/SCF.h"
#include "mlir/Transforms/DialectConversion.h"
diff --git a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
index 5619e08f4e9e..fa6930fe46c8 100644
--- a/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
+++ b/mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp
@@ -20,7 +20,7 @@
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVM.h"
#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVOps.h"
#include "mlir/IR/BuiltinOps.h"
diff --git a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
index 90a37dab12eb..652edfd8f727 100644
--- a/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
+++ b/mlir/lib/Conversion/VectorToGPU/VectorToGPU.cpp
@@ -18,7 +18,7 @@
#include "../PassDetail.h"
#include "mlir/Analysis/SliceAnalysis.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/NVGPU/NVGPUDialect.h"
#include "mlir/Dialect/SCF/SCF.h"
diff --git a/mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp b/mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp
index a8492466162b..1890c0ee1619 100644
--- a/mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp
+++ b/mlir/lib/Conversion/VectorToROCDL/VectorToROCDL.cpp
@@ -18,7 +18,7 @@
#include "mlir/Conversion/LLVMCommon/ConversionTarget.h"
#include "mlir/Conversion/LLVMCommon/Pattern.h"
#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/LLVMIR/ROCDLDialect.h"
#include "mlir/Dialect/Vector/IR/VectorOps.h"
diff --git a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index 8b83820a2f36..1b10085d994e 100644
--- a/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -10,7 +10,7 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
@@ -31,7 +31,7 @@
using namespace mlir;
using namespace mlir::gpu;
-#include "mlir/Dialect/GPU/GPUOpsDialect.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc"
//===----------------------------------------------------------------------===//
// MMAMatrixType
@@ -121,11 +121,11 @@ void GPUDialect::initialize() {
addTypes<MMAMatrixType>();
addOperations<
#define GET_OP_LIST
-#include "mlir/Dialect/GPU/GPUOps.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
>();
addAttributes<
#define GET_ATTRDEF_LIST
-#include "mlir/Dialect/GPU/GPUOpsAttributes.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
>();
addInterfaces<GPUInlinerInterface>();
}
@@ -1398,11 +1398,11 @@ void AllocOp::getCanonicalizationPatterns(RewritePatternSet &results,
results.add<SimplifyDimOfAllocOp>(context);
}
-#include "mlir/Dialect/GPU/GPUOpInterfaces.cpp.inc"
-#include "mlir/Dialect/GPU/GPUOpsEnums.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc"
#define GET_ATTRDEF_CLASSES
-#include "mlir/Dialect/GPU/GPUOpsAttributes.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc"
#define GET_OP_CLASSES
-#include "mlir/Dialect/GPU/GPUOps.cpp.inc"
+#include "mlir/Dialect/GPU/IR/GPUOps.cpp.inc"
diff --git a/mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp b/mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp
index f8f51059d8e9..85adf0013866 100644
--- a/mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp
@@ -13,8 +13,8 @@
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/IR/BlockAndValueMapping.h"
#include "mlir/IR/Builders.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp
index 0832fcb72b8b..59758b80357d 100644
--- a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp
@@ -13,9 +13,9 @@
#include "PassDetail.h"
#include "mlir/Dialect/Async/IR/Async.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
-#include "mlir/Dialect/GPU/Utils.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Dialect/GPU/Transforms/Utils.h"
#include "mlir/IR/BlockAndValueMapping.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/PatternMatch.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
index 12897d66678e..94c977f3eb88 100644
--- a/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
@@ -15,9 +15,9 @@
#include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
-#include "mlir/Dialect/GPU/Utils.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir/Dialect/GPU/Transforms/Utils.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/IR/BlockAndValueMapping.h"
#include "mlir/IR/Builders.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/MemoryPromotion.cpp b/mlir/lib/Dialect/GPU/Transforms/MemoryPromotion.cpp
index 62304fbb2cfa..72fa50c9d7b6 100644
--- a/mlir/lib/Dialect/GPU/Transforms/MemoryPromotion.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/MemoryPromotion.cpp
@@ -11,10 +11,11 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/MemoryPromotion.h"
+#include "mlir/Dialect/GPU/Transforms/MemoryPromotion.h"
+
#include "mlir/Dialect/Affine/LoopUtils.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SCF/SCF.h"
#include "mlir/IR/ImplicitLocOpBuilder.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
index d2e8ed671bee..ccdef7587790 100644
--- a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp
@@ -11,11 +11,11 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/ParallelLoopMapper.h"
+#include "mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h"
#include "PassDetail.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/SCF/SCF.h"
#include "mlir/IR/AffineMap.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/PassDetail.h b/mlir/lib/Dialect/GPU/Transforms/PassDetail.h
index 99faa71eb1a9..2f5c7dc803ca 100644
--- a/mlir/lib/Dialect/GPU/Transforms/PassDetail.h
+++ b/mlir/lib/Dialect/GPU/Transforms/PassDetail.h
@@ -12,13 +12,13 @@
#include "mlir/Dialect/Async/IR/Async.h"
#include "mlir/Dialect/DLTI/DLTI.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Pass/Pass.h"
namespace mlir {
#define GEN_PASS_CLASSES
-#include "mlir/Dialect/GPU/Passes.h.inc"
+#include "mlir/Dialect/GPU/Transforms/Passes.h.inc"
} // namespace mlir
diff --git a/mlir/lib/Dialect/GPU/Transforms/SerializeToBlob.cpp b/mlir/lib/Dialect/GPU/Transforms/SerializeToBlob.cpp
index 8dadb630f4a9..35bf5942f24d 100644
--- a/mlir/lib/Dialect/GPU/Transforms/SerializeToBlob.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/SerializeToBlob.cpp
@@ -12,7 +12,7 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/SerializeToCubin.cpp b/mlir/lib/Dialect/GPU/Transforms/SerializeToCubin.cpp
index 9742607440c4..b33db153c644 100644
--- a/mlir/lib/Dialect/GPU/Transforms/SerializeToCubin.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/SerializeToCubin.cpp
@@ -10,7 +10,8 @@
// adds that blob as a string attribute of the module.
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/Passes.h"
+
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#if MLIR_GPU_TO_CUBIN_PASS_ENABLE
#include "mlir/Pass/Pass.h"
diff --git a/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp b/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
index 15e1674a84d5..9c11a509c197 100644
--- a/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp
@@ -10,7 +10,8 @@
// adds that blob as a string attribute of the module.
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/Passes.h"
+
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/IR/Location.h"
#include "mlir/IR/MLIRContext.h"
diff --git a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
index ad5b6bd1db65..42659c1c3fd6 100644
--- a/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
+++ b/mlir/lib/Dialect/NVGPU/IR/NVGPUDialect.cpp
@@ -11,7 +11,7 @@
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/NVGPU/NVGPUDialect.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/IR/Builders.h"
#include "mlir/IR/DialectImplementation.h"
#include "mlir/IR/OpImplementation.h"
diff --git a/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToCubin.cpp b/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToCubin.cpp
index 7644acd45432..1c442b0147c8 100644
--- a/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToCubin.cpp
+++ b/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToCubin.cpp
@@ -6,8 +6,7 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/Passes.h"
-
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Target/LLVMIR/Dialect/NVVM/NVVMToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
diff --git a/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToHsaco.cpp b/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToHsaco.cpp
index 8f614d4b15b9..c204e86632ac 100644
--- a/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToHsaco.cpp
+++ b/mlir/test/lib/Dialect/GPU/TestConvertGPUKernelToHsaco.cpp
@@ -6,8 +6,7 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/Passes.h"
-
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Target/LLVMIR/Dialect/ROCDL/ROCDLToLLVMIRTranslation.h"
#include "mlir/Target/LLVMIR/Export.h"
diff --git a/mlir/test/lib/Dialect/GPU/TestGpuMemoryPromotion.cpp b/mlir/test/lib/Dialect/GPU/TestGpuMemoryPromotion.cpp
index 81e66f95270a..12dcec487aee 100644
--- a/mlir/test/lib/Dialect/GPU/TestGpuMemoryPromotion.cpp
+++ b/mlir/test/lib/Dialect/GPU/TestGpuMemoryPromotion.cpp
@@ -12,8 +12,8 @@
//===----------------------------------------------------------------------===//
#include "mlir/Dialect/Affine/IR/AffineOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/MemoryPromotion.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/MemoryPromotion.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SCF/SCF.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
diff --git a/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp b/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
index 3d81a26c0bc6..896488b66cc5 100644
--- a/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
+++ b/mlir/test/lib/Dialect/GPU/TestGpuRewrite.cpp
@@ -12,7 +12,7 @@
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Pass/Pass.h"
#include "mlir/Transforms/GreedyPatternRewriteDriver.h"
diff --git a/mlir/test/lib/Dialect/Linalg/TestLinalgCodegenStrategy.cpp b/mlir/test/lib/Dialect/Linalg/TestLinalgCodegenStrategy.cpp
index 74ca6b057291..883d9a3ad579 100644
--- a/mlir/test/lib/Dialect/Linalg/TestLinalgCodegenStrategy.cpp
+++ b/mlir/test/lib/Dialect/Linalg/TestLinalgCodegenStrategy.cpp
@@ -14,7 +14,7 @@
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
#include "mlir/Dialect/Linalg/Transforms/CodegenStrategy.h"
#include "mlir/Dialect/Linalg/Utils/Utils.h"
diff --git a/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp b/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp
index b82381625a78..4c8a9484eb4e 100644
--- a/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp
+++ b/mlir/test/lib/Dialect/Linalg/TestLinalgTransforms.cpp
@@ -13,7 +13,7 @@
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
#include "mlir/Dialect/Linalg/Passes.h"
#include "mlir/Dialect/Linalg/Transforms/HoistPadding.h"
diff --git a/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp b/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp
index 69aa4f64d504..e1b303f45f71 100644
--- a/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp
+++ b/mlir/test/lib/Dialect/SPIRV/TestEntryPointAbi.cpp
@@ -11,7 +11,7 @@
//
//===----------------------------------------------------------------------===//
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
#include "mlir/Dialect/SPIRV/IR/TargetAndABI.h"
#include "mlir/Pass/Pass.h"
diff --git a/mlir/test/lib/Dialect/Vector/TestVectorTransforms.cpp b/mlir/test/lib/Dialect/Vector/TestVectorTransforms.cpp
index 6901e51c22e8..f216ba8e80eb 100644
--- a/mlir/test/lib/Dialect/Vector/TestVectorTransforms.cpp
+++ b/mlir/test/lib/Dialect/Vector/TestVectorTransforms.cpp
@@ -11,7 +11,7 @@
#include "mlir/Analysis/SliceAnalysis.h"
#include "mlir/Dialect/Affine/IR/AffineOps.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/Linalg/IR/Linalg.h"
#include "mlir/Dialect/Linalg/Passes.h"
diff --git a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp b/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
index fb51af9002e1..9a3c1be9e8ea 100644
--- a/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
+++ b/mlir/tools/mlir-spirv-cpu-runner/mlir-spirv-cpu-runner.cpp
@@ -17,8 +17,8 @@
#include "mlir/Conversion/SPIRVToLLVM/SPIRVToLLVMPass.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/SPIRV/IR/SPIRVDialect.h"
diff --git a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
index 6d0edda7feb9..2c672794e3d2 100644
--- a/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
+++ b/mlir/tools/mlir-vulkan-runner/mlir-vulkan-runner.cpp
@@ -21,8 +21,8 @@
#include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h"
#include "mlir/Dialect/Arithmetic/IR/Arithmetic.h"
#include "mlir/Dialect/Func/IR/FuncOps.h"
-#include "mlir/Dialect/GPU/GPUDialect.h"
-#include "mlir/Dialect/GPU/Passes.h"
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
#include "mlir/Dialect/MemRef/IR/MemRef.h"
#include "mlir/Dialect/MemRef/Transforms/Passes.h"
diff --git a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
index 6c36e2022960..454d850f01b7 100644
--- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
+++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel
@@ -3451,9 +3451,9 @@ cc_library(
td_library(
name = "GPUOpsTdFiles",
srcs = [
- "include/mlir/Dialect/GPU/GPUBase.td",
- "include/mlir/Dialect/GPU/GPUOps.td",
- "include/mlir/Dialect/GPU/ParallelLoopMapperAttr.td",
+ "include/mlir/Dialect/GPU/IR/GPUBase.td",
+ "include/mlir/Dialect/GPU/IR/GPUOps.td",
+ "include/mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td",
],
includes = ["include"],
deps = [
@@ -3472,15 +3472,15 @@ gentbl_cc_library(
tbl_outs = [
(
["-gen-op-interface-decls"],
- "include/mlir/Dialect/GPU/GPUOpInterfaces.h.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpInterfaces.h.inc",
),
(
["-gen-op-interface-defs"],
- "include/mlir/Dialect/GPU/GPUOpInterfaces.cpp.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpInterfaces.cpp.inc",
),
],
tblgen = ":mlir-tblgen",
- td_file = "include/mlir/Dialect/GPU/GPUBase.td",
+ td_file = "include/mlir/Dialect/GPU/IR/GPUBase.td",
deps = [":OpBaseTdFiles"],
)
@@ -3493,42 +3493,42 @@ gentbl_cc_library(
"-gen-dialect-decls",
"-dialect=gpu",
],
- "include/mlir/Dialect/GPU/GPUOpsDialect.h.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpsDialect.h.inc",
),
(
[
"-gen-dialect-defs",
"-dialect=gpu",
],
- "include/mlir/Dialect/GPU/GPUOpsDialect.cpp.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpsDialect.cpp.inc",
),
(
["-gen-op-decls"],
- "include/mlir/Dialect/GPU/GPUOps.h.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOps.h.inc",
),
(
["-gen-op-defs"],
- "include/mlir/Dialect/GPU/GPUOps.cpp.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOps.cpp.inc",
),
(
["-gen-enum-decls"],
- "include/mlir/Dialect/GPU/GPUOpsEnums.h.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpsEnums.h.inc",
),
(
["-gen-enum-defs"],
- "include/mlir/Dialect/GPU/GPUOpsEnums.cpp.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpsEnums.cpp.inc",
),
(
["-gen-attrdef-decls"],
- "include/mlir/Dialect/GPU/GPUOpsAttributes.h.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpsAttributes.h.inc",
),
(
["-gen-attrdef-defs"],
- "include/mlir/Dialect/GPU/GPUOpsAttributes.cpp.inc",
+ "include/mlir/Dialect/GPU/IR/GPUOpsAttributes.cpp.inc",
),
],
tblgen = ":mlir-tblgen",
- td_file = "include/mlir/Dialect/GPU/GPUOps.td",
+ td_file = "include/mlir/Dialect/GPU/IR/GPUOps.td",
deps = [
":DLTIDialectTdFiles",
":GPUOpsTdFiles",
@@ -3543,9 +3543,7 @@ cc_library(
"lib/Dialect/GPU/IR/*.h",
],
),
- hdrs = [
- "include/mlir/Dialect/GPU/GPUDialect.h",
- ],
+ hdrs = glob(["include/mlir/Dialect/GPU/IR/*.h"]),
includes = ["include"],
deps = [
":ArithmeticDialect",
@@ -3570,25 +3568,25 @@ gentbl_cc_library(
"-gen-pass-decls",
"-name=GPU",
],
- "include/mlir/Dialect/GPU/Passes.h.inc",
+ "include/mlir/Dialect/GPU/Transforms/Passes.h.inc",
),
(
[
"-gen-pass-capi-header",
"--prefix=GPU",
],
- "include/mlir/Dialect/GPU/Passes.capi.h.inc",
+ "include/mlir/Dialect/GPU/Transforms/Passes.capi.h.inc",
),
(
[
"-gen-pass-capi-impl",
"--prefix=GPU",
],
- "include/mlir/Dialect/GPU/Passes.capi.cpp.inc",
+ "include/mlir/Dialect/GPU/Transforms/Passes.capi.cpp.inc",
),
],
tblgen = ":mlir-tblgen",
- td_file = "include/mlir/Dialect/GPU/Passes.td",
+ td_file = "include/mlir/Dialect/GPU/Transforms/Passes.td",
deps = [":PassBaseTdFiles"],
)
@@ -3600,12 +3598,7 @@ cc_library(
"lib/Dialect/GPU/Transforms/*.h",
],
),
- hdrs = [
- "include/mlir/Dialect/GPU/MemoryPromotion.h",
- "include/mlir/Dialect/GPU/ParallelLoopMapper.h",
- "include/mlir/Dialect/GPU/Passes.h",
- "include/mlir/Dialect/GPU/Utils.h",
- ],
+ hdrs = glob(["include/mlir/Dialect/GPU/Transforms/*.h"]),
defines = if_cuda_available(["MLIR_GPU_TO_CUBIN_PASS_ENABLE"]),
includes = ["include"],
deps = [
More information about the Mlir-commits
mailing list