MoveTritonGEN dialect into third_party/intel (#1024)

Signed-off-by: Tiotto, Ettore <ettore.tiotto@intel.com>
intel · May 2, 2024 · 9c515cf · 9c515cf
1 parent 86b302e
commit 9c515cf
Show file tree

Hide file tree

Showing 33 changed files with 59 additions and 36 deletions.
diff --git a/bin/RegisterTritonDialects.h b/bin/RegisterTritonDialects.h
@@ -1,7 +1,8 @@
 #pragma once
+#include "third_party/intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "third_party/nvidia/include/Dialect/NVGPU/IR/Dialect.h"
+
 #include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"

diff --git a/include/triton/Dialect/CMakeLists.txt b/include/triton/Dialect/CMakeLists.txt
@@ -1,5 +1,4 @@
 add_subdirectory(Triton)
-add_subdirectory(TritonGEN)
 add_subdirectory(TritonGPU)
 add_subdirectory(TritonIntelGPU)
 add_subdirectory(TritonNvidiaGPU)
diff --git a/include/triton/Dialect/TritonIntelGPU/Transforms/Passes.h b/include/triton/Dialect/TritonIntelGPU/Transforms/Passes.h
@@ -10,7 +10,8 @@
 #define TRITON_DIALECT_TRITON_INTEL_GPU_TRANSFORMS_PASSES_H
 
 #include "mlir/Pass/Pass.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
 
 namespace mlir {
 namespace triton {

diff --git a/lib/Dialect/CMakeLists.txt b/lib/Dialect/CMakeLists.txt
@@ -1,5 +1,4 @@
 add_subdirectory(Triton)
-add_subdirectory(TritonGEN)
 add_subdirectory(TritonGPU)
 add_subdirectory(TritonIntelGPU)
 add_subdirectory(TritonNvidiaGPU)
diff --git a/python/src/ir.cc b/python/src/ir.cc
@@ -19,7 +19,6 @@
 #include "triton/Dialect/Triton/IR/Dialect.h"
 #include "triton/Dialect/Triton/IR/Types.h"
 #include "triton/Dialect/Triton/IR/Utility.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Tools/Sys/GetEnv.hpp"
 #include <pybind11/functional.h>
 #include <pybind11/pybind11.h>

diff --git a/python/src/llvm.cc b/python/src/llvm.cc
@@ -1,6 +1,6 @@
-#include "mlir/IR/BuiltinOps.h" // mlir::ModuleOp
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+#include "mlir/IR/BuiltinOps.h" // mlir::ModuleOp
 #include "mlir/Target/LLVMIR/ModuleTranslation.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Target/SPIRV/SPIRVTranslation.h"
 #include "triton/Tools/Sys/GetEnv.hpp"
 #include "llvm/ADT/SmallVector.h"

diff --git a/third_party/intel/CMakeLists.txt b/third_party/intel/CMakeLists.txt
@@ -7,7 +7,7 @@ add_triton_plugin(TritonXPU
   ${CMAKE_CURRENT_SOURCE_DIR}/triton_xpu.cc
 
   LINK_LIBS
+  TritonGENToLLVM
   TritonIntelGPUToLLVM
   TritonIntelGPUTransforms
-  TritonGENToLLVM
 )
diff --git a/third_party/intel/include/CMakeLists.txt b/third_party/intel/include/CMakeLists.txt
@@ -1,3 +1,4 @@
+add_subdirectory(Dialect)
 add_subdirectory(GPUToTritonGEN)
 add_subdirectory(TritonGENToLLVM)
 add_subdirectory(TritonIntelGPUToLLVM)
diff --git a/third_party/intel/include/Dialect/CMakeLists.txt b/third_party/intel/include/Dialect/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(TritonGEN)
diff --git a/...e/triton/Dialect/TritonGEN/CMakeLists.txt → .../include/Dialect/TritonGEN/CMakeLists.txt b/...e/triton/Dialect/TritonGEN/CMakeLists.txt → .../include/Dialect/TritonGEN/CMakeLists.txt
diff --git a/...riton/Dialect/TritonGEN/IR/CMakeLists.txt → ...clude/Dialect/TritonGEN/IR/CMakeLists.txt b/...riton/Dialect/TritonGEN/IR/CMakeLists.txt → ...clude/Dialect/TritonGEN/IR/CMakeLists.txt
diff --git a/...Dialect/TritonGEN/IR/TritonGENAttrDefs.td → ...Dialect/TritonGEN/IR/TritonGENAttrDefs.td b/...Dialect/TritonGEN/IR/TritonGENAttrDefs.td → ...Dialect/TritonGEN/IR/TritonGENAttrDefs.td
diff --git a/...n/Dialect/TritonGEN/IR/TritonGENDialect.h → ...e/Dialect/TritonGEN/IR/TritonGENDialect.h b/...n/Dialect/TritonGEN/IR/TritonGENDialect.h → ...e/Dialect/TritonGEN/IR/TritonGENDialect.h
@@ -17,14 +17,15 @@
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Dialect.h"
 #include "mlir/IR/OpDefinition.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h.inc"
-#include "triton/Dialect/TritonGEN/IR/TritonGENOpsEnums.h.inc"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOpsEnums.h.inc"
 
 #define GET_ATTRDEF_CLASSES
-#include "triton/Dialect/TritonGEN/IR/TritonGENOpsAttrDefs.h.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOpsAttrDefs.h.inc"
 
 #define GET_OP_CLASSES
-#include "triton/Dialect/TritonGEN/IR/TritonGENOps.h.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOps.h.inc"
 
 namespace mlir {
 namespace triton {

diff --git a/.../Dialect/TritonGEN/IR/TritonGENDialect.td → .../Dialect/TritonGEN/IR/TritonGENDialect.td b/.../Dialect/TritonGEN/IR/TritonGENDialect.td → .../Dialect/TritonGEN/IR/TritonGENDialect.td
diff --git a/...iton/Dialect/TritonGEN/IR/TritonGENOps.td → ...lude/Dialect/TritonGEN/IR/TritonGENOps.td b/...iton/Dialect/TritonGEN/IR/TritonGENOps.td → ...lude/Dialect/TritonGEN/IR/TritonGENOps.td
@@ -13,8 +13,8 @@
 #ifndef TRITONGEN_OPS
 #define TRITONGEN_OPS
 
-include "triton/Dialect/TritonGEN/IR/TritonGENDialect.td"
-include "triton/Dialect/TritonGEN/IR/TritonGENAttrDefs.td"
+include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.td"
+include "intel/include/Dialect/TritonGEN/IR/TritonGENAttrDefs.td"
 include "mlir/IR/OpBase.td"
 include "mlir/IR/EnumAttr.td"
 include "mlir/Dialect/LLVMIR/LLVMOpBase.td"

diff --git a/third_party/intel/lib/CMakeLists.txt b/third_party/intel/lib/CMakeLists.txt
@@ -1,3 +1,4 @@
+add_subdirectory(Dialect)
 add_subdirectory(GPUToTritonGEN)
 add_subdirectory(TritonGENToLLVM)
 add_subdirectory(TritonIntelGPUToLLVM)

diff --git a/third_party/intel/lib/Dialect/CMakeLists.txt b/third_party/intel/lib/Dialect/CMakeLists.txt
@@ -0,0 +1 @@
+add_subdirectory(TritonGEN)
diff --git a/lib/Dialect/TritonGEN/CMakeLists.txt → ...ntel/lib/Dialect/TritonGEN/CMakeLists.txt b/lib/Dialect/TritonGEN/CMakeLists.txt → ...ntel/lib/Dialect/TritonGEN/CMakeLists.txt
diff --git a/lib/Dialect/TritonGEN/IR/CMakeLists.txt → ...l/lib/Dialect/TritonGEN/IR/CMakeLists.txt b/lib/Dialect/TritonGEN/IR/CMakeLists.txt → ...l/lib/Dialect/TritonGEN/IR/CMakeLists.txt
diff --git a/...Dialect/TritonGEN/IR/TritonGENDialect.cpp → ...Dialect/TritonGEN/IR/TritonGENDialect.cpp b/...Dialect/TritonGEN/IR/TritonGENDialect.cpp → ...Dialect/TritonGEN/IR/TritonGENDialect.cpp
@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
 
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/IR/Builders.h"
@@ -27,20 +27,20 @@ using namespace mlir::triton::TritonGEN;
 void TritonGENDialect::initialize() {
   addOperations<
 #define GET_OP_LIST
-#include "triton/Dialect/TritonGEN/IR/TritonGENOps.cpp.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOps.cpp.inc"
       >();
   addAttributes<
 #define GET_ATTRDEF_LIST
-#include "triton/Dialect/TritonGEN/IR/TritonGENOpsAttrDefs.cpp.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOpsAttrDefs.cpp.inc"
       >();
 
   // Support unknown operations because not all GEN operations are registered.
   allowUnknownOperations();
 }
 
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.cpp.inc"
-#include "triton/Dialect/TritonGEN/IR/TritonGENOpsEnums.cpp.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.cpp.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOpsEnums.cpp.inc"
 #define GET_ATTRDEF_CLASSES
-#include "triton/Dialect/TritonGEN/IR/TritonGENOpsAttrDefs.cpp.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOpsAttrDefs.cpp.inc"
 #define GET_OP_CLASSES
-#include "triton/Dialect/TritonGEN/IR/TritonGENOps.cpp.inc"
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENOps.cpp.inc"
diff --git a/lib/Dialect/TritonGEN/IR/TritonGENOps.cpp → ...lib/Dialect/TritonGEN/IR/TritonGENOps.cpp b/lib/Dialect/TritonGEN/IR/TritonGENOps.cpp → ...lib/Dialect/TritonGEN/IR/TritonGENOps.cpp
@@ -8,7 +8,9 @@
 
 #include "mlir/Dialect/Utils/StaticValueUtils.h"
 #include "mlir/IR/OpDefinition.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "llvm/ADT/STLExtras.h"
 #include <cstdint>
 

diff --git a/third_party/intel/lib/GPUToTritonGEN/GPUToTritonGEN.td b/third_party/intel/lib/GPUToTritonGEN/GPUToTritonGEN.td
@@ -15,7 +15,7 @@
 
 include "mlir/IR/PatternBase.td"
 include "mlir/Dialect/GPU/IR/GPUOps.td"
-include "triton/Dialect/TritonGEN/IR/TritonGENOps.td"
+include "intel/include/Dialect/TritonGEN/IR/TritonGENOps.td"
 
 def : Pat<(GPU_BarrierOp), (TritonGEN_BarrierOp)>;
 

diff --git a/third_party/intel/lib/GPUToTritonGEN/GPUToTritonGENPass.cpp b/third_party/intel/lib/GPUToTritonGEN/GPUToTritonGENPass.cpp
@@ -37,8 +37,9 @@
 #include "mlir/Transforms/GreedyPatternRewriteDriver.h"
 #include "llvm/Support/FormatVariadic.h"
 
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Conversion/TritonGPUToLLVM/Utility.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 
 #include "GPUOpsLowering.h"
 #include "IndexIntrinsicsOpLowering.h"

diff --git a/third_party/intel/lib/TritonGENToLLVM/TritonGENToLLVMPass.cpp b/third_party/intel/lib/TritonGENToLLVM/TritonGENToLLVMPass.cpp
@@ -30,8 +30,9 @@
 #include "llvm/Support/ErrorHandling.h"
 #include <type_traits>
 
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Conversion/TritonGPUToLLVM/Utility.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 
 namespace mlir {
 namespace triton {

diff --git a/third_party/intel/lib/TritonIntelGPUToLLVM/AllocateSharedMemory.cpp b/third_party/intel/lib/TritonIntelGPUToLLVM/AllocateSharedMemory.cpp
@@ -1,9 +1,11 @@
-#include "intel/include/TritonIntelGPUToLLVM/Passes.h"
 #include "mlir/Pass/Pass.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+#include "intel/include/TritonIntelGPUToLLVM/Passes.h"
+
 #include "triton/Analysis/Allocation.h"
 #include "triton/Analysis/Utility.h"
 #include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 
 using namespace mlir;

diff --git a/third_party/intel/lib/TritonIntelGPUToLLVM/DotOpToLLVM/DPAS.cpp b/third_party/intel/lib/TritonIntelGPUToLLVM/DotOpToLLVM/DPAS.cpp
@@ -2,12 +2,16 @@
 #include "../Utility.h"
 
 #include "mlir/IR/BuiltinTypes.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonIntelGPU/Transforms/Utility.h"
+
 #include "llvm/IR/DerivedTypes.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/raw_ostream.h"
+
 #include <optional>
 #include <thread>
 

diff --git a/third_party/intel/lib/TritonIntelGPUToLLVM/PrintOpToLLVM.cpp b/third_party/intel/lib/TritonIntelGPUToLLVM/PrintOpToLLVM.cpp
@@ -1,6 +1,7 @@
 #include "PatternTritonGPUOpToLLVM.h"
 #include "Utility.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
 
 namespace {
 

diff --git a/third_party/intel/lib/TritonIntelGPUToLLVM/TritonGPUToLLVM.cpp b/third_party/intel/lib/TritonIntelGPUToLLVM/TritonGPUToLLVM.cpp
@@ -8,14 +8,15 @@
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
 #include "mlir/Pass/Pass.h"
 
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "intel/include/GPUToTritonGEN/GPUToTritonGENPass.h"
 #include "intel/include/TritonGENToLLVM/TritonGENToLLVMPass.h"
 
 #include "triton/Analysis/Allocation.h"
 #include "triton/Analysis/AxisInfo.h"
 #include "triton/Analysis/Membar.h"
 #include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"
 

diff --git a/third_party/intel/lib/TritonIntelGPUToLLVM/TritonOpsToLLVM.cpp b/third_party/intel/lib/TritonIntelGPUToLLVM/TritonOpsToLLVM.cpp
@@ -1,8 +1,9 @@
 #include "PatternTritonGPUOpToLLVM.h"
-#include "triton/Analysis/Utility.h"
 
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
+#include "triton/Analysis/Utility.h"
 #include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 
 using namespace mlir;
 using namespace mlir::triton;

diff --git a/third_party/intel/lib/TritonIntelGPUToLLVM/Utility.h b/third_party/intel/lib/TritonIntelGPUToLLVM/Utility.h
@@ -10,9 +10,11 @@
 #define TRITON_CONVERSION_TRITONINTELGPU_TO_LLVM_UTILITY_H
 
 #include "mlir/Dialect/ControlFlow/IR/ControlFlowOps.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Conversion/TritonGPUToLLVM/Utility.h"
 #include "triton/Dialect/Triton/IR/Utility.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"
 
 #define DEBUG_TYPE "ttgpu_to_llvm"

diff --git a/third_party/intel/lib/TritonIntelGPUTransforms/DistributeToWarps.cpp b/third_party/intel/lib/TritonIntelGPUTransforms/DistributeToWarps.cpp
@@ -15,8 +15,9 @@
 #include "mlir/IR/BuiltinTypes.h"
 #include "mlir/IR/TypeUtilities.h"
 
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Dialect/Triton/IR/Utility.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonGPU/Transforms/Utility.h"
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"

diff --git a/third_party/intel/lib/TritonIntelGPUTransforms/PrefetchBlock.cpp b/third_party/intel/lib/TritonIntelGPUTransforms/PrefetchBlock.cpp
@@ -45,8 +45,9 @@
 #include "mlir/IR/PatternMatch.h"
 #include "mlir/Pass/Pass.h"
 
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Dialect/Triton/IR/Dialect.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonIntelGPU/Transforms/Passes.h"

diff --git a/third_party/intel/triton_xpu.cc b/third_party/intel/triton_xpu.cc
@@ -3,7 +3,9 @@
 #include "mlir/Pass/PassManager.h"
 #include "passes.h"
 #include "triton/Conversion/TritonToTritonGPU/Passes.h"
-#include "triton/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
+#include "intel/include/Dialect/TritonGEN/IR/TritonGENDialect.h"
+
 #include "triton/Dialect/TritonIntelGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonIntelGPU/Transforms/Passes.h"
 #include "triton/Dialect/TritonNvidiaGPU/IR/Dialect.h"