# TritonOps<!-- Autogenerated by mlir-tblgen; don't manually edit -->
### `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:

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

Traits: TensorSizeTrait, VerifyTensorLayoutsTrait

Interfaces: CallOpInterface, SymbolUserOpInterface

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>callee</code></td><td>::mlir::FlatSymbolRefAttr</td><td>flat symbol reference attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `operands` | variadic of any type

#### Results:

| Result | Description |
| :----: | ----------- |
&laquo;unnamed&raquo; | 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:

```mlir
// 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>sym_name</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>function_type</code></td><td>::mlir::TypeAttr</td><td>type attribute of function type</td></tr>
<tr><td><code>sym_visibility</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>arg_attrs</code></td><td>::mlir::ArrayAttr</td><td>Array of dictionary attributes</td></tr>
<tr><td><code>res_attrs</code></td><td>::mlir::ArrayAttr</td><td>Array of dictionary attributes</td></tr>
</table>

### `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:

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

Traits: AlwaysSpeculatableImplTrait, HasParent<FuncOp>, 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>message</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>file</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>func</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>line</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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:

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

#### 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:

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

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>allowTF32</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>maxNumImpreciseAcc</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>asm_string</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>constraints</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>pure</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>packed_element</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>axis</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>libname</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>libpath</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>symbol</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>pure</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>axis</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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:

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

#### 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:

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

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>start</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>end</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>order</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>prefix</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>axis</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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<ReduceOp>, 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>allow_reorder</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### 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:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>axis</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### 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<ScanOp>, 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:

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

#### 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