Skip to content

Commit

Permalink
[mlir][gpu] Move GPU headers into IR/ and Transforms/
Browse files Browse the repository at this point in the history 
Depends on D127350

Reviewed By: rriddle

Differential Revision: https://reviews.llvm.org/D127352
  • Loading branch information
@Mogball
Mogball committed Jun 9, 2022
1 parent 206f10d commit d7ef488
Show file tree
Hide file tree
Showing 57 changed files with 155 additions and 160 deletions.
2 changes: 1 addition & 1 deletion mlir/include/mlir-c/Dialect/GPU.h
View file
Expand Up @@ -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
37 changes: 2 additions & 35 deletions mlir/include/mlir/Dialect/GPU/CMakeLists.txt
View file
@@ -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)
27 changes: 27 additions & 0 deletions mlir/include/mlir/Dialect/GPU/IR/CMakeLists.txt
View file
@@ -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)
0 mlir/include/mlir/Dialect/GPU/GPUBase.td → mlir/include/mlir/Dialect/GPU/IR/GPUBase.td
View file
File renamed without changes.
16 changes: 8 additions & 8 deletions mlir/include/mlir/Dialect/GPU/GPUDialect.h → .../include/mlir/Dialect/GPU/IR/GPUDialect.h
View file
Expand Up @@ -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"
Expand Down Expand Up @@ -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
4 changes: 2 additions & 2 deletions mlir/include/mlir/Dialect/GPU/GPUOps.td → mlir/include/mlir/Dialect/GPU/IR/GPUOps.td
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion ...lir/Dialect/GPU/ParallelLoopMapperAttr.td → .../Dialect/GPU/IR/ParallelLoopMapperAttr.td
View file
Expand Up @@ -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">;
Expand Down
7 changes: 7 additions & 0 deletions mlir/include/mlir/Dialect/GPU/Transforms/CMakeLists.txt
View file
@@ -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)
6 changes: 3 additions & 3 deletions ...nclude/mlir/Dialect/GPU/MemoryPromotion.h → .../Dialect/GPU/Transforms/MemoryPromotion.h
View file
Expand Up @@ -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 {

Expand All @@ -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
8 changes: 4 additions & 4 deletions ...ude/mlir/Dialect/GPU/ParallelLoopMapper.h → ...alect/GPU/Transforms/ParallelLoopMapper.h
View file
Expand Up @@ -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"

Expand Down Expand Up @@ -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
10 changes: 5 additions & 5 deletions mlir/include/mlir/Dialect/GPU/Passes.h → ...lude/mlir/Dialect/GPU/Transforms/Passes.h
View file
Expand Up @@ -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 {
Expand Down Expand Up @@ -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_
0 mlir/include/mlir/Dialect/GPU/Passes.td → ...ude/mlir/Dialect/GPU/Transforms/Passes.td
View file
File renamed without changes.
6 changes: 3 additions & 3 deletions mlir/include/mlir/Dialect/GPU/Utils.h → ...clude/mlir/Dialect/GPU/Transforms/Utils.h
View file
Expand Up @@ -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"

Expand Down Expand Up @@ -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_
2 changes: 1 addition & 1 deletion mlir/include/mlir/InitAllDialects.h
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/include/mlir/InitAllPasses.h
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/CAPI/Dialect/GPU.cpp
View file
Expand Up @@ -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)
6 changes: 3 additions & 3 deletions mlir/lib/CAPI/Dialect/GPUPasses.cpp
View file
Expand Up @@ -7,19 +7,19 @@
//===----------------------------------------------------------------------===//

#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;

#ifdef __cplusplus
extern "C" {
#endif

#include "mlir/Dialect/GPU/Passes.capi.cpp.inc"
#include "mlir/Dialect/GPU/Transforms/Passes.capi.cpp.inc"

#ifdef __cplusplus
}
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUCommon/GPUOpsLowering.h
View file
Expand Up @@ -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 {
Expand Down
4 changes: 2 additions & 2 deletions mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUCommon/IndexIntrinsicsOpLowering.h
View file
Expand Up @@ -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"

Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUCommon/OpToFuncCallLowering.h
View file
Expand Up @@ -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"

Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUToNVVM/GPUToNVVM.td
View file
Expand Up @@ -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)>;
Expand Down
4 changes: 2 additions & 2 deletions mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUToROCDL/GPUToROCDL.td
View file
Expand Up @@ -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)>;
Expand Down
4 changes: 2 additions & 2 deletions mlir/lib/Conversion/GPUToROCDL/LowerGpuOpsToROCDLOps.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRVPass.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/GPUToVulkan/ConvertGPULaunchFuncToVulkanLaunchFunc.cpp
View file
Expand Up @@ -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"
Expand Down
2 changes: 1 addition & 1 deletion mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
View file
Expand Up @@ -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"

Expand Down
4 changes: 2 additions & 2 deletions mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp
View file
Expand Up @@ -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"
Expand Down

0 comments on commit d7ef488

Please sign in to comment.