TritonInstrumentOps

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:

Attribute	MLIR Type	Description
offsets	::mlir::DenseI32ArrayAttr	i32 dense array attribute

Results:

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