Synchronize bulk-async reductions before sC reuse in Rubin fused MoE finalize

Summary

Add bulk-async group synchronization after the fused MoE finalize block reductions in the Rubin SM107 contiguous grouped GEMM kernel.

Rationale

blk_reduce_bf16, blk_reduce_fp32, and blk_reduce_fp16 issue cp.reduce.async.bulk...bulk_group operations. The current epilogue synchronization does not wait for completion of that bulk-async group. A later reuse of the single-stage sC buffer can overwrite data while a reduction is still reading it.

This follow-up is intentionally separate from the merge-back because the current implementation matches the validated source branch.

Affected area

• tensorrt_llm/_torch/cute_dsl_kernels/rubin/moe/rubin_contiguous_grouped_gemm_finalize_fusion.py
• Fused finalize epilogue after the dtype-specific blk_reduce_bf16, blk_reduce_fp32, and blk_reduce_fp16 calls

Required change

After the dtype-specific reduction calls, commit and wait for the bulk-async group before sC can be reused and before tmem.free:

• cute.arch.cp_async_bulk_commit_group()
• cute.arch.cp_async_bulk_wait_group(0, read=True)

Keep the existing dtype dispatch unchanged.

Acceptance criteria

• The epilogue commits the issued cp.reduce.async.bulk operations.
• The epilogue waits for bulk-async completion before sC reuse and before tensor-memory release.
• The BF16, FP32, and FP16 reduction paths retain their existing dispatch behavior.
• Kernel validation covers repeated/multi-tile execution that reuses sC.

Backlinks

• Pull request: https://github.com/NVIDIA/TensorRT-LLM/pull/18311
• Review comment: https://github.com/NVIDIA/TensorRT-LLM/pull/18311#discussion_r3872091317
• Requested by: @farazkh80