fix OOB in cuda_ctc_decoder (#3754)

[gilbsgilbs]

first_matrix__bitonic_topk_kernel placed its top-k values array with:

int* block_topk_value = reinterpret_cast<int*>(block_topk_key + sizeof(float) * beam);

This is a classic pitfall with pointers arithmetic: block_topk_key is a float*, so + sizeof(float) * beam advances sizeof(float) * beam (= 4*4*beam bytes) instead of beam "elements" (= 4*beam bytes). Therefore, the value is written and read 8*beam..12*beam bytes past the end of the kernel's allocated smem.

On Hopper/Blackwell, the issue is systematic at the first synchronize after the kernel (probably because of tighter shared-memory bounds on newer architectures?). On Ampere/Ada, the issue only becomes apparent at higher beam sizes because it makes the overrun large enough.

Advance by beam elements instead of sizeof(float)*beam elements. Add a regression test that decodes with beam_size=128, so that the overrun is large enough to fault on any GPU.

Fixes #3754.

---

For future reference: compute-sanitizer --tool memcheck on the unpatched decoder yields the exact OOB line (at the queue.store):

========= Invalid __shared__ write of size 4 bytes
=========     at first_matrix__bitonic_topk_kernel<(int)256, (int)16>(...)+0x... in ctc_prefix_decoder_kernel_v2.cu:299
=========     by thread (0,0,0) in block (0,0,0)
=========     Access to 0x8a0 is out of bounds
========= ERROR SUMMARY: 41 errors

Also, linters are usually able to catch such pointer arithmetic bugs statically.

---

PLEASE NOTE THAT THE TORCHAUDIO REPOSITORY IS NO LONGER ACTIVELY MONITORED. You may not get a response. For open discussions, visit https://discuss.pytorch.org/.

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/audio/4199

• 📄 Preview Python docs built from this PR

Note: Links to docs will display an error until the docs builds have been completed.

This comment was automatically generated by Dr. CI and updates every 15 minutes.

[gilbsgilbs]

Just realized #4198 by @lanarkite99 has the same fix. Closing this one (although I prefer my test case as it is less verbose and reuses the patterns of the other tests).