tilelang-ascend-api

Use T.gemm_v0 with A and B allocated in shared/L1 memory and C in fragment/L0C memory, copying tiles with T.copy inside a loop. Set init=True on the first iteration to clear the accumulator, and wrap copies in T.Pipelined with num_stages=2 to overlap data movement with compute.

T.alloc_shared is Developer mode where the compiler automatically maps to L1 or UB, while T.alloc_ub is Expert mode that explicitly targets the Unified Buffer for vector computation. Use Developer mode for portability and Expert mode when you need precise memory hierarchy control.

Yes, T.tile.atomic_add accumulates a local UB or L0C tile into a GM destination, supporting int8, int16, float16, bfloat16, int32, and float32. The destination must be a GM buffer and the source a local tensor with matching dtype; it does not support return_prev or constant sources.

Use T.Pipelined with num_stages and enable the pass configs TL_ASCEND_AUTO_CV_COMBINE and tl.ascend_auto_cross_core_sync. Note that inter-core and intra-core pipelining cannot be enabled simultaneously, and synchronization uses T.set_cross_flag and T.wait_cross_flag.

These fast-path reductions support 1D buffers on axis 0, 2D buffers on axes 0/1/-1/-2, and 3D buffers only on trailing-tile axes. Invalid dim, real_shape, or output shape values raise frontend errors rather than silently passing to the backend.

T.Parallel with symbolic APIs like T.exp is the recommended cross-platform approach for element-wise computation. T.tile primitives like T.tile.exp directly trigger hardware instructions and suit Expert mode or mixed-mode kernels needing precise control.