GooseLLM — vLLM for NVIDIA V100 (SM70)

High-throughput LLM inference on Tesla V100 GPUs building on 1Cat-vLLM-0.0.2.

This fork adds:

• TileLang JIT-compiled FlashAttention — paged prefill kernel covering all attention layers (dense and MoE models). Compiled once per head dim during first request, no offline build step.
• TileLang-accelerated Gated DeltaNet / Flash Linear Attention — SM70 kernel paths for hybrid MoE layers (FLA_USE_TILELANG=1).
• Hybrid model support — automatically handles Qwen3.5-122B-A10B style mixed attention + GatedDeltaNet architectures (1056-token alignment pages split into 16-token sub-pages).
• Automatic prefix caching — safely reuses warm KV-cache prefixes without clearing shared cached blocks, including hybrid Qwen3.5 models.
• Expert parallelism — --enable-expert-parallel for MoE models, works with both AWQ and FP16.
• FP16 inference — supports Qwen3.6-35B-A3B and Qwen3.6-27B in FP16 mode (no quantization required).

Upstream FA2 kernel (csrc/flash_attention_v100/), SM70 decode kernel, AWQ SM70 autotune, and custom all-reduce remain from the original 1Cat work.

Acknowledgements

Thanks to 1CatAI for the upstream V100 vLLM builds and tile-ai/tilelang for the TileLang JIT compiler framework.

Quick Start

Docker Build

docker build \
  -f docker/Dockerfile.sm70-build \
  -t goosellm:sm70 \
  .

Local Build

# Save repo root for later reference
GPP=$(pwd)

# 1. Create and activate environment
conda create -n goosellm python=3.12 -y
conda activate goosellm

# 2. Install dependencies
python -m pip install --upgrade pip setuptools wheel
python -m pip install torch torchvision torchaudio \
    --index-url https://download.pytorch.org/whl/cu128
python -m pip install -r requirements/cuda.txt
python -m pip install 'setuptools>=77.0.3,<81.0.0' 'setuptools_scm>=8' grpcio-tools cmake build

# 3. Set build environment
export CUDA_HOME=/usr/local/cuda-12.8
export PATH=$CUDA_HOME/bin:$PATH
export LD_LIBRARY_PATH=$CUDA_HOME/lib64:${LD_LIBRARY_PATH:-}
export VLLM_TARGET_DEVICE=cuda
export VLLM_MAIN_CUDA_VERSION=12.8
export TORCH_CUDA_ARCH_LIST=7.0
export MAX_JOBS=$(nproc)
export NVCC_THREADS=4

# 4. Clone and install TileLang (JIT compiler, takes ~60s)
git clone https://github.com/tile-ai/tilelang.git ~/tilelang
cd ~/tilelang
git apply "$GPP/patches/tilelang-sm70-bf16.patch"
git apply "$GPP/patches/tilelang-nvcc-optlevel.patch"
cd 3rdparty/tvm
git apply "$GPP/patches/tilelang-tvm-padded-layout.patch"
cd ../..
pip install -e .
# Pin apache-tvm-ffi — 0.1.12+ breaks on TVM's DictAttrs.__dict__ field (Python won't setattr on dunder dunder dunder)
pip install "apache-tvm-ffi>=0.1.2,<0.1.12"
cd "$GPP"

# 5. Install TileLang FA-V100 kernels (pure Python, depends on tilelang)
pip install -e 3rdparty/tilelang-fa-v100/

# 6. Build SM70 FA2 kernel (CUDA extension)
cd csrc/flash_attention_v100
sed -i 's/if not torch.cuda.is_available():/if False: # if not torch.cuda.is_available():/' setup.py
python setup.py build_ext --inplace
cd "$GPP"

# 7. Build and install vLLM
rm -rf build vllm.egg-info .deps/*-build .deps/*-subbuild
SETUPTOOLS_SCM_PRETEND_VERSION=0.0.3.dev0 \
  python -m build --wheel --no-isolation --outdir dist-cu128-sm70
python -m pip install dist-cu128-sm70/*.whl --no-deps

Example Serving Commands

Before running any of the commands below, make sure you're in the conda environment:

conda activate goosellm

The examples enable automatic prefix caching explicitly. Warm repeated-prefix generation and tool calls have been validated on Qwen3.5-122B-A10B-AWQ with 4× V100-32GB. Use --no-enable-prefix-caching only when caching is not wanted.

MoE Model (Qwen3.6-35B-A3B-AWQ / Qwen3.5-122B-A10B-AWQ)

FLA_USE_TILELANG=1 NCCL_P2P_LEVEL=NVL \
python -m vllm.entrypoints.openai.api_server \
  --model QuantTrio/Qwen3.6-35B-A3B-AWQ \
  --tensor-parallel-size 4 \
  --dtype float16 \
  --gpu-memory-utilization 0.80 \
  --max-model-len 262144 \
  --max-num-seqs 1 \
  --max-num-batched-tokens 16384 \
  --enable-prefix-caching \
  --trust-remote-code \
  --attention-backend FLASH_ATTN_TILELANG_V100 \
  --enable-expert-parallel \
  --skip-mm-profiling \
  --limit-mm-per-prompt '{"image":0,"video":0}' \
  --compilation-config '{"cudagraph_mode":"full_and_piecewise"}' \
  --host 0.0.0.0 \
  --port 8000

Dense Model (Qwen3.6-27B-AWQ)

FLA_USE_TILELANG=1 NCCL_P2P_LEVEL=NVL \
python -m vllm.entrypoints.openai.api_server \
  --model QuantTrio/Qwen3.6-27B-AWQ \
  --tensor-parallel-size 4 \
  --dtype float16 \
  --gpu-memory-utilization 0.80 \
  --max-model-len 262144 \
  --max-num-seqs 1 \
  --max-num-batched-tokens 16384 \
  --enable-prefix-caching \
  --trust-remote-code \
  --attention-backend FLASH_ATTN_TILELANG_V100 \
  --skip-mm-profiling \
  --limit-mm-per-prompt '{"image":0,"video":0}' \
  --compilation-config '{"cudagraph_mode":"full_and_piecewise"}' \
  --host 0.0.0.0 \
  --port 8000

Docker MoE (Qwen3.6-35B-A3B-AWQ / Qwen3.5-122B-A10B-AWQ)

docker run --rm \
  --gpus all \
  --ipc=host \
  -p 8000:8000 \
  -v ~/.cache/huggingface:/root/.cache/huggingface \
  -e NCCL_P2P_LEVEL=NVL \
  -e FLA_USE_TILELANG=1 \
  goosellm:sm70 \
  python -m vllm.entrypoints.openai.api_server \
    --model QuantTrio/Qwen3.6-35B-A3B-AWQ \
    --quantization awq \
    --dtype float16 \
    --gpu-memory-utilization 0.8 \
    --max-model-len 262144 \
    --tensor-parallel-size 4 \
    --max-num-seqs 1 \
    --max-num-batched-tokens 16384 \
    --enable-prefix-caching \
    --trust-remote-code \
    --attention-backend FLASH_ATTN_TILELANG_V100 \
    --enable-expert-parallel \
    --skip-mm-profiling \
    --limit-mm-per-prompt '{"image":0,"video":0}' \
    --compilation-config '{"cudagraph_mode":"full_and_piecewise"}' 

Docker Dense (Qwen3.6-27B-AWQ)

docker run --rm \
  --gpus all \
  --ipc=host \
  -p 8000:8000 \
  -v ~/.cache/huggingface:/root/.cache/huggingface \
  -e NCCL_P2P_LEVEL=NVL \
  -e FLA_USE_TILELANG=1 \
  goosellm:sm70 \
  python -m vllm.entrypoints.openai.api_server \
    --model QuantTrio/Qwen3.6-27B-AWQ \
    --quantization awq \
    --dtype float16 \
    --gpu-memory-utilization 0.8 \
    --max-model-len 262144 \
    --tensor-parallel-size 4 \
    --max-num-seqs 1 \
    --max-num-batched-tokens 16384 \
    --enable-prefix-caching \
    --trust-remote-code \
    --attention-backend FLASH_ATTN_TILELANG_V100 \
    --skip-mm-profiling \
    --limit-mm-per-prompt '{"image":0,"video":0}' \
    --compilation-config '{"cudagraph_mode":"full_and_piecewise"}'

Results

Known Limitations

• MTP / DFlash — not working yet.
• Concurrent requests — more than 1 concurrent request is slow (single-kernel launch path needs optimization).
• Decode — single-token decode falls back to Triton on TileLang backend (slower than FA-V100 backend's custom CUDA decode kernel).
• Gemma 4 — not supported yet.

Environment Variables

Attention / Decode

Communication / NCCL

Dense FP16 Fastpath

MoE / AWQ

MoE FP16

GDN / FLA (Gated Delta Network / Flash Linear Attention)

Tuning knobs for SM70 linear attention kernel launch configurations.

Debug / Trace

References

• Original V100 kernel research: ai-bond/flash-attention-v100
• Upstream vLLM: 1CatAI/1Cat-vLLM