tilelang-ascend-attention

Structure it as Cube+Vector fusion: QK^T and PV GEMMs run on Cube cores, softmax on Vector cores, with workspace tensors declared via workspace_idx in @tilelang.jit carrying intermediate results between cores. Enable AUTO_CV_COMBINE and AUTO_CV_SYNC pass configs for automatic synchronization.

Online softmax maintains three running states: m_i (row max), sumexp (row exponential sum), and acc_o (output accumulator). Each KV block updates them using the correction factor exp(m_prev - m_new), which rescales historical sumexp and acc_o for numerical stability.

Enable TL_ASCEND_AUTO_CV_COMBINE, TL_ASCEND_AUTO_CV_SYNC, TL_ASCEND_AUTO_SYNC, and TL_ASCEND_MEMORY_PLANNING in pass_configs. These let Developer mode handle C/V fusion, synchronization, and memory planning automatically instead of manual set_flag/wait_flag.

This guide targets Atlas A2 and Atlas A3 hardware. It uses Ascend-specific features like L0C accumulators, UB memory, and Cube/Vector core partitioning that are specific to these NPU architectures.

L0C is an accumulator buffer that does not support reduce operations. You must T.copy the data from L0C to UB (unified buffer) first, then perform softmax reduction on the Vector core.

Use T.ceildiv to compute block counts instead of assuming SEQ_LEN divides evenly. Never assume exact divisibility when partitioning KV blocks across iterations.