TritonOps¶

`tt.call` (triton::CallOp)¶

Call operation

Syntax:

operation ::= `tt.call` $callee `(` $operands `)` attr-dict `:` functional-type($operands, results)

The tt.call operation represents a direct call to a function that is within the same symbol scope as the call. The operands and result types of the call must match the specified function type. The callee is encoded as a symbol reference attribute named “callee”.

Example:

%2 = tt.call @my_add(%0, %1) : (f32, f32) -> f32

Traits: TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: CallOpInterface, SymbolUserOpInterface

Attributes:¶

Attribute	MLIR Type	Description
`callee`	::mlir::FlatSymbolRefAttr	flat symbol reference attribute

Operands:¶

Operand	Description
`operands`	variadic of any type

Results:¶

Result	Description
«unnamed»	variadic of any type

`tt.func` (triton::FuncOp)¶

An operation with a name containing a single SSACFG region

Operations within the function cannot implicitly capture values defined outside of the function, i.e. Functions are IsolatedFromAbove. All external references must use function arguments or attributes that establish a symbolic connection (e.g. symbols referenced by name via a string attribute like SymbolRefAttr). An external function declaration (used when referring to a function declared in some other module) has no body. While the MLIR textual form provides a nice inline syntax for function arguments, they are internally represented as “block arguments” to the first block in the region.

Only dialect attribute names may be specified in the attribute dictionaries for function arguments, results, or the function itself.

Example:

// External function definitions.
tt.func @abort()
tt.func @scribble(i32, i64, memref<? x 128 x f32, #layout_map0>) -> f64

// A function that returns its argument twice:
tt.func @count(%x: i64) -> (i64, i64)
  attributes {fruit: "banana"} {
  return %x, %x: i64, i64
}

// A function with an argument attribute
tt.func @example_fn_arg(%x: i32 {swift.self = unit})

// A function with a result attribute
tt.func @example_fn_result() -> (f64 {dialectName.attrName = 0 : i64})

// A function with an attribute
tt.func @example_fn_attr() attributes {dialectName.attrName = false}

Traits: AffineScope, AutomaticAllocationScope, IsolatedFromAbove, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: CallableOpInterface, FunctionOpInterface, OpAsmOpInterface, Symbol

Attributes:¶

Attribute	MLIR Type	Description
`sym_name`	::mlir::StringAttr	string attribute
`function_type`	::mlir::TypeAttr	type attribute of function type
`sym_visibility`	::mlir::StringAttr	string attribute
`arg_attrs`	::mlir::ArrayAttr	Array of dictionary attributes
`res_attrs`	::mlir::ArrayAttr	Array of dictionary attributes

`tt.return` (triton::ReturnOp)¶

Function return operation

Syntax:

operation ::= `tt.return` attr-dict ($operands^ `:` type($operands))?

The tt.return operation represents a return operation within a function. The operation takes variable number of operands and produces no results. The operand number and types must match the signature of the function that contains the operation.

Example:

tt.func @foo() : (i32, f8) {
  ...
  tt.return %0, %1 : i32, f8
}

Traits: AlwaysSpeculatableImplTrait, HasParent, ReturnLike, TensorSizeTrait, Terminator, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`operands`	variadic of any type

`tt.addptr` (triton::AddPtrOp)¶

Syntax:

operation ::= `tt.addptr` $ptr `,` $offset attr-dict `:` type($result) `,` type($offset)

Traits: AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`ptr`	ptr or tensor of ptr values
`offset`	integer or tensor of integer values

Results:¶

Result	Description
`result`	ptr or tensor of ptr values

`tt.advance` (triton::AdvanceOp)¶

Advance a tensor pointer by offsets

Syntax:

operation ::= `tt.advance` $ptr `,` `[` $offsets `]` attr-dict `:` type($result)

Traits: AlwaysSpeculatableImplTrait, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`ptr`	ptr
`offsets`	variadic of 32-bit signless integer

Results:¶

Result	Description
`result`	ptr

`tt.assert` (triton::AssertOp)¶

Device-side assert, as in CUDA for correctness checking

Syntax:

operation ::= `tt.assert` $condition `,` $message `,` $file `,` $func `,` $line attr-dict `:` type($condition)

tt.assert takes a condition tensor, a message string, a file string, a function string, and a line number. If the condition is false, the message is printed, and the program is aborted.

Traits: TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Attributes:¶

Attribute	MLIR Type	Description
`message`	::mlir::StringAttr	string attribute
`file`	::mlir::StringAttr	string attribute
`func`	::mlir::StringAttr	string attribute
`line`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`condition`	tensor of floating-point values or tensor of integer values or tensor of ptr values

`tt.atomic_cas` (triton::AtomicCASOp)¶

Atomic cas

compare $cmp with data $old at location $ptr,

if $old == $cmp, store $val to $ptr,

else store $old to $ptr,

return $old

Traits: SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}, MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Attributes:¶

Attribute	MLIR Type	Description
`sem`	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4 {{% markdown %}}Enum cases: * relaxed (`RELAXED`) * acquire (`ACQUIRE`) * release (`RELEASE`) * acq_rel (`ACQUIRE_RELEASE`){{% /markdown %}}
`scope`	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3 {{% markdown %}}Enum cases: * gpu (`GPU`) * cta (`CTA`) * sys (`SYSTEM`){{% /markdown %}}

Operands:¶

Operand	Description
`ptr`	ptr or tensor of ptr values
`cmp`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr
`val`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.atomic_rmw` (triton::AtomicRMWOp)¶

Atomic rmw

load data at $ptr, do $rmw_op with $val, and store result to $ptr.

return old value at $ptr

Traits: SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}, MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Attributes:¶

Attribute	MLIR Type	Description
`atomic_rmw_op`	::mlir::triton::RMWOpAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7, 8, 9, 10 {{% markdown %}}Enum cases: * and (`AND`) * or (`OR`) * xor (`XOR`) * add (`ADD`) * fadd (`FADD`) * max (`MAX`) * min (`MIN`) * umax (`UMAX`) * umin (`UMIN`) * exch (`XCHG`){{% /markdown %}}
`sem`	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4 {{% markdown %}}Enum cases: * relaxed (`RELAXED`) * acquire (`ACQUIRE`) * release (`RELEASE`) * acq_rel (`ACQUIRE_RELEASE`){{% /markdown %}}
`scope`	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3 {{% markdown %}}Enum cases: * gpu (`GPU`) * cta (`CTA`) * sys (`SYSTEM`){{% /markdown %}}

Operands:¶

Operand	Description
`ptr`	ptr or tensor of ptr values
`val`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr
`mask`	1-bit signless integer or tensor of 1-bit signless integer values

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.bitcast` (triton::BitcastOp)¶

Cast between types of the same bitwidth

Syntax:

operation ::= `tt.bitcast` $from attr-dict `:` type($from) `->` type($result)

Traits: AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`from`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.broadcast` (triton::BroadcastOp)¶

Broadcast. No left-padding as of now.

Syntax:

operation ::= `tt.broadcast` $src attr-dict `:` functional-type(operands, results)

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, SameOperandsAndResultEncoding, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`src`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.cat` (triton::CatOp)¶

Concatenate 2 tensors

Syntax:

operation ::= `tt.cat` $lhs `,` $rhs attr-dict `:` functional-type(operands, results)

Traits: SameOperandsAndResultElementType, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`lhs`	tensor of floating-point values or tensor of integer values or tensor of ptr values
`rhs`	tensor of floating-point values or tensor of integer values or tensor of ptr values

Results:¶

Result	Description
`result`	tensor of floating-point values or tensor of integer values or tensor of ptr values

`tt.dot` (triton::DotOp)¶

Dot

Syntax:

operation ::= `tt.dot` $a`,` $b`,` $c attr-dict `:` type($a) `*` type($b) `->` type($d)

$d = matrix_multiply($a, $b) + $c

Traits: AlwaysSpeculatableImplTrait, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`allowTF32`	::mlir::BoolAttr	bool attribute
`maxNumImpreciseAcc`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`a`	tensor of floating-point values or tensor of integer values
`b`	tensor of floating-point values or tensor of integer values
`c`	tensor of floating-point values or tensor of integer values

Results:¶

Result	Description
`d`	tensor of floating-point values or tensor of integer values

`tt.elementwise_inline_asm` (triton::ElementwiseInlineAsmOp)¶

Inline assembly applying elementwise operation to a group of packed element.

Syntax:

operation ::= `tt.elementwise_inline_asm` $asm_string attr-dict ($args^ `:` type($args))? `->` type($result)

This will apply the given in inline assembly to packed_element number of elements of the inputs. The elements packed together is unknown and will depend on the backend implementation.

Traits: Elementwise, SameOperandsAndResultEncoding, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface

Attributes:¶

Attribute	MLIR Type	Description
`asm_string`	::mlir::StringAttr	string attribute
`constraints`	::mlir::StringAttr	string attribute
`pure`	::mlir::BoolAttr	bool attribute
`packed_element`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`args`	variadic of floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.expand_dims` (triton::ExpandDimsOp)¶

Expand_dims

Syntax:

operation ::= `tt.expand_dims` $src attr-dict `:` functional-type(operands, results)

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`axis`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`src`	tensor of floating-point values or tensor of integer values or tensor of ptr values

Results:¶

Result	Description
`result`	tensor of floating-point values or tensor of integer values or tensor of ptr values

`tt.extern_elementwise` (triton::ExternElementwiseOp)¶

Syntax:

operation ::= `tt.extern_elementwise` operands attr-dict `:` functional-type(operands, $result)

call an external function $symbol implemented in $libpath/$libname with $args return $libpath/$libname:$symbol($args…)

Traits: Elementwise, SameOperandsAndResultEncoding, SameVariadicOperandSize, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface

Attributes:¶

Attribute	MLIR Type	Description
`libname`	::mlir::StringAttr	string attribute
`libpath`	::mlir::StringAttr	string attribute
`symbol`	::mlir::StringAttr	string attribute
`pure`	::mlir::BoolAttr	bool attribute

Operands:¶

Operand	Description
`args`	variadic of floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.fp_to_fp` (triton::FpToFpOp)¶

Floating point casting for custom types

Syntax:

operation ::= `tt.fp_to_fp` $from attr-dict `:` type($from) `->` type($result)

Floating point casting for custom types (F8).

F8 <-> FP16, BF16, FP32, FP64

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`from`	floating-point or tensor of floating-point values

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values

`tt.get_num_programs` (triton::GetNumProgramsOp)¶

Syntax:

operation ::= `tt.get_num_programs` attr-dict `:` type($result)

Traits: AlwaysSpeculatableImplTrait, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`axis`	::mlir::IntegerAttr	32-bit signless integer attribute

Results:¶

Result	Description
`result`	32-bit signless integer

`tt.get_program_id` (triton::GetProgramIdOp)¶

Syntax:

operation ::= `tt.get_program_id` $axis attr-dict `:` type($result)

Traits: AlwaysSpeculatableImplTrait, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`axis`	::mlir::triton::ProgramIDDimAttr	allowed 32-bit signless integer cases: 0, 1, 2 {{% markdown %}}Enum cases: * x (`X`) * y (`Y`) * z (`Z`){{% /markdown %}}

Results:¶

Result	Description
`result`	32-bit signless integer

`tt.int_to_ptr` (triton::IntToPtrOp)¶

Cast int64 to pointer

Syntax:

operation ::= `tt.int_to_ptr` $from attr-dict `:` type($from) `->` type($result)

Traits: AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`from`	64-bit signless integer or tensor of 64-bit signless integer values

Results:¶

Result	Description
`result`	ptr or tensor of ptr values

`tt.load` (triton::LoadOp)¶

Load from a tensor of pointers or from a tensor pointer

Traits: AttrSizedOperandSegments, SameLoadStoreOperandsAndResultEncoding, SameLoadStoreOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface

Attributes:¶

Attribute	MLIR Type	Description
`boundaryCheck`	::mlir::DenseI32ArrayAttr	i32 dense array attribute
`padding`	::mlir::triton::PaddingOptionAttr	allowed 32-bit signless integer cases: 1, 2 {{% markdown %}}Enum cases: * zero (`PAD_ZERO`) * nan (`PAD_NAN`){{% /markdown %}}
`cache`	::mlir::triton::CacheModifierAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6 {{% markdown %}}Enum cases: * none (`NONE`) * ca (`CA`) * cg (`CG`) * wb (`WB`) * cs (`CS`) * wt (`WT`){{% /markdown %}}
`evict`	::mlir::triton::EvictionPolicyAttr	allowed 32-bit signless integer cases: 1, 2, 3 {{% markdown %}}Enum cases: * evict_normal (`NORMAL`) * evict_first (`EVICT_FIRST`) * evict_last (`EVICT_LAST`){{% /markdown %}}
`isVolatile`	::mlir::BoolAttr	bool attribute

Operands:¶

Operand	Description
`ptr`	ptr or tensor of ptr values or ptr
`mask`	1-bit signless integer or tensor of 1-bit signless integer values
`other`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.make_range` (triton::MakeRangeOp)¶

Make range

Syntax:

operation ::= `tt.make_range` attr-dict `:` type($result)

Returns an 1D int32 tensor.

Values span from $start to $end (exclusive), with step = 1

Traits: AlwaysSpeculatableImplTrait, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`start`	::mlir::IntegerAttr	32-bit signless integer attribute
`end`	::mlir::IntegerAttr	32-bit signless integer attribute

Results:¶

Result	Description
`result`	tensor of integer values

`tt.make_tensor_ptr` (triton::MakeTensorPtrOp)¶

Make a tensor pointer type with meta information of the parent tensor and the block specified

Syntax:

operation ::= `tt.make_tensor_ptr` $base `,` `[` $shape `]` `,` `[` $strides `]` `,` `[` $offsets `]` attr-dict `:` type($result)

tt.make_tensor_ptr takes both meta information of the parent tensor and the block tensor, then it returns a pointer to the block tensor, e.g. returns a type of tt.ptr<tensor<8x8xf16>>.

Traits: AlwaysSpeculatableImplTrait, SameVariadicOperandSize, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`order`	::mlir::DenseI32ArrayAttr	i32 dense array attribute

Operands:¶

Operand	Description
`base`	ptr
`shape`	variadic of 64-bit signless integer
`strides`	variadic of 64-bit signless integer
`offsets`	variadic of 32-bit signless integer

Results:¶

Result	Description
`result`	ptr

`tt.print` (triton::PrintOp)¶

Device-side print, as in CUDA for debugging

Syntax:

operation ::= `tt.print` $prefix attr-dict (`:` $args^ `:` type($args))?

tt.print takes a literal string prefix and an arbitrary number of scalar or tensor arguments that should be printed. format are generated automatically from the arguments.

Traits: TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Attributes:¶

Attribute	MLIR Type	Description
`prefix`	::mlir::StringAttr	string attribute

Operands:¶

Operand	Description
`args`	variadic of floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.ptr_to_int` (triton::PtrToIntOp)¶

Cast pointer to int64

Syntax:

operation ::= `tt.ptr_to_int` $from attr-dict `:` type($from) `->` type($result)

Traits: AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`from`	ptr or tensor of ptr values

Results:¶

Result	Description
`result`	64-bit signless integer or tensor of 64-bit signless integer values

`tt.reduce` (triton::ReduceOp)¶

Reduction using generic combination algorithm

Traits: AlwaysSpeculatableImplTrait, SameOperandsEncoding, SingleBlock, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`axis`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`operands`	variadic of tensor of floating-point values or tensor of integer values or tensor of ptr values

Results:¶

Result	Description
`result`	variadic of floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

`tt.reduce.return` (triton::ReduceReturnOp)¶

Terminator for reduce operator

Syntax:

operation ::= `tt.reduce.return` $result attr-dict `:` type($result)

Traits: AlwaysSpeculatableImplTrait, HasParent, ReturnLike, TensorSizeTrait, Terminator, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`result`	variadic of any type

`tt.reshape` (triton::ReshapeOp)¶

Reinterpret a tensor to a different shape. It may change elements order if the attribute is set.

Syntax:

operation ::= `tt.reshape` $src attr-dict `:` type($src) `->` type($result)

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`allow_reorder`	::mlir::BoolAttr	bool attribute

Operands:¶

Operand	Description
`src`	tensor of floating-point values or tensor of integer values or tensor of ptr values

Results:¶

Result	Description
`result`	tensor of floating-point values or tensor of integer values or tensor of ptr values

`tt.scan` (triton::ScanOp)¶

Reduction using generic combination algorithm

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, SameOperandsAndResultEncoding, SingleBlock, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:¶

Attribute	MLIR Type	Description
`axis`	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:¶

Operand	Description
`operands`	variadic of tensor of floating-point values or tensor of integer values or tensor of ptr values

Results:¶

Result	Description
`result`	variadic of tensor of floating-point values or tensor of integer values or tensor of ptr values

`tt.scan.return` (triton::ScanReturnOp)¶

Terminator for scan operator

Syntax:

operation ::= `tt.scan.return` $result attr-dict `:` type($result)

Traits: AlwaysSpeculatableImplTrait, HasParent, ReturnLike, TensorSizeTrait, Terminator, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface), RegionBranchTerminatorOpInterface

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`result`	variadic of any type

`tt.splat` (triton::SplatOp)¶

Splat

Syntax:

operation ::= `tt.splat` $src attr-dict `:` functional-type(operands, results)

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, SameOperandsAndResultEncoding, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`src`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr

Results:¶

Result	Description
`result`	tensor of floating-point values or tensor of integer values or tensor of ptr values

`tt.store` (triton::StoreOp)¶

Store by a tensor of pointers or by a tensor pointer

Traits: SameLoadStoreOperandsEncoding, SameLoadStoreOperandsShape, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::SideEffects::DefaultResource}

Attributes:¶

Attribute	MLIR Type	Description
`boundaryCheck`	::mlir::DenseI32ArrayAttr	i32 dense array attribute
`cache`	::mlir::triton::CacheModifierAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6 {{% markdown %}}Enum cases: * none (`NONE`) * ca (`CA`) * cg (`CG`) * wb (`WB`) * cs (`CS`) * wt (`WT`){{% /markdown %}}
`evict`	::mlir::triton::EvictionPolicyAttr	allowed 32-bit signless integer cases: 1, 2, 3 {{% markdown %}}Enum cases: * evict_normal (`NORMAL`) * evict_first (`EVICT_FIRST`) * evict_last (`EVICT_LAST`){{% /markdown %}}

Operands:¶

Operand	Description
`ptr`	ptr or tensor of ptr values or ptr
`value`	floating-point or tensor of floating-point values or integer or tensor of integer values or ptr or tensor of ptr values or ptr
`mask`	1-bit signless integer or tensor of 1-bit signless integer values

`tt.trans` (triton::TransOp)¶

Transpose a tensor

Syntax:

operation ::= `tt.trans` $src attr-dict `:` functional-type(operands, results)

Traits: AlwaysSpeculatableImplTrait, SameOperandsAndResultElementType, TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:¶

Operand	Description
`src`	tensor of floating-point values or tensor of integer values or tensor of ptr values

Results:¶

Result	Description
`result`	tensor of floating-point values or tensor of integer values or tensor of ptr values

TritonOps¶

tt.call (triton::CallOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.func (triton::FuncOp)¶

Attributes:¶

tt.return (triton::ReturnOp)¶

Operands:¶

tt.addptr (triton::AddPtrOp)¶

Operands:¶

Results:¶

tt.advance (triton::AdvanceOp)¶

Operands:¶

Results:¶

tt.assert (triton::AssertOp)¶

Attributes:¶

Operands:¶

tt.atomic_cas (triton::AtomicCASOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.atomic_rmw (triton::AtomicRMWOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.bitcast (triton::BitcastOp)¶

Operands:¶

Results:¶

tt.broadcast (triton::BroadcastOp)¶

Operands:¶

Results:¶

tt.cat (triton::CatOp)¶

Operands:¶

Results:¶

tt.dot (triton::DotOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.elementwise_inline_asm (triton::ElementwiseInlineAsmOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.expand_dims (triton::ExpandDimsOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.extern_elementwise (triton::ExternElementwiseOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.fp_to_fp (triton::FpToFpOp)¶

Operands:¶

Results:¶

tt.get_num_programs (triton::GetNumProgramsOp)¶

Attributes:¶

Results:¶

tt.get_program_id (triton::GetProgramIdOp)¶

Attributes:¶

Results:¶

tt.int_to_ptr (triton::IntToPtrOp)¶

Operands:¶

Results:¶

tt.load (triton::LoadOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.make_range (triton::MakeRangeOp)¶

Attributes:¶

Results:¶

tt.make_tensor_ptr (triton::MakeTensorPtrOp)¶

Attributes:¶

Operands:¶

Results:¶

tt.print (triton::PrintOp)¶

Attributes:¶

Operands:¶

tt.ptr_to_int (triton::PtrToIntOp)¶

Operands:¶

Results:¶

`tt.call` (triton::CallOp)¶

`tt.func` (triton::FuncOp)¶

`tt.return` (triton::ReturnOp)¶

`tt.addptr` (triton::AddPtrOp)¶

`tt.advance` (triton::AdvanceOp)¶

`tt.assert` (triton::AssertOp)¶

`tt.atomic_cas` (triton::AtomicCASOp)¶

`tt.atomic_rmw` (triton::AtomicRMWOp)¶

`tt.bitcast` (triton::BitcastOp)¶

`tt.broadcast` (triton::BroadcastOp)¶

`tt.cat` (triton::CatOp)¶

`tt.dot` (triton::DotOp)¶

`tt.elementwise_inline_asm` (triton::ElementwiseInlineAsmOp)¶

`tt.expand_dims` (triton::ExpandDimsOp)¶

`tt.extern_elementwise` (triton::ExternElementwiseOp)¶

`tt.fp_to_fp` (triton::FpToFpOp)¶

`tt.get_num_programs` (triton::GetNumProgramsOp)¶

`tt.get_program_id` (triton::GetProgramIdOp)¶

`tt.int_to_ptr` (triton::IntToPtrOp)¶

`tt.load` (triton::LoadOp)¶

`tt.make_range` (triton::MakeRangeOp)¶

`tt.make_tensor_ptr` (triton::MakeTensorPtrOp)¶

`tt.print` (triton::PrintOp)¶

`tt.ptr_to_int` (triton::PtrToIntOp)¶

`tt.reduce` (triton::ReduceOp)¶

`tt.reduce.return` (triton::ReduceReturnOp)¶

`tt.reshape` (triton::ReshapeOp)¶

`tt.scan` (triton::ScanOp)¶

`tt.scan.return` (triton::ScanReturnOp)¶

`tt.splat` (triton::SplatOp)¶

`tt.store` (triton::StoreOp)¶

`tt.trans` (triton::TransOp)¶