# NVGPUOps<!-- Autogenerated by mlir-tblgen; don't manually edit -->
### `nvgpu.cluster_arrive` (triton::nvgpu::ClusterArriveOp)

Syntax:

```
operation ::= `nvgpu.cluster_arrive` attr-dict
```


#### Attributes:

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

### `nvgpu.cluster_id` (triton::nvgpu::ClusterCTAIdOp)

Syntax:

```
operation ::= `nvgpu.cluster_id` attr-dict
```


Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Results:

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

### `nvgpu.cluster_wait` (triton::nvgpu::ClusterWaitOp)

Syntax:

```
operation ::= `nvgpu.cluster_wait` attr-dict
```


### `nvgpu.fence_async_shared` (triton::nvgpu::FenceAsyncSharedOp)

Syntax:

```
operation ::= `nvgpu.fence_async_shared` attr-dict
```


#### Attributes:

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

### `nvgpu.load_dsmem` (triton::nvgpu::LoadDSmemOp)

Syntax:

```
operation ::= `nvgpu.load_dsmem` operands attr-dict `:` functional-type(operands, results)
```


Interfaces: `MemoryEffectOpInterface (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{MemoryEffects::Read on ::mlir::SideEffects::DefaultResource}`

#### Attributes:

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `addr` | LLVM pointer type
| `ctaId` | 32-bit signless integer

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | LLVM type with size

### `nvgpu.store_dsmem` (triton::nvgpu::StoreDSmemOp)

Syntax:

```
operation ::= `nvgpu.store_dsmem` operands attr-dict `:` type(operands)
```


Interfaces: `MemoryEffectOpInterface (MemoryEffectOpInterface)`

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `addr` | LLVM pointer type
| `ctaId` | 32-bit signless integer
| `values` | variadic of LLVM type with size
| `pred` | 1-bit signless integer

### `nvgpu.stmatrix` (triton::nvgpu::StoreMatrixOp)

Syntax:

```
operation ::= `nvgpu.stmatrix` operands attr-dict `:` type(operands)
```


Interfaces: `MemoryEffectOpInterface (MemoryEffectOpInterface)`

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `addr` | LLVM pointer in address space 3
| `datas` | variadic of 32-bit signless integer

### `nvgpu.wgmma_commit_group` (triton::nvgpu::WGMMACommitGroupOp)

Syntax:

```
operation ::= `nvgpu.wgmma_commit_group` attr-dict
```


### `nvgpu.wgmma_fence` (triton::nvgpu::WGMMAFenceOp)

Syntax:

```
operation ::= `nvgpu.wgmma_fence` attr-dict
```


### `nvgpu.wgmma` (triton::nvgpu::WGMMAOp)

Syntax:

```
operation ::= `nvgpu.wgmma` $opA `,` $opB (`,` $opC^)? attr-dict `:` functional-type(operands, $res)
```


#### Attributes:

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>m</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>n</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>k</code></td><td>::mlir::IntegerAttr</td><td>32-bit signless integer attribute</td></tr>
<tr><td><code>eltTypeC</code></td><td>::mlir::triton::nvgpu::WGMMAEltTypeAttr</td><td><details><summary>wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'</summary>{{% markdown %}}Enum cases:
* s8 (`s8`)
* s32 (`s32`)
* e4m3 (`e4m3`)
* e5m2 (`e5m2`)
* f16 (`f16`)
* bf16 (`bf16`)
* tf32 (`tf32`)
* f32 (`f32`){{% /markdown %}}</details></td></tr>
<tr><td><code>eltTypeA</code></td><td>::mlir::triton::nvgpu::WGMMAEltTypeAttr</td><td><details><summary>wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'</summary>{{% markdown %}}Enum cases:
* s8 (`s8`)
* s32 (`s32`)
* e4m3 (`e4m3`)
* e5m2 (`e5m2`)
* f16 (`f16`)
* bf16 (`bf16`)
* tf32 (`tf32`)
* f32 (`f32`){{% /markdown %}}</details></td></tr>
<tr><td><code>eltTypeB</code></td><td>::mlir::triton::nvgpu::WGMMAEltTypeAttr</td><td><details><summary>wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'</summary>{{% markdown %}}Enum cases:
* s8 (`s8`)
* s32 (`s32`)
* e4m3 (`e4m3`)
* e5m2 (`e5m2`)
* f16 (`f16`)
* bf16 (`bf16`)
* tf32 (`tf32`)
* f32 (`f32`){{% /markdown %}}</details></td></tr>
<tr><td><code>layoutA</code></td><td>::mlir::triton::nvgpu::WGMMALayoutAttr</td><td><details><summary>wgmma layout, either 'row' or 'col'</summary>{{% markdown %}}Enum cases:
* row (`row`)
* col (`col`){{% /markdown %}}</details></td></tr>
<tr><td><code>layoutB</code></td><td>::mlir::triton::nvgpu::WGMMALayoutAttr</td><td><details><summary>wgmma layout, either 'row' or 'col'</summary>{{% markdown %}}Enum cases:
* row (`row`)
* col (`col`){{% /markdown %}}</details></td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `opA` | wgmma operand A/B type
| `opB` | wgmma operand A/B type
| `opC` | LLVM structure type

#### Results:

| Result | Description |
| :----: | ----------- |
| `res` | LLVM structure type

### `nvgpu.wgmma_wait_group` (triton::nvgpu::WGMMAWaitGroupOp)

Syntax:

```
operation ::= `nvgpu.wgmma_wait_group` $input attr-dict `:` type($input)
```


Interfaces: `InferTypeOpInterface`

#### Attributes:

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `input` | LLVM structure type

#### Results:

| Result | Description |
| :----: | ----------- |
| `output` | LLVM structure type