[Bugfix] Bounds-check moe_permute reverse-map write (#45492)

[waynehacking8]

Summary

Fixes #45492 — out-of-bounds write in expandInputRowsKernel
(csrc/libtorch_stable/moe/permute_unpermute_kernels/moe_permute_unpermute_kernel.inl).

expanded_source_row is read from a caller-provided map and then used directly
as a destination index:

expanded_source_row_to_expanded_dest_row[expanded_source_row] =
    static_cast<int>(expanded_dest_row);

The reverse map is sized num_rows * k, but that source-row map can contain
sentinel / padding entries that index past it (the reporter's case:
expanded_dest_row_to_expanded_source_row[0] == num_rows * k == 1032 for a
[129, 8] map), so the write lands one element past the allocation — memory
corruption on the fused-MoE permute path. Note the asymmetry: the neighboring
permuted_idx and 128-bit data-copy writes are both guarded
(!CHECK_SKIPPED || blockIdx.x < *num_dest_rows), but this reverse-map write
was not.

Fix

Bounds-check expanded_source_row before the write. Valid permutations are
unaffected (the index is already in range); only out-of-range sentinel/padding
slots are skipped. This preserves the consumer contract in
finalizeMoeRoutingKernel (skipped rows are still distinguished by their
>= num_valid dest value).

Why this is not a duplicate

gh pr list --repo vllm-project/vllm --state open --search "45492 in:body" and
--search "moe_permute expanded_source_row in:title,body" return no open PR.
The issue has 0 comments and no linked PR.

Test plan

The installed wheel is precompiled, so I validated the kernel's indexing with a
standalone compute-sanitizer harness that replicates lines 16–28 verbatim
(the cutlass data-copy is omitted — it is separately guarded and is not the OOB
under test), on an RTX PRO 6000 (SM120), CUDA 12.8:

nvcc -arch=sm_120 -DOLD -o harness_old harness.cu   # pre-fix
nvcc -arch=sm_120       -o harness_new harness.cu   # this PR
compute-sanitizer --tool memcheck ./harness_old
compute-sanitizer --tool memcheck ./harness_new

• Pre-fix: Invalid __global__ write of size 4 bytes at expandKernel … 1 bytes after the nearest allocation … of size 4128 bytes — ERROR SUMMARY: 2 errors (4128 = 1032 × 4 = the reverse map; index 1032 is one element past the end).
• This PR: ERROR SUMMARY: 0 errors. Both the reporter (sentinel) case and a fully-valid identity permutation run cleanly; the valid case produces the identical reverse map the old code would (checksum 531996), confirming the guard is a no-op on valid input.

Note

AI assistance (Claude) was used to investigate and draft this change; I have
reviewed every line and run the validation above.

[waynehacking8]

I re-audited this after #45393. Supported moe_permute callers construct token_expert_indices as arange(0, num_rows * k), so the sorted row map must remain a permutation in range; expert-parallel sentinels are sort keys, not row-map values. An out-of-range value therefore signals an upstream invariant failure, and skipping this write would leave a real inv_permuted_idx entry undefined for finalizeMoeRoutingKernel, so I’m closing the PR rather than masking that failure.