triton-ascend-a5-matmul-vector

Use a two-stage structure: a cube scope loop that runs tl.dot and pushes results to a UB buffer via al.fixpipe, and a vector scope loop that reads the buffer with bl.to_tensor, applies the post-processing, and stores to GM. Pair the stages with sync_block_set/wait events.

Use affinity APIs only for CV-fused kernels that combine tl.dot with fusable vector post-processing such as GELU, ReLU, bias-add, or quantization. Pure MatMul and pure Vector kernels must use native Triton, since affinity APIs add extra data movement and synchronization overhead there.

Deadlocks occur when sync_block_set/wait five-tuples (producer, consumer, event_id, src_pipe, dst_pipe) do not match exactly. Wrong outputs usually come from missing the buffer-free event (EVT1), which lets cube overwrite c_ub before vector finishes reading it.

Yes. With FixpipeDualDstMode.ROW_SPLIT each sub-vector core sees half the rows, so the output row offset must add sub_vec_id * (BLOCK_M // 2). Omitting this makes both sub-vectors write to the same GM region and corrupt results.

No. bl.alloc shapes must be tl.constexpr, and the same applies to bl.subview and bl.to_tensor target shapes. Passing runtime integers raises a get_buffer_ty() error at compile time.