[CuTeDSL] reduce example: replace inline PTX with cute.arch.store_async_dsmem

[zkyue]

Summary

The Blackwell reduce.py example still hand-rolls two cluster helpers, set_block_rank and store_shared_remote, that emit mapa.shared::cluster and st.async.shared::cluster.mbarrier::complete_tx::bytes.f32 via llvm.inline_asm. That was necessary when the example landed (#2917), but cute.arch now exposes the same functionality as built-ins:

• cute.arch.map_dsmem_ptr — mapa.shared::cluster, returns a dsmem Pointer
• cute.arch.store_async_dsmem — maps destination + mbarrier to the peer CTA and issues st.async with mbarrier::complete_tx::bytes

This PR removes the inline PTX and calls the built-in directly from cluster_reduce.

Changes

• Delete set_block_rank / store_shared_remote and the cutlass._mlir.dialects.llvm / T imports.
• cluster_reduce now calls
  cute.arch.store_async_dsmem(elem_pointer(...), val.bitcast(Int32), mbar_ptr, peer_cta_rank=lane_idx).
  The built-in stores raw b32 words, so the Float32 is reinterpreted as Int32 with Numeric.bitcast (an arith.bitcast, no data conversion). The transaction-byte accounting (num_warps * cluster_n * 4) and the mbarrier protocol are unchanged.
• Refresh the module docstring and the section comment; the Quack attribution is kept for the overall scheme.

Net: 1 file changed, 18 insertions(+), 96 deletions(-), and the example no longer contains any inline PTX.

Notes

• store_async_dsmem was added in the v4.8 dev update, so the example now requires a 4.8+ nvidia-cutlass-dsl for the cluster_n > 1 path. map_dsmem_ptr (4.6+) is not called directly because store_async_dsmem performs both mapa translations internally; keeping set_block_rank as a wrapper around it would have been dead code.
• elem_pointer is untouched.

Testing

Run on a B200 (sm_100a, CUDA 13.3) with nvidia-cutlass-dsl==4.8.0.dev0:

CUDA_VISIBLE_DEVICES=7 python -m pytest test/examples/CuTeDSL/sm_100a/test_rmsnorm.py -v
======================= 79 passed, 22 warnings in 6.48s ========================

This includes TestRMSNormClusterPath::test_cluster_path_correctness[32768] and [65536], which exercise exactly the changed path (cluster_n > 1 in the RMSNorm example, which imports row_reduce from this file). All non-cluster RMSNorm tests pass as well.

PTX of the cluster kernel (CUTE_DSL_KEEP=ptx) shows the same instruction sequence as the removed inline asm, with the built-in emitting a .b32 store instead of .f32 (same bytes moved):

mapa.shared::cluster.u32   %r321, %r320, %r9;
mapa.shared::cluster.u32   %r323, %r322, %r9;
st.async.shared::cluster.mbarrier::complete_tx::bytes.b32  [%r321], %r8, [%r323];

Also checked statically against cute.arch.store_async_dsmem's contract: both pointers are SMEM Pointers from SmemAllocator, the value is a single Int32 (4-byte alignment requirement satisfied by the Float32 buffer), and peer_cta_rank is the lane index as before.

The two follow-up commits are doc-only (mbarrier arrival count is 1, lanes 0..cluster_n-1 each store, the usage example now mirrors the caller's init-fence / cluster-sync sequence).

[zkyue]

Tested on B200 with nvidia-cutlass-dsl==4.8.0.dev0: test/examples/CuTeDSL/sm_100a/test_rmsnorm.py passes (79 passed), including both test_cluster_path_correctness cases that go through the changed cluster_reduce path. PTX dump of the cluster kernel shows mapa.shared::cluster.u32 x2 + st.async.shared::cluster.mbarrier::complete_tx::bytes.b32, i.e. the same sequence the inline asm used to emit. Updated the Testing section of the description accordingly.