[NFI]: Make createConvertTritonToTritonGPUWarp be in mlir::triton::in…

…tel namespace (#1276) Signed-off-by: Tiotto, Ettore <[email protected]>
intel · Jun 6, 2024 · 9a70f9b · 9a70f9b
1 parent 722784b
commit 9a70f9b
Show file tree

Hide file tree

Showing 5 changed files with 29 additions and 34 deletions.
diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h
@@ -49,7 +49,7 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   mlir::test::registerTestAllocationPass();
   mlir::test::registerTestMembarPass();
   mlir::triton::registerConvertTritonToTritonGPUPass();
-  mlir::triton::registerConvertTritonToTritonGPUWarpPass();
+  mlir::triton::intel::registerConvertTritonToTritonGPUWarpPass();
   mlir::triton::registerAllocateSharedMemoryPass();
   mlir::triton::registerConvertTritonGPUToLLVMPass();
   mlir::triton::registerConvertNVGPUToLLVMPass();

diff --git a/third_party/intel/include/TritonToTritonGPUWarp/Passes.h b/third_party/intel/include/TritonToTritonGPUWarp/Passes.h
@@ -3,14 +3,12 @@
 
 #include "intel/include/TritonToTritonGPUWarp/TritonToTritonGPUWarpPass.h"
 
-namespace mlir {
-namespace triton {
+namespace mlir::triton::intel {
 
 #define GEN_PASS_DECL
 #define GEN_PASS_REGISTRATION
 #include "intel/include/TritonToTritonGPUWarp/Passes.h.inc"
 
-} // namespace triton
-} // namespace mlir
+} // namespace mlir::triton::intel
 
 #endif
diff --git a/third_party/intel/include/TritonToTritonGPUWarp/TritonToTritonGPUWarpPass.h b/third_party/intel/include/TritonToTritonGPUWarp/TritonToTritonGPUWarpPass.h
@@ -2,22 +2,21 @@
 #define TRITON_CONVERSION_TRITONTOTRITONGPUWARP_TRITONTOTRITONGPUWARPPASS_H
 
 #include <memory>
-#include <optional>
 
 namespace mlir {
 
 class ModuleOp;
 template <typename T> class OperationPass;
 
-namespace triton {
+namespace triton::intel {
 
 std::unique_ptr<OperationPass<ModuleOp>>
 createConvertTritonToTritonGPUWarpPass();
 
 std::unique_ptr<OperationPass<ModuleOp>>
 createConvertTritonToTritonGPUWarpPass(unsigned numWarps);
 
-} // namespace triton
+} // namespace triton::intel
 } // namespace mlir
 
 #endif
diff --git a/third_party/intel/lib/TritonToTritonGPUWarp/TritonToTritonGPUWarpPass.cpp b/third_party/intel/lib/TritonToTritonGPUWarp/TritonToTritonGPUWarpPass.cpp
@@ -38,11 +38,11 @@ using namespace mlir;
 namespace tt = mlir::triton;
 namespace ttg = mlir::triton::gpu;
 
-namespace mlir::triton {
+namespace mlir::triton::intel {
 #define GEN_PASS_DECL_CONVERTTRITONTOTRITONGPUWARP
 #define GEN_PASS_DEF_CONVERTTRITONTOTRITONGPUWARP
 #include "intel/include/TritonToTritonGPUWarp/Passes.h.inc"
-} // namespace mlir::triton
+} // namespace mlir::triton::intel
 
 #define DEBUG_TYPE "convert-triton-to-tritongpu-warp"
 #define DBGS() (llvm::dbgs() << "[" DEBUG_TYPE "]: ")
@@ -126,7 +126,7 @@ struct LoopDotInfo {
 
 } // namespace
 
-namespace mlir::triton {
+namespace mlir::triton::intel {
 class ConvertTritonToTritonGPUWarp
     : public impl::ConvertTritonToTritonGPUWarpBase<
           ConvertTritonToTritonGPUWarp> {
@@ -465,4 +465,4 @@ class ConvertTritonToTritonGPUWarp
   }
 };
 
-} // namespace mlir::triton
+} // namespace mlir::triton::intel
diff --git a/third_party/intel/triton_xpu.cc b/third_party/intel/triton_xpu.cc
@@ -9,9 +9,7 @@
 #include "intel/include/Dialect/TritonIntelGPU/Transforms/Passes.h"
 #include "intel/include/Target/LLVMIR/Dialect/TritonGEN/TritonGENToLLVMIRTranslation.h"
 #include "intel/include/TritonIntelGPUToLLVM/Passes.h"
-
 #include "intel/include/TritonToTritonGPUWarp/Passes.h"
-#include "intel/include/TritonToTritonGPUWarp/TritonToTritonGPUWarpPass.h"
 
 #include "triton/Target/SPIRV/SPIRVTranslation.h"
 
@@ -22,7 +20,6 @@
 namespace py = pybind11;
 
 using namespace mlir::triton;
-using namespace mlir::triton::gpu;
 using ret = py::return_value_policy;
 
 // Macros to create a pass that takes pass options.
@@ -48,36 +45,37 @@ static uint32_t findKernels(llvm::Module &M,
 
 void init_triton_intel_passes_ttir(py::module &&m) {
   ADD_PASS_WRAPPER_OPT_1("add_convert_to_ttgpuir_warp",
-                         mlir::triton::createConvertTritonToTritonGPUWarp,
-                         unsigned);
+                         intel::createConvertTritonToTritonGPUWarp, unsigned);
 }
 
 void init_triton_intel_passes_ttgpuir(py::module &&m) {
-  py::enum_<intel::DeviceArch>(m, "DEVICE_ARCH", py::module_local())
-      .value("UNKNOWN", intel::DeviceArch::UNKNOWN)
-      .value("ATS", intel::DeviceArch::ATS)
-      .value("PVC", intel::DeviceArch::PVC)
+  py::enum_<gpu::intel::DeviceArch>(m, "DEVICE_ARCH", py::module_local())
+      .value("UNKNOWN", gpu::intel::DeviceArch::UNKNOWN)
+      .value("ATS", gpu::intel::DeviceArch::ATS)
+      .value("PVC", gpu::intel::DeviceArch::PVC)
       .export_values();
 
-  ADD_PASS_WRAPPER_0("add_to_llvmir", intel::createConvertTritonIntelGPUToLLVM);
+  ADD_PASS_WRAPPER_0("add_to_llvmir",
+                     gpu::intel::createConvertTritonIntelGPUToLLVM);
   ADD_PASS_WRAPPER_0("add_accelerate_matmul",
-                     intel::createTritonIntelGPUAccelerateMatmul);
+                     gpu::intel::createTritonIntelGPUAccelerateMatmul);
   ADD_PASS_WRAPPER_0("add_decompose_unsupported_conversions",
-                     intel::createIntelDecomposeUnsupportedConversions);
+                     gpu::intel::createIntelDecomposeUnsupportedConversions);
   ADD_PASS_WRAPPER_0("add_allocate_shared_memory",
-                     intel::createIntelAllocateSharedMemory);
-  ADD_PASS_WRAPPER_OPT_2("add_pipeline", intel::createTritonIntelGPUPipeline,
-                         int, bool);
+                     gpu::intel::createIntelAllocateSharedMemory);
+  ADD_PASS_WRAPPER_OPT_2("add_pipeline",
+                         gpu::intel::createTritonIntelGPUPipeline, int, bool);
   ADD_PASS_WRAPPER_0("add_remove_layout_conversions",
-                     intel::createTritonIntelGPURemoveLayoutConversions);
+                     gpu::intel::createTritonIntelGPURemoveLayoutConversions);
   ADD_PASS_WRAPPER_0("add_rewrite_tensor_pointer",
-                     intel::createTritonIntelGPURewriteTensorPointer);
+                     gpu::intel::createTritonIntelGPURewriteTensorPointer);
   ADD_PASS_WRAPPER_OPT_2("add_prefetch_block",
-                         intel::createTritonIntelGPUPrefetchBlock, int, bool);
+                         gpu::intel::createTritonIntelGPUPrefetchBlock, int,
+                         bool);
   ADD_PASS_WRAPPER_0("add_distribute_to_warps",
-                     intel::createTritonIntelGPUDistributeToWarps);
+                     gpu::intel::createTritonIntelGPUDistributeToWarps);
   ADD_PASS_WRAPPER_0("add_match_target_size",
-                     intel::createTritonIntelGPUMatchTargetSize);
+                     gpu::intel::createTritonIntelGPUMatchTargetSize);
 }
 
 void init_triton_intel(py::module &&m) {
@@ -88,8 +86,8 @@ void init_triton_intel(py::module &&m) {
   // load dialects
   m.def("load_dialects", [](mlir::MLIRContext &context) {
     mlir::DialectRegistry registry;
-    registry
-        .insert<TritonGEN::TritonGENDialect, intel::TritonIntelGPUDialect>();
+    registry.insert<TritonGEN::TritonGENDialect,
+                    gpu::intel::TritonIntelGPUDialect>();
     mlir::registerTritonGENDialectTranslation(registry);
     context.appendDialectRegistry(registry);
     context.loadAllAvailableDialects();