[BugFix] [Kernel] Fix GPU SEGV occuring in fused_moe kernel

[rasmith]

When running a large model (~500G), I encountered GPU SEGV. It was occurring in fused_moe_kernel:

Thread 98 "python" received signal SIGSEGV, Segmentation fault.
Warning: precise memory violation signal reporting is not enabled, reported
location may not be accurate.  See "show amdgpu precise-memory".
[Switching to thread 98, lane 0 (AMDGPU Lane 7:2:1:1/0 (0,0,0)[0,0,0])]
fused_moe_kernel () at fused_moe.py:142
142             b = tl.load(b_ptrs,

It seemed to me that this could be integer overflow. Changing offsets to use tl.int64 in fused_moe.py:fused_moe_kernel fixes the issue.

[github-actions]

👋 Hi! Thank you for contributing to the vLLM project.
Just a reminder: PRs would not trigger full CI run by default. Instead, it would only run fastcheck CI which starts running only a small and essential subset of CI tests to quickly catch errors. You can run other CI tests on top of those by going to your fastcheck build on Buildkite UI (linked in the PR checks section) and unblock them. If you do not have permission to unblock, ping simon-mo or khluu to add you in our Buildkite org.

Once the PR is approved and ready to go, your PR reviewer(s) can run CI to test the changes comprehensively before merging.

To run CI, PR reviewers can do one of these:

• Add ready label to the PR
• Enable auto-merge.

🚀

[rasmith]

/ready

[tlrmchlsmth]

Thanks for the fix! I noted a couple of spots where I think 32-bits should be fine, WDYT?

[rasmith]

@tlrmchlsmth Doing this causes SEGV when I tested.