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

12 files changed, 14 insertions, 14 deletions

diff --git a/mlir/lib/CMakeLists.txt b/mlir/lib/CMakeLists.txt
index e5cf39cf276..fece5cbb063 100644
--- 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 43af3662384..766377528a1 100644
--- 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 dd147fc586d..bf7577856db 100644
--- 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 cc0466cf3b1..813a3bee0ad 100644
--- 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 8d1ea54c906..e4a6f964f50 100644
--- 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 90273687cb2..6ca4cb39f83 100644
--- 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 2dfe2a81605..8898c43fc1d 100644
--- 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
index 9cc8c46ea1f..09da5cc16e9 100644
--- 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
index 8d000321e9a..af50d0270cf 100644
--- 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
index 06762ad232c..bda5979939c 100644
--- 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
index 75b98fdb9c5..01decce28ac 100644
--- 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 82727b48410..c670cbf8337 100644
--- 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"