Move GPU dialect to {lib,include/mlir}/Dialect

Per tacit agreement, individual dialects should now live in lib/Dialect/Name
with headers in include/mlir/Dialect/Name and tests in test/Dialect/Name.

PiperOrigin-RevId: 259896851

diff --git a/mlir/include/mlir/CMakeLists.txt b/mlir/include/mlir/CMakeLists.txt
index 0a07efb..55843c0 100644 (file)
--- a/mlir/include/mlir/CMakeLists.txt
+++ b/mlir/include/mlir/CMakeLists.txt
@@ -1,7 +1,6 @@
 add_subdirectory(AffineOps)
 add_subdirectory(Dialect)
 add_subdirectory(EDSC)
-add_subdirectory(GPU)
 add_subdirectory(Linalg)
 add_subdirectory(LLVMIR)
 add_subdirectory(StandardOps)

diff --git a/mlir/include/mlir/Dialect/CMakeLists.txt b/mlir/include/mlir/Dialect/CMakeLists.txt
index 8d93f89..5ae314a 100644 (file)
--- a/mlir/include/mlir/Dialect/CMakeLists.txt
+++ b/mlir/include/mlir/Dialect/CMakeLists.txt
@@ -1,4 +1,5 @@
 add_subdirectory(FxpMathOps)
+add_subdirectory(GPU)
 add_subdirectory(LoopOps)
 add_subdirectory(QuantOps)
 add_subdirectory(SPIRV)

diff --git a/mlir/include/mlir/GPU/CMakeLists.txt b/mlir/include/mlir/Dialect/GPU/CMakeLists.txt
similarity index 100%
rename from mlir/include/mlir/GPU/CMakeLists.txt
rename to mlir/include/mlir/Dialect/GPU/CMakeLists.txt

diff --git a/mlir/include/mlir/GPU/GPUDialect.h b/mlir/include/mlir/Dialect/GPU/GPUDialect.h
similarity index 97%
rename from mlir/include/mlir/GPU/GPUDialect.h
rename to mlir/include/mlir/Dialect/GPU/GPUDialect.h
index c1e8ce5..d034212 100644 (file)
--- a/mlir/include/mlir/GPU/GPUDialect.h
+++ b/mlir/include/mlir/Dialect/GPU/GPUDialect.h
@@ -20,8 +20,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_GPU_GPUDIALECT_H
-#define MLIR_GPU_GPUDIALECT_H
+#ifndef MLIR_DIALECT_GPU_GPUDIALECT_H
+#define MLIR_DIALECT_GPU_GPUDIALECT_H
 
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
@@ -166,9 +166,9 @@ private:
 };
 
 #define GET_OP_CLASSES
-#include "mlir/GPU/GPUOps.h.inc"
+#include "mlir/Dialect/GPU/GPUOps.h.inc"
 
 } // end namespace gpu
 } // end namespace mlir
 
-#endif // MLIR_GPUKERNEL_GPUDIALECT_H
+#endif // MLIR_DIALECT_GPU_GPUDIALECT_H

diff --git a/mlir/include/mlir/GPU/GPUOps.td b/mlir/include/mlir/Dialect/GPU/GPUOps.td
similarity index 100%
rename from mlir/include/mlir/GPU/GPUOps.td
rename to mlir/include/mlir/Dialect/GPU/GPUOps.td

diff --git a/mlir/include/mlir/GPU/Passes.h b/mlir/include/mlir/Dialect/GPU/Passes.h
similarity index 90%
rename from mlir/include/mlir/GPU/Passes.h
rename to mlir/include/mlir/Dialect/GPU/Passes.h
index 9dd4ca0..f9b569d 100644 (file)
--- a/mlir/include/mlir/GPU/Passes.h
+++ b/mlir/include/mlir/Dialect/GPU/Passes.h
@@ -19,8 +19,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-#ifndef MLIR_GPU_PASSES_H_
-#define MLIR_GPU_PASSES_H_
+#ifndef MLIR_DIALECT_GPU_PASSES_H_
+#define MLIR_DIALECT_GPU_PASSES_H_
 
 namespace mlir {
 
@@ -30,4 +30,4 @@ ModulePassBase *createGpuKernelOutliningPass();
 
 } // namespace mlir
 
-#endif // MLIR_GPU_PASSES_H_
+#endif // MLIR_DIALECT_GPU_PASSES_H_

diff --git a/mlir/lib/CMakeLists.txt b/mlir/lib/CMakeLists.txt
index e5cf39c..fece5cb 100644 (file)
--- a/mlir/lib/CMakeLists.txt
+++ b/mlir/lib/CMakeLists.txt
@@ -4,7 +4,6 @@ add_subdirectory(Conversion)
 add_subdirectory(Dialect)
 add_subdirectory(EDSC)
 add_subdirectory(ExecutionEngine)
-add_subdirectory(GPU)
 add_subdirectory(IR)
 add_subdirectory(LLVMIR)
 add_subdirectory(Linalg)

diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp
index 43af366..7663775 100644 (file)
--- a/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/ConvertKernelFuncToCubin.cpp
@@ -23,7 +23,7 @@
 
 #include "mlir/Conversion/GPUToCUDA/GPUToCUDAPass.h"
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"

diff --git a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
index dd147fc..bf75778 100644 (file)
--- a/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/ConvertLaunchFuncToCudaCalls.cpp
@@ -24,7 +24,7 @@
 
 #include "mlir/Conversion/GPUToCUDA/GPUToCUDAPass.h"
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"

diff --git a/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp b/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp
index cc0466c..813a3be 100644 (file)
--- a/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp
+++ b/mlir/lib/Conversion/GPUToCUDA/GenerateCubinAccessors.cpp
@@ -20,7 +20,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/IR/Attributes.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"

diff --git a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
index 8d1ea54..e4a6f96 100644 (file)
--- a/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
+++ b/mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp
@@ -20,7 +20,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/StandardTypes.h"
 #include "mlir/LLVMIR/LLVMDialect.h"

diff --git a/mlir/lib/Conversion/LoopsToGPU/LoopsToGPU.cpp b/mlir/lib/Conversion/LoopsToGPU/LoopsToGPU.cpp
index 9027368..6ca4cb3 100644 (file)
--- a/mlir/lib/Conversion/LoopsToGPU/LoopsToGPU.cpp
+++ b/mlir/lib/Conversion/LoopsToGPU/LoopsToGPU.cpp
@@ -23,8 +23,8 @@
 
 #include "mlir/Conversion/LoopsToGPU/LoopsToGPU.h"
 #include "mlir/AffineOps/AffineOps.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/Dialect/LoopOps/LoopOps.h"
-#include "mlir/GPU/GPUDialect.h"
 #include "mlir/IR/AffineExpr.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/StandardOps/Ops.h"

diff --git a/mlir/lib/Dialect/CMakeLists.txt b/mlir/lib/Dialect/CMakeLists.txt
index 2dfe2a8..8898c43 100644 (file)
--- a/mlir/lib/Dialect/CMakeLists.txt
+++ b/mlir/lib/Dialect/CMakeLists.txt
@@ -1,4 +1,5 @@
 add_subdirectory(FxpMathOps)
+add_subdirectory(GPU)
 add_subdirectory(LoopOps)
 add_subdirectory(QuantOps)
 add_subdirectory(SPIRV)

diff --git a/mlir/lib/GPU/CMakeLists.txt b/mlir/lib/Dialect/GPU/CMakeLists.txt
similarity index 85%
rename from mlir/lib/GPU/CMakeLists.txt
rename to mlir/lib/Dialect/GPU/CMakeLists.txt
index 9cc8c46..09da5cc 100644 (file)
--- a/mlir/lib/GPU/CMakeLists.txt
+++ b/mlir/lib/Dialect/GPU/CMakeLists.txt
@@ -4,7 +4,7 @@ add_llvm_library(MLIRGPU
   Transforms/KernelOutlining.cpp
 
   ADDITIONAL_HEADER_DIRS
-  ${MLIR_MAIN_INCLUDE_DIR}/mlir/GPU
+  ${MLIR_MAIN_INCLUDE_DIR}/mlir/Dialect/GPU
 )
 add_dependencies(MLIRGPU MLIRGPUOpsIncGen MLIRIR LLVMSupport)
 target_link_libraries(MLIRGPU MLIRIR MLIRStandardOps LLVMSupport)

diff --git a/mlir/lib/GPU/IR/DialectRegistration.cpp b/mlir/lib/Dialect/GPU/IR/DialectRegistration.cpp
similarity index 95%
rename from mlir/lib/GPU/IR/DialectRegistration.cpp
rename to mlir/lib/Dialect/GPU/IR/DialectRegistration.cpp
index 8d00032..af50d02 100644 (file)
--- a/mlir/lib/GPU/IR/DialectRegistration.cpp
+++ b/mlir/lib/Dialect/GPU/IR/DialectRegistration.cpp
@@ -15,7 +15,7 @@
 // limitations under the License.
 // =============================================================================
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 
 // Static initialization for GPU dialect registration.
 static mlir::DialectRegistration<mlir::gpu::GPUDialect> kernelDialect;

diff --git a/mlir/lib/GPU/IR/GPUDialect.cpp b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
similarity index 99%
rename from mlir/lib/GPU/IR/GPUDialect.cpp
rename to mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
index 06762ad..bda5979 100644 (file)
--- a/mlir/lib/GPU/IR/GPUDialect.cpp
+++ b/mlir/lib/Dialect/GPU/IR/GPUDialect.cpp
@@ -19,7 +19,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/Module.h"
@@ -43,12 +43,12 @@ GPUDialect::GPUDialect(MLIRContext *context)
     : Dialect(getDialectName(), context) {
   addOperations<LaunchOp, LaunchFuncOp,
 #define GET_OP_LIST
-#include "mlir/GPU/GPUOps.cpp.inc"
+#include "mlir/Dialect/GPU/GPUOps.cpp.inc"
                 >();
 }
 
 #define GET_OP_CLASSES
-#include "mlir/GPU/GPUOps.cpp.inc"
+#include "mlir/Dialect/GPU/GPUOps.cpp.inc"
 
 //===----------------------------------------------------------------------===//
 // LaunchOp

diff --git a/mlir/lib/GPU/Transforms/KernelOutlining.cpp b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
similarity index 98%
rename from mlir/lib/GPU/Transforms/KernelOutlining.cpp
rename to mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
index 75b98fd..01decce 100644 (file)
--- a/mlir/lib/GPU/Transforms/KernelOutlining.cpp
+++ b/mlir/lib/Dialect/GPU/Transforms/KernelOutlining.cpp
@@ -19,8 +19,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "mlir/GPU/GPUDialect.h"
-#include "mlir/GPU/Passes.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/Passes.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/Builders.h"
 #include "mlir/Pass/Pass.h"

diff --git a/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp b/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
index 82727b4..c670cbf 100644 (file)
--- a/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
+++ b/mlir/lib/Target/LLVMIR/ConvertToNVVMIR.cpp
@@ -22,7 +22,7 @@
 
 #include "mlir/Target/NVVMIR.h"
 
-#include "mlir/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/Module.h"
 #include "mlir/LLVMIR/NVVMDialect.h"

diff --git a/mlir/test/GPU/canonicalize.mlir b/mlir/test/Dialect/GPU/canonicalize.mlir
similarity index 100%
rename from mlir/test/GPU/canonicalize.mlir
rename to mlir/test/Dialect/GPU/canonicalize.mlir

diff --git a/mlir/test/GPU/invalid.mlir b/mlir/test/Dialect/GPU/invalid.mlir
similarity index 100%
rename from mlir/test/GPU/invalid.mlir
rename to mlir/test/Dialect/GPU/invalid.mlir

diff --git a/mlir/test/GPU/ops.mlir b/mlir/test/Dialect/GPU/ops.mlir
similarity index 100%
rename from mlir/test/GPU/ops.mlir
rename to mlir/test/Dialect/GPU/ops.mlir

diff --git a/mlir/test/GPU/outlining.mlir b/mlir/test/Dialect/GPU/outlining.mlir
similarity index 100%
rename from mlir/test/GPU/outlining.mlir
rename to mlir/test/Dialect/GPU/outlining.mlir

diff --git a/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp b/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
index 0077e95..edf6aea 100644 (file)
--- a/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
+++ b/mlir/tools/mlir-cuda-runner/mlir-cuda-runner.cpp
@@ -27,8 +27,8 @@
 #include "mlir/Conversion/GPUToNVVM/GPUToNVVMPass.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVM.h"
 #include "mlir/Conversion/StandardToLLVM/ConvertStandardToLLVMPass.h"
-#include "mlir/GPU/GPUDialect.h"
-#include "mlir/GPU/Passes.h"
+#include "mlir/Dialect/GPU/GPUDialect.h"
+#include "mlir/Dialect/GPU/Passes.h"
 #include "mlir/IR/Function.h"
 #include "mlir/IR/Module.h"
 #include "mlir/LLVMIR/LLVMDialect.h"