Skip to content

Allow ExLlamaV3 on Turing (sm_75), keep Volta rejected - #463

Draft
vcruz305 wants to merge 1 commit into
theroyallab:mainfrom
vcruz305:allow-turing
Draft

Allow ExLlamaV3 on Turing (sm_75), keep Volta rejected#463
vcruz305 wants to merge 1 commit into
theroyallab:mainfrom
vcruz305:allow-turing

Conversation

@vcruz305

@vcruz305 vcruz305 commented Sep 4, 2026

Copy link
Copy Markdown

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.py rejects anything below compute capability 8.0, which is correct as things stand: the ExLlamaV3 kernels use cp.async and mma.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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant