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

### `amdgpu.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 ::= `amdgpu.buffer_atomic_rmw` $atomic_rmw_op `,` $sem `,` $scope `,` $value `,` $ptr `[` $offsets `]` (`,` $mask^)?
              (`stride` `=` $stride^)?
              attr-dict `:` 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`

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


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

_Load from a scalar base pointer and a tensor offset_

Syntax:

```
operation ::= `amdgpu.buffer_load` $ptr `[` $offsets `]` (`,` $mask^)? (`,` $other^)?
              oilist(`cacheModifier` `=` $cache)
              (`stride` `=` $stride^)?
              attr-dict `:` 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.

Traits: `AttrSizedOperandSegments`, `SameLoadStoreOperandsAndResultEncoding`

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


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

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

Syntax:

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

AMD Buffer load operation. Similar to amdgpu.buffer_load op but directly wirtes to shared memory instead of into registers.

Traits: `AttrSizedOperandSegments`

Interfaces: `InferTypeOpInterface`

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


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

_Store into scalar base pointer and a tensor offset_

Syntax:

```
operation ::= `amdgpu.buffer_store` $value `,` $ptr `[` $offsets `]` (`,` $mask^)?
              oilist(`cacheModifier` `=` $cache)
              (`stride` `=` $stride^)?
              attr-dict `:` 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.

Traits: `AttrSizedOperandSegments`, `SameLoadStoreOperandsEncoding`

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


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

_Concat operation_

Syntax:

```
operation ::= `amdgpu.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 = amdgpu.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 = amdgpu.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 type values |


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

_Conditionally set barriers to synchronize partial threads in a block_

Syntax:

```
operation ::= `amdgpu.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 |


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

_Extract slice operation_

Syntax:

```
operation ::= `amdgpu.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 = amdgpu.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 type values |

#### Results:

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


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

_Perform transpose of register values belonging to each threads_

Syntax:

```
operation ::= `amdgpu.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 |


### `amdgpu.instruction_sched_hint` (triton::amdgpu::InstructionSchedHint)

_A placeholder op for instruction scheduling hints within a basic block_

Syntax:

```
operation ::= `amdgpu.instruction_sched_hint` attr-dict
```

A placeholder op for instruction scheduling hints applied to instructions within
a basic block where the placeholder op is located. This op is primarily intended
to be used to adjust instruction scheduling inside the resulting main loop
of a `tt.dot` operation. It's easier to identify dot ops at a high level and, thus,
to mark intended scheduling regions. The hint ops are eventually lowered
into LLVM AMDGPU instruction scheduling primitives, which are meant to control
how different kinds of instructions (valu/mfma, global/shared memory, etc.) should
interleave for better instruction level parallelism.

#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>variant</code></td><td>::mlir::triton::amdgpu::SchedHintAttr</td><td>Instruction Scheduling Hints for AMD GPUs</td></tr>
<tr><td><code>numDsReadsA</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>numDsReadsB</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>numDsWritesA</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>numDsWritesB</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>numGlobalLoadsA</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>numGlobalLoadsB</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>isBufferLoadsAEnabled</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>isBufferLoadsBEnabled</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>numMMAs</code></td><td>::mlir::triton::amdgpu::InstCounterAttr</td><td><details><summary>An instruction counter attribute.</summary>{{% markdown %}}
    The attribute holds the number of issued LLVM instructions of a specific kind as well as
    the data type.
  {{% /markdown %}}</details></td></tr>
</table>


### `amdgpu.upcast_mxfp` (triton::amdgpu::UpcastMXFPOp)

_Convert an mxfp tensor to bf16/fp16_

Syntax:

```
operation ::= `amdgpu.upcast_mxfp` $src `,` $scale  `fp_type` `=` $fp_type attr-dict `:` type($src) `,` type($scale) `->` type($result)
```

Compute the bf16 encoded in the given mxfp number as per
https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf

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>fp_type</code></td><td>::mlir::triton::ScaleDotElemTypeAttr</td><td>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6</td></tr>
<tr><td><code>fastMath</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### Operands:

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

#### Results:

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