Macqueen09/tile store tensor tt tensor 20260625

include/incubated/Conversion/TritonToLinalgIncubated/DescriptorConverter.h

@@ -35,6 +35,7 @@
 
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/TypeSwitch.h"
+#include "llvm/Config/llvm-config.h"
 #include "llvm/Support/Debug.h"
 
 namespace DescriptorConverter {
@@ -45,7 +46,9 @@ struct Descriptor {
   Value base;
   SmallVector<Value> shape;
   SmallVector<Value> strides;
+#if LLVM_VERSION_MAJOR >= 22
   triton::PaddingOptionAttr padding;
+#endif
 };
 
 bool hasATensorDescriptorType(mlir::TypeRange types);

include/incubated/Conversion/TritonToLinalgIncubated/TritonOpConverter.h

@@ -38,6 +38,7 @@
 
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/TypeSwitch.h"
+#include "llvm/Config/llvm-config.h"
 #include "llvm/Support/Debug.h"
 
 #define DEBUG_TYPE "triton-to-linalg"
@@ -190,6 +191,7 @@ class SplatConverter : public OpConversionPattern<triton::SplatOp> {
                   ConversionPatternRewriter &rewriter) const override;
 };
 
+#if LLVM_VERSION_MAJOR >= 22
 class UnsplatConverter : public OpConversionPattern<triton::UnsplatOp> {
 public:
   using OpConversionPattern<triton::UnsplatOp>::OpConversionPattern;
@@ -198,6 +200,7 @@ class UnsplatConverter : public OpConversionPattern<triton::UnsplatOp> {
   matchAndRewrite(triton::UnsplatOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;
 };
+#endif
 
 class ReshapeConverter : public OpConversionPattern<triton::ReshapeOp> {
 public:

include/mlir-ext/Dialect/CMakeLists.txt

@@ -1 +1,2 @@
 add_subdirectory(MathExt)
+add_subdirectory(TileIR)

include/mlir-ext/Dialect/TileIR/CMakeLists.txt

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

include/mlir-ext/Dialect/TileIR/IR/CMakeLists.txt

@@ -0,0 +1,25 @@
+set(MLIR_BINARY_DIR ${CMAKE_BINARY_DIR})
+
+# Ops tablegen
+set(LLVM_TARGET_DEFINITIONS TileIROps.td)
+mlir_tablegen(TileIRDialect.h.inc -gen-dialect-decls -dialect=tile)
+mlir_tablegen(TileIRDialect.cpp.inc -gen-dialect-defs -dialect=tile)
+mlir_tablegen(TileIROps.h.inc -gen-op-decls)
+mlir_tablegen(TileIROps.cpp.inc -gen-op-defs)
+add_mlir_doc(TileIRDialect TileIRDialect dialects/ -gen-dialect-doc)
+add_mlir_doc(TileIROps TileIROps dialects/ -gen-op-doc)
+add_public_tablegen_target(TileIRTableGen)
+
+# AttrDefs tablegen
+set(LLVM_TARGET_DEFINITIONS TileIRAttrDefs.td)
+mlir_tablegen(TileIROpsAttrDefs.h.inc -gen-attrdef-decls)
+mlir_tablegen(TileIROpsAttrDefs.cpp.inc -gen-attrdef-defs)
+mlir_tablegen(TileIREnums.h.inc -gen-enum-decls)
+mlir_tablegen(TileIREnums.cpp.inc -gen-enum-defs)
+add_public_tablegen_target(TileIRAttrDefsIncGen)
+
+# Types tablegen
+set(LLVM_TARGET_DEFINITIONS TileIRTypes.td)
+mlir_tablegen(TileIRTypes.h.inc -gen-typedef-decls)
+mlir_tablegen(TileIRTypes.cpp.inc -gen-typedef-defs)
+add_public_tablegen_target(TileIRTypesIncGen)

include/mlir-ext/Dialect/TileIR/IR/TileIRAttrDefs.td

@@ -0,0 +1,151 @@
+//===-- TileIRAttrDefs.td - TileIR attribute definitions -----*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines TileIR attributes: memory space, data layout, buffer policy,
+// pipe types, event IDs, and engine attributes. These are used as operands /
+// attributes on TileIR ops to encode explicit hardware control.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TILE_IR_ATTRDEFS
+#define TILE_IR_ATTRDEFS
+
+include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.td"
+
+include "mlir/IR/AttrTypeBase.td"
+include "mlir/IR/EnumAttr.td"
+
+//===----------------------------------------------------------------------===//
+// Attribute base class
+//===----------------------------------------------------------------------===//
+
+class TileIR_Attr<string name, string attrMnemonic, list<Trait> traits = []>
+    : AttrDef<TileIR_Dialect, name, traits> {
+  let mnemonic = attrMnemonic;
+}
+
+//===----------------------------------------------------------------------===//
+// Memory Space Enum
+//===----------------------------------------------------------------------===//
+
+def TileIR_MemorySpace : I64EnumAttr<
+  "MemorySpace", "TileIR memory space", [
+    I64EnumAttrCase<"GM",  0, "gm">,
+    I64EnumAttrCase<"L1",  1, "l1">,
+    I64EnumAttrCase<"L0A", 2, "l0a">,
+    I64EnumAttrCase<"L0B", 3, "l0b">,
+    I64EnumAttrCase<"L0C", 4, "l0c">,
+    I64EnumAttrCase<"UB",  5, "ub">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+//===----------------------------------------------------------------------===//
+// Data Layout Enum
+//===----------------------------------------------------------------------===//
+
+def TileIR_Layout : I64EnumAttr<
+  "Layout", "TileIR data layout format", [
+    I64EnumAttrCase<"ND",  0, "nd">,
+    I64EnumAttrCase<"NZ",  1, "nz">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+//===----------------------------------------------------------------------===//
+// Buffer Policy Enum
+//===----------------------------------------------------------------------===//
+
+def TileIR_Policy : I64EnumAttr<
+  "Policy", "TileIR buffer policy", [
+    I64EnumAttrCase<"SingleBuffer", 0, "single_buffer">,
+    I64EnumAttrCase<"DoubleBuffer", 1, "double_buffer">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+//===----------------------------------------------------------------------===//
+// Engine Enum
+//===----------------------------------------------------------------------===//
+
+def TileIR_Engine : I64EnumAttr<
+  "Engine", "TileIR engine type", [
+    I64EnumAttrCase<"Cube",    0, "cube">,
+    I64EnumAttrCase<"Vector",  1, "vector">,
+    I64EnumAttrCase<"MTE2",    2, "mte2">,
+    I64EnumAttrCase<"MTE3",    3, "mte3">,
+    I64EnumAttrCase<"MTE1",    4, "mte1">,
+    I64EnumAttrCase<"FixPipe", 5, "fixpipe">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+//===----------------------------------------------------------------------===//
+// Pipe Enum (for set_flag / wait_flag / pipe_barrier)
+//===----------------------------------------------------------------------===//
+
+def TileIR_Pipe : I64EnumAttr<
+  "Pipe", "TileIR pipe type for cross-engine synchronization", [
+    I64EnumAttrCase<"PIPE_M",    0, "PIPE_M">,
+    I64EnumAttrCase<"PIPE_V",    1, "PIPE_V">,
+    I64EnumAttrCase<"PIPE_MTE1", 2, "PIPE_MTE1">,
+    I64EnumAttrCase<"PIPE_MTE2", 3, "PIPE_MTE2">,
+    I64EnumAttrCase<"PIPE_MTE3", 4, "PIPE_MTE3">,
+    I64EnumAttrCase<"PIPE_FIX",  5, "PIPE_FIX">,
+    I64EnumAttrCase<"PIPE_S",    6, "PIPE_S">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+//===----------------------------------------------------------------------===//
+// Event ID Enum
+//===----------------------------------------------------------------------===//
+
+def TileIR_EventID : I64EnumAttr<
+  "EventID", "TileIR hardware event IDs for cross-engine synchronization", [
+    I64EnumAttrCase<"EVENT_ID0", 0, "EVENT_ID0">,
+    I64EnumAttrCase<"EVENT_ID1", 1, "EVENT_ID1">,
+    I64EnumAttrCase<"EVENT_ID2", 2, "EVENT_ID2">,
+    I64EnumAttrCase<"EVENT_ID3", 3, "EVENT_ID3">,
+    I64EnumAttrCase<"EVENT_ID4", 4, "EVENT_ID4">,
+    I64EnumAttrCase<"EVENT_ID5", 5, "EVENT_ID5">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+//===----------------------------------------------------------------------===//
+// NZ Fractal Layout Attribute (for Cube operands)
+//===----------------------------------------------------------------------===//
+
+def TileIR_NZLayoutAttr : TileIR_Attr<"NZLayout", "nz"> {
+  let summary = "NZ fractal layout attribute for Cube operands";
+  let description = [{
+    Specifies the NZ (Z-order) fractal layout used for Cube MMA operands.
+    The fractal_h and fractal_w parameters define the fractal tile size
+    (typically 16x16 for Ascend 910B).
+  }];
+  let parameters = (ins
+    "int64_t":$fractal_h,
+    "int64_t":$fractal_w
+  );
+  let assemblyFormat = "`<` $fractal_h `x` $fractal_w `>`";
+}
+
+//===----------------------------------------------------------------------===//
+// Alloc lifetime attribute
+//===----------------------------------------------------------------------===//
+
+def TileIR_Lifetime : I64EnumAttr<
+  "Lifetime", "TileIR buffer lifetime policy", [
+    I64EnumAttrCase<"Default",      0, "default">,
+    I64EnumAttrCase<"LoopCarried",  1, "loop_carried">,
+  ]> {
+  let cppNamespace = "::mlir::triton::tile";
+}
+
+#endif // TILE_IR_ATTRDEFS

include/mlir-ext/Dialect/TileIR/IR/TileIRDialect.h

@@ -0,0 +1,40 @@
+//===- TileIRDialect.h - MLIR TileIR dialect --------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines the TileIR dialect in MLIR, containing operations that
+// explicitly encode memory hierarchy, engine binding, data layout, and
+// synchronization for deterministic 1:1 lowering to HIVM.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TRITON_DIALECT_TILE_IR_DIALECT_H
+#define TRITON_DIALECT_TILE_IR_DIALECT_H
+
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+#include "mlir/IR/Dialect.h"
+#include "mlir/IR/OpDefinition.h"
+
+#include "triton/Dialect/Triton/IR/Dialect.h"
+
+#include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.h.inc"
+
+#include "mlir-ext/Dialect/TileIR/IR/TileIREnums.h.inc"
+
+#define GET_ATTRDEF_CLASSES
+#include "mlir-ext/Dialect/TileIR/IR/TileIROpsAttrDefs.h.inc"
+
+#define GET_TYPEDEF_CLASSES
+#include "mlir-ext/Dialect/TileIR/IR/TileIRTypes.h.inc"
+
+#define GET_OP_CLASSES
+#include "mlir-ext/Dialect/TileIR/IR/TileIROps.h.inc"
+
+namespace mlir::triton::tile {
+} // namespace mlir::triton::tile
+
+#endif // TRITON_DIALECT_TILE_IR_DIALECT_H

include/mlir-ext/Dialect/TileIR/IR/TileIRDialect.td

@@ -0,0 +1,48 @@
+//===-- TileIRDialect.td - TileIR dialect definition ---------*- tablegen -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This is the TileIR dialect definition file. TileIR is a hardware-oriented IR
+// that explicitly encodes memory hierarchy, engine binding, data layout, and
+// synchronization — enabling deterministic 1:1 lowering to HIVM (Ascend NPU).
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TILE_IR_DIALECT
+#define TILE_IR_DIALECT
+
+include "mlir/IR/OpBase.td"
+
+def TileIR_Dialect : Dialect {
+  let name = "tile";
+  let cppNamespace = "::mlir::triton::tile";
+  let summary = "The TileIR dialect for explicit hardware control on NPU.";
+
+  let description = [{
+    TileIR is a hardware-oriented intermediate representation that encodes:
+    - Memory hierarchy (GM/L1/L0A/L0B/L0C/UB)
+    - Engine binding (Cube, Vector, MTE2, MTE3)
+    - Data layout (ND, NZ fractal)
+    - Buffer policy (single/double buffer)
+    - Cross-engine synchronization (set_flag/wait_flag/pipe_barrier)
+    - Async compute (cube_launch)
+
+    TileIR ops map 1:1 to HIVM instructions, eliminating the need for
+    compiler-driven scheduling inference.
+  }];
+
+  let dependentDialects = [
+    "mlir::LLVM::LLVMDialect",
+    "triton::TritonDialect",
+  ];
+
+  let usePropertiesForAttributes = 1;
+  let useDefaultTypePrinterParser = 1;
+  let useDefaultAttributePrinterParser = 1;
+}
+
+#endif // TILE_IR_DIALECT