Environment
- GPU: RTX 5090 (SM120 / Blackwell)
- Python 3.11.13, PyTorch 2.10.0+cu130, CUDA 13.0
- triton-windows 3.8.0.post28
- Model: Qwen3.8-27B Uncensored GGUF (Deepy / prompt enhancer)
Bug
Selecting the vllm engine for Deepy/prompt enhancer crashes on the first prompt with:
File "shared/llm_engines/nanovllm/layers/attention_sm120.py", line 216, in q8_grouped_partials q8_grouped_async_kernel[...] triton.compiler.errors.CompilationError: at 44:8: async_copy.wait_group(0) gl.thread_barrier() ^ AttributeError("module 'triton.experimental.gluon.language' has no attribute 'thread_barrier'")
Likely cause
Triton's official gluon API doesn't have thread_barrier — the correct function is gl.barrier() (see https://triton-lang.org/main/gluon/api/generated/triton.experimental.gluon.language.barrier.html). This blocks vllm engine entirely on SM120 GPUs.
Workaround
Using the cg (Cuda Graph) engine instead works fine, no crash.
Separate but related finding
Upgrading llamacpp_gguf_cuda from 1.0.2 to 1.0.14 fixed a different issue: with 1.0.2, the Q8/FlashAttention decode path for Deepy was unavailable (falling back to slow PyTorch SDPA, ~6 tok/s). With 1.0.14 it works and speed jumped to ~37 tok/s on the same 5090. Might be worth calling out more clearly in INSTALLATION.md that older GGUF kernel wheels silently disable this path.
Environment
Bug
Selecting the
vllmengine for Deepy/prompt enhancer crashes on the first prompt with:
File "shared/llm_engines/nanovllm/layers/attention_sm120.py", line 216, in q8_grouped_partials q8_grouped_async_kernel[...] triton.compiler.errors.CompilationError: at 44:8: async_copy.wait_group(0) gl.thread_barrier() ^ AttributeError("module 'triton.experimental.gluon.language' has no attribute 'thread_barrier'") Likely cause
Triton's official gluon API doesn't have
thread_barrier— the correct function isgl.barrier()(see https://triton-lang.org/main/gluon/api/generated/triton.experimental.gluon.language.barrier.html). This blocksvllmengine entirely on SM120 GPUs.Workaround
Using the
cg(Cuda Graph) engine instead works fine, no crash.Separate but related finding
Upgrading
llamacpp_gguf_cudafrom 1.0.2 to 1.0.14 fixed a different issue: with 1.0.2, the Q8/FlashAttention decode path for Deepy was unavailable (falling back to slow PyTorch SDPA, ~6 tok/s). With 1.0.14 it works and speed jumped to ~37 tok/s on the same 5090. Might be worth calling out more clearly in INSTALLATION.md that older GGUF kernel wheels silently disable this path.