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.

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

[Cleanup][Kernel] Remove if-else with identical branches in marlin 2:4 #10687

Merged
merged 1 commit into from
Nov 27, 2024

Conversation

tlrmchlsmth
Copy link
Collaborator

@tlrmchlsmth tlrmchlsmth commented Nov 27, 2024

Clean up a useless if-else in marlin_24_cuda_kernel.cu:

if (group_blocks != -1) {
s_sh_rd = 8 * ((threadIdx.x / 32) % (thread_n_blocks / 4)) +
(threadIdx.x % 32) / 4;
} else {
s_sh_rd = 8 * ((threadIdx.x / 32) % (thread_n_blocks / 4)) +
(threadIdx.x % 32) / 4;
}

See the equivalent if-else in marlin_cuda_kernel.cu:

if (group_blocks != -1)
s_sh_rd = 8 * ((threadIdx.x / 32) % (thread_n_blocks / 4)) +
(threadIdx.x % 32) / 4;
else
s_sh_rd = 8 * ((threadIdx.x / 32) % (thread_n_blocks / 4)) +
(threadIdx.x % 32) % 4;

See discussion in #6030

From @alexm-neuralmagic in that issue:

this is a leftover from a copy-paste from the original (dense) marlin that had different branches. We have tests to verify all of these cases inside test_marlin_gemm.py that verify both dense and sparse and group_size == -1 and other group_sizes as well.

I'd love to leave a better explanation for the discrepancy between the equivalent if statement in marlin_cuda_kernel.cu, because this looks very suspicious.

(closes #6030)

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.

🚀

Copy link
Member

@mgoin mgoin 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 addressing :)

csrc/quantization/marlin/sparse/marlin_24_cuda_kernel.cu Outdated Show resolved Hide resolved
@mgoin mgoin added the ready ONLY add when PR is ready to merge/full CI is needed label Nov 27, 2024
@tlrmchlsmth tlrmchlsmth enabled auto-merge (squash) November 27, 2024 02:35
@youkaichao youkaichao disabled auto-merge November 27, 2024 06:55
@youkaichao youkaichao merged commit e225110 into vllm-project:main Nov 27, 2024
66 of 68 checks passed
@youkaichao youkaichao deleted the marlin_24_if_cleanup branch November 27, 2024 06:55
afeldman-nm pushed a commit to neuralmagic/vllm that referenced this pull request Dec 2, 2024
sleepwalker2017 pushed a commit to sleepwalker2017/vllm that referenced this pull request Dec 13, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
ready ONLY add when PR is ready to merge/full CI is needed
Projects
None yet
Development

Successfully merging this pull request may close these issues.

[Bug]: identical branches in csrc/quantization/marlin/sparse/marlin_24_cuda_kernel.cu
3 participants