dotllm 0.3.0-dev.2974
See the version list below for details.
dotnet tool install --global dotllm --version 0.3.0-dev.2974
dotnet new tool-manifest
dotnet tool install --local dotllm --version 0.3.0-dev.2974
#tool dotnet:?package=dotllm&version=0.3.0-dev.2974&prerelease
nuke :add-package dotllm --version 0.3.0-dev.2974
<div align="center">
dotLLM
High-performance LLM inference engine written natively in C#/.NET
Website · Documentation · Roadmap · Discussions
</div>
About
dotLLM is a ground-up LLM inference engine for .NET — not a wrapper around llama.cpp or Python libraries. All orchestration, model loading, tokenization, sampling, and CPU compute are implemented in pure C#, with CUDA GPU acceleration via PTX kernels loaded through the CUDA Driver API (no native shared library). It targets transformer and SSM/MoE model families including Llama, Mistral, Phi, Qwen, DeepSeek-V2/V3, Mixtral, Granite-MoE, Nemotron-H, and Mamba-3 with SIMD-optimized CPU plus CUDA and Vulkan GPU backends.
Status: Active post-Mamba-3 development — speculative decoding, paged KV-cache allocation, Native AOT (experimental), startup warm-up, constrained decoding (JSON/schema/regex/grammar), tool calling, prompt caching, logprobs, LoRA adapter loading, CUDA CPU/GPU hybrid offload, and native Vulkan real-weight parity for multiple architectures. See docs/SUPPORTED_MODELS.md for the current support matrix.
Key Features
Performance
- Zero-GC inference — unmanaged memory (
NativeMemory.AlignedAlloc, 64-byte aligned) for all tensor data; no managed heap allocations on the hot path - SIMD vectorization —
TensorPrimitives+ hand-tunedSystem.Runtime.Intrinsicsfor quantized matmul, RMSNorm, RoPE, softmax - Memory-mapped model loading — GGUF files loaded via
MemoryMappedFile; OS demand-paging means multi-GB models load in milliseconds - Quantized inference — FP16/BF16, Q8_0, K-quants, and other GGUF quantization formats; fused scale×int dot-product kernels operating directly on quantized blocks where implemented
Architecture Support
- Transformer/SSM models — dense transformers (incl. ternary BitNet b1.58), MLA/MoE DeepSeek variants, MoE families, hybrid Nemotron-H, and pure-SSM Mamba-3. See docs/SUPPORTED_MODELS.md for the full matrix.
- Attention mechanisms — MHA, MQA, GQA, sliding-window masks, and DeepSeek MLA via parameterized
ModelConfig, withIAttentionStrategyfor kernel selection - Position encoding — RoPE variants and Mamba-3 data-dependent RoPE plumbing, with a pluggable
IPositionEncodingsurface for additional encodings - Composable sampling —
ISamplerStepchain: repetition penalty → temperature → top-k → top-p → min-p → categorical sample
Serving
- OpenAI-compatible API —
/v1/chat/completions,/v1/completions, tool calling, streaming via ASP.NET - Paged KV-cache — PagedAttention with block-level allocation, prefix caching, and copy-on-write
- Speculative decoding — draft-verify-accept with KV-cache rollback (greedy mode today; non-greedy planned — see issue #121)
- Structured output — FSM/PDA-based constrained decoding guaranteeing valid JSON, JSON Schema, regex, and grammar
- (Planned) Continuous batching — iteration-level scheduling with preemption and priority queuing — Phase 9, see Roadmap
Extensibility
- Pluggable backends —
IBackendinterface with separate packages per backend (CPU, CUDA, Vulkan, ROCm/HIP foundation) - Diagnostic hooks — zero-cost
IInferenceHookpoints for activation capture, logit lens, SAE integration - LoRA adapters — runtime loading, no weight merging, CPU/Vulkan application paths, and concurrent multi-adapter serving
- OpenTelemetry observability —
System.Diagnostics.Metrics+Activityfor throughput, latency, and per-request tracing, OTLP exporter wired on the server side
Architecture Overview
dotLLM is organized as a layered architecture where each layer depends only on the layers below it:
┌─────────────────────────────────────────┐
│ DotLLM.Server │ ASP.NET OpenAI-compatible API
├─────────────────────────────────────────┤
│ DotLLM.Engine │ KV-cache, scheduler, samplers,
│ │ constraints, speculative decoding
├──────────┬──────────┬───────────────────┤
│ DotLLM. │ DotLLM. │ DotLLM.Cpu/Cuda │ GGUF/SafeTensors, BPE/SPM,
│ Models │Tokenizers│ (backends) │ SIMD kernels / CUDA kernels
├──────────┴──────────┴───────────────────┤
│ DotLLM.Core │ Interfaces, tensor types, config
└─────────────────────────────────────────┘
Each layer is its own project/package (the .nupkgs are attached to every GitHub release; see NuGet Packages for what is on nuget.org). DotLLM.Core defines all abstractions (ITensor, IBackend, IModel, ISamplerStep, etc.) while concrete implementations live in their respective projects.
Getting Started
There are two paths: grab a pre-built release and run it, or clone the repo and build from source.
Use a pre-built release
Pick one of three install options. Builds are published from the dev branch as GitHub prereleases (0.3.0-dev.<n>); there is no stable release yet, so the NuGet and dnx options need --prerelease. See docs/PACKAGING.md for the release channels.
Option A — run or install as a .NET tool (needs the .NET 10 SDK, not just the runtime: dnx and dotnet tool are SDK commands):
# Run without installing (downloads and caches the NuGet package `dotllm`)
dnx dotllm --prerelease -- run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" -n 64
dnx dotllm --prerelease -- serve QuantFactory/SmolLM-135M-GGUF # OpenAI-compatible API + chat UI
# ...or install the `dotllm` command globally
dotnet tool install -g dotllm --prerelease
dotllm run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" -n 64
dotllm serve QuantFactory/SmolLM-135M-GGUF
Notes on dnx (checked against the .NET 10 SDK):
- Arguments after the package id are forwarded to the tool, but
dnxclaims--help/-hand--versionfor itself:dnx dotllm --prerelease serve --helpprintsdnx's help, whilednx dotllm --prerelease -- serve --helpprintsdotllm serve's. Put--before the tool's arguments whenever they include those flags (it is harmless otherwise). --yes(-y) is accepted and skips the confirmationdnxmay ask for before running a package for the first time; use it in scripts. No prompt appeared when run from a non-interactive shell.- Pin a version with
dnx dotllm@0.3.0-dev.2470 -- ...for reproducibility. - The package id is
dotllm(lower case; the tool command is alsodotllm). Do not installDotLLM.Cli: that id belongs to the upstream project and is frozen at0.1.0-preview.3.
Option B — download a self-contained binary (no .NET install needed — the runtime is bundled):
Grab the archive for your platform from the fork's releases page (pick the newest dotLLM 0.3.0-dev.<n>; prereleases are never marked "latest", so releases/latest will not find them):
- Windows x64:
dotllm-<version>-win-x64.zip(also containsdotllm-tray.exe, the system tray) - Linux x64:
dotllm-<version>-linux-x64.tar.gz - macOS (Apple Silicon):
dotllm-<version>-osx-arm64.tar.gz
Every release also carries a SHA256SUMS file and a build-provenance attestation. The builds are not Authenticode-signed, so Windows SmartScreen will warn; verify the download instead:
sha256sum -c SHA256SUMS --ignore-missing
gh attestation verify dotllm-<version>-linux-x64.tar.gz --repo jamesburton/dotLLM
Unpack and run:
# Linux / macOS
tar -xzf dotllm-<version>-linux-x64.tar.gz
cd dotllm-<version>-linux-x64
./dotllm run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" -n 64
./dotllm serve QuantFactory/SmolLM-135M-GGUF # OpenAI-compatible API + chat UI
# Windows
Expand-Archive dotllm-<version>-win-x64.zip -DestinationPath .
cd dotllm-<version>-win-x64
.\dotllm.exe run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" -n 64
Tagged (v*) releases additionally attach experimental Native AOT builds (dotllm-<version>-aot-<rid>.{zip,tar.gz}) — smaller and faster to start, but please file an issue if you hit a crash.
Option C — reference the libraries from your .NET app — see NuGet Packages below.
Build from source
Clone the repository and build with the .NET 10 SDK.
Prerequisites:
- .NET 10 SDK
- Python 3.10+ with
pip install rich InquirerPy— only needed for the benchmark scripts underscripts/ - Optional: llama.cpp for comparison benchmarks (see llama.cpp setup)
git clone https://github.com/kkokosa/dotLLM.git
cd dotLLM
dotnet build -c Release
When built from source, replace dotllm <subcommand> in the Usage examples below with dotnet run --project src/DotLLM.Cli -c Release -- <subcommand>.
Usage
The dotllm CLI has these commands:
dotllm model— pull, add, list, search, inspect and delete GGUF models; create and manage named model profiles; import ollama modelsdotllm run— single-shot text generation with a performance summarydotllm chat— interactive multi-turn REPL with chat template formattingdotllm serve— OpenAI-compatible HTTP API (plus Anthropic- and ollama-compatible routes) with a built-in web chat UIdotllm ps/dotllm stop— list / unload the models of a running serverdotllm bench/dotllm perplexity— throughput and quality measurement
Naming and fetching models
One resolver turns a model name into a GGUF in run, chat, serve, bench, and in a request's model field:
| you write | resolves to |
|---|---|
C:\models\x.gguf |
that file |
owner/repo |
the repo's best local GGUF (Q4_K_M first, not the largest file) |
owner/repo:Q8_0 |
the file whose name contains the tag (same as --quant Q8_0) |
llama3.2:3b |
a model in your local ollama store, loaded in place (see ollama) |
my-profile |
a model profile |
Missing models are downloaded for you by dotllm run|chat|serve owner/repo[:tag] (resumable, with progress, into the Hugging Face hub cache). The server only downloads on demand if started with --auto-pull — otherwise a remote client could trigger multi-gigabyte downloads just by naming a repo. dotllm model pull fetches ahead of time. Models already in ~/.dotllm/models/ or the Hugging Face hub cache (HF_HUB_CACHE, ~/.cache/huggingface/hub) are found without re-downloading. Details: docs/SERVER.md.
Manage models
# Search HuggingFace for GGUF repos
dotllm model search llama --limit 5
# Download ahead of time (a tag picks the file: owner/repo:Q4_K_M)
dotllm model pull QuantFactory/SmolLM-135M-GGUF
# Multi-GGUF repo: name the file. Without --file (or a :tag) pull lists the files
# and needs an interactive terminal to choose one, so scripts/CI must pass it.
dotllm model pull TheBloke/Llama-2-7B-GGUF --file llama-2-7b.Q4_K_M.gguf
# List everything available locally
dotllm model list
# Show architecture, quantizations, and tokenizer info for a repo
dotllm model info QuantFactory/SmolLM-135M-GGUF
# Remove a downloaded repo (--quant narrows to one file) or a profile
dotllm model delete QuantFactory/SmolLM-135M-GGUF
Custom models and profiles
# Register your own GGUF (hard-linked, not copied) or download one from a URL; listed as local/<name>
dotllm model add C:\models\my-finetune.gguf --name my-finetune
# A profile (Modelfile equivalent): base model + system prompt + sampling defaults + device + keep-alive
dotllm model create terse --from bartowski/Llama-3.2-3B-Instruct-GGUF:Q4_K_M \
--system "Answer in one sentence." --temperature 0.2 --device vulkan --keep-alive 600
dotllm model show terse # merged settings and where the base model resolves
dotllm model cp terse terse-hot # copy a profile, or alias a model under a new name
Profiles live in ~/.dotllm/profiles/ (DOTLLM_PROFILES_DIR overrides). A profile's name is a model id for the server ("model": "terse"), which applies its system prompt and defaults to requests that omit them. dotllm run|chat resolve a profile to its base model but do not yet apply its system prompt or sampling. model create options: --from/-f, --system, --system-file, --temperature, --top-p, --top-k, --min-p, --repeat-penalty, --max-tokens, --seed, --stop (repeatable), --device/-d, --gpu-layers, --keep-alive, --description, --force.
Use your ollama models
dotllm run llama3.2:3b -p "Hello" # loaded in place from your ollama store, nothing copied
dotllm model import-ollama # turn the whole store into profiles (system prompt + parameters)
dotllm model pull ollama:llama3.2:3b # or pull from registry.ollama.ai (sha256-verified)
dotLLM only reads the ollama store, never writes to it; ollama's Go prompt templates are not carried over (the chat template embedded in the GGUF is used). See docs/SERVER.md.
Run — single-shot generation
Encodes a prompt, streams tokens to stdout, and prints a performance + memory summary.
# Greedy generation (default: temperature=0, max-tokens=128)
dotllm run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" -n 64
# Sampled generation
dotllm run QuantFactory/SmolLM-135M-GGUF -p "Once upon a time" -n 128 -t 0.7 --top-k 40 --top-p 0.95
# JSON output (for scripting / piping)
dotllm run QuantFactory/SmolLM-135M-GGUF -p "Hello" --json
# Long prompts: read from a file instead of the command line
dotllm run QuantFactory/SmolLM-135M-GGUF --prompt-file prompt.txt -n 256
# Select a specific quantization when a repo has multiple .gguf files
dotllm run QuantFactory/SmolLM-135M-GGUF -p "Test" -q Q8_0
# GPU inference: CUDA (NVIDIA GPU) ...
dotllm run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" --device gpu
# ... or Vulkan (AMD / Intel / NVIDIA; the GPU path on AMD iGPUs such as Strix Halo)
dotllm run QuantFactory/SmolLM-135M-GGUF -p "The capital of France is" --device vulkan
# NUMA / P-core aware CPU threading
dotllm run QuantFactory/SmolLM-135M-GGUF -p "Test" --threads 8 --decode-threads 4 --numa-pin
# KV-cache quantization (Q8_0 / Q4_0) to fit longer contexts in memory
dotllm run QuantFactory/SmolLM-135M-GGUF -p "..." --cache-type-k q8_0 --cache-type-v q8_0
# Constrained JSON output
dotllm run QuantFactory/SmolLM-135M-GGUF -p "List 3 colors as JSON." --response-format json_object
# Speculative decoding — target + draft must share the same vocabulary.
# Example pair (validated by scripts/test_models_speculative.py):
# target: Llama-3.2-3B-Instruct Q8_0 draft: Llama-3.2-1B-Instruct Q4_K_M
dotllm model pull bartowski/Llama-3.2-3B-Instruct-GGUF
dotllm model pull bartowski/Llama-3.2-1B-Instruct-GGUF
dotllm run bartowski/Llama-3.2-3B-Instruct-GGUF -q Q8_0 \
-p "Explain what a CPU cache is in one sentence." -n 48 \
--speculative-model bartowski/Llama-3.2-1B-Instruct-GGUF --speculative-k 5
Sample output:
── dotllm | Llama 30L/576H | Q8_0 | 16 threads | greedy ──────────────────
The capital of France is Paris. Paris is a city of romance and culture,
╭──────────────────────────────────────────────────────────────────────────╮
│ │
│ Generation Complete 163.27 tok/s │
│ │
│ Performance │
│ Prefill 12.3 ms 6 tokens 487.80 tok/s │
│ Decode 91.8 ms 15 tokens 163.40 tok/s │
│ Sampling 0.1 ms 15 tokens │
│ ────────────────────────────────────────────────────────── │
│ Total 104.2 ms 21 tokens 201.54 tok/s │
│ Load 456.7 ms │
│ │
│ Memory │
│ Weights 136.73 MiB (memory-mapped) │
│ Compute 2.25 MiB │
│ KV Cache 158.20 MiB (192 slots) │
│ ────────────────────────────────────────────────────────── │
│ Total 297.18 MiB │
│ │
│ length | 6 prompt, 15 generated │
╰──────────────────────────────────────────────────────────────────────────╯
Chat — interactive REPL
Multi-turn chat with persistent history, using the model's built-in chat template (falling back to ChatML). Prompt caching reuses KV-cache state across turns so subsequent turns skip redundant prefill.
# Basic chat
dotllm chat QuantFactory/SmolLM-135M-GGUF
# With a system prompt and sampling
dotllm chat QuantFactory/SmolLM-135M-GGUF --system "You are a helpful assistant." -t 0.8 --top-p 0.95
# GPU + KV-cache quantization for long contexts
dotllm chat bartowski/Llama-3.2-3B-Instruct-GGUF --device gpu --cache-type-k q8_0 --cache-type-v q8_0
chat takes the same --device values as run (auto default, cpu, vulkan, gpu/cuda); on Vulkan it uses the model's own device KV-cache (paged/quantized KV and the prompt prefix cache are off there, as in serve).
In-session commands: /exit or /quit to leave, /clear to reset history (keeps the system prompt), /system <text> to change the system prompt.
Sample session:
── dotllm chat | Llama 30L/576H | Q8_0 | 16 threads | greedy ─────────────
Type /exit to quit, /clear to reset history, /system <text> to set system prompt.
>>> Hello, how are you?
I'm doing well, thank you for asking! How can I help you today?
[42 prompt tokens, 18 generated tokens, 28 ms TTFT, 487.8 prefill tok/s, 163.4 decode tok/s]
>>> What is 2+2?
2 + 2 = 4.
[78 prompt tokens, 12 generated tokens, 45 ms TTFT, 312.5 prefill tok/s, 155.2 decode tok/s]
>>> /clear
History cleared.
>>> /exit
Serve — OpenAI-compatible API + chat UI
Starts a local HTTP server exposing an OpenAI-compatible API (/v1/chat/completions, /v1/completions, /v1/models, /v1/tokenize, streaming SSE, tool calling) plus a built-in single-page web chat UI. Paged KV-cache, prompt caching, and startup warm-up are on by default. The browser opens automatically unless --no-browser is set.
--device defaults to auto for serve: for each model it tries CUDA (if a GPU is present and the model fits), then Vulkan, then the CPU, and a load that fails on one device falls through to the next. Pass cpu, gpu:N/cuda:N (CUDA) or vulkan to force one (an explicit GPU that cannot load the model is an error, not a CPU server). dotllm run/chat/bench use the same auto default and the same loader. Idle models unload after --keep-alive seconds (default 300).
# Start the server with a loaded model and open the chat UI
dotllm serve QuantFactory/SmolLM-135M-GGUF
# Bind a public interface, custom port, API-only, no auto-browser
dotllm serve QuantFactory/SmolLM-135M-GGUF --host 0.0.0.0 --port 9000 --no-ui --no-browser
# Start without a model — pick one from the chat UI
dotllm serve
# Force Vulkan, and let clients name a not-yet-downloaded model ("owner/repo:Q4_K_M") to fetch it
dotllm serve --device vulkan --auto-pull
# CUDA with partial hybrid offload and more warm-up iterations
dotllm serve bartowski/Llama-3.2-3B-Instruct-GGUF --device gpu --gpu-layers 24 --warmup-iterations 5
# Speculative decoding — draft must share the target's vocabulary
dotllm serve bartowski/Llama-3.2-3B-Instruct-GGUF -q Q8_0 \
--speculative-model bartowski/Llama-3.2-1B-Instruct-GGUF --speculative-k 5
Thinking models (Qwen3.x, Bonsai, DeepSeek-R1 style) are handled for you: the <think> span is split into message.reasoning_content (OpenAI), a
thinking content block (/v1/messages) or message.thinking (/api/chat), and enable_thinking, chat_template_kwargs and reasoning_effort
(Anthropic thinking, ollama think) reach the chat template. --reasoning-format none|auto|deepseek (default auto) picks the behaviour —
none returns the raw output as before; details in docs/SERVER.md.
Any OpenAI-compatible client works against the running server:
curl -N http://localhost:8080/v1/chat/completions \
-H "Content-Type: application/json" \
-d '{
"model": "SmolLM-135M",
"messages": [{"role":"user","content":"Say hi in one word."}],
"stream": true
}'
Other API surfaces
| Route(s) | What |
|---|---|
POST /v1/messages, /v1/messages/count_tokens |
Anthropic Messages API (docs/ANTHROPIC_API.md) |
POST /v1/embeddings |
embeddings (docs/SERVER.md) |
/api/chat, /api/generate, /api/tags, /api/ps, /api/show, /api/version, /api/pull, /api/delete |
ollama-compatible API — point Open WebUI, Continue, ollama-python/js etc. at the server (OLLAMA_HOST=http://localhost:8080). /api/embed, /api/embeddings, /api/create, /api/copy, /api/push answer 501 with a pointer to the dotLLM equivalent; tools in /api/chat is 501 (use /v1/chat/completions) |
POST /v1/systemone |
Jev-compatible decision endpoint: typed questions in, per-option probabilities out, one prefill and no decode loop (built for Tev1-style decision models; tune with --decision-temperature / --decision-orderings) |
POST /v1/models/{unload,pull,enable,disable}, PUT /v1/settings, POST /v1/admin/shutdown |
management API, off by default; enabled by --allow-model-admin. dotllm ps, dotllm stop [model\|--all\|--server] and the Windows tray are clients of it. --allow-lora-admin separately gates runtime LoRA load/unload |
Reference: docs/SERVER.md. The Windows system tray app (starts/attaches to a server, autostart, update check) is documented in docs/TRAY.md.
To embed the same endpoints inside your own ASP.NET Core app, see Host the OpenAI API in your ASP.NET app below.
CLI option reference
Common options (shared by run, chat, and serve):
| Option | Short | Default | Description |
|---|---|---|---|
--device |
-d |
auto (run, chat, bench, serve); cpu (perplexity, so scores stay comparable) |
auto (CUDA if the model fits, else Vulkan, else CPU - a CPU result always prints a warning); cpu; vulkan; gpu/cuda, gpu:1/cuda:1 (CUDA). An explicit GPU that cannot be honoured is an error (what/why/model vs device memory, and --device cpu to opt in) - never a silent CPU run. Every command prints device: requested -> resolved. |
--gpu-layers |
(all if gpu, 0 if cpu) |
Transformer layers on GPU (hybrid offload) | |
--rope-scaling, --rope-freq-base, --rope-scale, --yarn-* |
(from GGUF) | RoPE / YaRN overrides (run, serve) |
|
--prefill-chunk-size |
0 | Max prompt tokens per prefill forward pass (0 = whole prompt) | |
--threads |
0 (auto) | CPU threads for inference | |
--decode-threads |
0 (auto) | Decode threads (capped at memory channels) | |
--numa-pin |
false | Pin workers to NUMA-local cores (multi-socket) | |
--pcore-only |
false | Pin workers to P-cores only (Intel hybrid) | |
--quant |
-q |
(auto) | Quant filter when a repo has multiple .gguf files (e.g., Q4_K_M) |
--cache-type-k |
f32 |
KV-cache key quant: f32, q8_0, q4_0 |
|
--cache-type-v |
f32 |
KV-cache value quant: f32, q8_0, q4_0 |
|
--kv-quant-k / --kv-quant-v |
(unset) | Synonyms of --cache-type-k/-v (f16, q8_0, q4_0); on run, bench and serve. The only way to select quantized KV for the CUDA hybrid-dense model (Bonsai-2 / qwen35), where they override DOTLLM_HYBRID_KV. See docs/KV_CACHE.md |
|
--speculative-model |
(none) | Draft model for speculative decoding (must share vocab) | |
--speculative-k |
3 | Draft tokens per speculative step | |
--mtp-head |
(auto) | serve: path of a separate MTP head GGUF (qwen4exp mtp-*.gguf); an auto-detected sibling is attached when it fits the device. See docs/SERVER.md |
|
--no-mtp |
off | serve: opt out of default-on MTP self-speculation (models with an embedded MTP head, expected concurrency < 5, no --speculative-model). --mtp is accepted for back-compat. See docs/SERVER.md |
Sampling & constraints (shared by run and chat):
| Option | Short | Default (run) |
Default (chat) |
Description |
|---|---|---|---|---|
--max-tokens |
-n |
128 | 512 | Max tokens per generation |
--temp |
-t |
0 (greedy) | 0 (greedy) | Sampling temperature |
--top-k |
0 (off) | 0 (off) | Top-K sampling | |
--top-p |
1.0 | 1.0 | Nucleus threshold | |
--min-p |
0 (off) | 0 (off) | Min-P threshold | |
--repeat-penalty |
1.0 | 1.0 | Repetition penalty | |
--repeat-last-n |
0 (full) | 0 (full) | Penalty lookback window | |
--seed |
-s (run only) |
(random) | (random) | Random seed for reproducibility |
--cache-window |
0 | 0 | Full-precision tail window for KV quant | |
--paged |
off | off | Use paged (block-based) KV-cache | |
--response-format |
text |
text |
text, json_object, json_schema, regex, grammar |
|
--schema |
— | — | JSON Schema (or @file.json) for json_schema |
|
--pattern |
— | — | Regex pattern for regex |
|
--grammar |
— | — | GBNF grammar (or @file.gbnf) for grammar |
|
--tools |
— | — | Tool definitions JSON (or @file.json) |
run-only:
--prompt/-p— input prompt (required unless--prompt-fileis given)--prompt-file— read the prompt from a file instead of--prompt. A single trailing newline is stripped. Mutually exclusive with--prompt--json— emit a single JSON result object (suppresses formatted output)
chat-only:
--system/-s— system prompt--system-file— read the system prompt from a file instead of--system. A single trailing newline is stripped. Mutually exclusive with--system--tool-choice—auto(default),none,required, or a function name--no-prompt-cache— disable KV-cache reuse across turns--prompt-cache-size— max cached sessions (default:1)--verbose/-v— debug output (finish reason, raw text, tool-call details)
serve-only:
| Option | Short | Default | Description |
|---|---|---|---|
--host |
localhost |
Bind address | |
--port |
-p |
8080 |
Port to listen on |
--no-ui |
false | Disable the built-in chat UI (API only) | |
--no-browser |
false | Don't auto-open the browser | |
--no-paged |
false | Disable paged KV-cache (paged is on by default for serve) |
|
--no-prompt-cache |
false | Disable KV-cache reuse across requests | |
--prompt-cache-size |
4 |
Max cached sessions | |
--no-warmup |
false | Disable startup warm-up passes | |
--warmup-iterations |
3 |
Warm-up iteration count | |
--auto-pull |
false | Download a missing Hugging Face / ollama model when a request names one (off by default so remote clients cannot start downloads) | |
--keep-alive |
300 |
Idle-unload after N seconds (0 = unload after each request, negative = never) | |
--max-resident-models |
1 |
Models resident at once (>1 keeps several loaded) | |
--resident-memory-budget |
0 (unlimited) |
Total byte budget across resident models | |
--allow-model-admin |
false | Enable the management API (/v1/models/*, /v1/settings, /v1/admin/shutdown, ollama /api/pull + /api/delete) |
|
--allow-lora-admin |
false | Enable POST /v1/lora/load and DELETE /v1/lora/{name} |
|
--mtp |
(default on) | Accepted for back-compat. serve auto-enables MTP self-speculation for models with an embedded MTP head when expected concurrency is < 5 and no --speculative-model is set; --no-mtp opts out |
|
--max-messages |
8192 |
Max messages per chat request on /v1/chat/completions and /v1/messages; 0 = unlimited; env DOTLLM_MAX_MESSAGES (flag wins) |
|
--expected-concurrency |
0 |
At 5+ concurrent requests Vulkan hybrid models serve through the continuous-batch scheduler | |
--decision-temperature, --decision-orderings |
0, 1 |
Calibration of POST /v1/systemone probabilities |
|
--rate-limit-rpm, --rate-limit-tpm, --rate-limit-concurrency |
0 (off) |
Per-API-key rate limits; setting any enables limiting |
bench, perplexity: dotllm bench <model> [-d cpu|vulkan|cuda] [-p 512] [-n 128] [-r 5] [--depth N] [--json] times prefill/decode (llama-bench equivalent); dotllm perplexity <model> --corpus wiki.test.raw [--context 512] [--stride N] [-d ...] scores a corpus (docs/PERPLEXITY.md). bench defaults to --device auto; perplexity defaults to --device cpu (pass -d auto|vulkan|cuda to score on a GPU).
ps, stop: clients of a running server (--url, or DOTLLM_URL; default http://localhost:8080). dotllm ps lists resident models; dotllm stop <model> / --all unloads, --server stops the whole server gracefully. Both need --allow-model-admin on the server for the write actions.
Short-flag gotcha:
-pis prompt underrunbut port underserve.-sis seed underrunbut system underchat. When in doubt, use the long form.
Development
Debug build
Building in Debug configuration (-c Debug) enables a debug command group with diagnostic tools for inspecting GGUF files and model internals. These commands are excluded from Release builds via #if DEBUG.
# Build in Debug mode
dotnet build src/DotLLM.Cli -c Debug
# Inspect GGUF file structure
dotnet run --project src/DotLLM.Cli -c Debug -- debug gguf-header model.gguf
dotnet run --project src/DotLLM.Cli -c Debug -- debug gguf-metadata model.gguf
dotnet run --project src/DotLLM.Cli -c Debug -- debug gguf-tensors model.gguf
dotnet run --project src/DotLLM.Cli -c Debug -- debug gguf-config model.gguf
# Tokenizer round-trip verification
dotnet run --project src/DotLLM.Cli -c Debug -- debug tokenize model.gguf --text "Hello world"
# Single forward pass with top-10 logit diagnostics
dotnet run --project src/DotLLM.Cli -c Debug -- debug forward-pass model.gguf --prompt "Hello"
# Inspect embedding vector for a token ID
dotnet run --project src/DotLLM.Cli -c Debug -- debug embed-lookup model.gguf --token-id 1
| Command | Description |
|---|---|
debug gguf-header |
GGUF header structure (magic, version, tensor/metadata counts) |
debug gguf-metadata |
All metadata key-value pairs |
debug gguf-tensors |
Tensor descriptors (name, shape, quantization type, offset) |
debug gguf-config |
Extracted ModelConfig (architecture, layers, dims, RoPE params) |
debug tokenize |
Encode text → token IDs → decode, verify round-trip fidelity |
debug forward-pass |
Single forward pass, top-10 predicted tokens with softmax probabilities |
debug embed-lookup |
Raw embedding vector for a given token ID |
Debug builds are significantly slower than Release (~2-10x) because JIT optimizations, inlining, and SIMD vectorization are reduced. Always use
-c Releasefor performance measurements.
Tests
Unit and integration tests:
dotnet test
Integration tests automatically download several GGUF models (~4.5 GB total) from HuggingFace to
~/.dotllm/test-cache/on first run. The firstdotnet testwill take a while; subsequent runs use the cache. To run only unit tests (no downloads):dotnet test tests/DotLLM.Tests.Unit.
Set DOTLLM_TEST_CACHE_DIR to relocate that cache — useful for pointing CI at a mounted cache volume so models aren't re-downloaded every job, or for moving it off a small system drive:
# Linux/macOS
export DOTLLM_TEST_CACHE_DIR=/mnt/ci-cache/dotllm-models
# Windows
$env:DOTLLM_TEST_CACHE_DIR = "D:\dotllm-test-cache"
The cache uses the same {owner}/{repo}/{file}.gguf layout as ~/.dotllm/models/, so you can point it at your CLI model directory to reuse models you already pulled:
$env:DOTLLM_TEST_CACHE_DIR = "$HOME\.dotllm\models"
Doing so means test models are downloaded into your
dotllm model pullinventory and will appear indotllm model list. That is intentional but opt-in — leave the variable unset to keep auto-downloaded test models separate from models you pulled explicitly. The value must be an absolute path (or a relative one, resolved against the working directory);~is not expanded.
GPU tests (tagged
Category=GPU) require an NVIDIA GPU and run full model inference — they can take 20-30 minutes. They are skipped automatically on machines without CUDA. To exclude them explicitly:dotnet test tests/DotLLM.Tests.Unit/ --filter "Category!=GPU"
Model correctness smoke tests (scripts/test_models.py) run dotLLM CLI with greedy decoding across architectures (Llama, Mistral, Phi, Qwen) and verify expected output:
# Build CLI first
dotnet build src/DotLLM.Cli -c Release
# List available test cases and which models are cached
python scripts/test_models.py --list
# Run tests for all cached models
python scripts/test_models.py
# Download missing models and run all tests
python scripts/test_models.py --download
# Run only specific architectures
python scripts/test_models.py --filter phi,qwen
Models are downloaded from HuggingFace to ~/.dotllm/models/ on first use and cached for subsequent runs.
Sample output:
Test Arch Result Time Details
=====================================================================================================
SmolLM-135M Llama PASS 2.1s Paris (163.3 tok/s)
Llama-3.2-1B-Instruct-Q4 Llama PASS 5.7s Paris (31.0 tok/s)
Qwen2.5-0.5B-Instruct Qwen PASS 3.2s Paris (78.5 tok/s)
Phi-3-mini-4k-instruct Phi PASS 12.4s Paris (14.2 tok/s)
=====================================================================================================
4/4 passed, 0 failed, 0 skipped
Benchmarks
Three scripts in scripts/ provide benchmarking at different levels:
bench_compare.py -- Single-point benchmark. Runs dotLLM (via BenchmarkDotNet) and optionally llama.cpp on one or more models, reports best-of-N throughput with CV (coefficient of variation):
# Benchmark dotLLM on SmolLM-135M (auto-downloads from HuggingFace)
python scripts/bench_compare.py --model QuantFactory/SmolLM-135M-GGUF --quant Q8_0
# Benchmark multiple models and quantizations
python scripts/bench_compare.py \
--model QuantFactory/SmolLM-135M-GGUF,bartowski/Llama-3.2-1B-Instruct-GGUF \
--quant Q4_K_M,Q8_0
# Compare dotLLM vs llama.cpp side-by-side
python scripts/bench_compare.py --model QuantFactory/SmolLM-135M-GGUF --dotllm --llamacpp
# Export results to JSON for later comparison
python scripts/bench_compare.py --model QuantFactory/SmolLM-135M-GGUF \
--export-json benchmarks/results/baseline.json --label baseline
Sample output:
=== dotLLM Benchmark Results ===
Model Prefill tok/s Decode tok/s Decode ms/tok Total tok/s CV
SmolLM-135M.Q8_0 229.2 182.7 5.47 175.3 14.7%
SmolLM-135M.Q4_K_M 165.0 230.1 4.35 198.2 20.5%
All values are best-of-N (max tok/s, min ms). CV is the coefficient of variation
across N iterations -- lower means more stable measurements.
bench_trend.py -- Interactive comparison of exported JSON results. Displays color-coded delta tables with noise-aware highlighting:
# Interactive mode: select runs and models to compare
python scripts/bench_trend.py
# Compare two specific result files
python scripts/bench_trend.py benchmarks/results/baseline.json benchmarks/results/optimized.json
# Show all results as a trend table
python scripts/bench_trend.py --all
Sample output (trend across three labeled runs):
Benchmark Trend
Label Date Model Prefill tok/s Decode tok/s CV
baseline 2026-03-11 SmolLM-135M.Q4_K_M 127.6 109.5 -
step29 2026-03-13 SmolLM-135M.Q4_K_M 142.2 127.0 -
step30 2026-03-13 SmolLM-135M.Q4_K_M 146.0 98.6 -
bench_history.py -- Benchmark across git commits. Creates worktrees for each commit, runs bench_compare in each, and displays trend tables with per-commit deltas:
# Benchmark last 5 commits on main
python scripts/bench_history.py myrun --last 5
# Benchmark from a specific commit to HEAD
python scripts/bench_history.py myrun --from f3d3bf8
# Show results from a previous run (no benchmarking)
python scripts/bench_history.py myrun --show
# Interactively select which commits to benchmark
python scripts/bench_history.py myrun --last 10 --select
Sample output:
Benchmark History -- Llama-3.2-3B-Instruct-Q8_0
Label Date Prefill tok/s %chg pf Decode tok/s %chg dc CV
test_run_0 (f3d3bf8) 2026-03-11 21.2 7.4 3.8%
test_run_1 (cdb5234) 2026-03-12 24.9 +17.5% 8.0 +8.1% 2.1%
test_run_2 (a062743) 2026-03-13 24.5 ~-1.6% 7.8 ~-2.5% 4.5%
test_run_3 (6c06fbf) 2026-03-14 24.6 ~+0.4% 7.8 ~+0.0% 3.2%
test_run_4 (d1978d2) 2026-03-15 25.4 +3.3% 7.0 -10.3% 5.1%
test_run_5 (572179d) 2026-03-16 25.5 ~+0.4% 7.8 +11.4% 4.2%
%chgcolumns show commit-to-commit deltas.~prefix means the change is within noise (CV threshold). CV requires multiple BenchmarkDotNet iterations (controlled by--runsin bench_compare).
Why best-of-N instead of median? On a non-isolated machine, run-to-run noise is typically 6--30%. The median includes runs degraded by OS scheduling jitter, thermal throttling, and background I/O. Best-of-N (maximum throughput) represents what the hardware can achieve and is more stable across sessions. CV is reported alongside so you can judge measurement quality -- if CV is high, the environment was noisy and even the best-of-N value should be taken with a grain of salt.
llama.cpp setup
To run comparison benchmarks against llama.cpp:
Get llama.cpp -- either download a prebuilt release or build from source:
git clone https://github.com/ggerganov/llama.cpp.git cd llama.cpp && cmake -B build -DCMAKE_BUILD_TYPE=Release && cmake --build build --config ReleasePoint bench_compare to the binary -- either:
- Set
LLAMACPP_BINenvironment variable to the path ofllama-completion(older builds name itllama-cli) - Or pass
--llamacpp-bin /path/to/llama-completionon each invocation
Without either, bench_compare looks for
llama-completiononPATH.- Set
Run comparison -- match the thread count:
python scripts/bench_compare.py --model QuantFactory/SmolLM-135M-GGUF --dotllm --llamacpp --threads 16
Always pass
--threadswhen comparing engines. Their defaults are different: dotLLM auto-selects all logical cores, llama.cpp auto-selects physical cores. On a 32-logical/16-physical machine that is 32 vs 16, which makes the comparison meaningless. bench_compare warns when you omit it.
Comparing .NET runtimes
--runtimes runs dotLLM on several target frameworks in one comparison, each as a separate BenchmarkDotNet job with its own result row:
python scripts/bench_compare.py --model QuantFactory/SmolLM-135M-GGUF \
--dotllm --llamacpp --threads 16 --runtimes net10.0,net11.0
Decode Throughput (tok/s) * = fastest (highest)
Model Quant llama.cpp dotLLM (net10.0) dotLLM (net11.0)
SmolLM-135M Q8_0 304.6 360.3 417.0*
Each requested framework needs an SDK able to target it -- benchmarking net11.0 requires the .NET 11 SDK installed, otherwise the job fails with NETSDK1045. global.json uses rollForward: latestMajor with allowPrerelease, so machines with only the .NET 10 SDK keep resolving to it and the default (host runtime only) is unaffected.
Do not set
DOTNET_ROLL_FORWARDto benchmark a newer runtime. BenchmarkDotNet generates its toolchain project to match the host runtime, so rolling the host forward makes it emit a project the pinned SDK cannot build.--runtimesis the supported mechanism.
llama.cpp is optional. All dotLLM benchmarks work without it. The
--llamacppflag simply adds a side-by-side comparison column.
NuGet Packages
The fork publishes one package to nuget.org:
| Package | Description |
|---|---|
dotllm |
dotnet tool — the dotllm command (model / run / chat / serve / bench / ...). Use dnx dotllm --prerelease or dotnet tool install -g dotllm --prerelease (see Getting Started) |
The library projects (DotLLM.Core, DotLLM.Cpu, DotLLM.Cuda, DotLLM.Vulkan, DotLLM.Hip, DotLLM.Models, DotLLM.Tokenizers, DotLLM.Engine, DotLLM.Server, DotLLM.HuggingFace, DotLLM.Diagnostics, DotLLM.Telemetry) are packed on every build and attached as .nupkg files to each GitHub release, but are not published to nuget.org by this fork. The DotLLM.* ids on nuget.org belong to the upstream project and hold only its frozen 0.1.0-preview.* builds, which do not contain this fork's work — do not use them. To consume the libraries, download the .nupkg files from a release into a local feed:
# after downloading the DotLLM.*.nupkg assets of one release into ./feed
dotnet nuget add source ./feed --name dotllm-local
dotnet add package DotLLM.Engine --version 0.3.0-dev.<n>
dotnet add package DotLLM.Cpu --version 0.3.0-dev.<n>
dotnet add package DotLLM.Models --version 0.3.0-dev.<n>
dotnet add package DotLLM.Tokenizers --version 0.3.0-dev.<n>
All packages in a release carry the same version. See docs/PACKAGING.md.
Host the OpenAI API in your ASP.NET app
DotLLM.Server is a library — reference it from your own ASP.NET Core host to expose dotLLM's OpenAI-compatible routes. Two patterns are supported.
Mode 1 — run a dedicated dotLLM WebApplication inside your process. Simplest; you hand off model loading and routing to dotLLM.
using DotLLM.Server;
var options = new ServerOptions
{
Model = "llama-3.2-3b.Q4_K_M.gguf",
Device = "gpu",
GpuLayers = 32,
Port = 8080,
PromptCacheEnabled = true,
UsePaged = true,
};
using var state = ServerStartup.LoadModel(options.Model, options);
var app = ServerStartup.BuildApp(state, args: [], serveUi: true);
await app.RunAsync($"http://localhost:{options.Port}");
Mode 2 — attach dotLLM routes to your own WebApplication. Use this when you want dotLLM's endpoints alongside your own routes, middleware, and services.
using DotLLM.Server;
var builder = WebApplication.CreateBuilder(args);
// Load the model up front so the state can be registered in DI
// before builder.Build() — dotLLM endpoints resolve ServerState from DI.
var modelPath = "llama-3.2-3b.Q4_K_M.gguf";
var state = ServerStartup.LoadModel(modelPath, new ServerOptions
{
Model = modelPath,
Device = "cpu",
PromptCacheEnabled = true,
ModelId = "llama-3.2-3b",
});
// Your own services...
builder.Services.AddAuthentication();
// Register dotLLM's state + JSON context in the same IServiceCollection:
builder.Services.AddDotLLM(state);
var app = builder.Build();
// Your own middleware and routes...
app.UseAuthentication();
app.MapGet("/hello", () => "hi");
// Mount dotLLM's OpenAI-compatible endpoints:
app.MapDotLLMEndpoints(serveUi: false);
app.Run();
Both modes transparently reuse the embedded chat UI assets if serveUi: true. The ServerState is IDisposable — dispose it on shutdown to release the model and KV-cache.
News
- 2026-08 — Mamba-3 CUDA host (
CudaMamba3TransformerModel) — closes the CPU/Vulkan-only gap. (#346) Four new F32 CUDA kernels (mamba3_data_rope_f32,mamba3_chunk_boundary_f32,mamba3_ssd_scan_siso_f32,mamba3_ssd_scan_mimo_f32) port the already-validated Vulkan GLSL compute shaders to CUDA C — one block per head, sequential in-kernel loop overtwith__syncthreads()(not a per-token host-loop like the GDN scan kernel), compiled-fmad=falsefor CPU bit-parity; per-token preprocessing (softplus/sigmoid/RMSNorm+bias/qk_pre_dot/scale) stays host-side C#, mirroring Vulkan's design decision.CudaMamba3StateCachemirrorsCudaGdnStateCache's allocation idiom with Mamba-3's four-buffer state (ssm_state/cum_angle/k_state/v_state);CudaModelLoader.LoadMamba3FromSafetensorsis the dedicated safetensors entry point, reusing the CPUMamba3WeightLoaderfor tensor resolution. No GGUF path on any backend — confirmed, not just unimplemented. SISO is now real-weight-verified againstib-ssm/mamba3-370M-10BTon an RTX 3060 (IbSsmMamba3CudaParityTests): prefill max_abs 1.502e-5, prefill+decode max_abs 2.074e-5, argmax matches — same standing as CPU/Vulkan. MIMO stays synthetic-fixture-only, same as CPU/Vulkan (no public MIMO checkpoint exists anywhere). F32 only — quantized/F16 CUDA paths are a follow-up. - 2026-08 — MTP (Multi-Token Prediction) self-speculative decoding — CPU, CUDA, CLI, and server. (#253) Extends the existing two-model speculative-decode infrastructure with single-model self-speculation: a small MTP head predicts several future tokens per step, verified in one batched forward against the base model, reusing the DeepSeek-V3 "NextN" GGUF tensor convention (
blk.{n}.nextn.{eh_proj,enorm,hnorm,embed_tokens}) — first seen at scale infroggeric/Qwen3.6-27B-MTP-GGUF, sharingArchitecture.Qwen3HybridDensewith PrismML's Bonsai-27B. CUDA (CudaMtpState) composes entirely from already-proven kernels (RmsNorm, RoPE, attention, SwiGLU, theGemmdispatcher) rather than new kernel work. CLI wiring:--no-mtpopt-out forrun/chat(auto-detects and enables by default); server wiring:--mtpopt-in forserve(disables the continuous-batch scheduler while active, mirroring--speculative-model's existing restriction), surfaced viamtp_activein/props. Along the way, fixed a correctness gap in GDN (Gated DeltaNet) recurrent architectures: unlike KV-cache, GDN's recurrent state has no position addressing, so a rejected draft token's effect couldn't be "un-reached" by truncation alone (#287) — newIModel.CheckpointRecurrentState()/RestoreRecurrentState()(CPU + CUDA) make rollback safe for bothSpeculativeDecoderandMtpSpeculativeDecoderon hybrid GDN/attention models. Real end-to-end validation viadotllm runagainst the realfroggeric/Qwen3.6-27B-MTP-GGUF(Q4_K_M) checkpoint. - 2026-08 — Mach-1 additive-codec load path (Phase B) —
SyzygyResearch/Mach-1-Additive-35Bloads via the existingQwen3MoeHybridforward pass, zero forward-pass changes. (#266) This checkpoint is the same Qwen3.6-35B-A3B model as Phase 10'sqwen35moeGGUF support, but compressed to ~1.7 bpw (7.53 GB) with a bespoke tail-biting-trellis + randomized-Hadamard codec that has no GGUF representation. Phase A (bit-exact codec decoder, previously shipped) is now wired into the standard loading surface:Qwen35MoeConfigExtractorparses the nestedtext_configHF config;Mach1PackedCheckpointreads thepacked/HF repo layout via the existing generic safetensors reader (no new binary parser needed — its composite|-delimited keys are just ordinary safetensors headers) and dispatches container kind per-file rather than assuming a fixed layout;Qwen3MoeHybridTransformerModel.LoadFromMach1Packeddecodes every tensor to dense fp32 and populates the exact same weight structures the GGUF loader builds. Every tier (routed experts, GDN spine, full-attention spine, LM headint5g64_packed, embeddings affine-int4-with-exceptions) and both layer kinds are independently validated against the real fixture — several bit-exact against the vendor's own golden tensor — plus the realchat_template.jinja's{% macro %}support (issue #273) verified end-to-end. Full 40-layer end-to-end generation and top-1 agreement against the base GGUF is gated on ~70-128 GB RAM for a fully dense-decoded model — a known hardware limitation the issue itself anticipated ("a validation vehicle, not a deliverable"), not a defect in the loader, which keeps all layers resident exactly like the GGUF/mmap path already does. Phase C (fused additive-domain GEMV keeping routed experts in their native ~6.2 GB packed form) remains pending on all three backends before the issue closes. - 2026-07 — PQ2_0 ternary quant type + dense Qwen3.5 hybrid architecture — PrismML Bonsai-27B runs end-to-end. New
QuantizationType.PQ2_0(#151) for PrismML's Bonsai-27B: unlike I2_S's single per-tensor scale, PQ2_0 packs a per-128-element-group FP16 scale ahead of its 2-bit ternary codes; CPU scalar + CUDA GEMV/dequant kernels ship for both. NewArchitecture.Qwen3HybridDense(qwen35, #157) — a dense-FFN sibling ofQwen3MoeHybridreusing the same GDN/full-attention token-mixing infrastructure, with a dense SwiGLU FFN in place of MoE routing. CPU and CUDA forward passes both validated end-to-end against the real 7.17 GBTernary-Bonsai-27B-Q2_0.gguf(64 layers: 48 Gated-DeltaNet + 16 full-attention, PQ2_0 throughout). Along the way, fixed a VRAM-oversubscription bug on 12 GB consumer cards: a full-table F32 embedding pre-dequant (~5 GB at Bonsai's 248320-token vocab) was silently tripping WDDM's host-RAM paging fallback, masquerading as a hang; replaced with a per-call host-side row dequant. Baselinedotllm bench --device cuda: prefill 9.86 tok/s, decode 1.71 tok/s on an RTX 3060 (unoptimized, correctness-first PQ2_0 kernels — a target for follow-up optimization). Follow-up same day: the PQ2_0 F16 decode GEMV (pq2_0_gemv.cu) was rewritten to mirror I2_S's proven v2 scheme — shared-memory x-staging (kept half-width so Bonsai's largest K=17408 fits under sm_86's 48 KB static cap) plus warp-per-row (16 rows/block) instead of one-warp-per-row-with-global-x-reads, and the launch grid is no longer capped at 256 blocks. Decode 1.71 → 7.09 tok/s (4.15×); prefill unchanged (untouched dequant-then-cuBLAS path). Validated against the CPU float reference across real Bonsai shapes plus tail-row-clamp edge cases (CudaPQ2_0GemvTest). A follow-up fused-projection round (pq2_0_gemv2_f16in, mirroringi2_s_gemv2_f16in) applied to dense-FFN gate+up, attention K+V, and GDN alpha+beta — correctness-validated (fused-vs-separate parity test) but measured no further decode gain (stayed at 7.09–7.10 tok/s across all three): compute now dominates the per-call shared-x-staging cost it targets, so GEMV is no longer decode's bottleneck. Next lever is likely kernel-launch-count reduction (the codebase already hasCudaDecodeGraph.csfor CUDA graph capture) rather than further GEMV tuning. Overnight follow-up (same PR): a lightweight category profiler found the attention Q+Gate de-interleave host loop (numHeads separate pairedcuMemcpyDtoDAsynccalls per layer) was ~12% of decode time on its own — replaced with a single gather-kernel launch (deinterleave_qgate_f32; same fix applied to GDN's Q/K/V split). Then found the GEMV kernel's v2 fix hadn't actually solved the coalescing problem it targeted — lanes still owned separate whole groups 34 bytes apart, just fewer of them; restructured so the warp cooperates on one group at a time (lane L reads code byte L — a genuinely coalesced 32-byte transaction). The same uncoalesced pattern turned out to also be present, untouched, in the prefill-path dequant kernel (dequant_pq2_0.cu) — fixing it there took prefill from 12.6 → 89.06 tok/s (7.07×). The identical defect was found and fixed in I2_S's (BitNet) sibling dequant kernel too. Final cumulative numbers from this session's original unoptimized baseline: prefill 9.86 → 89.06 tok/s (9.0×), decode 1.71 → 10.34 tok/s (6.0×) on an RTX 3060. APQ2_0_ROWS_PER_WARP=4tuning experiment was tried and measured worse (documented, reverted) — decode is no longer GEMV-bound at the margin the earlier fixes left it. Decode-gap follow-up: with GEMV no longer the bottleneck, the profiler was re-run to find what else remained. Fused the decode-time residual-copy+RmsNorm pattern (cuMemcpyDtoDAsyncimmediately followed byrmsnorm_f32, called 2×/layer × 64 layers — ~9% of decode time) into one kernel (copy_rmsnorm_f32): decode 10.34 → 10.51 tok/s (+1.6%), bit-exact validated against the separate-calls sequence. Then fused GDN's decay+sigmoid step (gdn_decay_f32+sigmoid_f32, called once each per GDN layer × 48 layers) intogdn_decay_sigmoid_f32— lives in the-fmad=falsebit-perfect kernel file since it sits in the GDN recurrence's precision-critical path: decode 10.51 → 10.59 tok/s (+0.8%), again bit-exact validated. Re-profiling afterward showed the GEMM-bound categories (dense FFN gate+up/down, GDN projections — already coalescing-optimized) now dominate decode time; the next non-GEMM candidate (GDN's causal-conv1d input staging, ~6% of decode) needs a riskier restructure of a-fmad=falsekernel for an estimated sub-1% gain and was left as documented future work rather than pursued for diminishing returns. Cumulative session decode: 1.71 → 10.59 tok/s (6.2×); prefill: 9.86 → 88.34 tok/s (9.0×). Occupancy follow-up (#159): a roofline check (RTX 3060's 360 GB/s bandwidth against Bonsai's 7.17 GB weights implies a ~50 tok/s decode ceiling; 10.59 tok/s was only ~21% of it) prompted a deeperncu(Nsight Compute) pass, which found the decode GEMV was latency-bound, not bandwidth-bound: a single-lane scatter store (fixed, decode 10.59 → 10.95 tok/s), then the real headline fix — the per-layer activation-staging loop issued up to 68 sequential single-halfloads before releasing compute; vectorizing touint4cut that ~8× (decode 10.74 → 12.72 tok/s). A second GEMV specialization split the kernel by shape: the small attention/GDN case (k≤5120) got its own tighter shared-memory buffer, lifting occupancy from 33%→100% theoretical (decode 12.72 → 14.68 tok/s); the larger FFN case (k=17408) got windowed activation staging instead, since it can't shrink the same way, reaching the same 83% occupancy ceiling (decode 14.68 → 15.79 tok/s). Two follow-on ideas (resizing the FFN kernel's launch grid to remove wave-quantization tail effects, and an explicit warp-shuffle broadcast for a redundant scalar read) were modeled, implemented, and measured — both regressed real throughput once the kernel was no longer latency-bound, and are reverted/documented inpq2_0_gemv.cu's history rather than shipped. Final cumulative numbers across both issues: decode 1.71 → 15.79 tok/s (9.2×), ~31.6% of the theoretical bandwidth roofline. Launch-count + correctness follow-up (#161): a freshncupass plus an advisor review found the largest remaining PQ2_0 GEMV kernels had shifted from latency-bound to genuinely ALU-bound (78%/68% compute throughput), and — via tracing real call sites rather than only the profiled kernel names — that ~65% of all PQ2_0-related kernel launches per decode token (~704 of ~1,088) were pureconvert_f32_to_f16/convert_f16_to_f32plumbing bridging the model's F32 activation pipeline to the F16-only GEMV kernels. New F32-native GEMV kernel variants convert inline during the existing vectorized stage/store steps, eliminating 864 of those launches per decode token entirely: decode ~15.7 → ~16.9 tok/s (+8%). A follow-up algebraic identity (Σ(code−1)·x = Σ(code·x) − Σx, letting the group-levelΣxbe computed once per warp and reused across rows instead of once per byte) removed the ternary decode loop's per-element subtract: decode ~17.0 → ~17.3 tok/s (+1.9%). Two other candidates were tried and reverted (a SwiGLU-epilogue fusion regressed 5.4% — interleaving two independent weight tensors broke a previously-sequential access pattern; CUDA graph capture was investigated and explicitly deferred — the KV-cache/attention path for this hybrid model needs a real device-position-aware redesign, not a quick port of the existing dense-F16 pattern). Along the way, found and fixed a real correctness bug (not a perf issue): the LM head and token-embedding tensors were uploaded in the raw on-disk PQ2_0 layout while every GEMV/dequant kernel assumed the split[scales][codes]layout every other weight gets — confirmed via CPU-vs-GPU logit comparison (sign-uncorrelated, ruling out a simple scale bug) and fixed by routing both tensors through the same repack path as everything else; a previously-failing real-weight smoke test (-∞logits) now passes clean. Final cumulative numbers across all three issues: decode 1.71 → 17.85 tok/s (10.4×), ~35.7% of the theoretical bandwidth roofline. - 2026-06 — Phase 9 Production Serving complete — advanced scheduling (Step 59). The continuous-batch scheduler gains the full Step 59 feature set. Batched forward across in-flight sequences via
IModel.ForwardBatch: steady-state decode (#348), multi-admit prefill (#349), and recurrent decode/prefill for Mamba-3 + Qwen3-MoE-Hybrid (#350) — the latter via a newIRecurrentSequenceStatecontainer the scheduler allocates and threads per sequence (IModel.CreateSequenceState/SupportsThreadedSequenceState), which also fixes the latent cross-sequence corruption of running >1 concurrent recurrent sequence against a shared default state;MaxRecurrentSequencesbounds per-seq state memory. Prefill/decode disaggregation (#351):Step()splits intoStepPrefill()/StepDecode()over the separate prefill queue + decode set — the seam a future disaggregated/multi-worker driver uses, while the in-process single-GPU path stays serial (the model forward is single-threaded). Fairness (#352): opt-in per-API-key start-time fair queuing in admission so a hammer-client can't starve light tenants (priority still dominates across tiers). Joins the already-shipped chunked prefill, priority admission, and recompute-on-resume preemption. Phase 9 (continuous batching, prefix sharing, advanced scheduling, rate limiting, metrics & tracing) is now Done (5/5). - 2026-06 — BitNet b1.58 support — Microsoft BitNet b1.58 2B4T runs natively via the official I2_S ternary GGUF (weights ∈ {-1, 0, +1}, 4 codes/byte + per-tensor scale). New
QuantizationType.I2_Sdequant + ternary GEMV/GEMM kernels, config-driven FFN activation with squared-ReLU (ReluSquared) gating, and attention/FFN Sub-LN wired into the parameterizedTransformerBlock(no separate model class). CPU correctness-first path (full-precision activations); a parallel/SIMD ternary kernel and GPU (W2A8) support are follow-ups (#77) - 2026-05 — Strix Halo Vulkan hardening + per-sequence recurrent-state isolation. Post-Phase-10 work landing across the Strix Halo dev host (AMD Radeon 8060S iGPU, gfx1151). Sliding-window mask off-by-one (
>vs>=) fixed across 5 GPU kernels (Vulkanattention_f32.comp+ 3 FA shaders, CUDAattention.cu+attention_f32.cu); CUDA PTX regen deferred to a CUDA host (acknowledged in commit). IQ3 family Vulkan support complete: IQ3_XXS (3.0625 bpw) + IQ3_S (3.4375 bpw) GEMV/GEMM kernels via SSBO codebooks, bit-perfect vsggml-quants.c(16 kernel parity tests), upload-path predicates + per-host matmul dispatch on all 4 Vulkan transformer hosts (dense / Mamba3 / NemotronH / Qwen3MoeHybrid), and 16 host-level IQ3 forward parity tests across all 4 hosts — audit-H3 trap-the-bug (kernel-level ✓ ≠ dispatch-level ✓) is now closed for IQ3 across the entire matrix. Vulkan Flash Attention shader (flash_attention_f32.comp) with online-softmax single-pass; Vulkan FA microbench landed in PERFORMANCE.md §6.4. ForwardBatch CPU Phase 5a/5b: lm_head GEMM fusion atseqLen = Σ N_i(commit479c23f), intra-block Q/K/V/O/gate/up/down fusion across the simple-subgroup (commit92c1345, ~2.09× speedup at 4× decode batch on SmolLM-135M Q8_0). Phase 5f Vulkan dense intra-block fusion (commit1c04887) mirrors the CPU pattern. Gemma 3 wiring (CPU + Vulkan):attn_logit_softcapping, per-layer sliding window,query_pre_attn_scalar,final_logit_softcappingplumbed through both backends with synthetic discriminative parity tests. Per-sequence recurrent-state isolation for hosts with non-attention state:IGdnStateinterface +SequenceForwardRequest.GdnStateslot lifts the multi-seq guard onVulkanQwen3MoeHybridTransformerModel.ForwardBatch;IMambaStatemirror lifts the multi-seq guard onVulkanMamba3TransformerModel.ForwardBatch. Multi-seq parity tests on both hosts compare against per-seq Forwards on fresh model instances (abs 5e-3 / rel 1e-3). LoRA Q8_0 stage-1 Path B3 disconfirmed empirically (commitdd2892f, ~840 µs vs ≤470 µs bar at canonical rank=16); residual −16% LoRA Q8_0 prefill gap is now bounded by what an upstreamAvx512Vnni.V512PR would buy (passive at https://github.com/dotnet/runtime/issues/86849). Real-GGUF GPU parity tests still pending (29.6 GiB Q6_K_XL needs a CUDA host with >12 GiB VRAM; llama.cpp Vulkan FA bench needs-DGGML_VULKAN=ONbuild). - 2026-05 — Qwen3MoeHybrid (
qwen35moe) end-to-end — Gated DeltaNet + sparse MoE. Phase 10 lands the CPU, CUDA, and Vulkan implementations of Alibaba's Qwen3.6-35B-A3B hybrid architecture (38 GDN linear-attention layers + 10 full GQA layers, each feeding a shared sparse MoE with 256 routed experts top-8 + a sigmoid-gated shared expert). CPU is bit-exact vs llama.cpp (top-1 token "Ta" matches the reference), and the per-forward expert dequant scratch (~30 GiB at Q6_K_XL scale) is eliminated by feeding the GGUF raw quant view (Q4_K / Q5_K / Q6_K / Q8_0 / Q5_0 / F32 / F16) directly into the GEMM path;--threads Nis honoured for parallel projections. CUDA ships per-layergdn_scan_step,gdn_decay,conv1d_causal,l2_normalize_heads, four fused elementwise kernels (sigmoid,silu,sigmoid_mul), Q6_K → F32 dequant, and an on-device MoE dispatcher viaCudaMoeFfn(no more D2H/H2D fallback) with a model-private F32 KV cache for the 10 full-attn layers; csproj fix preserves-fmad=falseacross PTX rebuilds. Vulkan ships seven new GDN-related compute shaders including a multi-token GDN scan (one dispatch per layer instead of O(seqLen) with 6 per-token D2D copies eliminated) and an on-devicegdn_decay_f32+sigmoid_inplace_f32pair replacing the host round-trip; resident routed-expert banks are opt-in viaDOTLLM_VK_MOE_RESIDENT=1. Pre-existingCudaMoeGroupedGemvQ5_K divergence fixed along the way (Q4_K-style memory layout was being used to decode Q5_K). 1981 unit tests pass. Real-GGUF end-to-end GPU parity tests remain pending pursuant to hardware (29.6 GiB Q6_K_XL exceeds local 12 GiB CUDA card; Vulkan SDK + Strix Halo SPV build needed). - 2026-04 — Vulkan MoE Strategy C grouped F16 coopmat path — routed MoE expert rows can now stay grouped by expert on-GPU for native F16 expert banks when
VK_KHR_cooperative_matrixis available. The path computes expert counts/offsets on device, groups routed rows, runs grouped F16 coopmat W1/W3/W2 projections, ungroups rows back to the existing weighted-scatter contract, and falls back to the prior indexed F32/Q8_0 path whenever the model has active MoE LoRA, non-F16 expert banks, incompatible dimensions, or no cooperative-matrix support. Covered by standalone group/offset/grouped-matmul kernel tests and an end-to-end synthetic Vulkan-vs-CPU MoE forward parity test. - 2026-04 — Vulkan: K-quants (Q4_K_M / Q5_K_M / Q6_K_M) + native F16 / BF16 matmul kernels, plus comprehensive real-weight Vulkan parity coverage. Vulkan's matmul library now ships GEMV (decode) + GEMM (prefill) shaders for every common GGUF quantisation: F32, F16, BF16, Q8_0, Q4_K_M, Q5_K_M, Q6_K_M — plus an opt-in F16 cooperative-matrix GEMM path on devices that enumerate F16xF16→F32 tiles (gfx1151 verified). K-quant + Q8_0 weights stay in source bytes on device and dequantise inside the shader inner loop, eliminating the 2× memory penalty of the previous F32-only upload path. F16 source weights also stay native; BF16 source preserves bytes via
uintBitsToFloat(bits<<16)(no F16 extension required). The kernel half lands atafb2272/15099b9/29a1459/c9c08c5; upload-path + dispatcher routing atb1ee6bc/83e0732/39b7646/6787492coversVulkanWeights(standard transformer),VulkanNemotronHWeights,VulkanMamba3Weightsand the MoE overlay slots. Real-weight CPU↔Vulkan parity at the integration tier now covers TinyLlama-1.1B, Phi-3.5-mini, Qwen2.5-0.5B, Granite-3-MoE, Mamba-3 ib-ssm 370M, DeepSeek-V2-Lite-Q4_K_M (MLA + multi-shared MoE, 19m 27s end-to-end) —RealHfSafetensorsEndToEndVulkanTests+RealGgufVulkanParityTests. Vulkan suite 264 → 360 (+96 unit tests). DeepSeek-V2-Lite via the BF16 SafeTensors path self-skips at 0.7 s with a clear actionable message (the F32-expanded weights would need ~58 GB; use the Q4_K_M GGUF instead — Phase 1 K-quants make it native). - 2026-04 — Long-context YaRN ramped-freq verification end-to-end —
DeepSeekV2Lite_LongContext_LogitsMatchPyTorchReferencewas previously a 50 ms self-skip while its PyTorch reference JSON was missing. New--prompt-file+--last-n-rowsflags ontests/scripts/compare_logits_py_reference.pyplus alogits_offsetfield on the JSON schema make >4K-token references practical (a full 4920 × 102400 reference would be ~5 GB JSON; subset-to-last-8-rows is ~7.4 MB). The test now actually runs (36m 55s on this iGPU box's CPU forward) and validates positions 4912–4920 — past DeepSeek-V2-Lite'soriginal_max_position_embeddings=4096where YaRN's rampedinv_freqactively diverges from plain RoPE — withinDriftTolerances.Tightagainst the PyTorch BF16 oracle. Validates the YaRN fix landed at1e0a930. The C# reference comparator was extended to sliceours[offset..offset+ref_seq_len, :]for the comparison; short-prompt references with offset=0 are unaffected. - 2026-04 — DeepSeek-V2/V3 default to Phase C latent KV-cache at the loader level (HF + GGUF) — the Phase A expanded cache scales as
O(numLayers × maxSeqLen × numHeads × (qkNope + v) × 4 bytes)which for V2-Lite atmax_position_embeddings=163840is ~68 GB and OOMs on most hosts. Phase C's hybrid (latent persistence + Phase A-equivalent prefill expand + absorbed decode) keeps the same numerics but storeskv_lora_rank=512per token instead of expanded heads — ~9 GB for the same model.HfConfigExtractorandGgufModelConfigExtractorboth now setUseHybridMlaCache=trueby default forArchitecture.DeepSeekV2/Architecture.DeepSeekV3; production code paths get the win automatically, unit tests that buildMlaConfigdirectly (bypassing the loader) still default to Phase A. - 2026-04 — Vulkan: fused rmsnorm + matmul_q8_0 (sub-tile design), 282 → 295 tok/s SmolLM-135M.Q8_0 decode on AMD Radeon 8060S iGPU (RDNA3.5; +4.6%, beats v0.2.0-alpha.2 baseline of 291). Eliminates 60 dispatches + 60 pipeline barriers per decode at the attn-norm+Q and ffn-norm+Gate fusion sites. Each WG outputs ROWS_PER_WG=8 matmul rows so the rmsnorm reduction is amortised 8× across the rows it feeds (72 vs 576 WGs for SmolLM Q proj); the matmul phase reads normalised values from on-chip shared memory, and a leader-WG slice of the dispatch writes the global normOutput so K/V/Up matmuls keep working unchanged. A first per-row variant was ~7% slower (576× redundant rmsnorm overwhelmed the dispatch+barrier savings); sub-tile-4 broke even; sub-tile-8 with K_MAX=1024 (4 KB sharedNorm) tipped the balance. Default-on; opt out via
DOTLLM_VULKAN_DISABLE_FUSED_RMSNORM_MATMUL=1. Falls back to the standalone (rmsnorm + matmul) pair for prefill (seqLen>1), non-Q8_0 weights, and hidden sizes >1024. Strict-argmax CPU↔Vulkan parity holds (8/9, unchanged). 8 unit parity tests use a CPU reference and exercise K=32 with M>1, sidestepping the standalone GEMV's knownrowUints*4over-stride bug (issue #1) — the fused shader uses the correctblocksPerRow*34byte stride (#136) - 2026-04 — Phi-3.5-mini-instruct + Granite-3.0-MoE end-to-end on real weights. Two loader extensions enable loading recent production checkpoints as-is from HuggingFace. Phi-3 (
Phi3ForCausalLM) stores attention as a fusedself_attn.qkv_proj [Q+K+V, hidden]and MLP as a fusedmlp.gate_up_proj [2*intermediate, hidden]; the newSplitFusedProjectionhelper probes for these fused names, allocates one 64-byte-aligned F32 buffer per split at the correct row offsets, and BF16/F16/F32-decodes in place. Fusion-absent models unchanged. Granite-MoE (GraniteMoeForCausalLM/model_type=granitemoe) packs all experts per-layer intoblock_sparse_moe.input_linear [E, 2*I, H]+block_sparse_moe.output_linear [E, H, I]+ arouter.layergate. NewArchitecture.GraniteMoevariant +LoadGraniteMoeLayerslices each expert's w1/w2/w3 from the fused tensors into per-expert F32 pointers; the existingMoeSwiGluMlpkernel handles top-8 routing natively. Real-weight verification:microsoft/Phi-3.5-mini-instruct(7.6 GB, 32-layer MHA) produces[3, 32064]logits 96192/96192 finite (stddev 7.00);ibm-granite/granite-3.0-3b-a800m-instruct(6.3 GB, 40 experts top-8, 24Q/8KV GQA) produces[3, 49155]logits 147465/147465 finite (stddev 16.13). Both gated by env vars (DOTLLM_PHI35_CHECKPOINT_PATH,DOTLLM_GRANITE3_CHECKPOINT_PATH) with auto-detect ofC:/temp/dotllm-{phi35-mini,granite3-moe}/(#136) - 2026-04 — DeepSeek-V2/V3 end-to-end MLA integration —
TransformerModel.Forwardgains a per-layer MLA branch (RMSNorm →MlaAttention.Execute→ residual → shared MoE/dense FFN dispatch),LoadDeepSeekMlaLayersafetensors weight loader resolves all MLA tensors (q_a_proj/q_a_layernorm/q_b_projor monolithicq_projwhenq_lora_rank=0,kv_a_proj_with_mqa/kv_a_layernorm/kv_b_proj,o_proj), andModelLoader.LoadFromSafetensorsnow dispatchesArchitecture.DeepSeekV2/DeepSeekV3(removedNotSupportedException). FFN routes to dense MLP (first-K prefix) or plural-shared-experts MoE (step 58c). GQA path for Llama/Mistral/Phi/Qwen unchanged. Verified against realyujiepan/deepseek-v2-tiny-random: 4-token prefill produces 409600/409600 finite logits (stddev 0.056). 3 new unit tests (LoRA-Q prefill, monolithic-Q prefill, LoRA-Q single-token); 1285 unit tests pass. Follow-ups: latent KV-cache, YaRN mscale, absorption optimisation, real DeepSeek-V2-Lite (16 GB) validation (#136) - 2026-04 — DeepSeek multi-shared-expert MoE (
n_shared_experts > 1) — extends the MoE shared-expert branch to run N shared experts in parallel (DeepSeek-V2/V3 convention; Qwen1.5-MoE at N=1 unchanged).MoeConfig.NumSharedExperts(default 1).MoeLayerWeights.SharedGateProj/UpProj/DownProjmigrated from singleninttonint[].MoeSwiGluMlp.ExecuteWithSharedExpertaccepts per-shared-projReadOnlySpan<nint>; for k=0 the GEMM writes directly into the down-buffer (bit-identical with prior single-shared path), subsequent experts compute into pooled scratch thenTensorPrimitives.Add. Plural HF namingmlp.shared_experts.{k}.{gate,up,down}_projsupported; singular Qwen1.5mlp.shared_expert.*still loads as length-1 array. This path is now used by the DeepSeek-V2/V3 MLA forward integration as well as Qwen-MoE shared-expert variants (#136) - 2026-04 — Mamba-3 MIMO streaming decode (kernel + block) — extends the SISO streaming SSD from step 60e to MIMO (rank > 1).
Mamba3State.k_staterank-extended to[R, H, N](SISO keeps[H, N]);ssm_state,cum_angle,v_stateremain rank-free per canonicalmamba3.pyallocate_inference_cache(). NewMamba3CanonicalSsd.ExecuteMimoStreamingmirrors SISO streaming with rank-aware K; boundary term derived from the tilelang MIMO forward:ssm[h,p,n] += v_state[h,p] · (Σ_r k_state[r,h,n]) · DT[0,h] · (1-trap[0,h])— the rank-sum analog of the SISO adjustment. Empty-span and single-chunk schedules are bit-identical to the existing one-shotExecuteMimo. Synthetic MIMO drift (R=3) across all split schedules: y max_abs ≤ 4.7e-15, ssm_state max_abs ≤ 3.4e-13 — pure F32-reorder noise. 18 new unit tests (8 kernel + 9 block + 1 state shape). Model-level wiring (Mamba3TransformerModelMIMO path) awaits the weight-loader extension for[H, R, N]biases +mimo_*projections (#136) - 2026-04 — MoE GroupedGEMM refactor (bit-identical) —
MoeSwiGluMlpmoves from per-token per-expert scalar loop to gather/batched-GEMM/scatter: bucket tokens by their top-k-assigned experts (ArrayPool-rented index lists, zero sustained allocations), run one batched SwiGLU per expert across all its assigned tokens (gather →w1/w3GEMM → silu*up →w2GEMM), scatter back preserving the original per-(token, slot) accumulation order for bit-identity. Public signatures unchanged; bothExecute(Mixtral) andExecuteWithSharedExpert(Qwen-MoE) route through the same grouped core. Verified bit-identical on all 8 existing MoE unit tests + 16 integration tests including the realyujiepan/qwen3-moe-tiny-randomend-to-end forward. Follow-up: fused GroupedGEMM with SIMD-batched SwiGLU for ~2-3× further (#136) - 2026-04 — MLA attention kernel (PoC) + DeepSeek-V2/V3 detection — new
MlaAttention.Executescalar kernel implements DeepSeek's Multi-head Latent Attention end-to-end: Q LoRA factorization (q_a_proj→q_a_layernorm→q_b_proj, or monolithic Q whenq_lora_rank=0), KV compression viakv_a_proj_with_mqa→kv_a_layernorm→kv_b_proj(split intoK_nope+ per-headV), decoupled RoPE onqk_rope_head_dimsub-dim with an MQA-sharedK_rope, per-head causal SDPA, ando_proj. NewArchitecture.{DeepSeekV2, DeepSeekV3}enum variants;HfConfigExtractordetectsDeepseekV2ForCausalLM/DeepseekV3ForCausalLMand populatesMlaConfig+ DeepSeek-flavouredMoeConfig(n_routed_experts,n_shared_experts,first_k_dense_replace). Verified againstyujiepan/deepseek-v2-tiny-random(config detection) and a reference implementation within 5e-4 absolute tolerance. Later work wired this intoTransformerModel.Forward, safetensors/GGUF loading, latent KV-cache defaults, and real DeepSeek-V2-Lite verification (#136) - 2026-04 — Mamba-3 streaming decode (chunk-boundary state) — Mamba-3 prefill+decode now reproduces one-shot prefill at chunk boundaries, closing the canonical
shifted_γ[t] = DT[t+1]·(1-trap[t+1])1-token lookahead gap.Mamba3Stategains two new persistent buffers per layer —k_state [n_head, d_state](previous chunk's last-token post-RoPE, pre-scale K) andv_state [n_head, head_dim](previous chunk's last-token V) — matching canonicalmamba3_ssm.modules.mamba3commit7438488(mamba3.py:142/mamba3_siso_fwd.py:341-352).Mamba3Block.ForwardgainskState/vStatespan parameters: at chunk start it folds the deferredssm += v_state · k_state · DT[0] · (1 - trap[0])term into the SSM state BEFORE the scan — exactly the contribution a one-shot forward would have produced viascale[T_prev-1] = γ + shifted_γ; at chunk end it persists this chunk's last-token post-RoPE K (frombHRN[T-1]) and V (fromxBuf[T-1]).Mamba3TransformerModel.Forward(..., Mamba3State)threads all four buffers. User-visible win: on the realib-ssm/mamba3-370M-10BT(48 layers, 32 heads,d_state=128) the last-token logit drift max_abs drops from ≈6.5 to ≈3.3 and — critically — the top-1 argmax now matches between one-shot and 2-prefill+1-decode (previously differed, so next-token sampling was inconsistent across modes). On tiny synthetic fixtures prefill+decode reproduces one-shot to F32-reorder noise (1e-8 on 2+2 splits, ~2e-6 on three-way splits). Opens the door to true O(N) streaming-decode generation (vs today's O(N²) growing-context prefill). Tightened assertions:PrefillThenDecode_ApproximatesOneShot_WithExpectedDriftrenamed toPrefillThenDecode_BitEqualsOneShot(abs_tol 5e-5→1e-5, three schedules);DecodeMatchesPrefillOnRealCheckpointceiling 100→10 with a newAssert.Equal(argmaxPrefill, argmaxSplit). NewMamba3BlockStreamingTests(7 split schedules × {y, ssm_state, k_state, v_state}) covers the kernel-level property. MIMO streaming deferred (canonical's tilelang MIMO kernel has no input-state signature; the MIMO Block path is unreachable today). 65 Mamba-3 unit tests + 20 integration tests all green (#136) - 2026-04 — Mamba-3 pooled forward scratch — eliminates per-call managed allocations on the Mamba-3 hot path. New
Mamba3ForwardScratch : IDisposable(mirrors the existingNemotronHForwardStatepattern) owns 13 namedNativeMemory.AlignedAlloc-backed 64-byte-aligned scratch buffers (Proj,X,Z,Dt,Adt,Trap,Gamma,Scale,AnglesRaw,B,C,QkPreDot,YScan) grown power-of-two on demand viaEnsureCapacity(seqLen)and reused across every layer of everyMamba3Block.Forward/ForwardMimocall.Mamba3TransformerModelallocates one scratch at construction and disposes it — on the ib-ssm 370M checkpoint this collapses ≈480 per-step managed float[] allocations (48 layers × ≈10 buffers) to zero after the first Forward. Numerics preserved bit-exact: every pre-existing Mamba-3 test (57 unit + 20 integration, including the canonical-Python SISO/MIMO block comparators and all 9 real-weight ib-ssm integration tests with end-to-end 5-token generation) still passes.ComputeMemoryBytesnow reports the scratch footprint. Scratch handles both SISO and MIMO B/C widths in one instance (backed bymax(1, mimoRank)·nHead·dState) (#136) - 2026-04 — Qwen-MoE support (Qwen1.5/2/3-MoE + shared experts) — extends the Mixtral MoE plumbing to the HF Qwen-MoE naming convention (
mlp.gate+mlp.experts.{j}.{gate_proj,up_proj,down_proj}instead of Mixtral'sblock_sparse_moe.gate+experts.{j}.w1/w2/w3), with optional shared-expert branch (Qwen1.5-MoE-A2.7B:mlp.shared_expert.*+ optionalmlp.shared_expert_gate.weightsigmoid scalar) and thenorm_topk_prob=falseraw-softmax gating used by Qwen1.5. NewArchitecture.QwenMoeenum variant dispatchesQwen{2,3}MoeForCausalLM/model_type=qwen{2,3}_moe.MoeConfiggainsNormTopKProb,SharedExpertIntermediateSize,HasSharedExpertGate,DecoderSparseStep, andMlpOnlyLayers— the last two let Qwen3-MoE interleave dense MLP and MoE layers in the same model (decoder_sparse_step=2→ layer 0 dense, layer 1 MoE). NewMoeSwiGluMlp.ExecuteWithSharedExpertoverload runs a parallel dense SwiGLU on every token and (optionally) multiplies it bysigmoid(hidden . shared_expert_gate)before adding to the routed top-k sum. Existing MixtralExecutecall-sites are untouched — the kernel change is additive. Verified end-to-end against the realyujiepan/qwen3-moe-tiny-randomHF checkpoint (~20 MB, 2 layers × 8 experts × top-2,decoder_sparse_step=2, no shared expert): detection → load → 3-token forward → finite logits. Synthetic-fixture coverage for the shared-expert + sigmoid-gate + raw-softmax path (Qwen1.5-MoE convention). DeepSeek-V2/V3 reuse the plural shared-expert loader path through the MLA integration (#136) - 2026-04 — Mamba-3 end-to-end text generation verified on ib-ssm 370M — full pipeline confirmation:
ModelLoader.LoadFromSafetensors+ModelLoader.LoadTokenizerFromHfDirectory+ iterativeMamba3TransformerModel.Forward→ argmax → decode. Test generates 5 tokens from"The capital of France is"→[263, 1407, 4100, 760, 310]→"The capital of France is a very important part of"(coherent English from a 10BT-trained base model). Uses growing-context one-shot prefill per step (O(N²) total, every forward canonical) to sidestep theshifted_γchunk-edge lookahead drift that affects streaming decode; ~1.4 s/token on CPU, 7 s end-to-end for 5 tokens. Streaming-SSD decode (O(N) per token, needsk_state+v_statethreading plus a streaming kernel) remains a follow-up. Gated integration testIbSsmMamba3GenerationTestsruns onDOTLLM_IBSSM_CHECKPOINT_PATH/ conventional paths (#136) - 2026-04 — Multi-shard safetensors loader (
model.safetensors.index.json) — HuggingFace sharded-checkpoint support for models above the default 5 GiB shard cap (e.g. Llama-3-8B-Instruct, Mistral-7B-Instruct). NewSafetensorsIndexparses theweight_map+metadata.total_sizesidecar;MultiShardSafetensorsFileopens every shard, resolves tensor-name → owning shard, and exposes the same lookup surface as the single-file reader via newISafetensorsTensorSourceinterface. Consumers (TransformerWeightsSafetensorsLoader,Mamba3WeightLoader,TransformerModel.LoadFromSafetensors,Mamba3TransformerModel.LoadFromSafetensors) accept the interface; existing single-file callers unchanged.ModelLoader.LoadFromSafetensorsauto-detects: directory path → probe formodel.safetensors.index.jsonfirst (multi-shard) else single.safetensors; index.json path → multi-shard; single-file path → single-shard, but prefers sibling index.json when present. Duplicate tensor-name handling favours the index's authoritative mapping and throws on conflicts in non-index shards. Verified end-to-end via reshardedhf-internal-testing/tiny-random-LlamaForCausalLM(2 shards, matching index.json, full forward pass:[3, 32000]finite logits, stddev=0.080) (#136) - 2026-04 — HF
tokenizer.jsoningest for Mamba-3 / Llama-family — JSON-driven BPE adapter bridging HuggingFacetokenizer.jsoninto the existingBpeTokenizer/SentencePieceEncodinginfrastructure.HfTokenizerJsonParserdecodes vocab, merges (both"a b"and modern["a","b"]forms), added tokens, Metaspace pretokenizer, ByteFallback decoder.HfBpeTokenizerFactorymaps merge rank → synthetic score (score = -rank) so the score-driven merge queue reproduces HF's earliest-rank-wins order. Special added tokens route through the control-token pre-split path; BOS/EOS auto-detected.ModelLoader.LoadTokenizerFromHfDirectorysurfaces it next toLoadFromSafetensors. Verified end-to-end on realib-ssm/mamba3-370M-10BT:"Hello world"→[15043, 3186](matches Llama-2 canonical IDs), non-ASCII round-trips via byte fallback,"The quick brown fox"→ forward through 48-layer 370M SSM → argmax decodes (#136) - 2026-04 — Mamba-3 persistent state buffers —
Mamba3State : IDisposableowns per-layerssm_state+cum_angle(NativeMemory.AlignedAlloc, 64-byte) and threads them acrossMamba3TransformerModel.Forward(..., Mamba3State)calls for stateful multi-call sequence advance. Building block for streaming decode (full streaming kernel —k_state+v_state+ chunk-boundaryshifted_γcontinuation — is the next stage). State threading is deterministic bit-for-bit on re-run; prefill+decode tracks one-shot up to the canonical chunk-edge lookahead drift, which is documented and pinned by a regression-ceiling assertion. Verified on the realib-ssm/mamba3-370M-10BTcheckpoint (gated byDOTLLM_IBSSM_CHECKPOINT_PATH) (#136) - 2026-04 — Mixtral-family MoE support — dense-routing top-k Mixture-of-Experts for Mixtral-convention models (Mixtral, Qwen*-MoE without shared experts, Phi-3.5-MoE). New
MoeConfigonModelConfig(NumExperts,NumExpertsPerTok,MoeIntermediateSize),Architecture.Mixtralenum variant,HfConfigExtractordetectsnum_local_experts/num_experts+num_experts_per_tokand surfaces Phi-3.5'smoe_intermediate_sizeoverride.MoeSwiGluMlpkernel: full softmax over experts → top-k partial max-scan (stable tiebreak: lower index wins, matchingtorch.topk) → renormalise by sum (Mixtral convention, NOT a second softmax) → per-expert SwiGLU MLP via existingFusedOps.SwiGLU→ weighted sum.TransformerModel.Forwardbranches onTransformerLayerWeights.Moe; safetensors loader resolvesblock_sparse_moe.gate+experts.{j}.w1/w2/w3, F16/BF16 → F32 upcast at load time. Verified against realyujiepan/mixtral-tiny-random(config + detection) and a synthetic 2-layer, 4-expert, top-2 fixture (full forward pass). Out of scope: shared experts (DeepSeek-V3), Qwen-MoEmlp.expertsnaming adapter, fused GroupedGEMM, expert parallelism, real Mixtral-8x7B validation (#136) - 2026-04 — Mamba-3 architecture (prefill) — pure-SSM Mamba-3 end-to-end loading and prefill forward pass.
Mamba3ConfigExtractorparses HFconfig.json(model_type=mamba3),Mamba3WeightLoaderingests F32 safetensors, andMamba3TransformerModel : IModelcomposesembed → N × (RMSNorm + Mamba3Block.Forward + residual) → final RMSNorm → LM head. Canonicalstate-spaces/mambasemantics: 8-slice in_proj (z, x, B, C, dd_dt, dd_A, trap, angles), per-token per-headA = -softplus(dd_A)clamp, data-dependent RoPE withtanh(angle)·π·DTcumsum, attention-style SSD scan, SISO and MIMO entry points.ModelLoader.LoadFromSafetensorsnow dispatchesArchitecture.Mamba3. Prefill only — persistent decode state, tokenizer ingest, and generation loop are follow-ups. Verified end-to-end against canonical Python fixtures, a synthetic tiny checkpoint, and the real 1.55 GBib-ssm/mamba3-370M-10BTcheckpoint (48 layers × 32 heads × d_state=128 × vocab=32000, gated integration test viaDOTLLM_IBSSM_CHECKPOINT_PATH) (#136) - 2026-04 — Safetensors loader for dense transformers —
ModelLoader.LoadFromSafetensors+TransformerModel.LoadFromSafetensorsingest HuggingFacemodel.safetensors+config.jsonfor Llama/Mistral/Phi/Qwen.HfConfigExtractormirrors the GGUF extractor pattern over HF JSON fields (hidden_size,num_hidden_layers,num_key_value_heads,rope_theta,tie_word_embeddings, …). bf16 tensors are upcast into 64-byte-aligned scratch at load time; F32 tensors are zero-copy mmap views.ModelLoader.Load(path)auto-detects.ggufvs.safetensors. Verified end-to-end onhf-internal-testing/tiny-random-LlamaForCausalLM - 2026-04 — First public release (v0.1.0-preview.1) — dotLLM goes public. NuGet packages for all 10 libraries +
DotLLM.Clias a globaldotnet tool. Self-contained single-file downloads for Windows / Linux / macOS (Apple Silicon) and experimental Native AOT builds for Linux / Windows attached to every GitHub Release. Companion website at dotllm.dev (#119) - 2026-04 — Wave 7: CPU performance cleanup pass —
TopKSamplerreplaces fullArray.Sortwith a hand-rolled size-K min-heap (O(N log K), stack-resident scratch);JsonSchemaConstraintadds first-char bucketing to skip the ~160 MB of struct clones per mask build when the tracker rejects most leading characters, plus LRU eviction instead of the previous full-flush cache overflow;Dequantize.Q5_0gains an AVX2 path matching Q8_0's throughput (reusesMatMulQ5_0.ExtractQ5HighBits/vpshufbbit-extraction);BpeTokenizerpre-splits special tokens via the existingTrie.TryMatchLongestinstead of the O(n × m) linear scan;ComputeThreadPoolnow pins the caller (inference) thread to the first candidate P-core on firstDispatch, eliminating the hybrid-CPU stall where pinned P-core workers idled at the barrier waiting for an E-core caller. New BenchmarkDotNet suites for TopK sampling, schema mask build, and special-token encode (#109) - 2026-04 — Phase 7 begins: Logprobs — OpenAI-compatible
logprobs: true+top_logprobs: N(0-20) on/v1/chat/completionsand/v1/completions. Per-token log-softmax captured before sampling, returned in both streaming SSE chunks and non-streaming responses. Chat UI gains opt-in logprobs visualization: color-coded token confidence (green/lime/yellow/orange/red), hover tooltips with top-K alternatives and probabilities, diagnostic cues for low confidence, ambiguity, and sampling effect.DotLLM.Sample.Logprobsconsole sample with ANSI-colored output (#101) - 2026-04 — Phase 6 complete: Speculative decoding — draft-verify-accept loop with modified rejection sampling. A small draft model proposes K candidate tokens; the target model verifies in one batched forward pass. Greedy fast-path for temperature=0.
IKvCache.Rollback()for KV-cache truncation on rejection.IDecodingConstraint.Clone()for constraint rollback. Serve UI gains draft model selector and K slider.--speculative-modeland--speculative-kCLI options. Speculative acceptance rate in response timings (#98) - 2026-04 — Paged KV-cache — block-based KV-cache memory management (the allocation half of PagedAttention): shared block pool, block tables, ref counting, copy-on-write. Foundation for advanced prefix sharing (step 37, hard requirement) and speculative decoding (step 43, cheap rollback/fork). Attention kernels still operate on contiguous buffers via staging-buffer gather — true paged attention kernels are a future step.
--paged(opt-in for CLI), on by default forserve.--no-uiflag for API-only hosting. See docs/KV_CACHE.md (#96) - 2026-04 — Native AOT (experimental) — opt-in
dotnet publish -p:PublishAot=trueproduces a single-filedotllmbinary with ~50ms startup (vs ~500ms JIT). Source-generated JSON serialization across all projects,CreateSlimBuilderfor ASP.NET,TrimmerRoots.xmlfor Spectre.Console.Cli type preservation. JIT remains default for best throughput (Dynamic PGO). See docs/AOT.md (#94) - 2026-04 — Phase 6 begins: Warm-up — configurable dummy inference passes at server startup trigger .NET Tier-1 JIT promotion (Dynamic PGO) and exercise CUDA/cuBLAS pipelines.
/readyprobe gates on warm-up completion.--no-warmupto disable,--warmup-iterations Nto configure (#92) - 2026-04 — Phase 5 complete: simple prompt caching completes the constrained decoding & API phase — all 7 steps done (JSON mode, JSON Schema, regex/CFG, tool calling, API server, chat UI, prompt caching)
- 2026-04 — Simple prompt caching —
PrefixCachekeeps live KV-cache instances across generation calls. On each turn, element-wise prefix match finds cached tokens and skips redundant prefill, processing only new suffix tokens. LRU eviction with configurable max sessions. Enabled by default inchatandservecommands (--no-prompt-cacheto disable). Cached token count reported in CLI stats, APItimings.cached_tokens, and Chat UI stats bar. Near-100% cache hit rate in multi-turn chat, dramatically reducing TTFT on subsequent turns (#90) - 2026-04 — Built-in web chat UI —
dotllm serve model.ggufstarts the API server and opens a browser to a bundled single-page chat UI (vanilla JS + TailwindCSS, embedded as resources in the DLL). Per-message inference stats (prefill/decode tok/s, TTFT), live sampling parameter control, model hot-swap from the UI, system prompt, verbose mode. New endpoints:/props,/v1/config,/v1/models/available,/v1/models/load. Streaming SSE now includesusage+timingsin the final chunk (#86) - 2026-04 — ASP.NET OpenAI-compatible API server —
DotLLM.Serverwith/v1/chat/completions(streaming SSE + non-streaming),/v1/completions,/v1/models,/v1/tokenize,/v1/detokenize, health/ready probes. Chat template formatting, tool calling withIToolCallParserdetection,response_formatconstrained decoding. Model loading at startup via--model/--deviceCLI args. Sequential request processing via semaphore (#84) - 2026-04 — Tool calling —
IToolCallParserimplementations for Llama 3.1+, Hermes/Qwen, Mistral, and generic fallback. Auto-detection factory selects parser from model architecture and chat template content.ToolCallSchemaBuildergenerates JSON Schema from tool definitions for constrained decoding (tool_choice=required).ToolCallDetectorfor post-generation detection,StreamingToolCallAccumulatorfor streaming.--toolsand--tool-choiceCLI options with multi-turn tool use in chat REPL. Parallel tool calls supported (#82) - 2026-04 — Regex + CFG constrained decoding —
RegexConstraintcompiles patterns to minimized DFA (Thompson NFA → subset construction → Hopcroft minimization) with equivalence-class compression.GrammarConstraintparses GBNF grammars into PDA with InlineArray-based call stack. Both use zero-alloc struct simulators and dictionary-cached token masks.--response-format regex --pattern <pattern>and--response-format grammar --grammar <gbnf|@file>CLI support (#80) - 2026-03 — JSON Schema constrained decoding —
JsonSchemaConstraintlayers schema tracking onJsonCharParserto enforce type constraints, required properties, enum values, nested structures. Schema compiled into flat node array with property-name tries. Zero-allocClone()via struct-copy.--response-format json_schema --schema <json|@file>CLI support (#78) - 2026-03 — Phase 5 begins: JSON mode constrained decoding —
JsonConstraintFSM guarantees syntactically valid JSON output via per-token vocabulary masking. Stack-based PDA (RFC 8259), AVX2-vectorized logit masking, state-keyed mask cache.--response-format json_objectCLI flag (#76) - 2026-03 — KV-cache quantization: Q8_0 and Q4_0 KV-cache compression on CPU and GPU (3.7–7.1× memory reduction). Separate
--cache-type-k/--cache-type-voptions, mixed-precision window--cache-window Nkeeps recent tokens in full precision. Dual-region storage with quantize-on-evict, per-tile dequantization in tiled attention (#74) - 2026-03 — CPU/GPU hybrid layer offloading:
--gpu-layers Nto run first N layers on GPU, remainder on CPU. Automatic FP16→FP32 hidden state transfer at boundary. Split KV-cache (GPU FP16 + CPU FP32). Partial VRAM usage proportional to offloaded layers (#72) - 2026-03 — CUDA GPU backend: PTX kernels via CUDA Driver API P/Invoke (no native shared library), cuBLAS HGEMM for prefill, custom quantized GEMV for decode (Q8_0, Q4_K, Q6_K), FP16 activation pipeline, on-the-fly weight dequantization, GPU KV-cache,
--device gpuCLI flag,--device bothbenchmarking (#70) - 2026-03 — NUMA-aware threading: adaptive spin-wait dispatch (generation counter with event fallback), NUMA topology detection (Windows/Linux), P-core/E-core awareness, CPU affinity pinning, auto-reduced decode thread count (#57)
- 2026-03 — Operator fusion: fused RMSNorm+quantize (decode-only, eliminates normOut intermediate buffer) and tiled SwiGLU (1KB L1-resident sigmoid buffer) reduce DRAM roundtrips on the decode hot path (#56)
- 2026-03 — Fast approximate exp/softmax: Schraudolph IEEE-754 bit-manipulation trick replaces polynomial exp (~3 SIMD ops vs ~12) in attention softmax. AVX2/AVX-512 fused shift+exp+sum pass eliminates 3 separate TensorPrimitives calls. Sampling softmax keeps full precision (#55)
- 2026-03 — Tiled attention with online softmax: O(N) memory flash-attention-style algorithm replaces O(N²) score matrix materialization, eliminates 64 MB/head allocations at ctx 4096, uses ~1 KB stack per head (#54)
- 2026-03 — Row-interleaved weight repacking: R4 layout stores 4 consecutive rows' blocks contiguously at model load time, improving cache/TLB locality for all quantized GEMV kernels (#52)
- 2026-03 — Q8_1 input quantization: precomputed block sums for Q5_0 kernels, 2-block loop unrolling, eliminates ~4 SIMD ops/block from Q5_0 vec_dot hot path (#51)
- 2026-03 — Fused decode dispatch: Q/K/V (3→1) and Gate/Up (2→1) projection fusion saves ~72 dispatches/layer, ~4% decode throughput improvement (#50)
- 2026-03 — Phase 2 complete: additional model architectures (Mistral, Phi, Qwen), sliding window attention, fused QKV support,
IModelinterface,ModelLoaderhelper (#34) - 2026-03 — Streaming token generation:
IAsyncEnumerable<GenerationToken>API with UTF-8-safe incremental text,CancellationTokensupport, and per-token finish reason/timings (#31) - 2026-03 — Chat template engine: Jinja2-subset interpreter (lexer→parser→evaluator),
IChatTemplateimplementation,GgufChatTemplateFactory,dotllm chatREPL command (#30) - 2026-03 — Mixed quantization + Q8_K: Q8_K input quantization (float32 scale, 256-element blocks, precomputed bsums), true 4-row fused K-quant kernels, re-enabled Q4_K×Q8_K/Q5_K×Q8_K/Q6_K×Q8_K fused GEMV/GEMM (#29)
- 2026-03 — Q4_K_M dequantization and vec_dot kernels: Q4_K, Q5_K, Q6_K scalar + AVX2 dequant and fused matmul kernels with full model-level dispatch (#28)
- 2026-03 — BDN inference benchmarks: end-to-end benchmarks with custom tok/s columns, auto model download, llama.cpp comparison script (#42)
- 2026-03 — Engine inference timings:
InferenceTimingsonInferenceResponse,onTokenGeneratedcallback, CLI refactored to useTextGenerator(#41) - 2026-03 — Multi-threaded CPU inference: zero-alloc
ComputeThreadPoolwithdelegate*dispatch, parallel GEMV/GEMM and head-parallel attention (#36) - 2026-03 — SIMD kernel tuning: FMA float accumulation, 4-row batched GEMV, AVX-512 paths, SIMD quantization (#26)
- 2026-03 — Phase 1 complete: sampling pipeline + stop conditions — first coherent multi-token generation (#24)
- 2026-03 — KV-cache: eval drops from 1091 ms/token to 227 ms/token (~4.8× speedup)
- 2026-03 — Llama forward pass: first token generation from embedding to logits
- 2026-02 — BPE Tokenizer with SentencePiece and tiktoken support (#16)
Roadmap
| Phase | Description | Status |
|---|---|---|
| 1 — End-to-End Generation | GGUF loading, dequantization, CPU ops, tokenizer, attention, forward pass, KV-cache, sampling | Done (9/9) |
| 2 — Practical Local Inference | Engine metrics, benchmarks, Q4_K_M, chat templates, streaming, multi-threading, more architectures, safetensors loader | Done (12/12) |
| 3 — CPU Performance | Decode dispatch, Q8_1 input, weight repacking, outer-product GEMM, tiled attention, fast exp, fusion, NUMA | In Progress (7/8) |
| 4 — GPU Acceleration | CUDA backend, CPU/GPU hybrid, KV-cache quantization | Done (3/3) |
| 5 — Constrained Decoding & API | JSON mode, JSON Schema, regex/CFG, tool calling, OpenAI API server, chat UI, prompt caching | Done (7/7) |
| 6 — Improved Serving | Warm-up, NativeAOT, paged KV-cache, speculative decoding | Done (4/4) |
| 7 — Diagnostics & Interpretability | Logprobs, hook system, logit lens, SAE integration, LoRA adapters | In Progress (2/5) |
| 8 — Model Expansion | MLA attention, ALiBi, SmolLM3, Gemma 4, Mixture of Experts, Mamba-3 | In Progress (4/6) |
| 9 — Production Serving | Continuous batching, prefix sharing, advanced scheduling, rate limiting, metrics & tracing | Done (5/5) |
| 10 — Qwen3MoeHybrid (Gated DeltaNet + MoE) | qwen35moe architecture: alternating GDN linear-attention / full GQA + shared sparse MoE FFN. CPU bit-exact vs llama.cpp; CUDA forward path (MoE GPU dispatcher live, F32 KV cache for full-attn layers); Vulkan compute shaders (multi-token GDN scan, on-device decay/sigmoid). Real-GGUF GPU parity tests pending. |
In Progress (5/5 impl, parity pending) |
See docs/ROADMAP.md for detailed steps, dependencies, and milestones.
Documentation
- Architecture & data flow
- GGUF binary format
- Quantization formats
- Attention mechanisms
- Position encoding
- Tokenizers & chat templates
- Sampling pipeline
- Constrained decoding
- Tool calling
- KV-cache management
- GPU inference
- CUDA backend architecture
- Batch scheduling
- Server & API (OpenAI / Anthropic / ollama routes, model naming, profiles)
- Vulkan backend
- Windows system tray
- Releases, dev channel, NuGet /
dnx - Native AOT deployment
- Full roadmap
Contributing
Contributions are welcome! dotLLM uses an issue-driven workflow — every change starts with a GitHub issue describing the work. Pick an existing issue or open a new one, then submit a PR targeting main.
Contact
Questions, ideas, or feedback? Open a thread in GitHub Discussions.
Author
Built by Konrad Kokosa — .NET MVP, author of Pro .NET Memory Management (2nd ed.), and AI/agents engineer at Nethermind. Over 20 years of .NET performance work.
- Website: dotllm.dev
- Personal: kokosa.dev
- GitHub: @kkokosa
License
dotLLM is licensed under the GNU General Public License v3.0.
Acknowledgments
- llama.cpp — reference for GGUF format, quantization kernels, and CUDA implementations
- Hugging Face — model ecosystem, transformers reference implementations, tokenizer specs
- .NET team —
TensorPrimitives,System.Runtime.Intrinsics,MemoryMappedFile, and the runtime that makes this possible
| Product | Versions Compatible and additional computed target framework versions. |
|---|---|
| .NET | net10.0 is compatible. net10.0-android was computed. net10.0-browser was computed. net10.0-ios was computed. net10.0-maccatalyst was computed. net10.0-macos was computed. net10.0-tvos was computed. net10.0-windows was computed. |
This package has no dependencies.
| Version | Downloads | Last Updated |
|---|---|---|
| 0.3.0-dev.2987 | 8 | 10/16/2026 |
| 0.3.0-dev.2979 | 11 | 10/16/2026 |
| 0.3.0-dev.2974 | 39 | 10/15/2026 |
| 0.3.0-dev.2973 | 48 | 10/14/2026 |
| 0.3.0-dev.2965 | 35 | 10/14/2026 |
| 0.3.0-dev.2959 | 38 | 10/14/2026 |
| 0.3.0-dev.2948 | 36 | 10/14/2026 |
| 0.3.0-dev.2945 | 30 | 10/14/2026 |
| 0.3.0-dev.2942 | 34 | 10/14/2026 |
| 0.3.0-dev.2930 | 33 | 10/14/2026 |
| 0.3.0-dev.2925 | 32 | 10/14/2026 |
| 0.3.0-dev.2922 | 32 | 10/14/2026 |
| 0.3.0-dev.2920 | 40 | 10/14/2026 |
| 0.3.0-dev.2918 | 33 | 10/14/2026 |
| 0.3.0-dev.2910 | 35 | 10/14/2026 |
| 0.3.0-dev.2908 | 30 | 10/14/2026 |
| 0.3.0-dev.2902 | 30 | 10/14/2026 |
| 0.3.0-dev.2899 | 37 | 10/14/2026 |
| 0.3.0-dev.2897 | 34 | 10/14/2026 |
| 0.3.0-dev.2895 | 48 | 10/14/2026 |