tilelang.ascend.language.sync¶

Ascend pipeline, cross-core, and pipe-buffer synchronization primitives.

Functions¶

ascend_pipe_barrier(pipe_t)

Insert an AscendC pipeline barrier.

ascend_set_flag(hard_event, event_id)

Insert an AscendC SetFlag for pipeline synchronization.

ascend_wait_flag(hard_event, event_id)

Insert an AscendC WaitFlag for pipeline synchronization.

ascend_sync_inter_arrive(pipe, flag_id)

Signal an Ascend inter-core synchronization event.

ascend_sync_inter_wait(pipe, flag_id)

Wait for an Ascend inter-core synchronization event.

ascend_threadfence()

Insert an AscendC thread fence.

ascend_cross_core_set_flag(mode_id, pipe, flag_id)

Insert an AscendC CrossCoreSetFlag for cross-core synchronization.

ascend_cross_core_wait_flag(mode_id, pipe, flag_id)

Insert an AscendC CrossCoreWaitFlag for cross-core synchronization.

ascend_get_buf(pipe, buf_id[, mode])

Acquire a pipe buffer for double-buffer management.

ascend_rls_buf(pipe, buf_id[, mode])

Release a pipe buffer for double-buffer management.

Module Contents¶

tilelang.ascend.language.sync.ascend_pipe_barrier(pipe_t)¶

Insert an AscendC pipeline barrier.

Generates: AscendC::PipeBarrier<pipe_t::XXX>();

Parameters:

pipe_t (str) – Pipeline type string, e.g., “PIPE_ALL”, “PIPE_V”, “PIPE_M”, “PIPE_MTE1”, “PIPE_MTE2”, “PIPE_MTE3”, “PIPE_S”.

Return type:

None

Example

>>> T.ascend_pipe_barrier("PIPE_ALL")
tilelang.ascend.language.sync.ascend_set_flag(hard_event, event_id)¶

Insert an AscendC SetFlag for pipeline synchronization.

Generates: AscendC::SetFlag<AscendC::HardEvent::XXX>(eventID);

Parameters:

  • hard_event (str) – HardEvent enum name, e.g., “MTE2_MTE1”, “S_MTE3”, “V_MTE2”.

  • event_id (int or PrimExpr) – Event ID for synchronization.

Return type:

None

Example

>>> T.ascend_set_flag("S_MTE3", event_id)
tilelang.ascend.language.sync.ascend_wait_flag(hard_event, event_id)¶

Insert an AscendC WaitFlag for pipeline synchronization.

Generates: AscendC::WaitFlag<AscendC::HardEvent::XXX>(eventID);

Parameters:

  • hard_event (str) – HardEvent enum name, e.g., “MTE2_MTE1”, “S_MTE3”, “V_MTE2”.

  • event_id (int or PrimExpr) – Event ID for synchronization.

Return type:

None

Example

>>> T.ascend_wait_flag("S_MTE3", event_id)
tilelang.ascend.language.sync.ascend_sync_inter_arrive(pipe, flag_id)¶

Signal an Ascend inter-core synchronization event.

Generates: AscendC::CrossCoreSetFlag<0, pipe>(flag_id);

Parameters:

  • pipe (str) – Pipeline type, e.g., “PIPE_MTE3”, “PIPE_FIX”.

  • flag_id (int or PrimExpr) – Inter-core synchronization flag ID.

Return type:

None

Example

>>> T.ascend_sync_inter_arrive("PIPE_FIX", flag_id)
tilelang.ascend.language.sync.ascend_sync_inter_wait(pipe, flag_id)¶

Wait for an Ascend inter-core synchronization event.

Generates: AscendC::CrossCoreWaitFlag<0, pipe>(flag_id);

Parameters:

  • pipe (str) – Pipeline type, e.g., “PIPE_MTE3”, “PIPE_FIX”.

  • flag_id (int or PrimExpr) – Inter-core synchronization flag ID.

Return type:

None

Example

>>> T.ascend_sync_inter_wait("PIPE_MTE3", flag_id)
tilelang.ascend.language.sync.ascend_threadfence()¶

Insert an AscendC thread fence.

Generates: asc_threadfence();

Example

>>> T.ascend_threadfence()
Return type:

None

tilelang.ascend.language.sync.ascend_cross_core_set_flag(mode_id, pipe, flag_id)¶

Insert an AscendC CrossCoreSetFlag for cross-core synchronization.

Generates: AscendC::CrossCoreSetFlag<modeId, pipe>(flagId);

Parameters:

  • mode_id (int) – Cross-core sync mode: 0, 1, 2, or 4.

  • pipe (str) – Pipeline type, e.g., “PIPE_MTE3”, “PIPE_FIX”.

  • flag_id (int or PrimExpr) – Cross-core synchronization flag ID.

Return type:

None

Example

>>> T.ascend_cross_core_set_flag(0, "PIPE_MTE3", 0x8)
tilelang.ascend.language.sync.ascend_cross_core_wait_flag(mode_id, pipe, flag_id)¶

Insert an AscendC CrossCoreWaitFlag for cross-core synchronization.

Generates: AscendC::CrossCoreWaitFlag<modeId, pipe>(flagId);

Parameters:

  • mode_id (int) – Cross-core sync mode: 0, 1, 2, or 4.

  • pipe (str) – Pipeline type, e.g., “PIPE_MTE3”, “PIPE_FIX”.

  • flag_id (int or PrimExpr) – Cross-core synchronization flag ID.

Return type:

None

Example

>>> T.ascend_cross_core_wait_flag(0, "PIPE_MTE3", 0x8)
tilelang.ascend.language.sync.ascend_get_buf(pipe, buf_id, mode=False)¶

Acquire a pipe buffer for double-buffer management.

Generates: asc_lock(pipe, buf_id, mode);

Parameters:

  • pipe (str) – Pipeline type, e.g., “PIPE_MTE1”, “PIPE_M”.

  • buf_id (int or PrimExpr) – Buffer slot index, typically (sk & 1) for double-buffering.

  • mode (bool) – Optional mode flag (default False).

Example

>>> T.ascend_get_buf("PIPE_MTE1", sk & 1)
tilelang.ascend.language.sync.ascend_rls_buf(pipe, buf_id, mode=False)¶

Release a pipe buffer for double-buffer management.

Generates: asc_unlock(pipe, buf_id, mode);

Parameters:

  • pipe (str) – Pipeline type, e.g., “PIPE_MTE1”, “PIPE_M”.

  • buf_id (int or PrimExpr) – Buffer slot index, typically (sk & 1) for double-buffering.

  • mode (bool) – Optional mode flag (default False).

Example

>>> T.ascend_rls_buf("PIPE_M", sk & 1)