GitHub social preview: img/logo_dark_theme_empty_background_1280_x_640.png (1280×640).
Local LLM inference with compressed weights, from GPU-resident models to 32B on a 12 GB Ampere card.
Method · Results · Limits · Install · Quickstart · Contribute · Русский
Status. The project is in active development. Published generate kernels target NVIDIA Ampere (sm_86, measured on RTX 3080 12 GB). Ada, other Ampere-family cards, and SM120 (GeForce RTX 50, first remote SKU RTX 5070 Ti) can generate experimentally; Turing, Hopper, and SM100 generate are refused.
Running a language model on a consumer GPU is a problem of both capacity and data movement. Weights must fit alongside the attention cache and working buffers; weights held in system memory must cross a slower link during inference. Quantizing the file is not enough if inference then expands every linear layer back to a full 16-bit matrix.
deep-fold is an experimental stack around compressed weight storage. A Go compressor writes NF4 weights to the CHR0 container. Ampere CUDA kernels reconstruct values as they multiply and discard them. Embeddings (PackedEmbed) decode only the requested rows. Models that fit stay GPU-resident and use a full greedy-step CUDA graph. Larger models keep a resident subset and stream packed weights from pinned host memory.
On one RTX 3080 12 GB, InternLM2.5-20B hard-12 is 39.0 tokens/s and 9/12 correct, versus 13.2 and 9/12 for Ollama Q4_K. That gap is mainly placement: deep-fold stays resident; Ollama offloads. On 3B and 14B, Ollama is faster. Qwen2.5-32B uses overflow CopyRing. A separate ignore-EOS 64-token plateau is the long-decode comparison with llama.cpp. The sections below give the method, evidence, and current limits.
Packed codes and scales remain in GPU memory. Linear kernels reconstruct the values needed for the current tile without writing a full dense weight matrix back to device memory.
Each group of 64 weights stores four-bit NF4 codes and one FP16 scale:
For aligned matrices, the code-and-scale payload is about 3.76× smaller than BF16. Metadata, alignment, norms, biases, activations, and KV cache add to the runtime footprint. Quantization is lossy.
| Workload | Execution path |
|---|---|
| Prompt processing (prefill) | Ampere tensor-core NF4 GEMM (mma.sync.m16n8k16); reconstruction into BF16 register fragments; prompt chunks of up to 32 positions |
| Resident generation (decode) | N=1 CUDA-core NF4 GEMV inside a full greedy-step CUDA graph |
| Models exceeding the GPU budget | TokenLoop with CopyRing: selected packed matrices stream from pinned host memory through two reusable GPU slots |
Decode V2 keeps token position and valid KV length in device buffers. Its attention kernel reads the valid cache prefix. Before selecting V2, the CLI estimates weights, KV, working buffers, and runtime reserves. Those kernels are written for Ampere tensor cores; other architectures need their own ports.
Figure 1. Two execution paths share the compressed representation. Resident mode keeps packed matrices on the GPU. Overflow adds pinned host storage and two reusable device slots, then the same GEMM kernel. Embeddings decode requested rows on a separate path.
The project connects a CPU compressor and inspectable file format to packed model loading, Ampere CUDA kernels, resident and overflow execution, a local chat CLI, and reproducible experiments. Chat reuses the KV prefix when the tokenized conversation prefix matches.
NF4 was introduced in QLoRA. Fused quantized inference also appears in systems such as Marlin. deep-fold's contribution is its implementation and integration; broader kernel rankings require matched measurements.
Implementation: CHR0 · NF4 layout · Ampere kernel · GEMM · GEMV · Decode V2 · CopyRing.
Reference system: one RTX 3080 12 GB (sm_86), Windows/WDDM. Author-reported measurements from this repository. Published results as of 18 September 2026.
Independent requests, a 2,048-token context limit, greedy generation, and up to 256 new tokens. Rates are arithmetic means of the per-request decode rates. Correctness is the number of answers accepted by the fixture's scorer.
| Model | deep-fold V2, tokens/s | Ollama Q4_K, tokens/s | Correct answers: V2 / Ollama |
|---|---|---|---|
| Qwen2.5-3B-Instruct | 190.9 | 197.2 | 8/12 · 7/12 |
| Qwen2.5-14B-Instruct | 54.8 | 66.2 | 11/12 · 10/12 |
| InternLM2.5-20B-Chat | 39.0 | 13.2 | 9/12 · 9/12 |
On 20B the published mean is about 3× Ollama on this fixture because the packed model stays on the card. Ollama is faster on the 3B and 14B rows. Twelve questions are a small regression check, not a general quality ranking of either quantization scheme.
The updated 20B run uses packed embedding rows: first-request time to first token is 434 ms, mean first-token latency 552 ms, after initialization. CHR loading takes 11.46 s. These are separate timings, not a cold-start guarantee. The 3B and 14B rows are earlier V2 measurements and have not been republished as packed-embedding reruns.
Figure 2. Same 12-item fixture on one RTX 3080. Decode V2 vs Ollama quality and mean decode rate. llama.cpp hard-12 is not published yet. Redraw: python -m gpu.lab.hard_v2_plate --redraw.
Evidence: V2 3B, 14B, 20B; Ollama 3B, 14B, 20B. See the lab log for run history.
A different protocol — ignore-EOS, 64 new tokens, max_seq=2048 — is the long-decode comparison with llama.cpp Q4_K on the same card. Decode V2 reports 197 / 57.5 / 40 tok/s on 3B / 14B / 20B. Exclusive Q4_K long (ctx 2048) is ~187 / 69.9 / 11.87. Those numbers are not the hard-12 table; V2 still attends the full buffer. Details: Decode V2 lab, 3080 comparison.
Figure 3. Resident Decode V2 (GEMV graph + MMA prefill-32) next to exclusive Q4_K long plates. Headline V2: 3B 197, 14B 57.5, 20B 40. Redraw: python -m gpu.lab.decodev2_plate --redraw.
Qwen2.5-32B uses about 16.2 GiB of packed weight payload. In the recorded configuration, roughly 9.5 GiB of weights stay on the GPU and 6.7 GiB are held in pinned host memory. CopyRing stages the host matrices for computation on the GPU.
On the separate 64-token generation run, deep-fold reports 2.49 decode steps/s, or 2.53 tokens/s when normalized to the server-style token counter. Ollama and llama.cpp auto-fit each report 2.54 tokens/s. On hard-12, the published means are 2.12 for deep-fold and 3.1 for Ollama, so the tie does not extend to every workload. 32B protocol.
Figure 4. Quote 32B long plateaus. Ollama 2.54 vs llama.cpp auto-fit 2.54 vs H2 2.53 is CPU-suffix vs CPU-suffix vs PCIe CopyRing, not a kernel ranking. 3B Decode V2 197 sits next to Q4_K ~187. Redraw: python -m gpu.lab.compare_plate --redraw.
- Measured platform: RTX 3080 12 GB, Windows. Ampere
sm_86is the published generate path.sm_80,sm_87, andsm_89are accepted experimentally and unbenchmarked. SM120 (GeForce RTX 50, first remote SKU RTX 5070 Ti) is experimental generate — no plate yet. Turing, Hopper, and datacenter Blackwell (SM100) generate are refused. Linux install works; Linux tokens/s have not been published. - Generation: greedy decoding. Context length consumes KV memory; the capacity estimate uses fixed reserves and cannot guarantee a fit under arbitrary competing GPU load.
- Model coverage: the public download allowlist contains Qwen2.5 3B, 14B, and 32B Instruct, plus InternLM2.5-20B-Chat. Other checkpoints require validation.
- Input format: compression starts from unquantized Hugging Face safetensors. GGUF import is not implemented. VQ 2-bit remains experimental and is excluded from automatic codec selection.
- Evidence: broad quality evaluation, full raw artifacts for every V2 run, repeated-session stability tests, and independent hardware replication remain open work.
Start with Python 3.11 or 3.12, Git, and an NVIDIA driver. The measured GPU capability is sm_86. CPU compression is separate and does not need a supported generate GPU.
| GPU / platform | Current status |
|---|---|
RTX 3080 12 GB, sm_86, Windows |
Measured configuration |
Other sm_86 hardware |
Accepted capability; no matching performance claim |
sm_80, sm_87, sm_89 |
Experimental generation |
| SM120 (GeForce RTX 50; first remote SKU RTX 5070 Ti) | Experimental generation; no named plate |
| Turing, Hopper, SM100 | Generation refused |
| AMD, macOS, CPU-only | No generate path; CPU compress is separate |
git clone https://github.com/wertick01/deep-fold.git
cd deep-fold
powershell -ExecutionPolicy Bypass -File scripts/setup.ps1
.\.venv\Scripts\Activate.ps1
deepfold doctorgit clone https://github.com/wertick01/deep-fold.git
cd deep-fold
bash scripts/setup.sh
source .venv/bin/activate
deepfold doctorSetup installs the CLI and CUDA PyTorch, prepares the Go compressor, and builds or locates the CUDA extension. A source build needs CUDA 12.4 (Ampere) or 12.8 (RTX 50) and a C++ toolchain; Windows setup can attempt their installation.
Full requirements and troubleshooting: English installation guide · Русская инструкция. If the command is unavailable, use python -m gpu.cli in the activated environment.
Qwen2.5-3B is the smallest supported starting point. Allow space for its approximately 6.2 GB source checkpoint plus the compressed file.
deepfold pull Qwen/Qwen2.5-3B-Instruct --dir ./models/Qwen2.5-3B-Instruct --yes
deepfold compress --in ./models/Qwen2.5-3B-Instruct --out ./models/qwen25-3b.nf4.chr --codec nf4
deepfold chat --model ./models/Qwen2.5-3B-Instruct --chr ./models/qwen25-3b.nf4.chr
For a single response:
deepfold run --model ./models/Qwen2.5-3B-Instruct --chr ./models/qwen25-3b.nf4.chr --prompt "Explain why the sky is blue." --max-new-tokens 128
--executor auto selects Decode V2 when the loaded model and estimated memory budget permit it; other supported cases use TokenLoop. --max-seq defaults to 512 for run and 2048 for chat. With --agent on a 12 GB resident 14B card the default is 4096 unless the flag is set.
Keep the model configuration, tokenizer, and any required remote-code files alongside the source model directory. The .chr supplies the packed weights. InternLM additionally requires python -m pip install -e ".[internlm]".
In chat, /help lists commands, /stats shows timings, /clear resets the conversation, and /chats resumes saved conversations. Later turns prefill only the new suffix when the chat-template prefix matches. Optional --agent mode adds workspace tools (grep, patch, tests, allowlisted argv) with an approval policy. See chat and agent documentation and web-search configuration.
More examples: Quickstart · Быстрый старт · Model notes.
The project is early. The occupancy, graphs, and MMA/GEMV paths are Ampere-specific. Contributors who can port those kernels to Ada, Hopper, or Blackwell, and people who can run the same plates on another GPU or on Linux, are the highest-leverage help right now. Independent Ampere replication is equally useful.
Open an issue or pull request on GitHub. A useful report includes deepfold doctor output, the checkout commit, GPU name and CUDA capability, OS, and (when generate ran) the full plate directory — not a single tokens/s number.
To record a local installation and smoke run:
# Windows
powershell -File scripts/plate.ps1 3b --out ./runs/local-3b# Linux
bash scripts/plate.sh 3b --out ./runs/local-3bThe runner can download and compress the model. It records the machine, checkout, and smoke results; it does not reproduce the hard-12 comparison by itself. Keep the entire output directory.
For controlled comparisons, record the commit, checkpoint revision, engine versions, prompt and output token counts, context limit, placement, warmup, and timing definitions. Report startup, first-token latency, and decode speed separately. Do not fold CUDA-graph capture into the warmup timer. Laboratory guide · Hard-12 evaluation · Committed runs.
| Area | Start here |
|---|---|
| Architecture ports and doctor gates | Ampere kernel, CLI probe |
| Compression and verification | Go CLI, CPU round trip |
| Resident inference | Decode V2, design |
| Overflow execution | CopyRing, recorded 32B data path |
| CUDA arithmetic | NF4 kernels and checks |
| CLI and interaction | CLI source, UX notes |
Numerical checks and end-to-end generation timings should accompany performance changes.
MIT. Downloaded models retain their own licenses.
Developed with Cursor.