triton.language.extra.cann.extension.sync_block_set

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

Sets a cross-core synchronization flag for producer-consumer sync between Cube and Vector cores.

Pairs with sync_block_wait() to coordinate execution between different core types. Each call increments a per-event counter that the corresponding sync_block_wait will decrement (semaphore-like behavior).

Must be used within an scope() context matching the sender'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]. Each ID maps to an independent counter.

• sender_pipe (PIPE) -- Sender-side pipeline type (e.g., PIPE_MTE1, PIPE_V). 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_cube_to_vector():
    with al.scope(core_mode="cube"):
        al.sync_block_set("cube", "vector", 0, al.PIPE.PIPE_MTE1, al.PIPE.PIPE_MTE3)
    with al.scope(core_mode="vector"):
        al.sync_block_wait("cube", "vector", 0, al.PIPE.PIPE_MTE1, al.PIPE.PIPE_MTE3)

Special Restrictions

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

• event_id: must be in range [0, 15].

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