# TritonGPUOps<!-- Autogenerated by mlir-tblgen; don't manually edit -->

### `ttg.async_commit_group` (triton::gpu::AsyncCommitGroupOp)

_Commit pending async copies into an async group that can be waited on_

Syntax:

```
operation ::= `ttg.async_commit_group` (`tokens` $inputTokens^)? attr-dict
```

Closes the current batch of async_copy_* operations
and allows for them to be waited on with `ttg.async_wait`.
This is required in order to ensure async copy operations can be waited on.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `InferTypeOpInterface`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `inputTokens` | variadic of async token type |

#### Results:

| Result | Description |
| :----: | ----------- |
| `asyncToken` | async token type |



### `ttg.async_copy_global_to_local` (triton::gpu::AsyncCopyGlobalToLocalOp)

_Copy data from global memory to local memory asynchronously_

Syntax:

```
operation ::= `ttg.async_copy_global_to_local` $src `,` $result (`mask` $mask^)? (`other` $other^)?
              attr-dict `:` type($src) `->` type($result)
```

This operation copies data from global memory to local memory asynchronously.
This is analogue to `tt.load` except the data are copied to local memory pointed
to by the memory descriptor instead of a distributed tensor. The rest of the
operands are the same as `tt.load`.
Contiguity is the maximum number of elements that can be loaded in a single vector with
the given layout and mask.
This allows op to use `async_copy_global_to_local` even if the alignment cannot be proven based on IR.

The data will only be available in local memory after `ttg.async_wait` is issued to wait on the
completion of `async_copy_global_to_local`. The async copy operations must be committed using
`ttg.async_commit_group` to close the batch and allow for them to be waited on.

Traits: `AttrSizedOperandSegments`, `GlobalToLocalCopyTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `InferTypeOpInterface`, `PredicatedOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cachePolicy</code></td><td>::mlir::Attribute</td><td>any attribute</td></tr>
<tr><td><code>isVolatile</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>contiguity</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | ranked tensor of ptr values |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `mask` | tensor of 1-bit signless integer values |
| `other` | floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte or ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `token` | async token type |



### `ttg.async_wait` (triton::gpu::AsyncWaitOp)

_Ensure all specified async_copy_* operations are complete._

Syntax:

```
operation ::= `ttg.async_wait` ($asyncToken^)? attr-dict
```

The `async_wait` op waits until at most "num" async copy groups are outstanding without synchronising CTA execution.
It takes zero or more `asyncToken` plus an integer `num` that specifies how many async copy groups can remain
outstanding after the `async_wait` op is completed. `num = 0` waits until all groups of async copies are complete.

This operation does not provide any synchronization in the CTA, if synchronization is needed use `ttg.local_barrier`
in addition to this operation.

Traits: `MemWaitOpTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `InferTypeOpInterface`

#### Attributes:

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `asyncToken` | variadic of async token type |

#### Results:

| Result | Description |
| :----: | ----------- |
| `retToken` | async token type |



### `ttg.barrier` (triton::gpu::BarrierOp)

_Synchronizes execution and memory accesses within a CTA or warp._

At the default CTA scope, the `barrier` op synchronises the execution and all operations between the selected address spaces for all
threads in the CTA. It is used to coordinate communication between threads in the CTA.

This operation waits until all threads in the CTA have reached a `barrier` (for synchronization) and operations
between the selected address spaces made by these threads prior to the op are visible to all threads in the CTA.

Data hazards between threads accessing the same memory can be avoided by synchronising the
specified scope in-between these accesses with a `barrier`.

A `barrier` operation only provides synchronization and memory guarantees on the selected address spaces in the CTA.

Warp scope currently supports only the `local` address space. It synchronizes
all lanes of each participating warp and orders their shared memory accesses.

The mandatory `addrspace` attribute is a bitmask describing which address spaces will be visible when the `barrier` completes:

* `none`         control-only synchronization (no memory ordering).
* `local`        shared-memory operations are complete and visible within the selected scope.
* `global_read`  global memory reads are complete and visible CTA-wide.
* `global_write` global memory writes are complete and visible CTA-wide.
* `tensor_read`  tensor memory read operations are complete and visible CTA-wide.
* `tensor_write` tensor memory write operations are complete and visible CTA-wide.
* `all`          convenience alias for `["local", "global_read", "global_write", "tensor_read", "tensor_write"]`.

Multiple address spaces can be combined (e.g. `local|tensor_write`). `none` cannot be combined with other address spaces.

Example:

```mlir
ttg.barrier local
ttg.barrier local|global_read|global_write
ttg.barrier warp local
```

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>addrSpace</code></td><td>::mlir::triton::gpu::AddrSpaceAttr</td><td></td></tr>
<tr><td><code>scope</code></td><td>::mlir::triton::gpu::BarrierScopeAttr</td><td>barrier synchronization scope</td></tr>
</table>



### `ttg.convert_layout` (triton::gpu::ConvertLayoutOp)

_Convert layout_

Syntax:

```
operation ::= `ttg.convert_layout` $src attr-dict `:` type($src) `->` type($result)
```

Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultElementType`, `SameOperandsAndResultShape`, `VerifyMemDescLayoutsTrait`, `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>force_warp_shuffle</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |



### `ttg.fp4_to_fp` (triton::gpu::Fp4ToFpOp)

_Upcast fp4 (e2m1) to fp_

Syntax:

```
operation ::= `ttg.fp4_to_fp` $src attr-dict `:` type($src) `->` type($result)
```

Upcast fp4 (e2m1) represented packed as i8s to fp.

The lower 4 bits of the i8s represent the first fp4 element, and the upper 4 bits
the second fp4 element.

The `axis` attribute specifies the axis along which the fp4 elements are packed.

Traits: `AlwaysSpeculatableImplTrait`, `VerifyMemDescLayoutsTrait`, `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>axis</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | ranked tensor of 8-bit signless integer or Two E2M1 floating-point values packed in one byte values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point values |



### `ttg.global_scratch_alloc` (triton::gpu::GlobalScratchAllocOp)

_Allocate a global memory buffer_

Syntax:

```
operation ::= `ttg.global_scratch_alloc` attr-dict `:` qualified(type($result))
```

This operation allocates a buffer in global memory that is private to the current program.
A custom third-party allocation can be marked using the optional
`third_party_allocation` unit attribute.
Concurrency sanitizer allocations that hold cluster-wide state can be marked
using the optional `shared_cluster_state` unit attribute.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>nbytes</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>alignment</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>third_party_allocation</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr>
<tr><td><code>shared_cluster_state</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr>
</table>

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ptr |



### `ttg.inline_asm` (triton::gpu::InlineAsmOp)

_Run inline assembly once per thread with explicit tensor layouts_

Syntax:

```
operation ::= `ttg.inline_asm` $asm_string attr-dict ($args^)? `:` functional-type($args, $result)
```

Each tensor operand is unpacked into all of its per-thread elements in
register order, including replicated register positions. Scalars and memory
descriptors contribute one operand; descriptors are passed as i32 addresses.
Results are packed according to their explicit types and layouts. There is
no elementwise broadcasting, chunking, or subword packing.

Assembly with memory descriptor operands must be impure. Descriptor accesses
are not modeled by the compiler. The caller must provide all synchronization
for them, including barriers, fences, and asynchronous completion before
storage reuse. Accesses must remain within the descriptor views.
Scalar results must be uniform and replicated tensor elements must agree.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `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>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `args` | variadic of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte or ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values or memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | variadic of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte or ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |



### `ttg.local_alloc` (triton::gpu::LocalAllocOp)

_Allocate tensor_

Syntax:

```
operation ::= `ttg.local_alloc` ($src^)? attr-dict `:` functional-type(operands, results)
```

This operation allocates buffer in shared memory and return a descriptor
containing the address and a view of the buffer.

Explicitly deallocating a buffer is optional; see local_dealloc.

The `src` operand is an optional initializer for the allocated buffer. It
must have the element type as the buffer. If `src` is not specified, the
returned buffer must be mutable.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `MemoryEffectOpInterface`

#### Attributes:

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.local_atomic_scatter_rmw` (triton::gpu::LocalAtomicScatterRMWOp)

_Atomically scatter RMW elements into shared memory along a specified axis_

Syntax:

```
operation ::= `ttg.local_atomic_scatter_rmw` $atomic_rmw_op `,` $dst `[` $indices `]` `,` $values (`,` $mask^)? attr-dict `:`
              functional-type(operands, results)
```

Atomically updates elements in a shared memory descriptor using an indices tensor
along a single specified axis. The values tensor has the same shape as the
indices tensor and the result returns the previous values observed at each
updated location.

For each input position I, the operation atomically updates dst where the
coordinate at the axis is replaced by indices[I]:
  dst[I[0], ..., indices[I], ..., I[n]] =
    rmw(dst[I[0], ..., indices[I], ..., I[n]], values[I])

If mask is present, the update is performed only for positions where
mask[I] is true. Masked-off positions do not update shared memory and return
an undefined old value.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `InferTypeOpInterface`

#### 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>allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7, 8, 9, 10</td></tr>
<tr><td><code>axis</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `dst` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `values` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |
| `indices` | ranked tensor of integer values |
| `mask` | tensor of 1-bit signless integer values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |



### `ttg.local_dealloc` (triton::gpu::LocalDeallocOp)

_Dealloc buffer_

Syntax:

```
operation ::= `ttg.local_dealloc` $src attr-dict `:` qualified(type($src))
```

This operation deallocates a buffer explicitly. Using the buffer after this
operation is undefined.

This operation is optional.  If you don't explicitly dealloc a buffer, the
compiler assumes it's deallocated at the first point that post-dominates all
uses of the alloc.

Because we assume a memdesc is dead at the first point that post-dominates
its uses, ops that wait for an async operation on a memdesc to complete
(such as ttng.warp_group_dot_wait) should also take the memdesc as an
operand.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.local_gather` (triton::gpu::LocalGatherOp)

_Gather elements from shared memory along a specified axis_

Syntax:

```
operation ::= `ttg.local_gather` $src `[` $indices `]` (`token` $token^)? attr-dict `:` qualified(type($src)) `,` type($indices) `->` type($result)
```

Gather elements from a shared memory descriptor using an indices tensor along a
single specified axis. The output tensor has the same shape as the indices tensor.

For each output position I, the operation reads from src where the coordinate at
the gather axis is replaced by indices[I]:
  result[I] = src[I[0], ..., indices[I], ..., I[n]]
where the axis dimension is replaced by the index value.

This matches the behavior of tt.gather but operates on shared memory descriptors.

Traits: `LocalLoadTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### 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` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `indices` | ranked tensor of integer values |
| `token` | async token type |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |



### `ttg.local_load` (triton::gpu::LocalLoadOp)

_Load a buffer from local memory into a distributed tensor_

Syntax:

```
operation ::= `ttg.local_load` $src (`token` $token^)? attr-dict `:` qualified(type($src)) `->` type($result)
```

Load a tensor from the local memory descriptor into a distributed tensor.

Traits: `LocalLoadTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `token` | async token type |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |



### `ttg.local_scatter` (triton::gpu::LocalScatterOp)

_Scatter elements to shared memory along a specified axis_

Syntax:

```
operation ::= `ttg.local_scatter` $dst `[` $indices `]` `,` $values (`token` $token^)? attr-dict `:` qualified(type($dst)) `,` type($indices) `,` type($values)
```

Scatter elements to a shared memory descriptor using an indices tensor along a
single specified axis. The values tensor has the same shape as the indices tensor.

For each input position I, the operation writes to dst where the coordinate at
the scatter axis is replaced by indices[I]:
  dst[I[0], ..., indices[I], ..., I[n]] = values[I]
where the axis dimension is replaced by the index value.

This is the inverse of local_gather and writes to shared memory at runtime-computed indices.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### 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 |
| :-----: | ----------- |
| `dst` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `values` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |
| `indices` | ranked tensor of integer values |
| `token` | async token type |



### `ttg.local_store` (triton::gpu::LocalStoreOp)

_Store a distributed tensor into a buffer in local memory_

Syntax:

```
operation ::= `ttg.local_store` $src `,` $dst attr-dict `:` type($src) `->` qualified(type($dst))
```

Store a distributed tensor into a buffer in local memory.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | ranked tensor of floating-point or integer or ptr or Two E2M1 floating-point values packed in one byte values |
| `dst` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.mask` (triton::gpu::MaskOp)

_Mask op for pipelining_

Traits: `SingleBlock`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `pred` | 1-bit signless integer |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | variadic of any non-token type |



### `ttg.mask.return` (triton::gpu::MaskReturnOp)

_Terminator for mask operator_

Syntax:

```
operation ::= `ttg.mask.return` $result attr-dict `:` type($result)
```

Traits: `AlwaysSpeculatableImplTrait`, `HasParent<MaskOp>`, `ReturnLike`, `Terminator`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`, `RegionBranchTerminatorOpInterface`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `result` | variadic of any non-token type |



### `ttg.memdesc_index` (triton::gpu::MemDescIndexOp)

_Take a subview of the descriptor._

Syntax:

```
operation ::= `ttg.memdesc_index` $src `[` $index `]` attr-dict `:` qualified(type($src)) `->` qualified(type($result))
```

This operation returns a new descriptor pointing to the `i`-th element of the
input descriptor along the 0-th dimension.

It doesn't affect the underlying memory.

For example, suppose that
 - the input shape is 2x4x16xf16,
 - the output shape is 4x16xf16, and
 - index = 1.
Then the output descriptor is equivalent to input[1], where input is the logical tensor.

Traits: `AlwaysSpeculatableImplTrait`, `MemDescViewTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `index` | 32-bit signless integer |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.memdesc_reinterpret` (triton::gpu::MemDescReinterpretOp)

_Reinterpret a memory descriptor as a different type and shape_

Syntax:

```
operation ::= `ttg.memdesc_reinterpret` $src attr-dict `:` qualified(type($src)) `->` qualified(type($result))
```

The `ttg.memdesc_reinterpret` operation reinterprets a memory descriptor
as one with a different shape and element type. Every physical byte,
tensor-memory row, and pipeline stage accessed by the result must belong to
the source view. The result must not itself be a layout subview.

Traits: `AlwaysSpeculatableImplTrait`, `MemDescViewTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.memdesc_reshape` (triton::gpu::MemDescReshapeOp)

_Creates a descriptor for the new shape_

Syntax:

```
operation ::= `ttg.memdesc_reshape` $src attr-dict `:` qualified(type($src)) `->` qualified(type($result))
```

This operation returns a new descriptor representing a reshaped view of the underlying buffer.
This doesn't affect the memory.

Traits: `AlwaysSpeculatableImplTrait`, `MemDescViewTrait`, `SameOperandsAndResultElementType`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.memdesc_subslice` (triton::gpu::MemDescSubsliceOp)

_Take a subview of the descriptor._

Syntax:

```
operation ::= `ttg.memdesc_subslice` $src `[` custom<Offsets>($offsets) `]` attr-dict `:` qualified(type($src))
              `->` qualified(type($result))
```

This operation returns a new descriptor representing a subview of the logical tensor.
It doesn't affect the underlying memory.

For example, suppose that
 - the input shape is 32x16xf16,
 - the output shape is 8x16xf16, and
 - offsets = [2, 1].
Then in Python syntax, the subview covers input[2:8+2, 1:16+1] where input is
the logical tensor.

The offsets must be larger or equal to the tile of the tensor (or zero).

Traits: `AlwaysSpeculatableImplTrait`, `MemDescViewTrait`, `VerifyMemDescLayoutsTrait`, `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>offsets</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.memdesc_trans` (triton::gpu::MemDescTransOp)

_Transpose the descriptor_

Syntax:

```
operation ::= `ttg.memdesc_trans` $src attr-dict `:` qualified(type($src)) `->` qualified(type($result))
```

This operation returns a new descriptor
representing a transposed view of the buffer.

Traits: `AlwaysSpeculatableImplTrait`, `InferTypeOpAdaptor`, `MemDescViewTrait`, `SameOperandsAndResultElementType`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `InferTypeOpInterface`, `NoMemoryEffect (MemoryEffectOpInterface)`, `TransposeOpInterface`

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 |
| :-----: | ----------- |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |



### `ttg.predicate_stage` (triton::gpu::PredicateStageOp)

_Pipeliner stage predicate_

Syntax:

```
operation ::= `ttg.predicate_stage` $iv `,` $ub `,` $step `maxStage` $maxStage `stage` $stage attr-dict `:` type($iv) `->` type($result)
```

Traits: `AlwaysSpeculatableImplTrait`, `VerifyMemDescLayoutsTrait`, `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>maxStage</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>stage</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `iv` | signless integer or index |
| `ub` | signless integer or index |
| `step` | signless integer or index |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | 1-bit signless integer |



### `ttg.thread_handoff` (triton::gpu::ThreadHandoffOp)

_Distribute a scalar at the boundary of a thread region_

Syntax:

```
operation ::= `ttg.thread_handoff` $src $source `->` $destination attr-dict `:` type($src)
```

The source scalar is available in `source` and is copied to `destination`.
Lane and warp coordinates are relative to the enclosing warp-specialized
partition; CTA coordinates are relative to the cluster.
The destination may only widen the source domain. Lane transfers execute
with every lane of the destination warps participating. Cross-warp transfers
execute with all lanes and warps in each participating CTA. Cross-CTA
transfers broadcast from CTA zero to the full enclosing domain.
This operation may synchronize and must not be speculated or eliminated.

Traits: `SameOperandsAndResultType`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `InferTypeOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>source</code></td><td>::mlir::triton::gpu::ThreadDomainAttr</td><td><details><summary></summary>{{% markdown %}}
TritonGPU tensors differ from usual tensors in that they contain a _layout_ attribute which determines
how the data should be partitioned across CUDA threads. Formally speaking, we define a layout as a function
\mathcal{L} that maps a multi-dimensional tensor index $i \in \mathbb{Z}^d$ to a set of integers T corresponding
to the indices of the CUDA threads allowed to access some data at index $i$.

For example, let us consider the layout function:
\mathcal{L}(0, 0) = {0, 4}
\mathcal{L}(0, 1) = {1, 5}
\mathcal{L}(1, 0) = {2, 6}
\mathcal{L}(1, 1) = {3, 7}

Then, attaching $\mathcal{L} to a tensor $T$ would mean that:
- T[0,0] is owned by both cuda thread 0 and 4
- T[0,1] is owned by both cuda thread 1 and 5
- T[1,0] is owned by both cuda thread 2 and 6
- T[1,1] is owned by both cuda thread 3 and 7

Right now, Triton implements two main classes of layouts: shared, and distributed.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>destination</code></td><td>::mlir::triton::gpu::ThreadDomainAttr</td><td><details><summary></summary>{{% markdown %}}
TritonGPU tensors differ from usual tensors in that they contain a _layout_ attribute which determines
how the data should be partitioned across CUDA threads. Formally speaking, we define a layout as a function
\mathcal{L} that maps a multi-dimensional tensor index $i \in \mathbb{Z}^d$ to a set of integers T corresponding
to the indices of the CUDA threads allowed to access some data at index $i$.

For example, let us consider the layout function:
\mathcal{L}(0, 0) = {0, 4}
\mathcal{L}(0, 1) = {1, 5}
\mathcal{L}(1, 0) = {2, 6}
\mathcal{L}(1, 1) = {3, 7}

Then, attaching $\mathcal{L} to a tensor $T$ would mean that:
- T[0,0] is owned by both cuda thread 0 and 4
- T[0,1] is owned by both cuda thread 1 and 5
- T[1,0] is owned by both cuda thread 2 and 6
- T[1,1] is owned by both cuda thread 3 and 7

Right now, Triton implements two main classes of layouts: shared, and distributed.
  {{% /markdown %}}</details></td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `src` | integer or floating-point or ptr |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | integer or floating-point or ptr |



### `ttg.thread_predicate` (triton::gpu::ThreadPredicateOp)

_Select the threads in a compiler-generated execution domain_

Syntax:

```
operation ::= `ttg.thread_predicate` $domain attr-dict
```

This operation marks entry into a restricted thread region and executes in
the enclosing domain before its branch. Its conservative effects let
synchronization analysis finish pending waits before entering the region.
It must not be speculated, merged, or eliminated before that analysis.

Traits: `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `InferTypeOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>domain</code></td><td>::mlir::triton::gpu::ThreadDomainAttr</td><td><details><summary></summary>{{% markdown %}}
TritonGPU tensors differ from usual tensors in that they contain a _layout_ attribute which determines
how the data should be partitioned across CUDA threads. Formally speaking, we define a layout as a function
\mathcal{L} that maps a multi-dimensional tensor index $i \in \mathbb{Z}^d$ to a set of integers T corresponding
to the indices of the CUDA threads allowed to access some data at index $i$.

For example, let us consider the layout function:
\mathcal{L}(0, 0) = {0, 4}
\mathcal{L}(0, 1) = {1, 5}
\mathcal{L}(1, 0) = {2, 6}
\mathcal{L}(1, 1) = {3, 7}

Then, attaching $\mathcal{L} to a tensor $T$ would mean that:
- T[0,0] is owned by both cuda thread 0 and 4
- T[0,1] is owned by both cuda thread 1 and 5
- T[1,0] is owned by both cuda thread 2 and 6
- T[1,1] is owned by both cuda thread 3 and 7

Right now, Triton implements two main classes of layouts: shared, and distributed.
  {{% /markdown %}}</details></td></tr>
</table>

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | 1-bit signless integer |



### `ttg.thread_scope` (triton::gpu::ThreadScopeOp)

_Preserve a compiler-generated single-thread scope during lowering_

Syntax:

```
operation ::= `ttg.thread_scope` (`->` `(` type($results)^ `)`)? $body attr-dict
```

The body executes unconditionally within an enclosing guard that selects
the canonical thread of each participating CTA. Entry/exit synchronization
remains outside the guard.
This wrapper survives control-flow lowering so memory operations can query
their enclosing scope, and is inlined after memory operations are lowered.

Traits: `RecursiveMemoryEffects`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `RegionBranchOpInterface`

#### Results:

| Result | Description |
| :----: | ----------- |
| `results` | variadic of any non-token type |



### `ttg.thread_yield` (triton::gpu::ThreadYieldOp)

_Yield results from a single-thread scope_

Syntax:

```
operation ::= `ttg.thread_yield` ($values^)? attr-dict (`:` type($values)^)?
```

Traits: `AlwaysSpeculatableImplTrait`, `HasParent<ThreadScopeOp>`, `ReturnLike`, `Terminator`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`, `RegionBranchTerminatorOpInterface`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `values` | variadic of any non-token type |



### `ttg.warp_id` (triton::gpu::WarpIdOp)

_Return the GPU warp ID_

Syntax:

```
operation ::= `ttg.warp_id` attr-dict
```

This operation returns the GPU warp ID. This can translate to reading
hardware registers if there are, or just thread ID divided by warp size.

The `omitUniformHint` attribute is indicating in NVIDIA backend whether to
omit emitting nvvm.shfl.sync idx 0 for LLVM.

Traits: `AlwaysSpeculatableImplTrait`, `VerifyMemDescLayoutsTrait`, `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>omitUniformHint</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr>
</table>

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | 32-bit signless integer |



### `ttg.warp_return` (triton::gpu::WarpReturnOp)

_Implicit terminator from partition regions_

Syntax:

```
operation ::= `ttg.warp_return` attr-dict
```

The `ttg.warp_return` operation is the implicit terminator that ends the
partition regions of a `ttg.warp_specialize` op. It has no operands as these
regions cannot return anything.

TODO: Support returning uniform values from partition regions.

Traits: `AlwaysSpeculatableImplTrait`, `HasParent<WarpSpecializePartitionsOp>`, `ReturnLike`, `Terminator`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`, `RegionBranchTerminatorOpInterface`

Effects: `MemoryEffects::Effect{}`



### `ttg.warp_specialize` (triton::gpu::WarpSpecializeOp)

_Asynchronously execute code on multiple warpgroups_

The `ttg.warp_specialize` op represents executing different code
simultaneously on different warp groups. A warp group is a group of
power-of-2 warps, which can be a different number of warps than in the
enclosing region.

The "default" region of the op represents the code executed by the currently
executing warp group. This region is allowed to implicitly capture. The op
contains a number of "partition" regions that are isolated from above. They
must be isolated because these regions represent different layout domains,
as the number of warps is different.

Semantically, execution of each region starts simultaneously for each warp
group, and all warp groups are joined at the end of the op.

Example:

```mlir
%0 = ttg.warp_specialize(%a, %b)
default {
  %out = some_operation(%a) // implicit capture of `%a`
  ttg.warp_yield %out : i32
}
partition0(%arg0: i32, %arg1: i32) num_warps(8) {
  some_async_dispatch(%arg0, %arg1)
  ttg.warp_return
}
partition1(%arg0: i32, %arg1: i32) num_warps(1) {
  some_async_dispatch(%arg0, %arg1)
  ttg.warp_return
} : (i32, i32) -> i32
```

Traits: `AsyncRegions`, `RecursiveMemoryEffects`, `RecursivelySpeculatableImplTrait`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `RegionBranchOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>partitionNumWarps</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
<tr><td><code>warpGroupStartIds</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
<tr><td><code>requestedRegisters</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
<tr><td><code>actualRegisters</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
</table>

#### Results:

| Result | Description |
| :----: | ----------- |
| `defaultPassthrough` | variadic of any non-token type |



### `ttg.warp_specialize.partitions` (triton::gpu::WarpSpecializePartitionsOp)

_Container op for `ttg.warp_specialize`_

Because MLIR requires entire operations be isolated from above, this op
contains the actual isolated from above regions of `ttg.warp_specialize`.

Traits: `HasParent<WarpSpecializeOp>`, `IsolatedFromAbove`, `RecursiveMemoryEffects`, `RecursivelySpeculatableImplTrait`, `Terminator`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `RegionBranchOpInterface`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `explicitCaptures` | variadic of any non-token type |



### `ttg.warp_yield` (triton::gpu::WarpYieldOp)

_Yield from the default region of `ttg.warp_specialize`_

Syntax:

```
operation ::= `ttg.warp_yield` ($values^)? attr-dict (`:` type($values)^)?
```

The `ttg.warp_yield` operation is the terminator for the "default" region of
a `ttg.warp_specialize` operation. The operands are passed transparently as
the SSA results of the `ttg.warp_specialize` operation.

Example:

```mlir
ttg.warp_yield %a, %b : i32, tensor<32xbf16, #blocked>
```

Traits: `AlwaysSpeculatableImplTrait`, `HasParent<WarpSpecializeOp>`, `ReturnLike`, `Terminator`, `VerifyMemDescLayoutsTrait`, `VerifyTensorLayoutsTrait`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`, `RegionBranchTerminatorOpInterface`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `values` | variadic of any non-token type |