coll/alltoall: guard block-scope warp quiet when no warps are unused

When SCOPE==BLOCK and PE_size < groupSize, block-scope alltoall delegates
the final transfer_quiet to whichever warp has index == first_unused_warp.
If all warps already took part in the transfer (num_warps ==
first_unused_warp), no warp matches and transfer_quiet is silently skipped.
The kernel returns without flushing the proxy channel, so in-flight puts
may land at the destination after the caller observes completion, which
manifests as corrupted alltoall output for small block launches (e.g. the
functional device alltoall test on B300 with num_threads = 32, PE_size =
16, giving num_warps = 1 and first_unused_warp = 1).

Add the num_warps > first_unused_warp guard so the else branch runs a
full-scope transfer_quiet in that case.

Author: Xuan-1998

src/include/non_abi/device/coll/alltoall.cuh

@@ -110,7 +110,10 @@ __device__ NVSHMEMI_DEVICE_ALWAYS_INLINE void nvshmemi_alltoall_allpush_threadgr
         nvshmemi_wait_until_greater_than_equals<uint64_t>((psync + next_rank), *pwrk,
                                                           NVSHMEMI_CALL_SITE_SIGNAL_WAIT_UNTIL_GE);
     }
-    if (SCOPE == NVSHMEMI_THREADGROUP_BLOCK && PE_size < groupSize) {
+    /* Delegate the quiet to a single spare warp when one exists; otherwise every warp was used
+     * above and we must fall back to a scope-wide quiet (num_warps == first_unused_warp would
+     * match no warp and silently skip the quiet). */
+    if (SCOPE == NVSHMEMI_THREADGROUP_BLOCK && PE_size < groupSize && num_warps > first_unused_warp) {
         if (my_warp_idx == first_unused_warp)
             nvshmemi_transfer_quiet<NVSHMEMI_THREADGROUP_WARP>(false);
     } else