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

### `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.ld_acquire` (triton::nvgpu::LoadAcquireOp)

Syntax:

```
operation ::= `nvgpu.ld_acquire` $sem `,` $scope `,` $addr (`,` $mask^)? attr-dict `:` functional-type($addr, $result)
```

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

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `addr` | LLVM pointer in address space 1 |
| `mask` | 1-bit signless integer |

#### Results:

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


### `nvgpu.tensor_memory_base` (triton::nvgpu::TensorMemoryBaseAddress)

Syntax:

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

Op to represent base address of tensor memory in a kernel.
This is used to simplify lowering from TritonGPU to LLVM.

Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Results:

| Result | Description |
| :----: | ----------- |
| `result` | LLVM pointer in address space 6 |


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

Syntax:

```
operation ::= `nvgpu.wgmma` $opA `,` $opB `,` $useC (`,` $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>wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'</td></tr>
<tr><td><code>eltTypeA</code></td><td>::mlir::triton::nvgpu::WGMMAEltTypeAttr</td><td>wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'</td></tr>
<tr><td><code>eltTypeB</code></td><td>::mlir::triton::nvgpu::WGMMAEltTypeAttr</td><td>wgmma operand type, either 's8', 's32', 'e4m3', 'e5m2', 'f16', 'bf16', 'tf32', or 'f32'</td></tr>
<tr><td><code>layoutA</code></td><td>::mlir::triton::nvgpu::WGMMALayoutAttr</td><td>wgmma layout, either 'row' or 'col'</td></tr>
<tr><td><code>layoutB</code></td><td>::mlir::triton::nvgpu::WGMMALayoutAttr</td><td>wgmma layout, either 'row' or 'col'</td></tr>
</table>

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `opA` | wgmma operand A/B type |
| `opB` | wgmma operand A/B type |
| `useC` | 1-bit signless integer |
| `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 |


### `nvgpu.warp_id` (triton::nvgpu::WarpIdOp)

Syntax:

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

Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Results:

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