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

### `tti.experimental_check_async_write_with_mbar_shared` (triton::instrument::ExperimentalCheckAsyncWriteWithMbarSharedOp)

_Check if writing to a buffer guarded by a mbar is valid_

Syntax:

```
operation ::= `tti.experimental_check_async_write_with_mbar_shared` $buffer `,` $mbar `{` $buffers `,` $states `,` $barriers `}` attr-dict `:` type($buffer) `,` type($mbar) `,` type($buffers) `,` type($states) `,` type($barriers) `->` type($outStates) `,` type($outBarriers)
```

Check if writing to a shared memory buffer guarded by a mbar is valid.
Update the buffer state and assert if the buffer is being read or written.

Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `buffer` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `mbar` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `buffers` | ranked tensor of floating-point or integer or ptr values |
| `states` | ranked tensor of floating-point or integer or ptr values |
| `barriers` | ranked tensor of floating-point or integer or ptr values |

#### Results:

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


### `tti.experimental_check_wait_mbar` (triton::instrument::ExperimentalCheckWaitMbarOp)

_Check if waiting on a mbar is valid and update the barrier state_

Syntax:

```
operation ::= `tti.experimental_check_wait_mbar` $mbar `{` $states `,` $barriers `}` attr-dict `:` type($mbar) `,` type($states) `,` type($barriers) `->` type($outStates) `,` type($outBarriers)
```

Check if waiting on a mbar is valid and update the barrier state.

Traits: `AlwaysSpeculatableImplTrait`

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

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `mbar` | memory descriptor type (`::mlir::triton::gpu::MemDescType`) in Triton IR type system |
| `barriers` | ranked tensor of floating-point or integer or ptr values |
| `states` | ranked tensor of floating-point or integer or ptr values |

#### Results:

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


### `tti.experimental_shared_buffer_pointers` (triton::instrument::ExperimentalSharedBufferPointersOp)

_Definte an array of pointers to shared memory buffers_

Syntax:

```
operation ::= `tti.experimental_shared_buffer_pointers` attr-dict `:` type($result)
```

Create a tensor of pointers to shared memory buffers.

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>offsets</code></td><td>::mlir::DenseI32ArrayAttr</td><td>i32 dense array attribute</td></tr>
</table>

#### Results:

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