triton.language.extra.cann.extension.debug_barrier

triton.language.extra.cann.extension.debug_barrier(sync_mode: SYNC_IN_VF, _semantic=None) → None

Inserts a synchronization barrier between vector/scalar load/store instructions.

Provides fine-grained control over which instruction types are blocked until prior instructions complete. Intended for use within an scope() context.

参数:

sync_mode (SYNC_IN_VF) -- Barrier type specifying which instruction classes to synchronize.

Example

@triton.jit
def triton_sub(in_ptr0, in_ptr1, out_ptr0, XBLOCK: tl.constexpr, XBLOCK_SUB: tl.constexpr):
    offset = tl.program_id(0) * XBLOCK
    base1 = tl.arange(0, XBLOCK_SUB)
    loops1: tl.constexpr = (XBLOCK + XBLOCK_SUB - 1) // XBLOCK_SUB
    for loop1 in range(loops1):
        x0 = offset + (loop1 * XBLOCK_SUB) + base1
        tmp0 = tl.load(in_ptr0 + (x0), None)
        tmp1 = tl.load(in_ptr1 + (x0), None)
        tmp2 = tmp0 - tmp1
        tl.debug_barrier()
        tl.store(out_ptr0 + (x0), tmp2, None)

Special Restrictions

• sync_mode: must be a SYNC_IN_VF enum value.

• Intended for use within an al.scope context.