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

### `amdg.arrive_barrier` (triton::amdgpu::ArriveBarrierOp)

_Perform the arrive operation on an mbarrier_

Syntax:

```
operation ::= `amdg.arrive_barrier` $alloc `,` $count attr-dict `:` qualified(type($alloc)) `->` type($result)
```

Performs the "arrive" operation on an mbarrier object in shared memory. The operation requires a `count` attribute
of at least 1, and decreases the pending arrival count of the mbarrier by the specific count. If the pending count reaches
zero, the phase changes (is decremented in a wraparound manner) and the pending count is reloaded with the init count value. Returns the phase
parity (0 for even, 1 for odd) of the mbarrier object prior to the "arrive" operation.

Example:

```mlir
ttag.arrive_barrier %barrier, 2 : !ttg.memdesc<1xi64, #shared, #smem, mutable>
```

Interfaces: `InferTypeOpInterface`, `MBarrierOpInterface`

#### Attributes:

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

#### Operands:

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

#### Results:

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



### `amdg.async_copy_local_to_global` (triton::amdgpu::AsyncCopyLocalToGlobalOp)

_Copy data from local memory to global memory asynchronously_

Syntax:

```
operation ::= `amdg.async_copy_local_to_global` $src `,` $dst (`mask` $mask^)?
              oilist(`cachePolicy` `=` $cachePolicy)
              attr-dict `:` qualified(type($src)) `->` type($dst)
```

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

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>contiguity</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 |
| `dst` | ranked tensor of ptr values |
| `mask` | tensor of 1-bit signless integer values |

#### Results:

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



### `amdg.async_copy_mbarrier_arrive` (triton::amdgpu::AsyncCopyMbarrierArriveOp)

_Arrive on mbarrier once all previously issued copies are completed_

Syntax:

```
operation ::= `amdg.async_copy_mbarrier_arrive` $barrier attr-dict `:` qualified(type($barrier))
```

Performs the "async arrive" operation by decrementing pending account by 1 when all previous async load to LDS (particularly, not TDM) have completed.
The instruction itself is asynchronous; it returns immediately. Decrements the barrier pending count. The update value for decrementing is fixed at 1.
If the pending count becomes zero, the phase changes (is decremented in a wraparound manner) and the pending count is reloaded with the init count value.

Interfaces: `MBarrierOpInterface`

#### Operands:

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



### `amdg.async_tdm_copy_global_to_local` (triton::amdgpu::AsyncTDMCopyGlobalToLocalOp)

_Copy data based on descriptor from global memory to local memory asynchronously_

Syntax:

```
operation ::= `amdg.async_tdm_copy_global_to_local` $desc `into` $result (`,` `barrier` `=` $barrier^)?
              attr-dict `:` qualified(type($desc)) (`,` qualified(type($barrier))^)? `->` qualified(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 by `result` instead of a distributed tensor. The data
copied depends on the global memory pointed to by `desc`. This operation
does not support shared memory swizzling.

The descriptor must be positioned beforehand by amdg.update_tensor_descriptor
(global offset, bounds, pred); this op applies only the per-warp distribution
and inherits `pred` from the descriptor.
The operation can also take an optional 64bit LDS barrier address, in which case
it sends an "LDS atomic arrive" to signal its completion.

`warp_used_hint` is an optional i32 bitmask performance hint for
partial TDM copy: bit `n` (LSB = warp 0) means warp `n` carries a
real descriptor; bits above `num_warps - 1` must be zero. Same data
lands in shared memory either way: `K = popcount(hint)` active warps
cover the block, while inactive warps still issue the TDM instruction
with a null descriptor.

For "all warps participate", omit the attribute (preferred) or pass
a full bitmask `K = num_warps`; an explicit `warp_used_hint = 0` is
rejected by the verifier.

Restriction: the hint must be *axis-aligned*.  Pin some warpId bit
positions to fixed values and leave the others free; the active
warps are exactly those with the pinned bits set accordingly.
K = `popcount(hint)` is therefore a power of two (= 2^(number of
free bits)).  Legal: `0x0F` (warps 0..3), `0xF0` (warps 4..7),
`0x55` / `0xAA` / `0xCC` (stride-by-`2^j`).  Rejected: non
axis-aligned patterns like `0x69` or `0x1B`.

Why this shape: power-of-two K keeps the per-warp tile recalc
(`block / K`) integer-clean, and the bit-pinning shape lets the
per-warp offset and active-mask predicate fall out of the same
`getFreeVariableMasks` LL machinery the no-hint path uses.

PartitionedSharedEncoding additionally requires `K = popcount(hint)`
to be a multiple of `numLogicalPieces` (the hinted path doesn't
support multi-instruction slicing).

Traits: `GlobalToLocalCopyTrait`

Interfaces: `InferTypeOpInterface`, `MBarrierOpInterface`, `TDMOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cache</code></td><td>::mlir::triton::CacheModifierAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7</td></tr>
<tr><td><code>warp_used_hint</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `desc` | Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |
| `result` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `barrier` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

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



### `amdg.async_tdm_copy_local_to_global` (triton::amdgpu::AsyncTDMCopyLocalToGlobalOp)

_Copy data based on descriptor from local memory to global memory asynchronously_

Syntax:

```
operation ::= `amdg.async_tdm_copy_local_to_global` $desc `from` $src (`,` `barrier` `=` $barrier^)?
              attr-dict `:` qualified(type($src)) (`,` qualified(type($barrier))^)? `->` qualified(type($desc))
```

This operation copies data from local memory to global memory
asynchronously. This is analogue to tt.store except the data are copied from
local memory pointed by `src` instead of a distributed tensor. The copy
destination depends on the global memory pointed to by `desc`. This
operation does not support shared memory padding or swizzling.

The descriptor must be positioned beforehand by
amdg.update_tensor_descriptor; this op applies only the per-warp distribution.
The operation can also take an optional 64bit LDS barrier address, in which case
it sends an "LDS atomic arrive" to signal its completion.

Interfaces: `InferTypeOpInterface`, `MBarrierOpInterface`, `TDMOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cache</code></td><td>::mlir::triton::CacheModifierAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7</td></tr>
</table>

#### Operands:

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

#### Results:

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



### `amdg.async_tdm_fused_copy_global_to_local` (triton::amdgpu::AsyncTDMFusedCopyGlobalToLocalOp)

_Fuse multiple TDM global-to-local copies into one async operation_

Syntax:

```
operation ::= `amdg.async_tdm_fused_copy_global_to_local` $descs `into` $dests attr-dict `:` qualified(type($descs)) `->` qualified(type($dests))
```

This operation represents an explicit 2-4 member fused TDM load.  Each
member provides a tensor descriptor, destination shared-memory descriptor,
and `warp_used_hint`.  Member hints must be legal and pairwise disjoint.
Descriptors carry tile position and predicate state; destinations remain
op-level operands.

Traits: `GlobalToLocalCopyTrait`, `SameVariadicOperandSize`

Interfaces: `InferTypeOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>warp_used_hints</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>allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `descs` | variadic of Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |
| `dests` | variadic of memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

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



### `amdg.async_tdm_gather` (triton::amdgpu::AsyncTDMGatherOp)

_Gather data from non-contiguous global memory rows to local memory asynchronously_

Syntax:

```
operation ::= `amdg.async_tdm_gather` $desc `[` $src_row_indices `]` `to` $dst (`,` `barrier` `=` $barrier^)?
              attr-dict `:` qualified(type($src_row_indices)) `,` qualified(type($dst)) (`,` qualified(type($barrier))^)? `->` qualified(type($desc))
```

This operation gathers data from non-contiguous rows in global memory to local
memory using TDM gather mode.
Unlike the regular async_tdm_copy_global_to_local which reads from contiguous memory,
this operation uses src_row_indices to specify which rows in global memory to read from.

The descriptor must be 2D. The src_row_indices specify which rows in global memory
to read from. The element type of src_row_indices determines the index size:
- I16: 16-bit indices, up to 16 rows per instruction
- I32: 32-bit indices, up to 8 rows per instruction
If more rows are needed, multiple TDM instructions will be issued.

The descriptor must be positioned beforehand by amdg.update_tensor_descriptor
(global offset / column via add_offsets, bounds, pred); this op only supplies the
per-instruction row slices (src_row_indices) and the LDS destination, and inherits
`pred` from the descriptor.

Traits: `GlobalToLocalCopyTrait`

Interfaces: `InferTypeOpInterface`, `MBarrierOpInterface`, `TDMOpInterface`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `desc` | Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |
| `src_row_indices` | tensor of 16-bit signless integer or 32-bit signless integer values |
| `dst` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `barrier` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

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



### `amdg.async_tdm_intrinsic_wait` (triton::amdgpu::AsyncTDMIntrinsicWait)

_Wait until there are less than or equal to the given number of outstanding TDM intrinsics_

Syntax:

```
operation ::= `amdg.async_tdm_intrinsic_wait` ($asyncToken^)? attr-dict
```

This operation waits until there are less than or equal to the given number
of outstanding TDM intrinsics (assembly instructions). This is necessary to
ensure that data is available in LDS (load, gather) or HBM (store, scatter)
before it is used.

Unlike AsyncTDMWait which counts IR operations, this operation counts the
actual number of TDM assembly instructions (e.g., tensor_load_to_lds,
tensor_store_from_lds) that are emitted. Which is required for the lowering
to LLVM.

Traits: `MemWaitOpTrait`

Interfaces: `InferTypeOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>count</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 |



### `amdg.async_tdm_scatter` (triton::amdgpu::AsyncTDMScatterOp)

_Scatter data from local memory to non-contiguous global memory rows asynchronously_

Syntax:

```
operation ::= `amdg.async_tdm_scatter` $desc `[` $dst_row_indices `]` `from` $src (`,` `barrier` `=` $barrier^)?
              attr-dict `:` qualified(type($dst_row_indices)) `,` qualified(type($src)) (`,` qualified(type($barrier))^)? `->` qualified(type($desc))
```

This operation scatters data from local memory to non-contiguous rows in global
memory using TDM scatter mode.
Unlike the regular async_tdm_copy_local_to_global which copies to contiguous memory,
this operation uses dst_row_indices to specify which rows in global memory to write to.

The descriptor must be 2D. The dst_row_indices specify which rows in global memory
to write to. The element type of dst_row_indices determines the index size:
- I16: 16-bit indices, up to 16 rows per instruction
- I32: 32-bit indices, up to 8 rows per instruction
If more rows are needed, multiple TDM instructions will be issued.

The descriptor must be positioned beforehand by amdg.update_tensor_descriptor
(global offset / column via add_offsets, bounds); this op only supplies the
per-instruction row slices (dst_row_indices) and the LDS source.

Interfaces: `InferTypeOpInterface`, `MBarrierOpInterface`, `TDMOpInterface`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `desc` | Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |
| `dst_row_indices` | tensor of 16-bit signless integer or 32-bit signless integer values |
| `src` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `barrier` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |

#### Results:

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



### `amdg.async_tdm_wait` (triton::amdgpu::AsyncTDMWait)

_Wait until there are less than or equal to the given number of outstanding TDM operations_

Syntax:

```
operation ::= `amdg.async_tdm_wait` $asyncToken attr-dict
```

This operation waits until there are less than or equal to the given number
of outstanding async TDM operations.
This is necessary to ensure that data is available in the LDS (load, gather)
or HBM (store, scatter) before it is used.

This operation counts the number of TDM IR operations
(AsyncTDMCopyGlobalToLocalOp, AsyncTDMFusedCopyGlobalToLocalOp,
AsyncTDMCopyLocalToGlobalOp, AsyncTDMScatterOp, AsyncTDMGatherOp).

Traits: `MemWaitOpTrait`

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 |



### `amdg.async_wait` (triton::amdgpu::AsyncWaitOp)

_Wait until there are less than or equal to the given number of outstanding async intrinsics_

Syntax:

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

Similar to ttg.async_wait but instead of waiting on oustanding ttg.async_commit_groups
this op waits on the number of outstanding async instructions/intrinsics as required for the
lowering to LLVM on the AMD backend.

Traits: `MemWaitOpTrait`

Interfaces: `InferTypeOpInterface`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>num_inst</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 |



### `amdg.buffer_atomic_cas` (triton::amdgpu::BufferAtomicCASOp)

_Atomic CAS op which does compare-exchange to a scalar base pointer and a tensor offset_

Syntax:

```
operation ::= `amdg.buffer_atomic_cas` $sem `,` $scope `,` $cmp `,` $val `,` $ptr `[` $offsets `]`
              (`stride` `=` $stride^)?
              attr-dict `:` qualified(type($ptr)) `->` type($result)
```

AMD Buffer Atomic CAS operation. Buffer atomics are similar to normal atomics, but access global memory via a
scalar base pointer and a tensor of offsets instead of a tensor of pointers.
Similar to TT_AtomicCASOp: Buffer atomic CAS op loads data at $ptr, and stores $val to $ptr atomically if value at $ptr equals $cmp, with
the specified memory semantics and scope. Atomic CAS ops return the pre-op value if used, otherwise the value is implicitly dropped.
Stride is the distance between the beginning of contiguous memory chunks. When performing a CAS, the `stride` is
the address difference between the first elements of each row in bytes. Compiler tries to obtain the `stride`
when it converts to the buffer ops because it is important for optimizing the cache memory access.

Traits: `SameLoadStoreOperandsAndResultEncoding`

Interfaces: `BufferOpInterface`

#### 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>allowed 32-bit signless integer cases: 1, 2, 3, 4</td></tr>
<tr><td><code>scope</code></td><td>::mlir::triton::MemSyncScopeAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `ptr` | ptr |
| `offsets` | tensor of 32-bit signless integer values |
| `cmp` | ranked tensor of floating-point or integer or ptr values |
| `val` | ranked tensor of floating-point or integer or ptr values |
| `stride` | 32-bit signless integer |

#### Results:

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



### `amdg.buffer_atomic_rmw` (triton::amdgpu::BufferAtomicRMWOp)

_Atomic RMW op which reads, modifies, and writes to a scalar base pointer and a tensor offset_

Syntax:

```
operation ::= `amdg.buffer_atomic_rmw` $atomic_rmw_op `,` $sem `,` $scope `,` $value `,` $ptr `[` $offsets `]` (`,` $mask^)?
              (`stride` `=` $stride^)?
              attr-dict `:` qualified(type($ptr)) `->` type($result)
```

AMD Buffer atomic RMW operation. Buffer atomics are similar to normal atomics, but access global memory via a
scalar base pointer and a tensor of offsets instead of a tensor of pointers.
Similar to other buffer ops, the `mask` is a boolean vector that determines if a given element should be processed with
the atomic RMW op. Elements with `mask[i] == 0` are dropped (i.e., the atomic is not executed).
Similar to TT_AtomicRMWOp: Buffer atomic RMW ops load data at $ptr, do $rmw_op with $val, and store result to $ptr with
the specified memory semantics and scope. Atomic RMW ops return the pre-op value if used, otherwise the value is implicitly dropped.
Stride is the distance between the beginning of contiguous memory chunks. When performing a RMW, the `stride` is
the address difference between the first elements of each row in bytes. Compiler tries to obtain the `stride`
when it converts to the buffer ops because it is important for optimizing the cache memory access.

Traits: `AttrSizedOperandSegments`, `SameLoadStoreOperandsAndResultEncoding`

Interfaces: `BufferOpInterface`, `PredicatedOpInterface`

#### 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>sem</code></td><td>::mlir::triton::MemSemanticAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3, 4</td></tr>
<tr><td><code>scope</code></td><td>::mlir::triton::MemSyncScopeAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `ptr` | ptr |
| `offsets` | tensor of 32-bit signless integer values |
| `value` | ranked tensor of floating-point or integer or ptr values |
| `stride` | 32-bit signless integer |
| `mask` | ranked tensor of 1-bit signless integer values |

#### Results:

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



### `amdg.buffer_load` (triton::amdgpu::BufferLoadOp)

_Load from a scalar base pointer and a tensor offset_

Syntax:

```
operation ::= `amdg.buffer_load` $ptr `[` $offsets `]` (`,` $mask^)? (`,` $other^)?
              (`stride` `=` $stride^)?
              attr-dict `:` qualified(type($ptr)) `->` type($result)
```

AMD Buffer load operation. Buffer store is similar to
a normal store but it accesses global memory via a scalar base pointer
and a tensor of offsets instead of a tensor of pointers. The other fields
are similar to a normal load, i.e., the `mask` is a boolean vector that
determines if a given element should be read from memory, and `other` is the
element that should be returned on lane `i` when `mask[i] == 0`.
Stride is the distance between the beginning of contiguous memory chunks.
When performing a load of a block, the `stride` is the address difference between
the first elements of each row in bytes. Compiler tries to obtain the `stride`
when it converts to the buffer ops because it is important for optimizing
the cache memory access.
Contiguity is the maximum number of elements that can be loaded in a single vector
with the given layout and mask.
This allows to use buffer_load even if the alignment cannot be proven based on IR.

Traits: `AttrSizedOperandSegments`, `SameLoadStoreOperandsAndResultEncoding`

Interfaces: `BufferOpInterface`, `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>contiguity</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `ptr` | ptr |
| `offsets` | tensor of 32-bit signless integer values |
| `stride` | 32-bit signless integer |
| `mask` | ranked tensor of 1-bit signless integer values |
| `other` | ranked tensor of floating-point or integer or ptr values |

#### Results:

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



### `amdg.buffer_load_to_local` (triton::amdgpu::BufferLoadToLocalOp)

_Load from a scalar base pointer and a tensor offset to shared memory_

Syntax:

```
operation ::= `amdg.buffer_load_to_local` $ptr `[` $offsets `]` (`mask` `=` $mask^)? (`other` `=` $other^)? (`stride` `=` $stride^)?
              oilist(`cachePolicy` `=` $cachePolicy) `into` $dest
              attr-dict `:` qualified(type($ptr)) `[` type($offsets) `]` type($other) `->` type($dest)
```

AMD Buffer load operation. Similar to amdg.buffer_load op but directly wirtes to shared memory instead of into registers.
Contiguity is the maximum number of elements that can be loaded in a single vector with the given layout and mask.
This allows to use buffer_load_to_local even if the alignment cannot be proven based on IR.

Traits: `AttrSizedOperandSegments`, `GlobalToLocalCopyTrait`

Interfaces: `BufferOpInterface`, `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>contiguity</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `dest` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `ptr` | ptr |
| `offsets` | tensor of 32-bit signless integer values |
| `mask` | ranked tensor of 1-bit signless integer values |
| `other` | ranked tensor of floating-point or integer or ptr values |
| `stride` | 32-bit signless integer |

#### Results:

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



### `amdg.buffer_store` (triton::amdgpu::BufferStoreOp)

_Store into scalar base pointer and a tensor offset_

Syntax:

```
operation ::= `amdg.buffer_store` $value `,` $ptr `[` $offsets `]` (`,` $mask^)?
              (`stride` `=` $stride^)?
              attr-dict `:` qualified(type($ptr)) `->` type($value)
```

AMD Buffer store operation. Buffer store is similar to
normal store but it accesses global memory via a scalar base pointer
and a tensor of offsets instead of a tensor of pointers. The other fields
are similar to a normal store , i.e., the `mask` is a boolean vector that
determines if a given element should be written to memory, and `value` is the
tensor of elements that should be written on lane `i` when `mask[i] == 1`.
Stride is the distance between the beginning of contiguous memory chunks.
When performing a block store, the `stride` is the address difference between
the first elements of each row in bytes. Compiler tries to obtain the `stride`
when it converts to the buffer ops because it is important for optimizing
the cache memory access.
Contiguity is the maximum number of elements that can be loaded in a single vector
with the given layout and mask.
This allows to use buffer_store even if the alignment cannot be proven based on IR.

Traits: `AttrSizedOperandSegments`, `SameLoadStoreOperandsEncoding`

Interfaces: `BufferOpInterface`, `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>contiguity</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `value` | ranked tensor of floating-point or integer or ptr values |
| `ptr` | ptr |
| `offsets` | tensor of 32-bit signless integer values |
| `stride` | 32-bit signless integer |
| `mask` | ranked tensor of 1-bit signless integer values |



### `amdg.cluster_barrier_arrive` (triton::amdgpu::ClusterBarrierArriveOp)

_Arrive at a cluster barrier_

Syntax:

```
operation ::= `amdg.cluster_barrier_arrive` attr-dict
```

Signals that the cluster has arrived at a barrier, used to synchronizing CTAs within a cluster.

See ClusterBarrierWaitOp for how to wait on the arrived cluster barrier.



### `amdg.cluster_barrier_wait` (triton::amdgpu::ClusterBarrierWaitOp)

_Wait on a cluster barrier_

Syntax:

```
operation ::= `amdg.cluster_barrier_wait` attr-dict
```

Waits for all CTAs of the same cluster to have arrived at a cluster barrier.
Arrive and wait operations must come in pairs. Waiting before arriving or arriving
more than once without a corresponding wait will result in undefined behavior.



### `amdg.concat` (triton::amdgpu::ConcatOp)

_Concat operation_

Syntax:

```
operation ::= `amdg.concat` $sources attr-dict `:` type($sources) `->` type($result)
```

The "concat" operation combines a list of source n-dimensional tensors into a single larger destination tensor.

All source tensors must have the same shape, element type, and encoding.
The concatenation dimension is inferred from the source and destination shapes provided by the user.
For example, two tensors of shape 64x128 can produce a destination shape of 128x128,
indicating concatenation along dimension 0; or 64x256, indicating concatenation along dimension 1.

Generally, source tensors passed as op arguments can be arranged into the resulting shape in multiple ways.
For example, given four tensors of shape 64x64:
  concat s0<64x64>, s1<64x64>, s2<64x64>, s3<64x64> -> <128x128>

They can be laid out in different configurations within the result tensor:
  1) s0 s1     2) s0 s2
     s2 s3        s1 s3

From a logical tensor perspective, the source tensors are treated as elements of a tensor of tensors.
In other words, the 1-D array of input tensors is conceptually reshaped into an n-D grid.
The semantics of this op assume a row-major order (or its n-D generalization),
meaning the fastest-varying dimension is filled first, and the slowest-varying dimension is filled last.
In the example above, this corresponds to layout 1).

The source and destination tensors must have identical linear layouts at the CTA tile level.
That is, all base vectors for input dimensions must match, except for the register input dimension.
The register basis must align on the subset that defines the logical tensor shape of a single CTA tile.

This ensures that the concatenation is a no-op, meaning no data rearrangement among threads is required
to assemble the destination tensor with the given shape and layout.
However, the order of CTA tiles within the layout does not need to match between source and destination layouts.
It is the responsibility of the op's lowering logic to handle this correctly.

This op is designed to work on logical tensors directly, avoiding the need for complex layout reinterpretation or reshaping.
For example, the `tt.join` operation only supports concatenation along the innermost dimension,
and requires that the resulting innermost dimension provide 2 elements per thread, distributed across registers.
In contrast, this `concat` op imposes no constraints on the concatenation dimension or the size of dimensions.

* sources: a list of the input tensors.

Example 1:

```mlir
#blocked = #ttg.blocked<{sizePerThread = [1, 8],
    threadsPerWarp = [8, 8], warpsPerCTA = [4, 1], order = [1, 0]}>
%0 = amdg.concat %arg0, %arg1: tensor<32x64xf32, #blocked>,tensor<32x64xf32, #blocked>,
  -> tensor<64x64xf32, #blocked>
```

Example 2:
```mlir
#src_layout = #ttg.linear<{register=[[0, 1], [0, 2], [0, 8], [0, 16], [0, 64], [64, 0]], lane=[[1, 0], [2, 0], [4, 0], [8, 0], [16, 0], [0, 4]], warp=[[0, 32], [32, 0]], block=[]}>
#dst_layout = #ttg.linear<{register=[[0, 1], [0, 2], [0, 8], [0, 16], [0, 64], [0, 128], [64, 0], [128, 0]], lane=[[1, 0], [2, 0], [4, 0], [8, 0], [16, 0], [0, 4]], warp=[[0, 32], [32, 0]], block=[]}>
%0 = amdg.concat %arg0, %arg1, %arg2, %arg3 : tensor<128x128xf16, #src_layout>, tensor<128x128xf16, #src_layout>, tensor<128x128xf16, #src_layout>,
                                                tensor<128x128xf16, #src_layout> -> tensor<256x256xf16, #dst_layout>
```


Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `sources` | variadic of ranked tensor of floating-point or integer or ptr values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of any non-token type values |



### `amdg.cond_barrier` (triton::amdgpu::CondBarrierOp)

_Conditionally set barriers to synchronize partial threads in a block_

Syntax:

```
operation ::= `amdg.cond_barrier` $pred attr-dict
```

condBarrierOp sets barrier instruction only when the given argument is true.
This provides a way to synchronize partial threads in a block, deliberately
diverges the execution sequences. However, user should guarantee all threads
converge at the end by calling condBarrierOp(true) with the remaining threads.
Conceptually, this is similar to having an execution barrier inside an if statement.
This op allows us to avoid blocking the whole block when suitable to help scheduling.
NB. This doesn't set any memory fence.

#### Operands:

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



### `amdg.extract_slice` (triton::amdgpu::ExtractSliceOp)

_Extract slice operation_

Syntax:

```
operation ::= `amdg.extract_slice` $source $static_offsets attr-dict `:` type($source) `to` type($result)
```

The "extract_slice" operation enables extracting a slice of a tensor in
registers.

The "extract_slice" operation supports the following arguments:

* source: the base tensor on which to create a view tensor
* offsets: offsets into the base tensor at which to create the view

In distributed layouts, tensors are divided into CTA tiles.
A CTA tile represents the smallest contiguous portion of a tensor that is
distributed across all threads and warps within a workgroup.
The ExtractSlice operation extracts a portion of the tensor that is a
multiple of CTA tiles.

The source and destination must have matching linear layouts at the CTA
tile level. This ensures that the extract_slice is a no-op, meaning no data
rearrangement between threads is required to extract the destination tensor
with the given shape and layout.

  +-------+-------+
  |  W0   |  W1   |
  |       |       |
  |   +   |   +   |
  |  W2   |  W3   |  <-- Single CTA tile (distributed across warps W0-W3)
  |       |       |
  |   +   |   +   |
  |       |       |
  +-------+-------+
  |          Source Tensor                    Extracted Slice
  |             .                           +--------------+
  |             .                           |  W0  |  W1   |
  |             .                           |      |       |
  |                                         |  +   |   +   |
  |                                         |  W2  |  W3   |
  |                                         |      |       |
  |                                         |  +   |   +   |
  |                                         |      |       |
  |                                         +-------+------+
  |                                         |  W0  |   W1  |
  |                                         |      |       |
  |                                         |  +   |   +   |
  |                                         |  W2     W3   |
  |                                         |      |       |
  |                                         |  +   |   +   |
  |                                         |      |       |
  |                                         +--------------+


This op is designed to work on logical tensors directly, avoiding the need
for complex layout reinterpretation or reshaping. For example, the tt.split
operation only supports splitting along the innermost dimension,
and requires that the resulting innermost dimension provide 2 elements per thread,
distributed across registers. In contrast, extract_slice op imposes no constraints
on the extraction dimension or the size of dimensions.

Example 1:

```mlir
#blocked = #ttg.blocked<{sizePerThread = [1, 8],
    threadsPerWarp = [4, 16], warpsPerCTA = [4, 1], order = [0, 1]}>
#blocked1 = #ttg.blocked<{sizePerThread = [1, 8],
    threadsPerWarp = [16, 4], warpsPerCTA = [4, 1], order = [0, 1]}>
%1 = ttg.convert_layout %0 : tensor<128x128xf16, #blocked>
    -> tensor<128x128xf16, #blocked1>
// create a slice of base tensor %1 with static offsets
%2 = amdg.extract_slice %0 [0, 0] :
  tensor<128x128xf16, #blocked1> to tensor<128x32xf16, #blocked1>
```

Example 1 shows how "extract_slice" operation may be used. In this example a
new slice of 128x32 is created. "extract_slice" works on tensors
where the desired slice has the same layout on a CTA tile as the source tensor.
"%0" cannot be sliced directly as the resulting slice does not satisfy this condition.
Therefore it needs to be converted to a layout suitable for slicing.
"#blocked1" layout is appropriate for this as it keeps the
sizePerThread the same thus keeping coalescing properties the same.
In order to utilize all threads in a warp, "threadsPerWarp" is set to
[16,4] for this new layout. This layout conversion carried out before
using "extract_slice" ensures slicing still uses all threads efficiently. The
size of the slice is determined by the result type.

Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>static_offsets</code></td><td>::mlir::DenseI64ArrayAttr</td><td>i64 dense array attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `source` | ranked tensor of any non-token type values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | ranked tensor of any non-token type values |



### `amdg.in_thread_transpose` (triton::amdgpu::InThreadTransposeOp)

_Perform transpose of register values belonging to each threads_

Syntax:

```
operation ::= `amdg.in_thread_transpose` $src attr-dict `:` type($src) `->` type($result)
```

This operation performs a layout transpose over values in registers per thread.
Specifically, given the input layout's blocked layout, it transposes the two last dimensions(rank-1 and rank-2)
along the register dimension of the underlying linear layout.

Conversion example:
* input layout: blocked layout with sizePerThread=[2, 2], order=[0, 1]. It's linear layout register bases = [[1, 0], [2, 0], [0, 1], [0, 2]]
* output layout: same thread and warp bases as in input, register bases = [[0, 1], [0, 2], [1, 0], [2, 0]]

This operation enables efficient coalesced loading from HBM with following vectorized writing to shared memory
in cases when HBM and shared memory order differ and target AMD hardware does not natively support this transposition.
This is a specific variant of ttg.convert_layout and will be converted to ttg.convert_layout when lowering to llvm.
We do not want this conversion to be optimized out, because we need to explicitly materialize instructions
to transpose within each thread after loading from HBM and before writing to shared memory.

Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Operands:

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

#### Results:

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



### `amdg.init_barrier` (triton::amdgpu::InitBarrierOp)

_Initialize a barrier in the given shared memory allocation._

Syntax:

```
operation ::= `amdg.init_barrier` $alloc `,` $count attr-dict `:` qualified(type($alloc))
```

Initializes a shared memory allocation with mbarrier information.
`alloc` is a descriptor to the shared memory allocation. `count` is the
number of arrives expected by the barrier.

Interfaces: `MBarrierOpInterface`

#### Attributes:

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

#### Operands:

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



### `amdg.local_load_packed_transposed` (triton::amdgpu::LocalLoadPackedTransposedOp)

_Load a transposed packed tensor from shared memory into a distributed tensor_

Syntax:

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

Requires a M/N packed and M/N contiguous tensor in shared memory and will yield a K packed K contiguous tensor in registers.
The packing change will change the shape of the tensor by doubling the M/N dimension and halving the K dimension.
For example if A is 16x64 in shared memory, the result of this operation will be 32x32.

Traits: `LocalLoadTrait`

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



### `amdg.masked_load` (triton::amdgpu::MaskedLoadOp)

_Masked load operation_

Syntax:

```
operation ::= `amdg.masked_load` $ptr `,` $mask `,` $falseVal (`,` $multicastMask^)?
              oilist(`cacheModifier` `=` $cache)
              (`volatile` $isVolatile^)?
              (`forceNoAlias` $forceNoAlias^)?
              attr-dict `:` functional-type(operands, results)
```

Load operation with masking and multicast support. If the mask is true, loads from the given pointer. Works with LLVM types as a utility op for making LLVM conversion easier.
On architectures supporting multicast, the `multicastMask`specifies which CTAs in the cluster request the same data. This allows the hardware to efficiently broadcast the
data to multiple CTAs in the cluster.

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cache</code></td><td>::mlir::triton::CacheModifierAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7</td></tr>
<tr><td><code>isVolatile</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>forceNoAlias</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `ptr` | LLVM pointer type |
| `mask` | 1-bit signless integer |
| `falseVal` | LLVM dialect-compatible type |
| `multicastMask` | 32-bit signless integer |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | LLVM dialect-compatible type |



### `amdg.masked_store` (triton::amdgpu::MaskedStoreOp)

_Masked Store operation_

Syntax:

```
operation ::= `amdg.masked_store` $ptr `,` $value `,` $mask
              oilist(`cacheModifier` `=` $cache)
              (`forceNoAlias` $forceNoAlias^)?
              attr-dict `:` type(operands)
```

Store operation with masking support. If the mask is true, Store from the given pointer. Works with LLVM types as a utility op for making LLVM conversion easier.

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cache</code></td><td>::mlir::triton::CacheModifierAttr</td><td>allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7</td></tr>
<tr><td><code>forceNoAlias</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `ptr` | LLVM pointer type |
| `value` | LLVM dialect-compatible type |
| `mask` | 1-bit signless integer |



### `amdg.memory_counter_wait` (triton::amdgpu::MemoryCounterWaitOp)

_Wait for specified hardware counters_

Syntax:

```
operation ::= `amdg.memory_counter_wait` oilist( `load` `(` $load `)` | `store` `(` $store `)` | `ds` `(` $ds `)` ) attr-dict
```

Wait for the specified counters to be less-than or equal-to the provided
values before continuing.

Counters can lower to different instructions on different architectires,
including clamping to the some HW supported max value or combining multiple
counters into one.

#### Attributes:

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



### `amdg.scaled_downcast_fp4` (triton::amdgpu::ScaledDowncastFp4Op)

_Scale and downcast floating-point values to packed FP4_

Syntax:

```
operation ::= `amdg.scaled_downcast_fp4` $input `scale` $scale attr-dict
              `:` type($input) `,` type($scale) `->` type($output)
```

Divide FP16, BF16, or FP32 values by an E8M0 scale and convert them to
packed e2m1 values. The result packs two consecutive fp4 values per i8
along `axis`: the element at even position provides the low nibble and the
following element the high nibble, so the result extent along `axis` is
half the input extent (the inverse of `scaled_upcast_fp4`).

The `axis` attribute identifies the dimension along which the fp4 elements
are packed and along which scales are shared. The scale must be compact
along `axis`: one raw E8M0 i8 payload covers multiple consecutive packed
bytes along that axis (unlike :func:`scaled_upcast_fp4`, expanded scales
with one scale per output element are not supported). Each scale block along
`axis` must cover a multiple of 8 fp4 values.

This will be hardware accelerated on architectures supporting
`v_cvt_scalef32_pk_*` instructions (e.g. CDNA4 and CDNA5).

Traits: `AlwaysSpeculatableImplTrait`

Interfaces: `CastFpOpInterface`, `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 |
| :-----: | ----------- |
| `input` | ranked tensor of 16-bit float or bfloat16 type or 32-bit float values |
| `scale` | ranked tensor of 8-bit signless integer values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 8-bit signless integer values |



### `amdg.scaled_downcast_fp8` (triton::amdgpu::ScaledDowncastFp8Op)

_Scale and downcast floating-point values to fp8_

Syntax:

```
operation ::= `amdg.scaled_downcast_fp8` $input `scale` $scale attr-dict
              `:` type($input) `,` type($scale) `->` type($output)
```

Divide FP16, BF16, or FP32 values by an E8M0 scale and convert them to
fp8 (e4m3) or bf8 (e5m2) values. This is the elementwise inverse of
`scaled_upcast_fp8`: input and result share the same shape and encoding,
and each fp8 result occupies its own byte (no packing).

The `axis` attribute identifies the dimension along which scales are
shared. The scale must be compact along `axis`: one raw E8M0 i8 payload
covers multiple consecutive output elements along that axis (unlike
:func:`scaled_upcast_fp8`, expanded scales with one scale per output element
are not supported). Each scale block along `axis` must cover a multiple of 8
output elements.

This will be hardware accelerated on architectures supporting
`v_cvt_scalef32_pk_{fp8,bf8}_*` instructions (e.g. CDNA4 and CDNA5).

Traits: `AlwaysSpeculatableImplTrait`

Interfaces: `CastFpOpInterface`, `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 |
| :-----: | ----------- |
| `input` | ranked tensor of 16-bit float or bfloat16 type or 32-bit float values |
| `scale` | ranked tensor of 8-bit signless integer values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of f8E4M3FN type or f8E5M2 type values |



### `amdg.scaled_upcast_fp4` (triton::amdgpu::ScaledUpcastFp4Op)

_Upcast fp4 and then multiply scale_

Syntax:

```
operation ::= `amdg.scaled_upcast_fp4` $input `scale` $scale attr-dict
              `:` type($input) `,` type($scale) `->` type($output)
```

Upcast fp4 (e2m1) values packed as i8 values and multiply with the given
E8M0 scale encoded as BF16. This maps to `v_cvt_scalef32_*` intrinsics
on the AMD CDNA4 architecture.

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`

Interfaces: `CastFpOpInterface`, `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 |
| :-----: | ----------- |
| `input` | ranked tensor of 8-bit signless integer values |
| `scale` | ranked tensor of bfloat16 type or 8-bit signless integer values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 16-bit float or bfloat16 type values |



### `amdg.scaled_upcast_fp8` (triton::amdgpu::ScaledUpcastFp8Op)

_Upcast Fp8 and then multiply scale_

Syntax:

```
operation ::= `amdg.scaled_upcast_fp8` $input `scale` $scale attr-dict
              `:` type($input) `,` type($scale) `->` type($output)
```

Upcast fp8 (e4m3/e5m2) values and multiply with the given E8M0 scale
encoded as BF16. This maps to `v_cvt_scalef32_*` intrinsics
on the AMD CDNA4 architecture.

Traits: `AlwaysSpeculatableImplTrait`, `Elementwise`, `SameOperandsAndResultEncoding`, `SameOperandsAndResultShape`

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

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of f8E4M3FN type or f8E5M2 type values |
| `scale` | ranked tensor of bfloat16 type or 8-bit signless integer values |

#### Results:

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 16-bit float or bfloat16 type values |



### `amdg.tdm_prefetch` (triton::amdgpu::TDMPrefetchOp)

_Prefetch data based on a TDM descriptor from global memory to L2._

Syntax:

```
operation ::= `amdg.tdm_prefetch` $desc `[` $indices `]` `,` $pred `,` `speculative` `=` $speculative
              (`returnOffsets` $returnOffsets^)?
              attr-dict `:` qualified(type($desc))
              (`->` type($maybeOffsets)^)?
```

This operation prefetches data from global memory to L2. It is analogous to the AsyncTDMCopyGlobalToLocalOp,
but it does not copy the data to local memory and instead only prefetches the data into the L2 cache.
Speculative prefetches can generate more efficient assembly because they do not require out of bounds checks.
However, they are dropped by the hardware in case the virtual address translation is not already cached at CU level.

Interfaces: `InferTypeOpInterface`, `MemoryEffectOpInterface (MemoryEffectOpInterface)`, `PredicatedOpInterface`

Effects: `MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::amd::L2Cache}`

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>speculative</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>returnOffsets</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `desc` | Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |
| `indices` | variadic of 32-bit signless integer |
| `pred` | 1-bit signless integer |

#### Results:

| Result | Description |
| :----: | ----------- |
| `maybeOffsets` | ranked tensor of floating-point or integer or ptr values |



### `amdg.update_tensor_descriptor` (triton::amdgpu::UpdateTensorDescriptorOp)

_Update selected fields of a TDM descriptor; produce a new descriptor SSA value_

Syntax:

```
operation ::= `amdg.update_tensor_descriptor` $desc
              oilist(
              `add_offsets` `=` `[` $add_offsets `]`
              | `set_bounds` `=` `[` $set_bounds `]`
              | `pred` `=` $pred
              )
              attr-dict `:` qualified(type($desc))
```

Returns a new TDM descriptor with selected fields rewritten.  Each
parameter is independently optional; only the fields the caller names
are written.  Everything else is inherited from the input descriptor.

| parameter      | kind        | effect on descriptor                                            |
|----------------|-------------|-----------------------------------------------------------------|
| `add_offsets`  | incremental | `global_addr += sum(add_offsets[i] * stride[i] * elem_size)`    |
| `set_bounds`   | rewrite     | `tensor_dim[i] = set_bounds[i]` (absolute)                      |
| `pred`         | rewrite     | `pred = pred`                                                   |
| `clamp_bounds` | derive      | `tensor_dim[i] = max(0, tensor_dim[i] - add_offsets[i])`        |

The LDS destination and the TDM barrier (mbarrier) are op-level operands
on the async_tdm_copy ops, not descriptor fields, so they are not settable
here.

`add_offsets` is incremental (deltas in element units) and does not
touch `tensor_dim` — for OOB-correct loops, set `set_bounds` once
outside the loop (interior tiles see a loop-invariant tensor_dim
that hoists), or set `set_bounds` per iteration when OOB must fire.

`clamp_bounds` is a unit attribute that, when present, additionally
shrinks `tensor_dim` by `add_offsets` so the descriptor carries the
OOB extent of the advanced tile. It requires `add_offsets` and is
mutually exclusive with `set_bounds`. It lets a descriptor derive its
OOB bounds from the advance without a separate `set_bounds`, and works
for host-built descriptors whose shape is only known at runtime.

Examples:
```
// K-loop interior: bump tile position only
%d1 = amdg.update_tensor_descriptor %d, add_offsets=[%c0, %BLOCK_K]
        : !tt.tensordesc<64x64xf16, #shared>

// Prologue: position at first tile + set the predicate
%d1 = amdg.update_tensor_descriptor %d,
        add_offsets=[%pid_m_off, %c0],
        pred=%p
        : !tt.tensordesc<64x64xf16, #shared>

// Peel epilogue: install real OOB extent for the partial last tile
%d1 = amdg.update_tensor_descriptor %d, set_bounds=[%m_remain, %k_remain]
        : !tt.tensordesc<64x64xf16, #shared>
```

Traits: `AlwaysSpeculatableImplTrait`, `AttrSizedOperandSegments`

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>clamp_bounds</code></td><td>::mlir::UnitAttr</td><td>unit attribute</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `desc` | Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |
| `add_offsets` | variadic of 32-bit signless integer |
| `set_bounds` | variadic of 32-bit signless integer |
| `pred` | 32-bit signless integer |

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | Tensor descriptor type (`::mlir::triton::TensorDescType`) in Triton IR type system |



### `amdg.wait_barrier` (triton::amdgpu::WaitBarrierOp)

_Wait until the mbarrier phase completes._

Syntax:

```
operation ::= `amdg.wait_barrier` $alloc `,` $phase attr-dict `:` qualified(type($alloc))
```

Blocks the program progress until the mbarrier object in `alloc` completes
its current phase.

Interfaces: `MBarrierOpInterface`

#### Operands:

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