[Denverjin]

Add Ascend sparse-attention prefill and decode examples with a fully manual pipeline and a persistent T.Stage pipeline. The staged implementation uses compiler-managed versions and synchronization, native pair gathers, high-level QK/PV copies and GEMMs, and FP8/FP4 dequantization.

QK and PV use two Buffer views of one four-version L0B allocation. Their MajorK and MajorMN layouts remain distinct while shared .data provides one version clock and a 64 KiB ring without T.assume_conflict. This requires the Cube-view support in fork PR #11.

This PR includes these compiler changes:

• Denverjin/tilelang#5: Support native NZ output for L0C-to-UB copies.
• Denverjin/tilelang#7: Select AIV streams outside scheduled pipelines.
• Denverjin/tilelang#12: Preserve manual source order across physical phase ties.

Remaining compiler prerequisites:

• Denverjin/tilelang#6: Preserve Cube configuration phases across manual stages.
• tile-ai/tilelang#3606, from fork #8: Preserve complete periodic buffer lifetimes.
• tile-ai/tilelang#3607, from fork #9: Support exact placement of versioned on-chip buffers.
• Denverjin/tilelang#10: Preserve physical views and pointer units during layout lowering.
• Denverjin/tilelang#11: Preserve independent Cube view layouts on shared storage rings.
• tile-ai/tilelang#3530: Partial UB-to-L1 NZ copies and access footprints.

The placement prerequisite now supports UB, L1, L0A/L0B/L0C, and SS through T.annotate_buffer_offsets, with an explicit fixed version count required for each placed buffer. The staged kernel still uses the earlier T.annotate_l1_offsets API; integrating the new prerequisite also requires migrating that call and adding an explicit fixed q_l1 version count on the prefetched-query path.

The included performance records identify the historical integrated compiler. They are not measurements of this partially integrated branch.