‘tt’ Dialect

The Triton IR in MLIR

Triton Dialect.

Dependent Dialects:

  • Arith:

    • addf, addi, andi, cmpf, cmpi, divf, fptosi, …

  • Math:

    • exp, sin, cos, log, …

  • StructuredControlFlow:

    • for, if, while, yield, condition

  • ControlFlow:

    • br, cond_br

[TOC]

Type constraints

ptr

Types

PackedFloat4E2M1Type

Two E2M1 floating-point values packed in one byte

Syntax: !tt.fp4e2m1x2

The low nibble holds the first value and the high nibble the second. Tensor dimensions and pointer offsets count packed pairs.

PointerType

Pointer type (::mlir::triton::PointerType) in Triton IR type system

Pointer type in Triton IR type system. Pointer types may only point to scalar element types.

Parameters:

Parameter

C++ type

Description

pointeeType

Type

addressSpace

PtrAddrSpace

TensorDescType

Tensor descriptor type (::mlir::triton::TensorDescType) in Triton IR type system

A portable abstraction for TMA descriptors. This is the base tensor descriptor type for tiled tensor memory access.

Shape and elementType describe the block dimensions and data type. The optional sharedLayout attribute carries the shared memory encoding (e.g. swizzle pattern) that is assigned during lowering.

For specialized access patterns like im2col, see TensorDescIm2ColType in the TritonNvidiaGPU dialect.

Parameters:

Parameter

C++ type

Description

shape

::llvm::ArrayRef<int64_t>

elementType

Type

sharedLayout

Attribute