triton.language.extra.cann.extension.sync_block_all

triton.language.extra.cann.extension.sync_block_all(mode, event_id, _semantic=None)

Performs global synchronization across all cores of a specified type.

Inserts a sync barrier to resolve RAW, WAR, and WAW data hazards on shared global memory across Cube cores, Vector cores, or both. Also supports sub-vector-level synchronization within Vector cores.

参数:

• mode (str) -- Synchronization scope. One of: "all_cube" — sync all Cube cores; "all_vector" — sync all Vector cores; "all" — sync all Cube and Vector cores; "all_sub_vector" — sync between Vector sub-blocks.

• event_id (int) -- Event marker ID in range [0, 15].

Example

@triton.jit
def test_sync_block_all():
    al.sync_block_all("all_cube", 8)
    al.sync_block_all("all_vector", 9)
    al.sync_block_all("all", 10)
    al.sync_block_all("all_sub_vector", 11)

Special Restrictions

• mode: must be one of 'all_cube', 'all_vector', 'all', 'all_sub_vector'.

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