tilelang-ascend-example-reduction

Use T.alloc_ub to allocate UB buffers, T.reduce_min or T.reduce_sum for the reduction along the last dimension, and T.copy for data movement between global memory and UB. In Expert mode, wrap compute in T.Scope("V") and synchronize with T.barrier_all().

Allocate UB tensors with a leading stages dimension set to 2, then alternate buffers with cur = mm % stages and nxt = (mm + 1) % stages. Prefetch the next block with T.copy while reducing the current block so data movement overlaps computation.

Yes, the same pipeline structure works with T.reduce_max and T.reduce_sum by swapping the reduction primitive. The example uses T.reduce_min over dim=-1, but the double-buffering and synchronization pattern is identical for other reduction operations.

M must be an integer multiple of block_M and N must be an integer multiple of block_N. For non-divisible shapes, apply padding beforehand or switch to Developer mode using T.ceildiv for boundary handling.

The example targets Atlas A2 and Atlas A3 hardware using the ascendc JIT target. It compiles through tilelang.jit with target="ascendc" and runs on the vector cores with VEC_NUM set to 2.