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

[coderabbitai]

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: [None][feat] Add Rubin SM107 CuTe DSL foundation and BF16 kernels #18311
• Review comment: [None][feat] Add Rubin SM107 CuTe DSL foundation and BF16 kernels #18311 (comment)
• Requested by: @farazkh80

[farazkh80]

Tested on SM107. The race does not reproduce: with fp32 output the kernel passes its reference check deterministically across repeated runs, and inserting cp_async_bulk_commit_group + cp_async_bulk_wait_group(0, read=True) before the epilogue barrier changes nothing. The bf16 mismatches seen in the harness are nondeterministic scatter-add rounding order vs an fp32 reference with atol 0.1 (max diff equals one bf16 ulp at the output magnitude), not corruption. Downgrading this to defensive hardening; no code change planned for the merge-back.