diff --git a/include/incubated/Conversion/TritonToLinalgIncubated/DescriptorConverter.h b/include/incubated/Conversion/TritonToLinalgIncubated/DescriptorConverter.h index 4982f005..3373a105 100644 --- a/include/incubated/Conversion/TritonToLinalgIncubated/DescriptorConverter.h +++ b/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 shape; SmallVector strides; +#if LLVM_VERSION_MAJOR >= 22 triton::PaddingOptionAttr padding; +#endif }; bool hasATensorDescriptorType(mlir::TypeRange types); diff --git a/include/incubated/Conversion/TritonToLinalgIncubated/TritonOpConverter.h b/include/incubated/Conversion/TritonToLinalgIncubated/TritonOpConverter.h index f3871f13..ef4dc0d6 100644 --- a/include/incubated/Conversion/TritonToLinalgIncubated/TritonOpConverter.h +++ b/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 { ConversionPatternRewriter &rewriter) const override; }; +#if LLVM_VERSION_MAJOR >= 22 class UnsplatConverter : public OpConversionPattern { public: using OpConversionPattern::OpConversionPattern; @@ -198,6 +200,7 @@ class UnsplatConverter : public OpConversionPattern { matchAndRewrite(triton::UnsplatOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const override; }; +#endif class ReshapeConverter : public OpConversionPattern { public: diff --git a/include/mlir-ext/Dialect/CMakeLists.txt b/include/mlir-ext/Dialect/CMakeLists.txt index 35948d69..6a305ce6 100644 --- a/include/mlir-ext/Dialect/CMakeLists.txt +++ b/include/mlir-ext/Dialect/CMakeLists.txt @@ -1 +1,2 @@ add_subdirectory(MathExt) +add_subdirectory(TileIR) diff --git a/include/mlir-ext/Dialect/TileIR/CMakeLists.txt b/include/mlir-ext/Dialect/TileIR/CMakeLists.txt new file mode 100644 index 00000000..7d59dce8 --- /dev/null +++ b/include/mlir-ext/Dialect/TileIR/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(IR) \ No newline at end of file diff --git a/include/mlir-ext/Dialect/TileIR/IR/CMakeLists.txt b/include/mlir-ext/Dialect/TileIR/IR/CMakeLists.txt new file mode 100644 index 00000000..a3751b63 --- /dev/null +++ b/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) \ No newline at end of file diff --git a/include/mlir-ext/Dialect/TileIR/IR/TileIRAttrDefs.td b/include/mlir-ext/Dialect/TileIR/IR/TileIRAttrDefs.td new file mode 100644 index 00000000..3a5d8348 --- /dev/null +++ b/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 traits = []> + : AttrDef { + 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 \ No newline at end of file diff --git a/include/mlir-ext/Dialect/TileIR/IR/TileIRDialect.h b/include/mlir-ext/Dialect/TileIR/IR/TileIRDialect.h new file mode 100644 index 00000000..81cf8b6e --- /dev/null +++ b/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 \ No newline at end of file diff --git a/include/mlir-ext/Dialect/TileIR/IR/TileIRDialect.td b/include/mlir-ext/Dialect/TileIR/IR/TileIRDialect.td new file mode 100644 index 00000000..993d35b5 --- /dev/null +++ b/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 \ No newline at end of file diff --git a/include/mlir-ext/Dialect/TileIR/IR/TileIROps.td b/include/mlir-ext/Dialect/TileIR/IR/TileIROps.td new file mode 100644 index 00000000..0d508da7 --- /dev/null +++ b/include/mlir-ext/Dialect/TileIR/IR/TileIROps.td @@ -0,0 +1,553 @@ +//===-- TileIROps.td - TileIR operation 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 all TileIR operations. TileIR ops explicitly encode memory +// hierarchy, engine binding, data layout, synchronization, and async compute — +// enabling deterministic 1:1 lowering to HIVM (Ascend NPU). +// +//===----------------------------------------------------------------------===// + +#ifndef TILE_IR_OPS +#define TILE_IR_OPS + +include "mlir/IR/OpBase.td" +include "mlir/IR/EnumAttr.td" +include "mlir/Interfaces/SideEffectInterfaces.td" +include "mlir/Interfaces/InferTypeOpInterface.td" + +include "triton/Dialect/Triton/IR/TritonTypes.td" +include "triton/Dialect/Triton/IR/TritonInterfaces.td" + +include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.td" +include "mlir-ext/Dialect/TileIR/IR/TileIRAttrDefs.td" +include "mlir-ext/Dialect/TileIR/IR/TileIRTypes.td" + +//===----------------------------------------------------------------------===// +// Op Base +//===----------------------------------------------------------------------===// + +class TileIR_Op traits = []> : + Op { + let hasVerifier = 1; +} + +//===----------------------------------------------------------------------===// +// Memory Operations +//===----------------------------------------------------------------------===// + +// --------------------------------------------------------------------------- +// tile.alloc — Allocate a buffer in a specific memory space +// --------------------------------------------------------------------------- +def TileIR_AllocOp : TileIR_Op<"alloc", [Pure]> { + let summary = "Allocate a buffer in a specific memory space"; + let description = [{ + Allocates a buffer with given shape, dtype, and memory space attributes. + Returns a !tile.buf value that can be used by copy/subview/cube_launch ops. + + Example: + ```mlir + %buf = tile.alloc {space=#tile.UB, shape=[128,128], dtype=f32, + policy=#tile.single_buffer, layout=#tile.nd, + lifetime="loop_carried"} + : !tile.buf<[128,128], f32, #tile.UB> + ``` + }]; + + let arguments = (ins + TileIR_MemorySpace:$space, + OptionalAttr:$shape, + OptionalAttr:$dtype, + OptionalAttr:$policy, + OptionalAttr:$layout, + OptionalAttr:$lifetime, + OptionalAttr:$comment + ); + + let results = (outs TileIR_BufType:$result); + + let assemblyFormat = "attr-dict `:` type($result)"; +} + +// --------------------------------------------------------------------------- +// tile.subview — Extract a sub-view (slice) from a buffer +// --------------------------------------------------------------------------- +def TileIR_SubViewOp : TileIR_Op<"subview", [Pure]> { + let summary = "Extract a sub-view (slice) from a buffer"; + let description = [{ + Creates a zero-copy view into a portion of a buffer. Used for slot + selection from multi-buffered buffers (e.g., ping-pong slot extraction). + + Example: + ```mlir + %slot = tile.subview %buf[%off, %c0, %c0] [1, 128, 128] [1, 1, 1] + : !tile.buf<[2,128,128],f16,#tile.L0A> -> !tile.buf<[128,128],f16,#tile.L0A> + ``` + }]; + + let arguments = (ins + TileIR_BufType:$source, + Variadic:$offsets, + I64ArrayAttr:$sizes, + I64ArrayAttr:$strides + ); + + let results = (outs TileIR_BufType:$result); + + let assemblyFormat = [{ + $source `[` $offsets `]` `[` $sizes `]` `[` $strides `]` + attr-dict `:` type($source) `->` type($result) + }]; +} + +// --------------------------------------------------------------------------- +// tile.gm_offset — Compute global memory address with strides +// --------------------------------------------------------------------------- +def TileIR_GmOffsetOp : TileIR_Op<"gm_offset", [Pure, SameVariadicOperandSize]> { + let summary = "Compute global memory offset with multi-dimensional strides"; + let description = [{ + Computes a GM pointer from a base pointer and multi-dimensional indices + with explicit strides. Replaces raw pointer arithmetic. Each index is + paired with a stride, so `indices` and `strides` must have the same arity. + + Example: + ```mlir + %off = tile.gm_offset %Q [%bz, %by, %row, %c0] strides [%sQb, %sQh, %sQs, %sQd] + : !tt.ptr -> !tt.ptr + ``` + }]; + + let arguments = (ins + TT_Ptr:$base, + Variadic:$indices, + Variadic:$strides + ); + + let results = (outs TT_Ptr:$result); + + let assemblyFormat = [{ + $base `[` $indices `]` `strides` `[` $strides `]` attr-dict + `:` type($base) `->` type($result) + }]; +} + +// --------------------------------------------------------------------------- +// tile.copy — Copy data between memory regions with layout conversion +// --------------------------------------------------------------------------- +def TileIR_CopyOp : TileIR_Op<"copy", [Pure, MemoryEffects<[MemWrite]>]> { + let summary = "Copy data between memory regions with optional layout conversion"; + let description = [{ + Copies data from source to destination, optionally performing layout + conversion (e.g., ND→NZ for Cube operands) and transpose. + + Example: + ```mlir + tile.copy %src → %dst {engine=#tile.mte2, src_layout=#tile.nd, + dst_layout=#tile.nz<16,16>, transpose=true} + ``` + }]; + + let arguments = (ins + AnyType:$src, + AnyType:$dst, + OptionalAttr:$engine, + OptionalAttr:$src_layout, + OptionalAttr:$dst_nz_layout, + OptionalAttr:$transpose, + OptionalAttr:$comment + ); + + let assemblyFormat = [{ + $src `->` $dst attr-dict `:` type($src) `,` type($dst) + }]; +} + +// --------------------------------------------------------------------------- +// tile.load — Load data from GM to local memory +// --------------------------------------------------------------------------- +def TileIR_LoadOp : TileIR_Op<"load", [Pure, MemoryEffects<[MemRead]>]> { + let summary = "Load data from global memory to local memory"; + let description = [{ + Loads a tensor from a GM pointer into the specified destination memory + space (typically UB for Vector compute). + + Example: + ```mlir + %t = tile.load %ptr {engine=#tile.mte2, dst_space=#tile.UB} + : !tt.ptr -> !tile.tensor<[128,128],f32,#tile.UB> + ``` + }]; + + let arguments = (ins + TT_Ptr:$src, + OptionalAttr:$engine, + OptionalAttr:$dst_space + ); + + let results = (outs TileIR_TensorType:$result); + + let assemblyFormat = "$src attr-dict `:` type($src) `->` type($result)"; +} + +// --------------------------------------------------------------------------- +// tile.store — Store data from local memory to GM +// --------------------------------------------------------------------------- +def TileIR_StoreOp : TileIR_Op<"store", [Pure, MemoryEffects<[MemWrite]>]> { + let summary = "Store data from local memory to global memory"; + let description = [{ + Stores a tensor from local memory (e.g., UB) to a GM pointer destination. + + Example: + ```mlir + tile.store %tensor → %ptr {engine=#tile.mte3, src_layout=#tile.nd, + dst_layout=#tile.nd} + ``` + }]; + + let arguments = (ins + AnyType:$src, + TT_Ptr:$dst, + OptionalAttr:$engine, + OptionalAttr:$src_layout, + OptionalAttr:$dst_layout, + OptionalAttr:$comment + ); + + let assemblyFormat = [{ + $src `->` $dst attr-dict `:` type($src) `,` type($dst) + }]; +} + +// --------------------------------------------------------------------------- +// tile.vec_fill — Fill a buffer with a scalar value (Vector engine) +// --------------------------------------------------------------------------- +def TileIR_VecFillOp : TileIR_Op<"vec_fill", [Pure, MemoryEffects<[MemWrite]>]> { + let summary = "Fill a buffer with a scalar value on the Vector engine"; + let description = [{ + Initializes a buffer by broadcasting a scalar value. Typically used to + reset running max/sum states or zero-fill accumulators. + + Example: + ```mlir + tile.vec_fill %buf, %ZERO_F32 {compute_unit=#tile.vector_core} + ``` + }]; + + let arguments = (ins + AnyType:$buf, + AnyType:$value, + OptionalAttr:$compute_unit + ); + + let assemblyFormat = "$buf `,` $value attr-dict `:` type($buf) `,` type($value)"; +} + +//===----------------------------------------------------------------------===// +// Compute Operations +//===----------------------------------------------------------------------===// + +// --------------------------------------------------------------------------- +// tile.cube_launch — Async matrix multiply on Cube engine +// --------------------------------------------------------------------------- +def TileIR_CubeLaunchOp : TileIR_Op<"cube_launch", [MemoryEffects<[MemWrite]>]> { + let summary = "Launch async matrix multiply on the Cube engine"; + let description = [{ + Asynchronously launches a matrix multiplication on the Cube engine. + The operands a/b are from L1, acc from L0C, stage_a/stage_b from L0A/L0B + (data staging), and dst is the GM output pointer. + + Returns a token that can be used with tile.cube_wait to synchronize. + + Example: + ```mlir + tile.cube_launch a(%q_l1) b(%k_l1) acc(%l0c0) + stage_a(%l0a0) stage_b(%l0b0) dst(%m1_slot) + {transpose_b=true, init=true, mma="128x128x128_f32_f16_f16"} + : (!tile.buf<...>, !tile.buf<...>, ..., !tt.ptr) + ``` + }]; + + let arguments = (ins + AnyType:$a, + AnyType:$b, + AnyType:$acc, + AnyType:$stage_a, + AnyType:$stage_b, + TT_Ptr:$dst, + OptionalAttr:$transpose_a, + OptionalAttr:$transpose_b, + OptionalAttr:$init, + OptionalAttr:$mma, + OptionalAttr:$comment + ); + + let assemblyFormat = [{ + `a` `(` $a `)` `b` `(` $b `)` `acc` `(` $acc `)` + `stage_a` `(` $stage_a `)` `stage_b` `(` $stage_b `)` `dst` `(` $dst `)` + attr-dict `:` `(` type($a) `,` type($b) `,` type($acc) `,` + type($stage_a) `,` type($stage_b) `,` type($dst) `)` + }]; +} + +// --------------------------------------------------------------------------- +// tile.cube_wait — Wait for Cube async operation completion +// --------------------------------------------------------------------------- +def TileIR_CubeWaitOp : TileIR_Op<"cube_wait", []> { + let summary = "Wait for a Cube async operation to complete"; + let description = [{ + Blocks until the Cube async operation identified by the token completes. + Used to drain the pipeline after cube_launch. + + Example: + ```mlir + tile.cube_wait %token + ``` + }]; + + let arguments = (ins); + + let assemblyFormat = "attr-dict"; +} + +// --------------------------------------------------------------------------- +// tile.reduce — Reduce a tensor along an axis +// --------------------------------------------------------------------------- +def TileIR_ReduceOp : TileIR_Op<"reduce", [Pure]> { + let summary = "Reduce a tensor along a specified axis"; + let description = [{ + Performs a reduction along a given axis using the specified arithmetic + operation. The scope attribute specifies which engine performs the reduce. + + Example: + ```mlir + %res = tile.reduce %tensor {axis=1, op=@arith.maximumf, scope=#tile.vector} + : !tile.tensor<[128,128],f32,#tile.UB> -> !tile.tensor<[128,1],f32,#tile.UB> + ``` + }]; + + let arguments = (ins + TileIR_TensorType:$src, + I64Attr:$axis, + FlatSymbolRefAttr:$op, + OptionalAttr:$scope + ); + + let results = (outs TileIR_TensorType:$result); + + let assemblyFormat = [{ + $src attr-dict `:` type($src) `->` type($result) + }]; +} + +// --------------------------------------------------------------------------- +// tile.elemwise — Element-wise operation +// --------------------------------------------------------------------------- +def TileIR_ElemwiseOp : TileIR_Op<"elemwise", [Pure]> { + let summary = "Element-wise operation (math function applied to tensor)"; + let description = [{ + Applies an element-wise operation (e.g., @math.exp) to input tensors. + The function symbol references an MLIR standard function. + + Example: + ```mlir + %res = tile.elemwise @math.exp(%arg) + ``` + }]; + + let arguments = (ins + FlatSymbolRefAttr:$func, + Variadic:$inputs + ); + + let results = (outs TileIR_TensorType:$result); + + let assemblyFormat = "$func `(` $inputs `)` attr-dict `:` functional-type($inputs, $result)"; +} + +// --------------------------------------------------------------------------- +// tile.broadcast — Broadcast a tensor to a larger shape +// --------------------------------------------------------------------------- +def TileIR_BroadcastOp : TileIR_Op<"broadcast", [Pure]> { + let summary = "Broadcast a tensor to a larger shape along an axis"; + let description = [{ + Broadcasts a tensor along a specified axis to match a target shape. + Used to expand (e.g., [128,1] → [128,128]) for element-wise operations. + + Example: + ```mlir + %bc = tile.broadcast %tensor {axis=1, shape=[128,128]} + : !tile.tensor<[128,1],f32,#tile.UB> -> !tile.tensor<[128,128],f32,#tile.UB> + ``` + }]; + + let arguments = (ins + TileIR_TensorType:$src, + I64Attr:$axis, + I64ArrayAttr:$shape + ); + + let results = (outs TileIR_TensorType:$result); + + let assemblyFormat = "$src attr-dict `:` type($src) `->` type($result)"; +} + +// --------------------------------------------------------------------------- +// tile.cast — Type conversion (e.g., f32 → f16) +// --------------------------------------------------------------------------- +def TileIR_CastOp : TileIR_Op<"cast", [Pure]> { + let summary = "Type conversion between element types"; + let description = [{ + Converts element types of a tensor (e.g., f32 → f16 for GM write). + + Example: + ```mlir + %f16 = tile.cast %f32 : !tile.tensor<[128,128],f32,#tile.UB> + -> !tile.tensor<[128,128],f16,#tile.UB> + ``` + }]; + + let arguments = (ins + TileIR_TensorType:$src + ); + + let results = (outs TileIR_TensorType:$result); + + let assemblyFormat = "$src attr-dict `:` type($src) `->` type($result)"; +} + +//===----------------------------------------------------------------------===// +// Tensor Conversion Operations +//===----------------------------------------------------------------------===// + +// --------------------------------------------------------------------------- +// tile.to_tensor — Convert a buffer to an immutable tensor view +// --------------------------------------------------------------------------- +def TileIR_ToTensorOp : TileIR_Op<"to_tensor", [Pure]> { + let summary = "Convert a buffer to a ranked tensor value"; + let description = [{ + Materializes a buffer as a standard ranked tensor so it can feed Triton + compute ops (tt.dot, arith.*, math.*, reduce, ...). The result is a builtin + RankedTensorType (not a !tile.* type), which is what those ops require. + + Example: + ```mlir + %t = tile.to_tensor %buf : !tile.buf<[128,128],f16,l0a> -> tensor<128x128xf16> + ``` + }]; + + let arguments = (ins + TileIR_BufType:$src + ); + + let results = (outs TT_Tensor:$result); + + let assemblyFormat = "$src attr-dict `:` type($src) `->` type($result)"; +} + +// --------------------------------------------------------------------------- +// tile.store_tensor — Store a standard tensor into a TileIR buffer +// --------------------------------------------------------------------------- +def TileIR_StoreTensorOp : TileIR_Op<"store_tensor", [Pure, MemoryEffects<[MemWrite]>]> { + let summary = "Store a tensor value into a buffer"; + let description = [{ + Writes a standard ranked tensor value back into a mutable TileIR buffer. + This is the inverse boundary of tile.to_tensor, whose result intentionally + stays a standard tensor so Triton/arith/math ops can consume it directly. + + Example: + ```mlir + tile.store_tensor %tensor → %buf : tensor<128x128xf32>, !tile.buf<[128,128],f32,#tile.UB> + ``` + }]; + + let arguments = (ins + TT_Tensor:$src, + TileIR_BufType:$dst + ); + + let assemblyFormat = "$src `->` $dst attr-dict `:` type($src) `,` type($dst)"; +} + +//===----------------------------------------------------------------------===// +// Synchronization Operations +//===----------------------------------------------------------------------===// + +// --------------------------------------------------------------------------- +// tile.set_flag — Set a cross-engine synchronization flag +// --------------------------------------------------------------------------- +def TileIR_SetFlagOp : TileIR_Op<"set_flag", []> { + let summary = "Set a cross-engine synchronization flag"; + let description = [{ + Signals that a producer pipe has completed its work, enabling the + consumer pipe to proceed. Maps to AscendC SetFlag(). + + Example: + ```mlir + tile.set_flag [producer=PIPE_MTE3, consumer=PIPE_MTE2, event=EVENT_ID2] + ``` + }]; + + let arguments = (ins + TileIR_Pipe:$producer, + TileIR_Pipe:$consumer, + TileIR_EventID:$event + ); + + let assemblyFormat = [{ + `[` $producer `,` $consumer `,` $event `]` attr-dict + }]; +} + +// --------------------------------------------------------------------------- +// tile.wait_flag — Wait for a cross-engine synchronization flag +// --------------------------------------------------------------------------- +def TileIR_WaitFlagOp : TileIR_Op<"wait_flag", []> { + let summary = "Wait for a cross-engine synchronization flag"; + let description = [{ + Blocks until the specified producer pipe has set the event flag. + Maps to AscendC WaitFlag(). + + Example: + ```mlir + tile.wait_flag [producer=PIPE_MTE2, consumer=PIPE_V, event=EVENT_ID3] + ``` + }]; + + let arguments = (ins + TileIR_Pipe:$producer, + TileIR_Pipe:$consumer, + TileIR_EventID:$event + ); + + let assemblyFormat = [{ + `[` $producer `,` $consumer `,` $event `]` attr-dict + }]; +} + +// --------------------------------------------------------------------------- +// tile.pipe_barrier — Intra-engine pipeline barrier +// --------------------------------------------------------------------------- +def TileIR_PipeBarrierOp : TileIR_Op<"pipe_barrier", []> { + let summary = "Intra-engine pipeline barrier (RAW hazard)"; + let description = [{ + Inserts a barrier within a single pipe/engine to ensure read-after-write + hazards are resolved. Maps to AscendC PipeBarrier(). + + Example: + ```mlir + tile.pipe_barrier [PIPE_V] + ``` + }]; + + let arguments = (ins + TileIR_Pipe:$pipe + ); + + let assemblyFormat = "`[` $pipe `]` attr-dict"; +} + +#endif // TILE_IR_OPS diff --git a/include/mlir-ext/Dialect/TileIR/IR/TileIRTypes.td b/include/mlir-ext/Dialect/TileIR/IR/TileIRTypes.td new file mode 100644 index 00000000..1f13e6cb --- /dev/null +++ b/include/mlir-ext/Dialect/TileIR/IR/TileIRTypes.td @@ -0,0 +1,87 @@ +//===-- TileIRTypes.td - TileIR type 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 types: !tile.buf (mutable buffer) and +// !tile.tensor (immutable tensor view), both parameterized by shape, +// element type, and memory space. +// +//===----------------------------------------------------------------------===// + +#ifndef TILE_IR_TYPES +#define TILE_IR_TYPES + +include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.td" +include "mlir-ext/Dialect/TileIR/IR/TileIRAttrDefs.td" + +include "mlir/IR/AttrTypeBase.td" +include "mlir/IR/EnumAttr.td" + +//===----------------------------------------------------------------------===// +// Type base class +//===----------------------------------------------------------------------===// + +class TileIR_Type traits = []> + : TypeDef { + let mnemonic = typeMnemonic; +} + +//===----------------------------------------------------------------------===// +// !tile.buf<[shape], dtype, space> — mutable buffer type +//===----------------------------------------------------------------------===// + +def TileIR_BufType : TileIR_Type<"Buf", "buf"> { + let summary = "TileIR mutable buffer type"; + let description = [{ + A !tile.buf represents a mutable buffer allocated in a specific memory space + (GM, L1, L0A, L0B, L0C, or UB). It carries shape, element type, and + memory space — enabling the IR to track physical buffer locations. + + Example: + ```mlir + !tile.buf<[128,128], f16, #tile.L1> + !tile.buf<[2,128,128], f16, #tile.L0A> + ``` + }]; + + let parameters = (ins + ArrayRefParameter<"int64_t">:$shape, + "mlir::Type":$elementType, + EnumParameter:$memorySpace + ); + + let assemblyFormat = "`<` `[` $shape `]` `,` $elementType `,` $memorySpace `>`"; +} + +//===----------------------------------------------------------------------===// +// !tile.tensor<[shape], dtype, space> — immutable tensor view type +//===----------------------------------------------------------------------===// + +def TileIR_TensorType : TileIR_Type<"Tensor", "tensor"> { + let summary = "TileIR immutable tensor view type"; + let description = [{ + A !tile.tensor represents an immutable view of data in a specific memory + space, typically used as the result of tile.to_tensor or tile.load. + Unlike !tile.buf, it is read-only and is used as input to element-wise + and reduce operations. + + Example: + ```mlir + !tile.tensor<[128,128], f32, #tile.UB> + ``` + }]; + + let parameters = (ins + ArrayRefParameter<"int64_t">:$shape, + "mlir::Type":$elementType, + EnumParameter:$memorySpace + ); + + let assemblyFormat = "`<` `[` $shape `]` `,` $elementType `,` $memorySpace `>`"; +} + +#endif // TILE_IR_TYPES \ No newline at end of file diff --git a/lib/Conversion/TritonToLinalgIncubated/BlockPtrAnalysis.cpp b/lib/Conversion/TritonToLinalgIncubated/BlockPtrAnalysis.cpp index 47520124..ff321993 100644 --- a/lib/Conversion/TritonToLinalgIncubated/BlockPtrAnalysis.cpp +++ b/lib/Conversion/TritonToLinalgIncubated/BlockPtrAnalysis.cpp @@ -55,6 +55,7 @@ #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" #include "llvm/Support/FormatVariadic.h" +#include "llvm/Config/llvm-config.h" #include #include @@ -1350,8 +1351,15 @@ void BlockDataParser::rewriteCustomOp( resultTypes.emplace_back(ty); } auto newCustomOp = +#if LLVM_VERSION_MAJOR >= 22 rewriter.create(loc, resultTypes, op.getName(), newInputs, newOutputs, adaptor.getTempBuffers()); +#else + rewriter.create(loc, resultTypes, + rewriter.getStringAttr(op.getName()), + newInputs, + newOutputs); +#endif auto operandSegmentSizesAttr = newCustomOp->getAttr("operandSegmentSizes"); newCustomOp->setAttrs(op->getAttrs()); newCustomOp->setAttr("operandSegmentSizes", operandSegmentSizesAttr); diff --git a/lib/Conversion/TritonToLinalgIncubated/DescriptorConverter.cpp b/lib/Conversion/TritonToLinalgIncubated/DescriptorConverter.cpp index bbaced7a..8852087a 100644 --- a/lib/Conversion/TritonToLinalgIncubated/DescriptorConverter.cpp +++ b/lib/Conversion/TritonToLinalgIncubated/DescriptorConverter.cpp @@ -45,6 +45,7 @@ #include "mlir/IR/OpDefinition.h" #include "mlir/IR/ValueRange.h" #include "mlir/Transforms/DialectConversion.h" +#include "llvm/Config/llvm-config.h" namespace DescriptorConverter { using namespace mlir; @@ -72,7 +73,9 @@ Descriptor unpackDescriptor(TensorDescType type, Value desc, res.strides.push_back(rewriter.createOrFold( makeDescOp.getLoc(), rewriter.getI64Type(), st)); } + #if LLVM_VERSION_MAJOR >= 22 res.padding = makeDescOp.getPaddingAttr(); +#endif return res; } @@ -125,7 +128,9 @@ LogicalResult DescriptorLoadConverter::matchAndRewrite( ); // 3. replace tt.load auto boundaryCheck = getFullBoundaryCheckAttr(rewriter, blockShape); +#if LLVM_VERSION_MAJOR >= 22 triton::PaddingOptionAttr padding = desc.padding; +#endif auto cache = triton::CacheModifierAttr::get(rewriter.getContext(), triton::CacheModifier::NONE); auto evict = triton::EvictionPolicyAttr::get(rewriter.getContext(), @@ -143,7 +148,13 @@ LogicalResult DescriptorLoadConverter::matchAndRewrite( loc, descTy.getSignlessBlockType(), tensorPtr, Value(), // mask Value(), // other - boundaryCheck, padding, cache, evict, isVolatile); + boundaryCheck, +#if LLVM_VERSION_MAJOR >= 22 + padding, +#else + triton::PaddingOptionAttr(), +#endif + cache, evict, isVolatile); rewriter.replaceOp(op, newLoad.getResult()); diff --git a/lib/Conversion/TritonToLinalgIncubated/TritonOpConverter.cpp b/lib/Conversion/TritonToLinalgIncubated/TritonOpConverter.cpp index 512874d6..9384a8b8 100644 --- a/lib/Conversion/TritonToLinalgIncubated/TritonOpConverter.cpp +++ b/lib/Conversion/TritonToLinalgIncubated/TritonOpConverter.cpp @@ -59,6 +59,8 @@ #include "bishengir/Dialect/HIVM/IR/HIVM.h" #endif +#include "llvm/Config/llvm-config.h" + namespace TTOpConverters { using namespace mlir; using namespace triton; @@ -875,6 +877,7 @@ SplatConverter::matchAndRewrite(triton::SplatOp op, OpAdaptor adaptor, return success(); } +#if LLVM_VERSION_MAJOR >= 22 LogicalResult UnsplatConverter::matchAndRewrite(triton::UnsplatOp op, OpAdaptor adaptor, ConversionPatternRewriter &rewriter) const { @@ -898,6 +901,7 @@ UnsplatConverter::matchAndRewrite(triton::UnsplatOp op, OpAdaptor adaptor, rewriter.replaceOp(op, extractOp.getResult()); return success(); } +#endif LogicalResult ReshapeConverter::matchAndRewrite(triton::ReshapeOp op, OpAdaptor adaptor, @@ -2344,6 +2348,7 @@ DotScaledConverter::matchAndRewrite(triton::DotScaledOp op, OpAdaptor adaptor, #endif RankedTensorType dstType = cast(op.getType()); +#if LLVM_VERSION_MAJOR >= 22 auto lhsElemType = op.getAElemType(); auto rhsElemType = op.getBElemType(); @@ -2404,6 +2409,7 @@ DotScaledConverter::matchAndRewrite(triton::DotScaledOp op, OpAdaptor adaptor, rewriter.replaceOp(op, finalResult); return success(); } +#endif if (!lhsScale) { return op.emitError("lhsScale is required for non-FP8 input"); @@ -2425,7 +2431,11 @@ DotScaledConverter::matchAndRewrite(triton::DotScaledOp op, OpAdaptor adaptor, Type bf16Ty = rewriter.getBF16Type(); Type fp16Ty = rewriter.getF16Type(); Type fp32Ty = rewriter.getF32Type(); + #if LLVM_VERSION_MAJOR >= 22 bool fastMath = op.getFastMath(); +#else + bool fastMath = false; +#endif auto createNanSplat = [&](RankedTensorType tensorTy) -> Value { auto floatTy = cast(tensorTy.getElementType()); diff --git a/lib/Conversion/TritonToLinalgIncubated/TritonToLinalgPass.cpp b/lib/Conversion/TritonToLinalgIncubated/TritonToLinalgPass.cpp index e109349a..36d0d6b5 100644 --- a/lib/Conversion/TritonToLinalgIncubated/TritonToLinalgPass.cpp +++ b/lib/Conversion/TritonToLinalgIncubated/TritonToLinalgPass.cpp @@ -77,6 +77,7 @@ #include "llvm/Support/Casting.h" #include "llvm/Support/Debug.h" #include "llvm/Support/ErrorHandling.h" +#include "llvm/Config/llvm-config.h" #include "llvm/Support/LogicalResult.h" #ifdef __TLE_STRUCT__ @@ -626,7 +627,9 @@ void TritonToLinalgIncubatedPass::populateTritonToLinalgConversionPatterns( patterns.add(patterns.getContext()); patterns.add(patterns.getContext()); patterns.add(patterns.getContext()); +#if LLVM_VERSION_MAJOR >= 22 patterns.add(patterns.getContext()); +#endif patterns.add(patterns.getContext()); patterns.add(patterns.getContext()); // reduce converters @@ -1130,10 +1133,17 @@ void TritonToLinalgIncubatedPass::runOnOperation() { MemRefType syncBlockLockArgType = MemRefType::get(SmallVector(1, ShapedType::kDynamic), IntegerType::get(context, 8)); + #if LLVM_VERSION_MAJOR >= 22 llvm::LogicalResult syncBlockLockArg = func.insertArgument(syncBlockLockArgIdx, // argIndex syncBlockLockArgType, // argType nullptr, func->getLoc()); // dicAttr + (void)syncBlockLockArg; +#else + func.insertArgument(syncBlockLockArgIdx, // argIndex + syncBlockLockArgType, // argType + nullptr, func->getLoc()); // dicAttr +#endif func->setAttr("SyncBlockLockArgIdx", IntegerAttr::get(IntegerType::get(&getContext(), 64), 0)); // 64: 64位整型 @@ -1145,10 +1155,17 @@ void TritonToLinalgIncubatedPass::runOnOperation() { NamedAttribute workspaceArgAttr(StringAttr::get(context, "workspace"), UnitAttr::get(context)); + #if LLVM_VERSION_MAJOR >= 22 llvm::LogicalResult workspaceArg = func.insertArgument(/*argIndex*/ workspaceArgIdx, /*argType*/ workspaceArgType, /*dicAttr*/ nullptr, func->getLoc()); + (void)workspaceArg; +#else + func.insertArgument(/*argIndex*/ workspaceArgIdx, + /*argType*/ workspaceArgType, + /*dicAttr*/ nullptr, func->getLoc()); +#endif func->setAttr("WorkspaceArgIdx", IntegerAttr::get(IntegerType::get(&getContext(), 64), 1)); // 64: 64位整型 diff --git a/lib/Dialect/CMakeLists.txt b/lib/Dialect/CMakeLists.txt index 822eeefa..6e147475 100644 --- a/lib/Dialect/CMakeLists.txt +++ b/lib/Dialect/CMakeLists.txt @@ -2,6 +2,7 @@ add_subdirectory(TritonTilingExt) add_subdirectory(TPtr) add_subdirectory(MathExt) add_subdirectory(TritonStructured) +add_subdirectory(TileIR) if (FLIR_BUILD_INCUBATED) add_subdirectory(TritonStructuredIncubated) endif() diff --git a/lib/Dialect/TileIR/CMakeLists.txt b/lib/Dialect/TileIR/CMakeLists.txt new file mode 100644 index 00000000..7d59dce8 --- /dev/null +++ b/lib/Dialect/TileIR/CMakeLists.txt @@ -0,0 +1 @@ +add_subdirectory(IR) \ No newline at end of file diff --git a/lib/Dialect/TileIR/IR/CMakeLists.txt b/lib/Dialect/TileIR/IR/CMakeLists.txt new file mode 100644 index 00000000..9849c841 --- /dev/null +++ b/lib/Dialect/TileIR/IR/CMakeLists.txt @@ -0,0 +1,15 @@ +add_triton_library(TileIRIR + TileIRAttrs.cpp + TileIRDialect.cpp + TileIROps.cpp + + DEPENDS + TileIRTableGen + TileIRAttrDefsIncGen + TileIRTypesIncGen + + LINK_LIBS PUBLIC + MLIRIR + MLIRLLVMDialect + TritonIR +) \ No newline at end of file diff --git a/lib/Dialect/TileIR/IR/TileIRAttrs.cpp b/lib/Dialect/TileIR/IR/TileIRAttrs.cpp new file mode 100644 index 00000000..e372d199 --- /dev/null +++ b/lib/Dialect/TileIR/IR/TileIRAttrs.cpp @@ -0,0 +1,16 @@ +//===- TileIRAttrs.cpp - TileIR attribute implementations -----------------===// +// +// 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 implements TileIR attribute custom methods. +// +//===----------------------------------------------------------------------===// + +#include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.h" + +using namespace mlir; +using namespace mlir::triton::tile; \ No newline at end of file diff --git a/lib/Dialect/TileIR/IR/TileIRDialect.cpp b/lib/Dialect/TileIR/IR/TileIRDialect.cpp new file mode 100644 index 00000000..10888dfb --- /dev/null +++ b/lib/Dialect/TileIR/IR/TileIRDialect.cpp @@ -0,0 +1,52 @@ +//===- TileIRDialect.cpp - TileIR Dialect registration --------------------===// +// +// 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 implements the TileIR dialect, registering all operations, +// attributes, and types. +// +//===----------------------------------------------------------------------===// + +#include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.h" + +#include "mlir/Dialect/LLVMIR/LLVMDialect.h" +#include "mlir/IR/Builders.h" +#include "mlir/IR/BuiltinTypes.h" +#include "mlir/IR/DialectImplementation.h" +#include "mlir/IR/MLIRContext.h" +#include "mlir/IR/Operation.h" + +#include "llvm/ADT/StringExtras.h" +#include "llvm/ADT/TypeSwitch.h" +#include "llvm/Support/SourceMgr.h" + +using namespace mlir; +using namespace mlir::triton::tile; + +void TileIRDialect::initialize() { + addOperations< +#define GET_OP_LIST +#include "mlir-ext/Dialect/TileIR/IR/TileIROps.cpp.inc" + >(); + addAttributes< +#define GET_ATTRDEF_LIST +#include "mlir-ext/Dialect/TileIR/IR/TileIROpsAttrDefs.cpp.inc" + >(); + addTypes< +#define GET_TYPEDEF_LIST +#include "mlir-ext/Dialect/TileIR/IR/TileIRTypes.cpp.inc" + >(); +} + +#include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.cpp.inc" +#include "mlir-ext/Dialect/TileIR/IR/TileIREnums.cpp.inc" +#define GET_ATTRDEF_CLASSES +#include "mlir-ext/Dialect/TileIR/IR/TileIROpsAttrDefs.cpp.inc" +#define GET_TYPEDEF_CLASSES +#include "mlir-ext/Dialect/TileIR/IR/TileIRTypes.cpp.inc" +#define GET_OP_CLASSES +#include "mlir-ext/Dialect/TileIR/IR/TileIROps.cpp.inc" \ No newline at end of file diff --git a/lib/Dialect/TileIR/IR/TileIROps.cpp b/lib/Dialect/TileIR/IR/TileIROps.cpp new file mode 100644 index 00000000..1aea5478 --- /dev/null +++ b/lib/Dialect/TileIR/IR/TileIROps.cpp @@ -0,0 +1,169 @@ +//===- TileIROps.cpp - TileIR operation implementations -------------------===// +// +// 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 implements TileIR operations with their verification logic. +// +//===----------------------------------------------------------------------===// + +#include "mlir-ext/Dialect/TileIR/IR/TileIRDialect.h" + +#include "mlir/IR/Builders.h" +#include "mlir/IR/BuiltinTypes.h" +#include "mlir/IR/OpImplementation.h" +#include "mlir/Interfaces/SideEffectInterfaces.h" + +#include "triton/Tools/Sys/GetEnv.hpp" +#include "llvm/ADT/STLExtras.h" +#include + +using namespace mlir; +using namespace mlir::triton::tile; + +//===----------------------------------------------------------------------===// +// AllocOp +//===----------------------------------------------------------------------===// + +LogicalResult AllocOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// SubViewOp +//===----------------------------------------------------------------------===// + +LogicalResult SubViewOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// GmOffsetOp +//===----------------------------------------------------------------------===// + +LogicalResult GmOffsetOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// CopyOp +//===----------------------------------------------------------------------===// + +LogicalResult CopyOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// LoadOp +//===----------------------------------------------------------------------===// + +LogicalResult LoadOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// StoreOp +//===----------------------------------------------------------------------===// + +LogicalResult StoreOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// VecFillOp +//===----------------------------------------------------------------------===// + +LogicalResult VecFillOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// CubeLaunchOp +//===----------------------------------------------------------------------===// + +LogicalResult CubeLaunchOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// CubeWaitOp +//===----------------------------------------------------------------------===// + +LogicalResult CubeWaitOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// ReduceOp +//===----------------------------------------------------------------------===// + +LogicalResult ReduceOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// ElemwiseOp +//===----------------------------------------------------------------------===// + +LogicalResult ElemwiseOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// BroadcastOp +//===----------------------------------------------------------------------===// + +LogicalResult BroadcastOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// CastOp +//===----------------------------------------------------------------------===// + +LogicalResult CastOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// ToTensorOp +//===----------------------------------------------------------------------===// + +LogicalResult ToTensorOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// StoreTensorOp +//===----------------------------------------------------------------------===// + +LogicalResult StoreTensorOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// SetFlagOp +//===----------------------------------------------------------------------===// + +LogicalResult SetFlagOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// WaitFlagOp +//===----------------------------------------------------------------------===// + +LogicalResult WaitFlagOp::verify() { + return success(); +} + +//===----------------------------------------------------------------------===// +// PipeBarrierOp +//===----------------------------------------------------------------------===// + +LogicalResult PipeBarrierOp::verify() { + return success(); +} \ No newline at end of file