Barrier between mark reads and compaction writes in moe_topk_softmax_mi300

[titoatwork]

Description of changes:

The compaction loop in moe_topk_softmax_mi300_task_impl reads active_expert_ids[i] and atomicAdd-compacts into the same array with no ordering between them (moe_topk_softmax_mi300.cuh, 'Compact active expert marks' block). Compaction writes land on indices 0..k-1 while an inactive expert below the compacted count may not have been read yet, so a cleared slot can be made to look active: a phantom expert inflates the active set nondeterministically. This change stages every mark thread-locally, issues __syncthreads(), then compacts.

Scope notes, stated precisely:

• The only corruption mode is phantom inclusion. Dropped experts are impossible by construction here: the mark is used only as a >=0 predicate and each thread stores its own index, so overwriting a live mark cannot un-count it. I verified this with a CPU ordering litmus (below) rather than assuming it.
• No behavior change at current sizes: at 64 experts with blockDim 256 only threads 0..63 run, a single wavefront, which issues all loads before any store. The race is live once active counts span wavefronts, so this is hardening before it can fire (e.g. larger expert counts).
• The added barrier sits under a block-uniform condition (null check on a kernel pointer), matching the existing barriers in this function. Register/stack cost is two small per-thread arrays sized (NUM_EXPERTS+255)/256.

Evidence (CPU ordering model of the loop, N=512 experts, std::atomic visibility, relaxed ordering):

• Adversarial interleaving, phantom config (299 active, target inactive): old corrupts deterministically (dense=300 vs 299), fixed version correct.
• Adversarial interleaving, dropped config: correct in both versions, consistent with the construction argument above.
• 300 fully-concurrent trials each: old corrupts 1/300 (phantom config) and 0/300 (dropped config); fixed version 0/300 both. The litmus models ordering only, not the kernel; happy to share the source.

Verification: diff is +21/-5 in one file, no new clang-format v15 violations introduced (file has pre-existing ones, verified identical before/after outside this hunk), change compiles by inspection (constant-expression array bound, uniform barrier). No MI300X access here, so device execution is unverified. The litmus and the code-structure argument above are the evidence.

Related Issues:

Linked Issues: none filed (reporting directly).