ROCm split-phase random sample: phase B launches 1024 threads onto 128-entry shared arrays

[localai-org-maint-bot]

Row: BACKEND-ROCM

RandomSampleSplitBK in src/vt/rocm/rocm_sample.hip declares its reduction
scratch at the partial count:

__shared__ float sh_score[kSampleSplitBlocks];    // 128
__shared__ int64_t sh_idx[kSampleSplitBlocks];    // 128

and is launched at the vocab block width:

RandomSampleSplitBK<<<static_cast<unsigned>(n), kVocabBlock, 0, s>>>(   // 1024

Every thread then executes the stores unconditionally:

sh_score[threadIdx.x] = best_v;
sh_idx[threadIdx.x] = best_j;

The if (threadIdx.x < blocks_per_row) above them guards only the global read
of the partials, not the shared write. Threads 128-1023 write past both arrays.
sh_score is 128 floats, so sh_score[128] lands on sh_idx[0] -- the element
the kernel writes to out[row]. The overflowing threads carry kNegInf and
kArgSentinel, so the store races the real sh_idx[0] and can leave a float
bit pattern where a token id belongs.

The path is default-ON (VT_SAMPLE_SPLIT defaults to enabled) and is selected
for any v >= 4096 && n <= 64, which is every production vocab at low
concurrency. The consequence is a wrong sampled token, not a crash.

Why the gate did not catch it

The split path has no test. ROCm random_sample agrees with CPU on the vast majority of rows in tests/vt/test_ops_sample.cpp runs V = 128, below the
v >= 4096 bar, so it takes the single-block kernel. The distribution case runs
V = 4. #3010 reports test_ops_sample 29/29 and the full suite 225/225 green,
and both numbers are true and neither touches the kernel the change adds.

Fix

Launch phase B with kSampleSplitBlocks threads, which is the width its shared
arrays are sized for and the width its reduction tree assumes, plus a test at
V = 8192, N = 8 that selects the split path.

Found while reviewing #3010. Fixed in the same flow.