Merge OpenAI Triton commit f4a9d54 (#4688)

The main desire is to include commit
f4a9d54, which will fix CI.

This PR change the Triton base from
673ca35
to f4a9d54 (Jul 11).
Pass rate: 97.78%

Please do not squash and merge this PR.

---------
Signed-off-by: Anatoly Myachev <[email protected]>

Author: anmyachev

.github/workflows/build-macos.yml

@@ -97,6 +97,8 @@ jobs:
           # macos-latest has 3 vcpus and 7GB DRAM, to save memory we limit the number of jobs to 3
           # https://docs.github.com/en/actions/using-github-hosted-runners/about-github-hosted-runners/about-github-hosted-runners#standard-github-hosted-runners-for-public-repositories
           MAX_JOBS: 3
+          # Add elapsed time in seconds to ninja status to monitor where build stalls
+          NINJA_STATUS: "[%f/%t, %es elapsed] "
         run: |
           source ~/.venv/bin/activate
           echo "PATH is '$PATH'"

bin/RegisterTritonDialects.h

@@ -15,6 +15,7 @@
 #include "third_party/nvidia/include/Dialect/NVGPU/IR/Dialect.h"
 #include "third_party/nvidia/include/Dialect/NVWS/IR/Dialect.h"
 #include "third_party/proton/dialect/include/Dialect/Proton/IR/Dialect.h"
+#include "triton/Dialect/Gluon/Transforms/Passes.h"
 #include "triton/Dialect/Triton/IR/Dialect.h"
 #include "triton/Dialect/TritonGPU/IR/Dialect.h"
 #include "triton/Dialect/TritonInstrument/IR/Dialect.h"
@@ -67,6 +68,7 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
   mlir::triton::nvidia_gpu::registerTritonNvidiaGPUPasses();
   mlir::test::intel::registerTestAxisInfoPass();
   mlir::triton::instrument::registerTritonInstrumentPasses();
+  mlir::triton::gluon::registerGluonPasses();
   mlir::test::registerTestAliasPass();
   mlir::test::registerTestAlignmentPass();
   mlir::test::registerAMDTestAlignmentPass();
@@ -135,5 +137,6 @@ inline void registerTritonDialects(mlir::DialectRegistry &registry) {
       mlir::triton::amdgpu::TritonAMDGPUDialect,
       mlir::triton::proton::ProtonDialect, mlir::ROCDL::ROCDLDialect,
       mlir::triton::gpu::intel::TritonIntelGPUDialect,
-      mlir::triton::TritonGEN::TritonGENDialect>();
+      mlir::triton::TritonGEN::TritonGENDialect,
+      mlir::triton::gluon::GluonDialect>();
 }

include/triton/Dialect/CMakeLists.txt

@@ -2,3 +2,4 @@ add_subdirectory(Triton)
 add_subdirectory(TritonGPU)
 add_subdirectory(TritonNvidiaGPU)
 add_subdirectory(TritonInstrument)
+add_subdirectory(Gluon)

include/triton/Dialect/Gluon/CMakeCache.txt

@@ -0,0 +1,2 @@
+add_subdirectory(IR)
+add_subdirectory(Transforms)

include/triton/Dialect/Gluon/CMakeLists.txt

@@ -0,0 +1,2 @@
+add_subdirectory(IR)
+add_subdirectory(Transforms)

include/triton/Dialect/Gluon/IR/CMakeLists.txt

@@ -0,0 +1,16 @@
+set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR})
+
+set(LLVM_TARGET_DEFINITIONS GluonDialect.td)
+mlir_tablegen(Dialect.h.inc -gen-dialect-decls -dialect=gluon)
+mlir_tablegen(Dialect.cpp.inc -gen-dialect-defs -dialect=gluon)
+mlir_tablegen(Ops.h.inc -gen-op-decls)
+mlir_tablegen(Ops.cpp.inc -gen-op-defs)
+add_mlir_doc(GluonDialect GluonDialect dialects/ -gen-dialect-doc)
+add_public_tablegen_target(GluonTableGen)
+
+set(LLVM_TARGET_DEFINITIONS GluonAttrDefs.td)
+mlir_tablegen(GluonAttrDefs.h.inc -gen-attrdef-decls)
+mlir_tablegen(GluonAttrDefs.cpp.inc -gen-attrdef-defs)
+mlir_tablegen(OpsEnums.h.inc -gen-enum-decls)
+mlir_tablegen(OpsEnums.cpp.inc -gen-enum-defs)
+add_public_tablegen_target(GluonAttrDefsIncGen)

include/triton/Dialect/Gluon/IR/Dialect.h

@@ -0,0 +1,8 @@
+#pragma once
+#include "triton/Dialect/Triton/IR/Dialect.h"
+#include "triton/Dialect/TritonGPU/IR/Dialect.h"
+
+#include "triton/Dialect/Gluon/IR/Dialect.h.inc"
+
+#define GET_ATTRDEF_CLASSES
+#include "triton/Dialect/Gluon/IR/GluonAttrDefs.h.inc"

include/triton/Dialect/Gluon/IR/GluonAttrDefs.td

@@ -0,0 +1,15 @@
+#ifndef GLUON_ATTRDEFS
+#define GLUON_ATTRDEFS
+
+include "mlir/IR/AttrTypeBase.td"
+include "triton/Dialect/Gluon/IR/GluonDialect.td"
+
+def Gluon_AutoEncodingAttr : AttrDef<Gluon_Dialect, "AutoEncoding"> {
+  let mnemonic = "auto_encoding";
+  let attrName = "gluon.auto_encoding";
+  let description = [{
+    An encoding that is inferred from neighboring ops in the graph.
+  }];
+}
+
+#endif

include/triton/Dialect/Gluon/IR/GluonDialect.td

@@ -0,0 +1,22 @@
+#ifndef GLUON_DIALECT
+#define GLUON_DIALECT
+
+include "mlir/IR/OpBase.td"
+
+def Gluon_Dialect : Dialect {
+  let name = "gluon";
+  let cppNamespace = "::mlir::triton::gluon";
+  let description = [{
+    Gluon dialect.
+  }];
+
+  let dependentDialects = [
+    "triton::TritonDialect",
+    "triton::gpu::TritonGPUDialect",
+    "mlir::gpu::GPUDialect",
+  ];
+  let useDefaultAttributePrinterParser = 1;
+  let usePropertiesForAttributes = 1;
+}
+
+#endif

include/triton/Dialect/Gluon/Transforms/CMakeLists.txt

@@ -0,0 +1,3 @@
+set(LLVM_TARGET_DEFINITIONS Passes.td)
+mlir_tablegen(Passes.h.inc -gen-pass-decls -name Gluon)
+add_public_tablegen_target(GluonTransformsIncGen)