diff --git a/mlir/include/mlir-c/Dialect/GPU.h b/mlir/include/mlir-c/Dialect/GPU.h --- a/mlir/include/mlir-c/Dialect/GPU.h +++ b/mlir/include/mlir-c/Dialect/GPU.h @@ -23,6 +23,6 @@ } #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 --- 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/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt copy from mlir/include/mlir/Dialect/GPU/CMakeLists.txt copy to mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt --- a/mlir/include/mlir/Dialect/GPU/CMakeLists.txt +++ b/mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt @@ -16,12 +16,6 @@ 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) @@ -31,5 +25,3 @@ 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) diff --git a/mlir/include/mlir/Dialect/GPU/GPUBase.td b/mlir/include/mlir/Dialect/GPU/IR/GPUBase.td 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 rename from mlir/include/mlir/Dialect/GPU/GPUDialect.h rename to mlir/include/mlir/Dialect/GPU/IR/GPUDialect.h --- 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 @@ } // 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 rename from mlir/include/mlir/Dialect/GPU/GPUOps.td rename to mlir/include/mlir/Dialect/GPU/IR/GPUOps.td --- 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 rename from mlir/include/mlir/Dialect/GPU/ParallelLoopMapperAttr.td rename to mlir/include/mlir/Dialect/GPU/IR/ParallelLoopMapperAttr.td --- 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 --- /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 rename from mlir/include/mlir/Dialect/GPU/MemoryPromotion.h rename to mlir/include/mlir/Dialect/GPU/Transforms/MemoryPromotion.h --- 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 @@ } // 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 rename from mlir/include/mlir/Dialect/GPU/ParallelLoopMapper.h rename to mlir/include/mlir/Dialect/GPU/Transforms/ParallelLoopMapper.h --- 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 @@ ArrayRef 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 rename from mlir/include/mlir/Dialect/GPU/Passes.h rename to mlir/include/mlir/Dialect/GPU/Transforms/Passes.h --- 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 @@ /// 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 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 rename from mlir/include/mlir/Dialect/GPU/Utils.h rename to mlir/include/mlir/Dialect/GPU/Transforms/Utils.h --- 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 @@ llvm::function_ref 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 --- 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 --- 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 --- 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 --- 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 @@ 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 @@ addTypes(); 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(); } @@ -1398,11 +1398,11 @@ results.add(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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- 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 --- a/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel +++ b/utils/bazel/llvm-project-overlay/mlir/BUILD.bazel @@ -3451,9 +3451,9 @@ 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 @@ 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 @@ "-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 @@ "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 @@ "-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 @@ "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 = [