triton.language.extra.cann.extension.sync_block_wait

triton.language.extra.cann.extension.sync_block_wait(sender, receiver, event_id, sender_pipe=None, receiver_pipe=None, _semantic=None)

Waits on a cross-core synchronization flag set by sync_block_set().

Blocks execution until the corresponding event counter is positive (signaling the producer has completed), then decrements it by 1. Pairs with sync_block_set for producer-consumer synchronization between Cube and Vector cores.

Must be used within an scope() context matching the receiver's core type.

参数:

  • sender (str) -- Sending core type. Must be "cube" or "vector" (must differ from receiver).

  • receiver (str) -- Receiving core type. Must be "cube" or "vector".

  • event_id (int) -- Sync flag identifier in range [0, 15]. Must match the ID used by the corresponding sync_block_set.

  • sender_pipe (PIPE) -- Sender-side pipeline type. Defaults to PIPE_FIX if sender is cube, PIPE_MTE3 if sender is vector.

  • receiver_pipe (PIPE) -- Receiver-side pipeline type. Defaults to PIPE_MTE2.

Example

@triton.jit
def kernel_sync_vector_to_cube():
    with al.scope(core_mode="vector"):
        al.sync_block_set("vector", "cube", 1, al.PIPE.PIPE_V, al.PIPE.PIPE_FIX)
    with al.scope(core_mode="cube"):
        al.sync_block_wait("vector", "cube", 1, al.PIPE.PIPE_V, al.PIPE.PIPE_FIX)

Special Restrictions

  • sender/receiver: must be 'cube' or 'vector', and must differ from each other.

  • event_id: must match the ID used by the corresponding sync_block_set.

  • Must be used within an al.scope context matching the receiver's core type.