imp

KO | EN

gitlite — search

#python #react #llm #machine-learning #rust #deep-learning #docker #macos #windows #vue #security #ios #mcp #open-source

(★ 35)

From-scratch C++/CUDA inference engine for the NVIDIA RTX 5090 (sm_120a) — the best single-GPU backend for agentic AI: tool calling, long-context loops, reasoning and concurrent sub-agents on top of the fastest single-stream decode on the 5090 (beats llama.cpp, at-or-ahead of vLLM on NVFP4). 100% written by Claude Code.

imp 최신버젼 다운로드

최종 버전 다운로드 (.zip)

// repository documentation

Was this content helpful?

★ 0 (0 ratings)

From-scratch CUDA inference engine for the NVIDIA RTX 5090 (<code>sm_120a</code>). 
 The best single-GPU backend for agentic AI: tool calling, long-context loops, reasoning and concurrent sub-agents, 
 on top of the fastest single-user inference on the 5090: faster decode than llama.cpp (+42–48% dense GGUF at b9976, 2026-07-12), 
 at-or-ahead of vLLM on NVFP4, and the only engine running native NVFP4 on consumer Blackwell. 
 ~97k lines, 100% written by <a href="https://claude.ai/claude-code">Claude Code</a>.

---

## What is imp

A C++23/CUDA inference engine that targets exactly one architecture: the NVIDIA RTX 5090 / RTX PRO 6000 (GB202, `sm_120a`). The build emits raw `sm_120a` SASS via direct gencode: no portability layer, no FP16 dequant fallback in the hot path, no wrapper around llama.cpp or vLLM. A `compute_120f` PTX fallback covers the other Blackwell consumer SKUs (RTX 5080 / 5070 Ti). imp ships its own GGUF and SafeTensors loaders, BPE tokenizer, paged KV cache, attention kernels, MoE routing, Gated DeltaNet + Mamba2 scans, CUDA Graphs, and an OpenAI/Anthropic-compatible HTTP server built for **agentic workloads**: tool calling, `json_schema` constrained decoding, separable reasoning/thinking channels, prefix-cached long-context multi-turn, and concurrent sub-agent fan-out. In total ~97k lines of C++/CUDA (plus ~50k lines of tooling and tests), every line generated by an AI coding agent.

## Why this exists

The RTX 5090 shipped with native FP4 tensor cores, 32 GB GDDR7 at 1.8 TB/s, and a new ISA (`sm_120a`) that no existing inference engine fully exploits. llama.cpp targets broad hardware compatibility and runs GGUF through dequant-to-FP16 paths: fast everywhere, but leaving consumer-Blackwell-specific features (NVFP4 block-scaled `mma.sync` MMA, FP8 `f8f6f4` scores) on the table. vLLM targets datacenter Blackwell (B200/B300, `sm_100`) and gates key backends like FlashInfer MoE on `tcgen05`, an opcode family that consumer Blackwell (`sm_120`) doesn't have, so vLLM's NVFP4 path falls back to Marlin or fails outright on the 5090.

Consumer Blackwell is **not** a smaller datacenter Blackwell: `sm_120a` has no `tcgen05`, no TMEM, no `wgmma`, and no TMA warp-specialized grouped GEMM. imp is built against that reality: the FP4 path is `mma.sync mxf4nvf4` with FlashAttention-2-style block-scaling, not the Hopper/B200 kernel designs.

That leaves a gap: 30B-class MoE models (Qwen3-Coder, Qwen3.6, Gemma-4) fit in 32 GB as NVFP4 prequant and should decode at 150–270 tok/s on this hardware, but neither engine delivers that today. imp exists to close that gap: a single-target engine that treats the RTX 5090 as the only GPU that matters and optimizes accordingly.

NVFP4 prequant (from [NVIDIA Model Optimizer](https://github.com/NVIDIA/Model-Optimizer) / llm-compressor) is the native weight format for Blackwell's FP4 tensor cores. imp loads these SafeTensors directly into its NVFP4 decode cache and CUTLASS v4.5.2 GEMM kernels with no re-quantization (there are no FP4 cuBLASLt kernels on `sm_120`, so CUTLASS is the primary GEMM path). For GGUF models, imp builds an NVFP4 decode cache that converts Q8_0/Q6_K weights to FP4 at init time, getting the bandwidth benefit of sub-byte weights on the decode hot path while keeping full-precision prefill via source dequant.

## Performance

**On a single RTX 5090, imp is the fastest engine for single-user and agentic
inference**: it leads llama.cpp on decode and is at-or-ahead of vLLM on NVFP4
(prefill *and* decode), while being the only engine that runs native NVFP4 on
consumer Blackwell. That single-stream speed is the foundation the agentic
surface (tool calling, long-context loops, concurrent sub-agents) is built on.
Every number below is measured, dated, and commit-anchored in
[BENCHMARKS.md](docs/BENCHMARKS.md) (single 5090, greedy, CUDA 13.3, CUDA Graphs on;
decode is the reliable A/B signal):

- **GGUF dense decode:** Qwen3-8B Q8_0 at **~270 tok/s** (CI-gated baseline),
  **+42–48% over llama.cpp** (re-swept 2026-07-12 vs build 9976 `e3546c794`,
  full offload, flash attention); hybrid Qwen3.6-35B UD-Q4_K_M **+18%**,
  Gemma-4-26B MoE **+21%** same sweep.
- **NVFP4 SafeTensors decode:** 30B-class MoE at **~271–390 tok/s**
  (2026-07-12 scoreboard: Qwen3-30B-A3B 390, Qwen3-Coder-30B 389, Qwen3.6-35B
  hybrid 311, Gemma-4-26B 271): **uncontested on `sm_120`**: vLLM's NVFP4 path
  needs `tcgen05` (absent on consumer Blackwell) and llama.cpp has no native
  NVFP4 at all.
- **NVFP4 long-context prefill:** at-or-ahead of vLLM after #687 (FP16-QK FA2 as
  the primary hd=128 prefill): MoE pp4096 **leads vLLM (+4%)**, MoE pp2048 +27%,
  dense pp2048 a tie, and TTFT wins everywhere.

**Honest losses & ties** (narrow, and documented): NVFP4 dense pp4096 still
trails vLLM by ~4% (1.04×). llama.cpp's MXFP4/MoE-GGUF decode caught up
substantially by build 9976: the non-hero Qwen3-30B Q4_K_M margin is within
noise (+1.7%). The brief gpt-oss-20b tie from that sweep was reversed by the
FP8 attention-projection decode sidecar (#984/PR #990: imp ~391 vs llama
329–345, +13–19%). The former 35B-hybrid GGUF loss was closed by extending
the FP8 SSM-projection sidecar to UD quants (#962, 224 → 267–272 tok/s, +18%
ahead of llama.cpp).

Every number, with date, commit SHA, CUDA version, quant and the exact
command: **[BENCHMARKS.md](docs/BENCHMARKS.md)**. Methodology details:
[`docs/performance.md`](docs/performance.md).

## Should I use this?

**Use [llama.cpp](https://github.com/ggerganov/llama.cpp) if** you want a stable, mature engine with broad model/hardware support and CPU fallback. It runs everywhere and supports everything.

**Use [vLLM](https://github.com/vllm-project/vllm) if** you serve high-concurrency batched workloads on datacenter GPUs (H100/B200). It has continuous batching, PagedAttention at scale, and production-grade serving.

**Use imp if** you run a Blackwell consumer/workstation card (RTX 5090, RTX PRO 6000) and want the fastest backend for **agentic AI** (coding agents, tool-using assistants and reasoning loops) on one card. It is the fastest engine on the 5090 for both decode and NVFP4 prefill (native NVFP4/FP4, no dequant overhead), with a first-class agentic surface: tool calling, `json_schema` constrained decoding, separable reasoning channels, prefix-cached long-context multi-turn, and moderate concurrency for sub-agent fan-out. The trade-off is scope, not speed: it's single-GPU by design, with no multi-GPU and no datacenter-scale batching.

## Known limitations

- **Single GPU only.** No tensor parallelism, no multi-GPU.
- **Consumer Blackwell only.** `sm_120a` SASS + `compute_120f` PTX fallback. Other GB202 SKUs like the China-market RTX 5090D and 5090D v2 share compute capability 12.0 and run the native SASS, but only the 32 GB RTX 5090 is tested. Cut-down SKUs scale down accordingly: the 24 GB / ~1.34 TB/s 5090D v2 fits a smaller model set (the 30B-class NVFP4 heroes are sized against 32 GB) and decodes roughly 25% slower, since decode is bandwidth-bound. No Hopper, Ada, Ampere, datacenter Blackwell. No AMD, Intel, Apple, or CPU paths.
- **GGUF prefill: largely fixed 2026-06-07.** The INT8-IMMA prefill family (#612–#617) puts Qwen3-30B-A3B and Qwen3-14B-Q6_K AHEAD of llama.cpp; Q8_0 dense sits at 1.13×, gemma-4 at 1.20×, Qwen3.6-35B at 1.55× (GDN share is quality-locked). NVFP4 prefill: imp WINS MoE pp4096 (+4% vs vLLM), MoE pp2048 (+27%), and TTFT everywhere; only dense pp4096 still trails, at ~1.04× (2026-06-13, post #687 FP16-QK FA2 primary hd=128 prefill; see docs/audit/prefill_gap_2026_06_07.md for the kernel-level decomposition).
- **Legacy GGUF hybrid decode trails the NVFP4 path.** The FP16 GDN-projection tax on Qwen3.6-35B hybrid decode was closed on the NVFP4 checkpoint path by the #949 FP8 SSM-projection sidecar (35B NVFP4 decode 257 → ~320 tok/s, ahead of llama.cpp). The Q4_K GGUF path is a no-op for that sidecar and decodes slower (~213 tok/s); GGUF is the legacy path — NVFP4 SafeTensors is the priority.
- **Only tested models work reliably.** Anything not on the [supported list](docs/supported-models.md) may load but hasn't been verified.
- **Prefill numbers are noisy.** cuBLAS autotuning causes up to 2.6× variance across container restarts.
- **Moderate concurrency, not datacenter scale.** Concurrent sub-agent fan-out (tens of requests) is a tuned, supported workload; batch=64+ continuous-batching throughput on a rack is not; that's vLLM/SGLang territory.
- **Single-user / agentic scope.** Built and tuned for single-user and agentic inference on a 5090, not high-concurrency datacenter serving.

## Quickstart

Everything runs in Docker, no local CUDA toolkit needed. Prebuilt images are on
[GHCR](https://github.com/kekzl/imp/pkgs/container/imp) (built per release for
x86-64 + sm_120a):

```bash
# Drop a GGUF or SafeTensors model into ./models/
mkdir -p models

# Run the server from the prebuilt image. The cache volume is optional but
# pays for itself on the second start: it holds the transformed weights
# (Qwen3-14B-NVFP4 init 7.9 s -> 2.1 s) and the measured library reserve.
docker run --gpus all -v ./models:/models -v imp-cache:/home/imp/.cache/imp \
  -p 8080:8080 \
  ghcr.io/kekzl/imp:latest --model /models/your-model.gguf

# Hit the OpenAI-compatible endpoint (the model id is the file/dir basename;
# GET /v1/models lists it, and the `model` field is required).
curl -s http://localhost:8080/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"your-model.gguf","messages":[{"role":"user","content":"Hello!"}],"max_tokens":64}'
```

Or open <http://localhost:8080> in a browser: the server ships a small chat UI
that streams the answer and plots inter-token latency while it is written.

Or build from source (tracks `main` instead of the latest release):

```bash
git clone https://github.com/kekzl/imp.git && cd imp
docker compose build imp-server
docker run --gpus all -v ./models:/models -p 8080:8080 \
  imp:latest --model /models/your-model.gguf
```

CLI reference, server flags, config, and C API: [`docs/usage.md`](docs/usage.md).

## Supported models

| Family | Variants | Quantizations |
|---|---|---|
| Qwen3 / Qwen3-MoE | dense + MoE (Coder-30B-A3B) | Q4_K_M, Q6_K, Q8_0, NVFP4 |
| Qwen3-VL | text + vision, tower in the checkpoint | SafeTensors (bf16) |
| Qwen3.5 / Qwen3.6 | GDN + attention (+ MoE) | Q4_K_M, Q8_0, NVFP4 |
| Gemma-4 | 26B-A4B MoE, 31B dense | Q4_K_M, Q5_K_M, Q8_0, NVFP4 |
| Gemma-3 | text + vision (SigLIP) | GGUF |
| Phi-4-reasoning-plus | dense, fused projections | NVFP4 |
| gpt-oss-20b | MoE (32 experts, top-4), Harmony | SafeTensors MXFP4 (native) |
| Nemotron-H | Mamba2 + Attention + MoE | NVFP4, GGUF |
| Llama / Mistral / DeepSeek | dense + MoE | GGUF (Q*_K, Q8_0) |
| DeepSeek-V2 (MLA) | Multi-head Latent Attention + MoE | SafeTensors (bf16) |

**Vision** is Qwen3-VL (tower in the checkpoint, no extra flag) and Gemma-3/Gemma-4
through GGUF + `--mmproj`. Those are the only ones: a multimodal SafeTensors
checkpoint from another family — Gemma-4 NVFP4, the Qwen3.5/3.6 MoE checkpoints —
loads **text-only**, and a request that sends it an image is refused with
`400 vision_unavailable` rather than answered from the text. Details:
[supported-models.md](docs/supported-models.md#vision).

VRAM, decode tok/s, and per-model notes: [`docs/supported-models.md`](docs/supported-models.md).

## Features

| | |
|---|---|
| **Quantization** | GGUF Q4_K_M/Q5_K_M/Q6_K/Q8_0 + IQ4_NL/IQ4_XS, SafeTensors NVFP4 (prequant), MXFP4. NVFP4 KV cache (`--kv-nvfp4`) for 4× context compression at decode parity. SWA-aware KV sizing (`kv_cache.swa_sizing=auto`): sliding-window layers allocate only their window (gemma-3/4 ~5-6× fewer KV tokens, gpt-oss ~2×) — auto-active when prefix caching is off. Auto context ceiling 128K, VRAM- and model-bounded. |
| **LoRA** | PEFT adapter hot-swap (`--lora name=path`, per-request `"lora"` field): runtime low-rank deltas, no weight patching, works with every quant path. |
| **Attention** | Prefill: FP16 cuBLAS below the auto `fmha_prefill_threshold` (the largest chunk whose S-matrix fits, ~2.5k tokens), then the FMHA family above it, an `mma.sync` `m16n8k32` FP8-E4M3 score kernel and a register-resident FlashAttention-2 kernel (head_dim 128 and 256, both default-on). Decode: paged attention (block_size 16) switching on KV dtype (FP16/FP8/INT8/INT4/NVFP4/MXFP4). Auto-dispatch per phase × dtype × layer, see [`docs/attention-dispatch.md`](docs/attention-dispatch.md). |
| **Architectures** | Dense transformers, Mixture-of-Experts (top-k grouped GEMM), Multi-head Latent Attention (DeepSeek-V2; materialized + opt-in absorbed latent-KV-cache decode), Gated DeltaNet (fused recurrent scan), Mamba2 (SSM), SigLIP/Gemma-4v vision encoders, and the Qwen3-VL tower — dynamic resolution (no fixed image size), DeepStack features injected at the LM's first layers, three-axis M-RoPE. |
| **`sm_120a` kernels** | NVFP4 block-scaled `mma.sync mxf4nvf4` GEMM/GEMV (CUTLASS v4.6.1), FP8 `f8f6f4` attention scores, FA2 block-scaling, packed `cvt.e2m1`/`cvt.e4m3x2` dequant, PDL, Green Contexts. **No** `tcgen05`/TMEM/`wgmma`/TMA-WS; those are datacenter Blackwell only. |
| **Server** | OpenAI `/v1/chat/completions` + `/v1/responses` (Responses API, Agents SDK / Codex dialect, native SSE events with incremental tool-call argument deltas) + `/v1/completions` + `/v1/embeddings` + `/v1/rerank` (Cohere/Jina-shaped cross-encoder reranking, also served at `/rerank`) + `/tokenize` + `/detokenize`; Anthropic `/v1/messages` with real per-token SSE streaming (`ping` keepalives), `cache_control` prompt caching (prefix-cache pinning + `cache_read`/`cache_creation_input_tokens` usage reporting; prefix cache default-on since #538, model-fingerprint-gated on disk so a cache file from a different model is never replayed); Prometheus `/metrics` with TTFT / inter-token-latency histograms, cancellation counters, and a memory breakdown that separates capacity from occupancy (per-tier `imp_memory_reserved_bytes` / `imp_memory_live_bytes`, `imp_kv_blocks_total` / `imp_kv_blocks_used`, `imp_vram_budget_bytes` / `imp_vram_own_bytes` / `imp_vram_own_peak_bytes`). Image input on `/v1/chat/completions` (`image_url` content parts, data-URI or fetched over HTTP; the prefix cache is keyed on the picture, not just the token ids). Tool/function calling, constrained decoding in four flavours — `json_object`, `json_schema` (both whole-token validated), `response_format: regex` / `guided_regex`, and GBNF grammars via `response_format: grammar` / llama.cpp's `grammar` / vLLM's `guided_grammar` (a pushdown simulator, so recursive and bracket-balanced formats are enforceable, not just regular ones) — `reasoning_content` separation (DeepSeek format) + think budget. Per-request speculative-decode override (`"speculative": true/false`) and an opt-in deterministic mode (`--set runtime.deterministic=true`, ordered MoE reduction). Client-disconnect cancellation reclaims the slot + KV within one scheduler tick. Model swapping on request (`server.model_swap`, default on): a request naming another model in the models directory swaps to it — in-flight generations drain first and are never cancelled, and a failed load restores the previous model, so the server is never left empty. `/v1/models` lists the directory with `loaded: true|false`. Serial by nature (32 GB holds one model; the requesting call pays one load, warmed by the on-disk weight cache); `server.model_swap=false` restores strict single-model semantics. Suspend-to-RAM: `POST /admin/suspend` parks the weights in host RAM and frees the GPU completely (CUDA context reset — ~0 MiB for the process), `POST /admin/resume` restores them warm and serves again in seconds (inference answers 503 while suspended). API-key auth, rate limiting, JSONL request logging. Agent-shaped load harnesses: `tools/agent_bench.py` (TTFT/ITL p50/p99 under concurrency, warm-vs-cold cache) and `tools/agent_replay_bench.py` (per-turn TTFT on a growing multi-turn transcript, prefix cache on/off); `make bench-agentic`. |
| **Runtime** | CUDA Graphs (auto per model), `imp.conf` + CLI config, Jinja2 chat templates with macro support. `degen_suite.py` is the coherence quality-gate after hot-path changes. |

## How was this built?

imp is written entirely by [Claude Code](https://claude.ai/claude-code) (Opus 4.6 → 4.8): 1200+ commits, ~97k lines of C++/CUDA in the engine itself (plus ~50k more in tooling and 140+ test files / 1500+ tests). A human ([@kekzl](https://github.com/kekzl)) directs architecture decisions, picks models, and runs benchmarks. The AI agent writes all code: GGUF parser, BPE tokenizer, paged KV cache, CUTLASS NVFP4 GEMM dispatch, FP8/FA2 attention kernels, MoE routing, GDN/Mamba2 scan kernels, HTTP server, CI, the build system; nothing is copy-pasted from other engines.

This is a running experiment in how far AI-generated systems code scales. The [architecture doc](docs/architecture.md) covers the end-to-end pipeline if you want to see how the pieces fit together.

## Documentation

| Document | |
|---|---|
| [Architecture](docs/architecture.md) ([diagram](docs/architecture.svg)) | Load → engine init → prefill → decode pipeline |
| [Memory architecture](docs/MEMORY_ARCHITECTURE.md) | Lifetime tiers, allocators, ownership invariants, VRAM planning |
| [Supported models](docs/supported-models.md) | Tested families with VRAM + tok/s |
| [Performance](docs/performance.md) | Throughput numbers + methodology |
| [Usage & reference](docs/usage.md) | Build, server, CLI, C API |
| [Quantization](docs/quantization.md) | GGUF, NVFP4, MXFP4, FP8 KV: formats and trade-offs |
| [Attention dispatch](docs/attention-dispatch.md) | Per-(phase x dtype x layer) kernel selection |
| [sm_120a kernels](docs/sm120.md) | Kernel optimization notes |
| [Determinism](docs/determinism.md) | Reproducibility guarantees + known limits |
| [imp.conf reference](imp.conf.example) | All runtime configuration keys |
| [Roadmap](docs/roadmap.md) | Open issues and in-flight work |
| [Changelog](CHANGELOG.md) | Per-release notes |

## Building from source

```bash
# With CUDA 13.3+ on the host:
cmake -B build -DCMAKE_BUILD_TYPE=Release
cmake --build build -j$(nproc)

# Or via Docker (canonical):
make build          # → imp:test image
make verify-fast    # build + tests + perf gate + peak-VRAM gate
```

Full build options and test commands: [`docs/usage.md`](docs/usage.md). Contributing: [`CONTRIBUTING.md`](CONTRIBUTING.md).

## License

MIT, see [`LICENSE`](LICENSE).

## Acknowledgements

Stands on the shoulders of [llama.cpp](https://github.com/ggerganov/llama.cpp). The GGUF format, GGML quantization schemes, and most practical conventions for local LLM inference were established there.

Built by [@kekzl](https://github.com/kekzl) with [Claude Code](https://claude.ai/claude-code). Heavy use of [CUTLASS](https://github.com/NVIDIA/cutlass) v4.5.2 for `sm_120a` NVFP4/MXFP4 GEMM and grouped MoE kernels; the FMHA/FA2 attention kernels are hand-written `mma.sync`. Other references: Flash Attention 2, NVIDIA Model Optimizer, [llm-compressor](https://github.com/vllm-project/llm-compressor).

imp

imp 최신버젼 다운로드

오늘의 추천 프로젝트

Was this content helpful?

Recent Feedback

imp

imp 최신버젼 다운로드

오늘의 추천 프로젝트

Was this content helpful?

Recent Feedback

Download README