Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

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

Closed
wants to merge 3 commits into from
Closed

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

wants to merge 3 commits into from

Conversation

rasmith
Copy link
Contributor

@rasmith rasmith commented Nov 15, 2024 •
edited by github-actions bot
Loading

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.

@rasmith
Fix segv occuring in fused_moe.py kernel
789365b 
Signed-off-by: Randall Smith <Randall.Smith@amd.com>
@github-actions GitHub Actions
Copy link

👋 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 rasmith changed the title Fix segv occuring in fused_moe.py kernel [BugFix] [Kernel] Fix GPU SEGV occuring in fused_moe kernel Nov 15, 2024
@rasmith
Copy link
Contributor Author

rasmith commented Nov 15, 2024

/ready

tlrmchlsmth
tlrmchlsmth reviewed Nov 16, 2024
View reviewed changes
Copy link
Collaborator

@tlrmchlsmth tlrmchlsmth left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

vllm/model_executor/layers/fused_moe/fused_moe.py Outdated Show resolved Hide resolved
vllm/model_executor/layers/fused_moe/fused_moe.py Outdated Show resolved Hide resolved
vllm/model_executor/layers/fused_moe/fused_moe.py Show resolved Hide resolved
@rasmith
remove unncessary cast to tl.int64
0eb50d5 
Signed-off-by: Randall Smith <Randall.Smith@amd.com>
@tlrmchlsmth tlrmchlsmth added the ready ONLY add when PR is ready to merge/full CI is needed label Nov 16, 2024
@tlrmchlsmth tlrmchlsmth enabled auto-merge (squash) November 16, 2024 00:31
@tlrmchlsmth tlrmchlsmth enabled auto-merge (squash) November 16, 2024 00:34
rasmith
rasmith commented Nov 16, 2024
View reviewed changes
vllm/model_executor/layers/fused_moe/fused_moe.py Outdated
@@ -116,7 +116,7 @@ def fused_moe_kernel(
a_ptrs = a_ptr + (offs_token[:, None] // top_k * stride_am +
offs_k[None, :] * stride_ak)

off_experts = tl.load(expert_ids_ptr + pid_m).to(tl.int64)
off_experts = tl.load(expert_ids_ptr + pid_m)
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@tlrmchlsmth Doing this causes SEGV when I tested.

@tlrmchlsmth tlrmchlsmth enabled auto-merge (squash) November 16, 2024 00:47
@rasmith
Merge and drop two commits
23e7dfe 
Signed-off-by: Randall Smith <Randall.Smith@amd.com>
auto-merge was automatically disabled November 16, 2024 01:38

Head branch was pushed to by a user without write access

@mergify mergify bot added documentation Improvements or additions to documentation ci/build labels Nov 16, 2024
@rasmith rasmith closed this Nov 16, 2024
@rasmith
Copy link
Contributor Author

rasmith commented Nov 16, 2024

DCO issues, need to close and reopen.

@rasmith rasmith deleted the ransmith_fused_moe_segv_fix branch November 16, 2024 01:41
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@tlrmchlsmth tlrmchlsmth tlrmchlsmth approved these changes

Assignees
No one assigned
Labels
ci/build documentation Improvements or additions to documentation ready ONLY add when PR is ready to merge/full CI is needed
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
2 participants
@rasmith @tlrmchlsmth