Allow ExLlamaV3 on Turing (sm_75), keep Volta rejected - #463
Draft
vcruz305 wants to merge 1 commit into
Draft
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This is blocked on turboderp-org/exllamav3#325 and should not merge before that one does. I am opening it as a draft now so the tabbyAPI side is ready when it lands, and because I would like an opinion on the caveat at the bottom.
The check in
common/hardware.pyrejects anything below compute capability 8.0, which is correct as things stand: the ExLlamaV3 kernels usecp.asyncandmma.m16n8k16, and both need sm_80 or higher. exllamav3#325 adds sm_75 fallbacks for those. It has been validated on a T4 and on dual 2080 Ti, and separately I have been running Qwen3.8-27B at 4.00bpw on a Quadro RTX 6000 with my own port of the same three changes, at 31 to 33 tok/s decode with a 262k context.This accepts 7.5 and up. Volta stays rejected: it reports 7.0 and does not have the instructions the fallbacks assume. That means comparing the full
(major, minor)tuple, since the current code takes[0]and cannot tell 7.5 from 7.0.I also updated the failure message in
backends/exllamav3/model.py, which currently names ampere and the 30 series.The part I am unsure about
Compute capability is necessary here but not sufficient. The published exllamav3 wheels contain no sm_75 cubin, so even after #325 merges, a Turing user doing a normal install would pass this check and then die at the first matmul with "no kernel image is available for execution on the device". Getting sm_75 kernels means building exllamav3 from source with
TORCH_CUDA_ARCH_LIST=7.5.So this function arguably wants to ask whether the installed ExLlamaV3 has kernels for the device, rather than whether the device is capable in principle. That is a different shape of check and I did not want to change it uninvited, so this PR stays minimal. I am happy to rework it as a real capability probe if you would prefer that, and it would fail more clearly for the 8.0-and-up case too.