TritonInstrumentOps

tti.dot_i8 (triton::instrument::DotI8Op)

Non-saturating NVIDIA MMAv2 i8 dot

Syntax:

operation ::= `tti.dot_i8` $a `,` $b `,` $c `,` `aSigned` `=` $aSigned `,` `bSigned` `=` $bSigned
              attr-dict `:` type($a) `*` type($b) `->` type($d)

Performs a wrapping i8 matrix multiplication into an i32 accumulator using NVIDIA MMAv2. The A and B operands have independent signedness.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, DotOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
aSigned	::mlir::BoolAttr	bool attribute
bSigned	::mlir::BoolAttr	bool attribute

Operands:

Operand	Description
a	ranked tensor of 8-bit signless integer values
b	ranked tensor of 8-bit signless integer values
c	ranked tensor of 32-bit signless integer values

Results:

Result	Description
d	ranked tensor of 32-bit signless integer values

tti.experimental_assert_uniform (triton::instrument::ExperimentalAssertUniformOp)

Assert the uniform condition

Syntax:

operation ::= `tti.experimental_assert_uniform` $condition `,` $message attr-dict-with-keyword

Assert that the condition is true given all threads in the warp group have the same value, so only one thread needs to evaluate the assert and print the message.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:

Attribute	MLIR Type	Description
message	::mlir::StringAttr	string attribute

Operands:

Operand	Description
condition	1-bit signless integer

tti.experimental_buffer_descriptors (triton::instrument::ExperimentalBufferDescriptorsOp)

Define an array of buffer descriptors

Syntax:

operation ::= `tti.experimental_buffer_descriptors` $offsets `,` $lengths `,` $memType attr-dict `:` type($result)

Create a tensor of buffer descriptors packing 32-bit pointer offsets and 32-bit lengths into 64-bit elements.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
offsets	::mlir::DenseI32ArrayAttr	i32 dense array attribute
lengths	::mlir::DenseI32ArrayAttr	i32 dense array attribute
memType	::mlir::triton::instrument::MemTypeAttr	allowed 32-bit signless integer cases: 0, 1

Results:

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

tti.experimental_cluster_cta_id (triton::instrument::ExperimentalClusterCTAIdOp)

Get the CTA id within the current cluster

Syntax:

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

Return the cluster-local CTA id used to index ConSan’s multi-CTA scratch slabs. For single-CTA kernels this is always zero.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:

Result	Description
result	32-bit signless integer

tti.experimental_fpsan_embed (triton::instrument::ExperimentalFPSanEmbedOp)

Embed float value in the fpsan integer ring

Syntax:

operation ::= `tti.experimental_fpsan_embed` $val attr-dict `:` functional-type(operands, $result)

Traits: AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
val	floating-point or ranked tensor of floating-point values

Results:

Result	Description
result	integer or ranked tensor of integer values

tti.experimental_fpsan_unembed (triton::instrument::ExperimentalFPSanUnembedOp)

Unembed fpsan integer payload as a float value

Syntax:

operation ::= `tti.experimental_fpsan_unembed` $val attr-dict `:` functional-type(operands, $result)

Traits: AlwaysSpeculatableImplTrait, Elementwise, SameOperandsAndResultEncoding, SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
val	integer or ranked tensor of integer values

Results:

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

tti.experimental_gsan_atomic_cas (triton::instrument::ExperimentalGSanAtomicCASOp)

Lower a GSan-instrumented atomic cas

Syntax:

operation ::= `tti.experimental_gsan_atomic_cas` $sem `,` $scope `,` $ptr `,` $cmp `,` $val attr-dict `:`
              functional-type(operands, $result)

Traits: SameOperandsAndResultEncoding, SameOperandsAndResultShape

Interfaces: AtomicOpInterface

Attributes:

Attribute	MLIR Type	Description
sem	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

Operands:

Operand	Description
ptr	ptr or ranked tensor of ptr values
cmp	floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values
val	floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values

Results:

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

tti.experimental_gsan_atomic_load (triton::instrument::ExperimentalGSanAtomicLoadOp)

Lower a GSan-instrumented atomic load

Syntax:

operation ::= `tti.experimental_gsan_atomic_load` $sem `,` $scope `,` $ptr (`,` $mask^)? attr-dict `:`
              functional-type(operands, $result)

Traits: SameLoadStoreOperandsAndResultEncoding, SameLoadStoreOperandsAndResultShape

Interfaces: AtomicOpInterface, InferTypeOpInterface

Attributes:

Attribute	MLIR Type	Description
sem	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

Operands:

Operand	Description
ptr	ptr or ranked tensor of ptr values
mask	1-bit signless integer or ranked tensor of 1-bit signless integer values

Results:

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

tti.experimental_gsan_atomic_poll (triton::instrument::ExperimentalGSanAtomicPollOp)

Instrument a completed atomic poll for GSan

Syntax:

operation ::= `tti.experimental_gsan_atomic_poll` $sem `,` $scope `,` $ptr `,` $matched attr-dict `:` qualified(type($ptr))

Records the successful read performed by an atomic poll after the hardware polling loop has completed.

Traits: SameOperandsEncoding

Attributes:

Attribute	MLIR Type	Description
sem	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

Operands:

Operand	Description
ptr	ptr or ranked tensor of ptr values
matched	1-bit signless integer or ranked tensor of 1-bit signless integer values

tti.experimental_gsan_atomic_rmw (triton::instrument::ExperimentalGSanAtomicRMWOp)

Lower a GSan-instrumented atomic rmw

Syntax:

operation ::= `tti.experimental_gsan_atomic_rmw` $atomic_rmw_op `,` $sem `,` $scope `,` $ptr `,` $val (`,` $mask^)? attr-dict `:`
              functional-type(operands, $result)

Traits: SameOperandsAndResultEncoding, SameOperandsAndResultShape

Interfaces: AtomicOpInterface

Attributes:

Attribute	MLIR Type	Description
atomic_rmw_op	::mlir::triton::RMWOpAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4, 5, 6, 7, 8, 9, 10
sem	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

Operands:

Operand	Description
ptr	ptr or ranked tensor of ptr values
val	floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values
mask	1-bit signless integer or ranked tensor of 1-bit signless integer values

Results:

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

tti.experimental_gsan_atomic_store (triton::instrument::ExperimentalGSanAtomicStoreOp)

Lower a GSan-instrumented atomic store

Syntax:

operation ::= `tti.experimental_gsan_atomic_store` $sem `,` $scope `,` $ptr `,` $value (`,` $mask^)? attr-dict `:`
              type($ptr)

Traits: SameLoadStoreOperandsEncoding, SameLoadStoreOperandsShape

Interfaces: AtomicOpInterface

Attributes:

Attribute	MLIR Type	Description
sem	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

Operands:

Operand	Description
ptr	ptr or ranked tensor of ptr values
value	floating-point or ranked tensor of floating-point values or integer or ranked tensor of integer values or ptr or ranked tensor of ptr values
mask	1-bit signless integer or ranked tensor of 1-bit signless integer values

tti.experimental_gsan_atomic_tensordesc_access (triton::instrument::ExperimentalGSanAtomicTensorDescAccessOp)

Instrument an atomic TMA descriptor access for GSan

Syntax:

operation ::= `tti.experimental_gsan_atomic_tensordesc_access` $desc `[` $coords `]` `,` $sem `,` $scope attr-dict `:`
              qualified(type($desc))

Passes a native tensor descriptor and coordinates to the GSan runtime for atomic instrumentation with the specified memory semantics and scope.

Attributes:

Attribute	MLIR Type	Description
sem	::mlir::triton::MemSemanticAttr	allowed 32-bit signless integer cases: 1, 2, 3, 4
scope	::mlir::triton::MemSyncScopeAttr	allowed 32-bit signless integer cases: 1, 2, 3

Operands:

Operand	Description
desc	tensor descriptor type
coords	variadic of 32-bit signless integer

tti.experimental_gsan_cluster_barrier_init (triton::instrument::ExperimentalGSanClusterBarrierInitOp)

Initialize a GSan cluster-barrier rendezvous

Syntax:

operation ::= `tti.experimental_gsan_cluster_barrier_init` $scratch attr-dict `:` type($scratch)

Initializes cluster-shared profile scratch used to synchronize GSan vector clocks. This operation is internal to GSan instrumentation.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Operands:

Operand	Description
scratch	ptr

tti.experimental_gsan_cluster_barrier_sync (triton::instrument::ExperimentalGSanClusterBarrierSyncOp)

Synchronize GSan clocks at a cluster barrier

Syntax:

operation ::= `tti.experimental_gsan_cluster_barrier_sync` $scratch attr-dict `:` type($scratch)

Joins the vector clocks of all CTAs participating in a non-relaxed cluster barrier. This operation is internal to GSan instrumentation.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory, MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Operands:

Operand	Description
scratch	ptr

tti.experimental_gsan_grid_dependency_wait (triton::instrument::ExperimentalGSanGridDependencyWaitOp)

Acquire completed predecessor clocks after a grid dependency wait

Syntax:

operation ::= `tti.experimental_gsan_grid_dependency_wait` attr-dict

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory, MemoryEffects::Write on ::mlir::triton::GlobalMemory}

tti.experimental_gsan_indexed_tensordesc_access (triton::instrument::ExperimentalGSanIndexedTensorDescAccessOp)

Instrument an indexed TMA gather or scatter for GSan

Syntax:

operation ::= `tti.experimental_gsan_indexed_tensordesc_access` $desc `[` $indices `,` $offset `]` `,` $isStore (`,` $pred^)? attr-dict
              `:` qualified(type($desc)) `,` type($indices)

Passes a native tensor descriptor and its row indices directly to the GSan runtime, which checks row and column bounds before updating shadow memory.

Attributes:

Attribute	MLIR Type	Description
isStore	::mlir::BoolAttr	bool attribute

Operands:

Operand	Description
desc	tensor descriptor type
indices	ranked tensor of 32-bit signless integer values
offset	32-bit signless integer
pred	1-bit signless integer

tti.experimental_gsan_init (triton::instrument::ExperimentalGSanInitOp)

Initialize GSan thread

Syntax:

operation ::= `tti.experimental_gsan_init` $acquireStreamClock attr-dict

Attributes:

Attribute	MLIR Type	Description
acquireStreamClock	::mlir::BoolAttr	bool attribute

tti.experimental_gsan_kernel_exit (triton::instrument::ExperimentalGSanKernelExitOp)

Publish a completed GSan thread clock to its launch stream

Syntax:

operation ::= `tti.experimental_gsan_kernel_exit` attr-dict

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

tti.experimental_gsan_mbarrier_arrive (triton::instrument::ExperimentalGSanMBarrierArriveOp)

Record an mbarrier arrival in GSan

Syntax:

operation ::= `tti.experimental_gsan_mbarrier_arrive` $scratch `,` $barrier `,` $pred `,` $count attr-dict `:` type($scratch) `,`
              qualified(type($barrier))

Records an arrival for every mbarrier targeted by the corresponding hardware operation. When publishClock is true, the issuing CTA’s vector clock is also published without advancing its local epoch. Waiters acquire the issuing CTA’s completed epochs and its imported clocks, but not its current epoch. A false value records only phase progress for asynchronous completion signals that provide no generic memory ordering. multicastMasks contains CTA-id broadcast masks whose recipient bitsets are unioned. When multicast is false, the arrival targets the leader selected by the barrier’s CGA layout. This operation is internal to GSan instrumentation.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory, MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:

Attribute	MLIR Type	Description
count	::mlir::IntegerAttr	32-bit signless integer attribute
multicastMasks	::mlir::DenseI32ArrayAttr	i32 dense array attribute
multicast	::mlir::BoolAttr	bool attribute
sourceBroadcastMask	::mlir::IntegerAttr	32-bit signless integer attribute
publishClock	::mlir::BoolAttr	bool attribute

Operands:

Operand	Description
scratch	ptr
barrier	memory descriptor type (::mlir::triton::gpu::MemDescType) in Triton IR type system
pred	1-bit signless integer

tti.experimental_gsan_mbarrier_init (triton::instrument::ExperimentalGSanMBarrierInitOp)

Initialize an mbarrier in GSan

Syntax:

operation ::= `tti.experimental_gsan_mbarrier_init` $scratch `,` $barrier `,` $expectedCount attr-dict `:` type($scratch) `,`
              qualified(type($barrier))

Associates an mbarrier with its expected arrival count. This operation is internal to GSan instrumentation.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory, MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:

Attribute	MLIR Type	Description
expectedCount	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:

Operand	Description
scratch	ptr
barrier	memory descriptor type (::mlir::triton::gpu::MemDescType) in Triton IR type system

tti.experimental_gsan_mbarrier_table_init (triton::instrument::ExperimentalGSanMBarrierTableInitOp)

Initialize GSan mbarrier release state

Syntax:

operation ::= `tti.experimental_gsan_mbarrier_table_init` $scratch `,` $capacity attr-dict `:` type($scratch)

Initializes GSan’s internal mbarrier table, used to sequence memory accesses through cluster mbarriers. This operation is internal to GSan instrumentation.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Attributes:

Attribute	MLIR Type	Description
capacity	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:

Operand	Description
scratch	ptr

tti.experimental_gsan_mbarrier_wait (triton::instrument::ExperimentalGSanMBarrierWaitOp)

Acquire GSan clocks after an mbarrier wait

Syntax:

operation ::= `tti.experimental_gsan_mbarrier_wait` $scratch `,` $barrier `,` $phase `,` $pred attr-dict `:` type($scratch) `,`
              qualified(type($barrier))

Acquires the clocks published to the completed mbarrier phase, excluding each publisher’s current epoch. Global-memory handoffs that rely solely on an mbarrier may therefore be conservatively reported as races. This operation is internal to GSan instrumentation.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Read on ::mlir::triton::GlobalMemory, MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Operands:

Operand	Description
scratch	ptr
barrier	memory descriptor type (::mlir::triton::gpu::MemDescType) in Triton IR type system
phase	32-bit signless integer
pred	1-bit signless integer

tti.experimental_gsan_tensor_access (triton::instrument::ExperimentalGSanTensorAccessOp)

Instrument a tensor load/store access for GSan

Syntax:

operation ::= `tti.experimental_gsan_tensor_access` $ptr `,` $isStore (`,` $mask^)? attr-dict `:` type($ptr)

Emits runtime instrumentation for a tensor pointer access. The pointer and optional mask are consumed by the GSan runtime.

Attributes:

Attribute	MLIR Type	Description
isStore	::mlir::BoolAttr	bool attribute

Operands:

Operand	Description
ptr	ptr or ranked tensor of ptr values
mask	1-bit signless integer or ranked tensor of 1-bit signless integer values

tti.experimental_gsan_tensordesc_access (triton::instrument::ExperimentalGSanTensorDescAccessOp)

Instrument a TMA tensor descriptor access for GSan

Syntax:

operation ::= `tti.experimental_gsan_tensordesc_access` $desc `[` $coords `]` `,` $isStore (`,` $pred^)? attr-dict `:`
              qualified(type($desc))

Passes the native tensor descriptor and its coordinates directly to the GSan runtime, which decodes the descriptor and iterates over the tile.

Traits: AttrSizedOperandSegments

Attributes:

Attribute	MLIR Type	Description
isStore	::mlir::BoolAttr	bool attribute

Operands:

Operand	Description
desc	tensor descriptor type
coords	variadic of 32-bit signless integer
pred	1-bit signless integer

tti.experimental_local_gather (triton::instrument::ExperimentalLocalGatherOp)

Gather elements from shared memory with logical base offsets

Syntax:

operation ::= `tti.experimental_local_gather` $src `[` $indices `]` `offsets` `=` `[` $offsets `]`
              attr-dict `:` qualified(type($src)) `,` type($indices) `->` type($result)

Gather elements from a shared memory descriptor using an index tensor along one axis, after shifting the logical source coordinates by rank-sized scalar offsets. This is intentionally private to instrumentation passes.

Attributes:

Attribute	MLIR Type	Description
axis	::mlir::IntegerAttr	32-bit signless integer attribute

Operands:

Operand	Description
src	memory descriptor type (::mlir::triton::gpu::MemDescType) in Triton IR type system
indices	ranked tensor of integer values
offsets	variadic of 32-bit signless integer

Results:

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

tti.experimental_lock_acquire (triton::instrument::ExperimentalLockAcquireOp)

Acquire a lock.

Syntax:

operation ::= `tti.experimental_lock_acquire` $lock (`,` $pred^)? attr-dict `:` type($lock)

Enter a critical section by acquiring a lock with single thread.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Operands:

Operand	Description
lock	ptr or ranked tensor of ptr values
pred	1-bit signless integer

tti.experimental_lock_release (triton::instrument::ExperimentalLockReleaseOp)

Release a lock.

Syntax:

operation ::= `tti.experimental_lock_release` $lock (`,` $pred^)? attr-dict `:` type($lock)

Leave a critical section by releasing a lock with single thread.

Interfaces: MemoryEffectOpInterface (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{MemoryEffects::Write on ::mlir::triton::GlobalMemory}

Operands:

Operand	Description
lock	ptr or ranked tensor of ptr values
pred	1-bit signless integer

tti.experimental_memdesc_to_i32 (triton::instrument::ExperimentalMemDescToI32Op)

Convert a memdesc into its base pointer as i32

Syntax:

operation ::= `tti.experimental_memdesc_to_i32` $memdesc attr-dict `:` type($memdesc)

Extract the base pointer from the given memdesc and return it as a 32-bit integer. This can be used to compare the memdesc against tensors of barrier pointers maintained by the concurrency sanitizer.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand	Description
memdesc	memory descriptor type (::mlir::triton::gpu::MemDescType) in Triton IR type system

Results:

Result	Description
result	32-bit signless integer

tti.experimental_memory_offset_to_i32 (triton::instrument::ExperimentalMemoryOffsetToI32Op)

Materialize a memory-space offset as an i32 address

Syntax:

operation ::= `tti.experimental_memory_offset_to_i32` $offset `,` $memType attr-dict

Add a static byte or encoded tensor-memory offset to the current function’s memory-space base and return the 32-bit address used by ConSan. This provides the runtime identity corresponding to a compile-time BufferRegion key without materializing a descriptor table.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable, InferTypeOpInterface, NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Attribute	MLIR Type	Description
offset	::mlir::IntegerAttr	32-bit signless integer attribute
memType	::mlir::triton::instrument::MemTypeAttr	allowed 32-bit signless integer cases: 0, 1

Results:

Result	Description
result	32-bit signless integer