deepseek-ai/DeepSeek-V3
github.com/deepseek-ai/DeepSeek-V3 ↗ · 18 files · 102 KB
Paper Link 👁️ Table of Contents 1. Introduction 2.
Overview
DeepSeek-V3 is the official reference implementation ('DeepSeek-Infer Demo') for the DeepSeek-V3 671B-parameter Mixture-of-Experts language model (37B activated per token). It ships a minimal PyTorch/Triton inference stack — model definition, FP8 quantization kernels, a torchrun-based chat/batch generation CLI, and checkpoint-conversion scripts — intended as a correctness reference rather than a production serving framework (the README points users to SGLang, vLLM, LMDeploy, TensorRT-LLM, and LightLLM for real deployments). Distinctive features are the Multi-head Latent Attention (MLA) architecture with an optional low-rank 'absorb' KV-cache path, auxiliary-loss-free MoE routing with grouped top-k gating, and native FP8 (e4m3, 128x128 block-scaled) weight support.
System map
The repo is a small, flat reference codebase: all executable code lives under inference/ (five Python files plus a configs/ directory of JSON model presets). The primary entrypoint is generate.py, launched via torchrun for multi-node, tensor-parallel inference; convert.py and fp8_cast_bf16.py are standalone CLI utilities for checkpoint preparation. model.py contains the entire model definition (Transformer, Block, MLA, MoE, Gate, parallel linears) and kernel.py contains the Triton kernels for FP8 quantization and GEMM. There is no build system, package manifest, or test suite — the only CI is a GitHub Actions stale-issue bot. Weights are safetensors files on disk, sharded per rank by convert.py, and runtime state is GPU-resident KV-cache buffers.
- All runnable code lives in inference/: model.py (model definition, 808 lines), kernel.py (Triton kernels), generate.py (CLI entrypoint), convert.py (HF->demo checkpoint converter), fp8_cast_bf16.py (FP8->BF16 weight converter), plus configs/*.json presets.
- generate.py is the primary entrypoint, designed to be launched by torchrun; main() reads WORLD_SIZE, RANK, and LOCAL_RANK environment variables and calls dist.init_process_group('nccl') when world_size > 1.
- model.py defines the full architecture in one file: Transformer wraps ParallelEmbedding, a ModuleList of Block (MLA attention + MLP or MoE FFN), final RMSNorm, and a ColumnParallelLinear head.
- kernel.py provides three Triton kernels — act_quant_kernel (per-128-block activation quantization to float8_e4m3fn), weight_dequant_kernel (blockwise FP8->BF16 dequant), and an autotuned fp8_gemm_kernel accumulating in fp32.
- Checkpoints are stored as per-rank safetensors shards named model{rank}-mp{world_size}.safetensors, produced by convert.py and consumed by generate.py via safetensors.torch.load_model.
- Model hyperparameters are supplied as JSON files under inference/configs/ (config_16B.json, config_236B.json, config_671B.json, config_v3.1.json) deserialized directly into the ModelArgs dataclass in main().
- Dependencies are pinned in inference/requirements.txt: torch==2.4.1, triton==3.0.0, transformers==4.46.3, safetensors==0.4.5; the README states Linux + Python 3.10 only, with Mac and Windows unsupported.
- There is no package manifest (no setup.py/pyproject.toml) and imports are bare sibling imports ('from kernel import ...', 'from model import ...'), so scripts must be run from inside the inference/ directory.
- The only CI/CD automation is .github/workflows/stale.yml, a daily actions/stale@v9 job that marks issues stale after 30 days and closes them after 14 more; there is no build, lint, or test workflow.
- Runtime state is GPU-resident: MLA registers non-persistent buffers kv_cache/pe_cache (absorb impl) or k_cache/v_cache (naive impl) sized [max_batch_size, max_seq_len, ...] at construction.
- Tensor parallelism is implemented via ColumnParallelLinear/RowParallelLinear, ParallelEmbedding, and expert-sharded MoE, synchronized with dist.all_reduce / dist.all_gather at layer boundaries.
- The tokenizer is loaded from the checkpoint directory with transformers AutoTokenizer.from_pretrained(ckpt_path), so the converted output dir must contain HF tokenizer files (convert.py copies files matching '*token*').
Feature catalog
User-facing surface is deliberately small: a torchrun-launched CLI supporting interactive multi-turn chat and batch file inference, two checkpoint-conversion utilities, and a set of compile-time implementation switches (attention impl, GEMM impl, gate score function, quantization scale format). Config presets cover 16B/236B/671B model sizes plus the V3.1 variant with ue8m0 scale format. The README documents the model itself (benchmarks, weights layout, deployment ecosystem) more than the code, which is positioned as a demo.
- Interactive chat mode (--interactive) presents a '>>> ' prompt loop with two built-in slash commands: /exit to quit and /clear to reset message history.
- Batch mode (--input-file FILE) reads one prompt per line, applies the chat template to each as a single user message, and prints Prompt/Completion pairs; prompt count is asserted <= args.max_batch_size.
- Sampling supports temperature-scaled multinomial sampling (implemented as argmax of probs/exponential(1), i.e. the Gumbel-max trick) and greedy decoding when temperature <= 0 via logits.argmax.
- Multi-turn context is maintained in a messages list and rendered each turn by tokenizer.apply_chat_template(messages, add_generation_prompt=True).
- convert.py converts a Hugging Face checkpoint directory into per-rank demo shards, with required CLI args --hf-ckpt-path, --save-path, --n-experts, --model-parallel.
- fp8_cast_bf16.py converts an FP8 HF checkpoint to BF16 by dequantizing every 1-byte weight with its *_scale_inv tensor and rewriting model.safetensors.index.json.
- Two attention implementations are selectable via the module-level attn_impl global: 'naive' (materializes per-head k/v caches) and 'absorb' (default; caches latent kv + rope pe vectors).
- Two GEMM paths are selectable via the gemm_impl global: 'bf16' dequantizes FP8 weights then calls F.linear; 'fp8' quantizes activations with act_quant and runs fp8_gemm.
- MoE gating supports grouped top-k expert routing with either softmax or sigmoid scoring (score_func), an optional learned bias for auxiliary-loss-free balancing, and route_scale weight scaling.
- YaRN-style RoPE extension is implemented in precompute_freqs_cis using original_seq_len, rope_factor, beta_fast/beta_slow correction ranges, and an mscale-adjusted softmax_scale for long context.
- Four shipped config presets parameterize the same code for different models: config_16B (DeepSeek-V2-Lite-like, q_lora_rank 0, softmax gate), config_236B, config_671B (61 layers, 256 experts, sigmoid gate, fp8), and config_v3.1 (adds scale_fmt ue8m0).
- model.py doubles as a smoke-test script: its __main__ block builds a default ModelArgs Transformer on CUDA and prints the output logits shape.
Data flow
Data moves along three pipelines: (1) offline checkpoint conversion — HF safetensors are renamed via the convert.py mapping table, column/row-sharded across mp ranks, and experts are partitioned contiguously; (2) startup — JSON config -> ModelArgs -> Transformer construction -> safetensors load_model into each rank's shard; (3) inference — token IDs flow through vocab-parallel embedding, N Blocks of RMSNorm->MLA->residual and RMSNorm->(MLP|MoE)->residual, a final norm + head producing logits that are all-gathered across ranks, then sampled and appended back into the token buffer. FP8 support threads through as blockwise scale tensors attached to Linear weights.
- Config data flows: --config JSON -> json.load -> ModelArgs(**dict); every architecture knob (dim, n_layers, n_routed_experts, dtype, scale_fmt, rope params) is a ModelArgs field with defaults matching the 16B model.
- Weight-name conversion is driven by the convert.py `mapping` dict, which renames HF keys (e.g. self_attn.q_a_proj -> attn.wq_a, mlp.gate_proj -> ffn.w1, weight_scale_inv -> scale, e_score_correction_bias -> bias) and records the shard axis (0 or 1) per tensor type.
- Expert sharding in convert.py assigns expert i to rank floor(i / n_local_experts); non-expert tensors with a shard axis are narrowed into mp contiguous slices, asserted divisible by mp.
- At startup each rank loads only its shard file model{rank}-mp{world_size}.safetensors after a 2-token warmup generate call that exercises the model before weights are loaded.
- Token flow: input IDs -> ParallelEmbedding (each rank embeds only its vocab slice, masks others to zero, then all_reduce) -> h [bsz, seqlen, dim].
- MLA forward computes q via wq or wq_a->q_norm->wq_b, splits q into q_nope/q_pe, applies RoPE to q_pe and to the k_pe slice of wkv_a output, then either writes k/v caches (naive) or caches normalized kv latent + pe (absorb).
- Attention scores are softmaxed in float32 then cast back to the input dtype; the causal mask is an additive -inf upper-triangular tensor built only when seqlen > 1.
- MoE forward flattens tokens to [tokens, dim], routes via Gate (scores -> optional group masking -> top-k indices -> gather original scores -> normalize -> route_scale), then loops over local experts scattering weighted outputs, adds shared_experts output, and all_reduces across ranks.
- Final logits are produced by ColumnParallelLinear head and reassembled across ranks with dist.all_gather + concat along the vocab dimension.
- The decode loop writes prompt tokens into a -1-padded buffer, iterates cur_pos from min(prompt_lens) to total_len calling model.forward on the slice [prev_pos:cur_pos], and forces prompt positions via prompt_mask so prefilled tokens override sampled ones.
- FP8 quantization data flow: act_quant produces (float8_e4m3fn tensor, fp32 scales) per 128-element block; fp8_gemm multiplies scaled blocks accumulating in fp32; weight scales live as a .scale Parameter on each Linear whose weight element_size() == 1.
- fp8_cast_bf16 reads model.safetensors.index.json weight_map to locate tensors across shard files, keeps a 2-file LRU cache of loaded safetensors on cuda, dequantizes each FP8 weight with weight_dequant, and strips *_scale_inv entries from the output index.
Behavior traces
Three flows were traced end-to-end. Startup: torchrun env -> process group -> config -> model construction -> warmup -> weight load. Interactive turn: input -> broadcast -> chat template -> generate loop -> decode -> history append. Per-decode-step forward: embed -> freqs_cis slice -> per-Block (attn_norm -> MLA -> residual; ffn_norm -> MLP/MoE -> residual) -> final norm -> head -> all-gather -> sample. Offline conversion runs glob -> safe_open -> rename -> shard/expert-partition -> save_file -> tokenizer copy.
- Startup trace: torchrun sets WORLD_SIZE/RANK/LOCAL_RANK -> main() reads them via os.getenv -> dist.init_process_group('nccl') if world_size>1 -> non-rank-0 rebinds print to a no-op -> torch.cuda.set_device(local_rank) -> set_default_dtype(bfloat16) -> ModelArgs loaded from --config JSON.
- Startup continues: Transformer(args) constructor sets globals world_size/rank from torch.distributed and Linear.dtype (float8_e4m3fn if args.dtype=='fp8' else bfloat16), then builds embed/layers/norm/head and precomputes freqs_cis.
- Startup finishes with a warmup: tokenizer loaded via AutoTokenizer, generate() called on the tokenized string 'DeepSeek' for 2 tokens, then safetensors load_model applies the rank's shard into the already-allocated module tree.
- Interactive turn trace: input('>>> ') on rank 0 -> dist.broadcast_object_list propagates to other ranks -> '/exit' breaks, '/clear' empties messages -> messages.append(user) -> apply_chat_template -> generate() -> tokenizer.decode(skip_special_tokens=True) -> print -> messages.append(assistant).
- Decode-step trace inside generate(): for each cur_pos, model.forward(tokens[:, prev_pos:cur_pos], prev_pos) -> Transformer.forward embeds tokens, slices freqs_cis[start_pos:start_pos+seqlen], builds causal mask if seqlen>1, runs every Block, applies final norm on the last position only (h[:, -1]), projects to logits, all_gathers across ranks.
- MLA 'absorb' trace: wq_a -> q_norm -> wq_b -> view to heads -> split q_nope/q_pe -> apply_rotary_emb(q_pe) -> wkv_a -> split kv/k_pe -> apply_rotary_emb(k_pe) -> kv_norm(kv) written to kv_cache, k_pe to pe_cache -> scores = einsum(q_nope,kv_cache)+einsum(q_pe,pe_cache) -> softmax fp32 -> einsum with kv_cache then wkv_b value slice -> wo.
- Gate trace: linear(x, gate.weight) -> softmax-or-sigmoid scores -> add bias if present -> if n_groups>1, compute group scores (amax without bias, topk(2) sum with bias), keep top topk_groups groups via masked_fill(-inf) -> topk(n_activated_experts) indices -> gather weights from pre-bias scores -> sigmoid weights normalized -> *= route_scale.
- MoE trace: gate(x) -> bincount per-expert counts -> for each local expert index with count>0, torch.where(indices==i) selects tokens and y[idx] += expert(x[idx]) * weight -> shared_experts MLP on all tokens -> dist.all_reduce(y) -> return (y+z).view(shape).
- Conversion trace: glob(hf_ckpt_path/*.safetensors) -> safe_open each file -> per tensor: skip 'model.layers.61', strip 'model.' prefix, apply name replacements, look up (new_key, shard_dim) in mapping -> for each rank i: filter experts to that rank's contiguous range or narrow the tensor along shard_dim -> save_file per rank -> copy *token* files.
- fp8->bf16 trace: load index.json weight_map -> sort *.safetensors -> per file: load to cuda, drop *_scale_inv entries, weight_dequant each 1-byte tensor with its scale_inv (warning + passthrough on KeyError) -> save_file -> prune scale_inv keys from weight_map -> write new index.json.
Testing signals
There is effectively no test infrastructure: no test directory, no test framework in requirements.txt, and no CI workflow that builds or runs code — the sole GitHub Actions job is a stale-issue bot. Verification relies on a __main__ smoke block in model.py, a hardcoded warmup generation at startup, and liberal use of runtime asserts for shape/divisibility/contiguity preconditions. The README's benchmark tables are external evaluation results, not runnable tests.
- No test files or directories exist in the repo; requirements.txt contains only runtime deps (torch, triton, transformers, safetensors) with no pytest, and the only workflow is stale.yml, which runs actions/stale and no code.
- model.py's __main__ block is the only executable self-check: it builds a default Transformer on CUDA with a random [2,128] token batch and prints logits size — it asserts nothing and requires a GPU to run.
- generate.py performs a warmup call generate(model, [tokenizer.encode('DeepSeek')], 2, -1, 1.) before weights are loaded, which smoke-tests allocation and the forward path but is not a correctness assertion (eos_id=-1 is a placeholder).
- Defensive asserts are used as the main correctness mechanism: vocab/out/in features and expert counts must divide by world_size, tensors must be contiguous, and scale/weight dims must be 2-D.
- fp8_gemm_kernel is wrapped in triton.autotune over 36 configs (BLOCK_SIZE_M x BLOCK_SIZE_N x num_stages) keyed on N,K — a performance-regression mechanism, not a correctness test.
- The .gitignore is a stock Python template (references pytest_cache, coverage, tox, Django/Flask artifacts) that does not reflect any tooling actually present — e.g., no safetensors or checkpoint patterns are ignored.
- Benchmark tables in README.md (MMLU, HumanEval, AIME, etc.) are published results from external evaluation harnesses; nothing in the repo reproduces them.
- Issue hygiene is semi-automated: stale.yml marks issues stale after 30 days and closes after 14, exempting 'pinned' and 'security' labels; PRs are exempt (days-before-pr-stale: -1).
- Untested critical surface includes the entire convert.py shard/rename pipeline and the Gate routing math — a silent mismatch there yields wrong weights or wrong routing with no failure signal. (inferred)
Risk map
The codebase is a demo, and its risks match: correctness depends on runtime asserts (stripped under -O), module-level globals that must be set before model construction, a hardcoded dim==7168 special case in Gate, and silent behaviors like dropping all layer-61 (MTP) weights during conversion. The distributed path uses pickle-based broadcast_object_list and suppresses stdout on non-zero ranks. Memory is pre-allocated aggressively (full-length KV caches). There is no input validation on paths and no test coverage anywhere.
- Gate creates its correction bias only when self.dim == 7168 — a magic-number coupling to the 671B config; any other model dim that needs e_score_correction_bias (e.g., V3.1-style variants on different dims) would silently load-mismatch.
- world_size, rank, gemm_impl, attn_impl, and block_size are module-level globals in model.py; world_size/rank are only (re)assigned inside Transformer.__init__, so constructing modules before the Transformer or changing impl flags post-hoc gives stale values.
- gemm_impl and attn_impl are not exposed as CLI flags in generate.py — changing them requires editing model.py, and nothing validates the combination against the loaded checkpoint dtype.
- main() rebinds the builtin name 'print' to a no-op lambda on non-zero ranks via 'global print' — any library code or later edits relying on print on those ranks are silently swallowed.
- convert.py unconditionally skips every tensor containing 'model.layers.61', silently discarding the MTP module weights without warning.
- Interactive mode synchronizes ranks with dist.broadcast_object_list, which pickles Python objects over the wire — acceptable for a trusted torchrun cluster but unsafe if any peer is untrusted. (inferred)
- User-supplied paths (--ckpt-path, --config, --input-file, --hf-ckpt-path, --save-path) are used directly in open()/glob()/os.makedirs with no validation; this is a CLI tool so impact is limited to the invoking user's privileges. (inferred)
- Validation is assert-based throughout (CLI arg checks, batch-size limit, divisibility, contiguity); running under python -O strips all of it, turning config errors into shape mismatches deep in the forward pass.
- KV caches are allocated at construction as zeros(max_batch_size, max_seq_len, ...) — with defaults max_seq_len=16384 that's a large always-on GPU allocation per layer regardless of actual prompt lengths.
- The MoE expert dispatch is a Python loop over experts with torch.where gather per expert — O(n_routed_experts) kernel launches per layer per step; correct but far from production performance (README itself steers users to SGLang/vLLM).
- weight_dequant inside the absorb attention path runs every forward step when wkv_b is FP8 (wkv_b.scale is not None), dequantizing the full kv_b weight per decode step — a hidden per-token cost.
- sample() divides probs in-place by an exponential_ tensor (probs.div_); logits are also divided by max(temperature,1e-5) — correct Gumbel-max sampling, but in-place mutation of freshly computed probs is fragile if reused. (inferred)
Extension points
This is a reference implementation, not a plugin architecture — there are no registries or hook interfaces. Extension happens through four narrow seams: the ModelArgs/configs JSON surface for new model sizes, the convert.py `mapping` table for new HF weight names, the module-level Literal switches (gemm_impl, attn_impl, score_func, scale_fmt) for implementation variants, and the fp8_gemm_configs autotune list for kernel tuning. New chat commands are string-matched literals in the interactive loop.
- ModelArgs is the single configuration surface: adding support for a new checkpoint size means adding a JSON file under inference/configs/ whose keys match ModelArgs field names — unknown keys would raise TypeError in ModelArgs(**json.load(f)).
- convert.py's `mapping` dict is the extension table for weight renames: to support a new HF parameter name, add an entry mapping the HF key segment to (demo_name, shard_axis_or_None); unmapped keys hit 'assert key in mapping'.
- attn_impl is a Literal['naive','absorb'] global switch; a third attention backend would slot into MLA.__init__'s cache-buffer branch and MLA.forward's score/output branches.
- gemm_impl is a Literal['bf16','fp8'] switch consulted in the linear() helper; a new GEMM path plugs in there and applies globally to every Linear subclass.
- Gate.score_func is a Literal['softmax','sigmoid'] switch with per-branch logic in forward (scoring and weight normalization); adding a scorer means touching both places.
- scale_fmt is threaded from ModelArgs -> Linear.scale_fmt -> linear() -> act_quant -> act_quant_kernel constexpr; 'ue8m0' is the only implemented format (power-of-two scales via ceil(log2)), so new scale formats extend kernel.py's branch.
- fp8_gemm_configs is an explicit list of triton.Config tile parameters; new shapes/tunings are added by editing this list, and @triton.autotune keys on N,K.
- Chat slash commands are hardcoded string comparisons in the interactive loop ('/exit', '/clear'); a new command is another elif before messages.append.
- Block selects dense-vs-MoE FFN purely by layer_id < args.n_dense_layers; architectures with interleaved or partially-MoE layouts require changing this branch, not just config.
- Experts are constructed per-rank with None placeholders in the ModuleList for non-local indices — extending to expert-parallel offload or remote experts would hook at this construction and the dispatch loop.
- There is no plugin/registry system; all behavior changes are source edits. The repo's own recommendation for anything beyond demo use is to switch frameworks (SGLang/vLLM/LMDeploy/TRT-LLM/LightLLM). (inferred)
Contributor handoff
Small codebase, high blast radius: five Python files where every edit is near the numerics. Read README.md section 6 first for the intended workflow, then generate.py (orchestration), model.py (all architecture), kernel.py (Triton numerics), and convert.py (weight naming contract). Validation is manual — there are no tests — so the smoke path is `python model.py` (needs CUDA) or a short torchrun generate. Known traps: module globals set only in Transformer.__init__, the dim==7168 gate-bias hack, silent layer-61 drops in convert.py, assert-based validation, and the requirement to run everything from inside inference/ with per-rank shard naming.
- Read order for onboarding: README.md (workflow + ecosystem) -> inference/generate.py (CLI + distributed setup) -> inference/model.py (ModelArgs + all modules) -> inference/kernel.py (FP8 kernels) -> inference/convert.py (weight-name contract) -> README_WEIGHTS.md (checkpoint layout).
- Environment setup is exactly: cd inference && pip install -r requirements.txt (torch 2.4.1, triton 3.0.0, transformers 4.46.3, safetensors 0.4.5); Linux + Python 3.10 + CUDA only — Triton kernels and torch.set_default_device('cuda') mean no CPU/macOS path.
- Canonical run commands: convert with 'python convert.py --hf-ckpt-path <hf> --save-path <out> --n-experts 256 --model-parallel 16', then generate with 'torchrun --nnodes N --nproc-per-node 8 generate.py --ckpt-path <out> --config configs/config_671B.json --interactive'.
- Checkpoint shard naming is a strict contract between convert.py and generate.py: files must be named model{rank}-mp{world_size}.safetensors in --ckpt-path; a mismatch in world_size between conversion and torchrun silently loads nothing-compatible.
- Always run scripts from inside inference/ — imports are bare ('from kernel import ...', 'from model import ...') with no package structure.
- Trap: Gate.bias exists iff args.dim == 7168. When adding a config or changing dim, also check whether the checkpoint has e_score_correction_bias tensors, or load_model will fail on unexpected keys.
- Trap: world_size/rank globals and Linear.dtype are only assigned inside Transformer.__init__; never instantiate Block/MLA/Linear standalone expecting distributed or fp8 behavior.
- Trap: convert.py drops all 'model.layers.61' tensors (MTP module). If asked to support MTP/speculative decoding, this skip plus the missing eh_proj/enorm/hnorm mapping entries are the first things to change.
- Smoke-test procedure after edits: 'python model.py' should print a logits shape (requires CUDA); for a real check, convert a small config (config_16B.json) and run a 2-token generate — there is no automated test to run instead.
- Code conventions: every public class/function carries a docstring with Args/Returns; type hints use typing.Literal for mode switches; distributed code checks 'if world_size > 1' before every collective; keep new code consistent.
- Safe edit boundaries: kernel.py Triton code and Gate routing math are numerics-critical and untested — changes there need manual numerical comparison; generate.py CLI plumbing and convert.py key renames are lower-risk. (inferred)
- Do not commit model weights or generated shards — nothing in .gitignore excludes *.safetensors or checkpoint dirs, so a careless 'git add -A' can stage hundreds of GB.
Languages
| Language | Files | Share of source |
|---|---|---|
| Python | 5 | 100% |
| JSON | 4 | docs/config |
| Markdown | 4 | docs/config |
| YAML | 1 | docs/config |
Start reading here
Probable source entrypoint: inference/ — inferred from the directory tree and manifest, not the upstream docs.
Source roots
inference/ (10 files)
Source layout — inference/
| Path | Files |
|---|---|
inference/configs/ | 4 |
inference/convert.py | 1 |
inference/fp8_cast_bf16.py | 1 |
inference/generate.py | 1 |
inference/kernel.py | 1 |
inference/model.py | 1 |
inference/requirements.txt | 1 |
Repository structure
| Top-level path | Files |
|---|---|
inference/ | 10 |
.github/ | 3 |
Docs found: README.md
CI & config entrypoints
README.md— Repository readme.github/workflows/stale.yml— GitHub Actions workflow
Structural health signals
Average score 7/10 across 5 analyzed source files; 1 hotspot surfaced.
inference/model.py— 808 lines, score 6.2 (Large file, Deep nesting)inference/kernel.py— 196 lines, score 6.9 (Deep nesting, High churn)inference/generate.py— 185 lines, score 6.9 (Deep nesting, High churn)
These are deterministic review leads — size, churn, and structural signals — not proof of a bug in this project.
Recent history
9b4e9788e4a32025-08-27 — Merge pull request #969 from youkaichao/rmsnormadecc0efbe2f2025-08-27 — fix rmsnorm and act_quant_kernel82f6008c8c6a2025-08-27 — fix act_quant_kernel (#968)b15f0dbbbe6a2025-08-27 — support scale_fmt=ue8m0 (#964)4592be48c07f2025-08-26 — fp32 gate bias