diff --git a/.gitattributes b/.gitattributes index a6344aac8c09253b3b630fb776ae94478aa0275b..69019d5e7f95da0b93350a06377b14bad3b25dea 100644 --- a/.gitattributes +++ b/.gitattributes @@ -33,3 +33,5 @@ saved_model/**/* filter=lfs diff=lfs merge=lfs -text *.zip filter=lfs diff=lfs merge=lfs -text *.zst filter=lfs diff=lfs merge=lfs -text *tfevents* filter=lfs diff=lfs merge=lfs -text +docs/evidence/libqi35metal.dylib filter=lfs diff=lfs merge=lfs -text +docs/evidence/qi35_model filter=lfs diff=lfs merge=lfs -text diff --git a/config/qi35-calibration-prompts-v1.json b/config/qi35-calibration-prompts-v1.json new file mode 100644 index 0000000000000000000000000000000000000000..ddcd65195f9302acd958fee9d7132da9be718606 --- /dev/null +++ b/config/qi35-calibration-prompts-v1.json @@ -0,0 +1,66 @@ +{ + "schema": "misaka.qi35.calibration-corpus.v1", + "tokenizer_sha256": "87a7830d63fcf43bf241c3c5242e96e62dd3fdc29224ca26fed8ea333db72de4", + "prompts": [ + { + "name": "hello_world", + "text": "Hello, world", + "ids": [9419, 11, 1814] + }, + { + "name": "paris", + "text": "The capital of France is", + "ids": [760, 6511, 314, 9338, 369] + }, + { + "name": "arithmetic", + "text": "2 + 2 =", + "ids": [17, 478, 220, 17, 283] + }, + { + "name": "freezing", + "text": "Water freezes at", + "ids": [27336, 90162, 506] + }, + { + "name": "python_comprehension", + "text": "In Python, a list comprehension", + "ids": [623, 12654, 11, 264, 1103, 59036] + }, + { + "name": "japanese_capital", + "text": "\u65e5\u672c\u306e\u9996\u90fd\u306f", + "ids": [161607, 110161, 14876] + }, + { + "name": "quick_fox", + "text": "The quick brown fox", + "ids": [760, 3841, 13477, 37550] + }, + { + "name": "gravity", + "text": "Explain gravity in one sentence:", + "ids": [814, 20139, 22525, 303, 799, 11316, 25] + }, + { + "name": "rust_ownership", + "text": "Rust ownership prevents", + "ids": [49, 575, 14834, 27028] + }, + { + "name": "logic_cat", + "text": "If all cats are mammals and Luna is a cat, then Luna is", + "ids": [2592, 660, 18833, 513, 53702, 321, 48130, 369, 264, 7993, 11, 1179, 48130, 369] + }, + { + "name": "fibonacci", + "text": "def fibonacci(n):", + "ids": [727, 73111, 1393, 1590] + }, + { + "name": "french_translation", + "text": "Translate good morning to French:", + "ids": [26583, 1603, 6353, 310, 8323, 25] + } + ] +} diff --git a/docs/adr-receipt-v3-node-ownership.md b/docs/adr-receipt-v3-node-ownership.md new file mode 100644 index 0000000000000000000000000000000000000000..c79f6b8be468b7af978489e0df7f725d2a3ef5fe --- /dev/null +++ b/docs/adr-receipt-v3-node-ownership.md @@ -0,0 +1,109 @@ +# ADR — Receipt v3 統合: 規約所有権の凍結と一括破壊変更(node ↔ runtime) + +**日付**: 2026-07-20 **状態**: Accepted(node 側設計者と合意) **適用**: runtime-palw の次期一括変更 +**背景**: 最近の設計決定(PQ-only・Hash64・単一プール・supersession)は全て **node 側**に住んでおり、 +runtime の現行レシート(final-v7 系)は「Q4_K_M 較正期の使い捨て」と割り切り済み。互換 shim なしの +**クリーンブレイク**で node 規約へ寄せる。「写した瞬間に不整合が固定される」を防ぐため、**写す前に**本 +所有権表を凍結する。 + +--- + +## 0. 所有権表(凍結) + +| 規約 | 原本 | 帰結 | +|---|---|---| +| 署名方式・鍵体系 | **node**(ML-DSA-87 のみ) | runtime の Ed25519 経路は排除 | +| ハッシュ関数・幅・keying | **node**(keyed BLAKE2b-512 / Hash64) | 合意が再計算する全値が対象 | +| 射影・match 述語 | **node**(§8 exact-match 集合) | `runtime_class_id` は非合意 telemetry | +| 署名コンテキスト台帳 | **node** の `signature_domains.rs` | Receipt v3 の active context を core const と alignment test で固定 | +| Receipt v3 hash-domain / wire / 正準直列化 | **node** の `mil/palw/src/receipt_v3.rs` | fixed-width・宣言順・LE。runtime は golden vector へ byte 一致させる | +| 実行内容(trace 語彙・各 root の生成規則) | **runtime** | node 所有の Receipt v3 field へ格納し、署名対象 wire は変更しない | + +**所有権の一本化が肝** — 署名コンテキスト registry と Receipt v3 の byte-hash domain 定数は用途が +異なるためファイルは分かれるが、どちらも node が原本であり runtime は独自 wire 規約を持たない。 +Receipt v3 context `misaka-palw-v3/receipt/mldsa87` は既に active registry 行であり、PENDING ではない。 +runtime-local の trace/root domain は実行内容の内部規則として残るが、最終値は node-owned field へ入る。 + +## 1. 衝突 1(署名)— ML-DSA-87、鍵は credential、Ed25519 は「証明可能に不在」へ + +- consensus 経路(mint に効く receipt / leaf / vote)の署名は **すべて ML-DSA-87**。 +- **署名鍵 = 登録 credential 鍵** を v1 とする。delegation(credential が署名したセッション鍵)は運用 + 要求が出てから。導入する場合もセッション鍵は 87 のまま(パラメタセットを 2 つにしない)。 +- runtime の Ed25519 は EVM の secp256k1 と同じ手筋: **feature-gate で「provably-Ed25519-free な + mint 経路バイナリ」**にする(悪い状態を表現不能に)。 +- 既存 final-v7 系レシートは**移行せず破棄**(較正期の使い捨て、既定方針どおり)。 + +## 2. 衝突 2(ハッシュ)— Hash64 keyed BLAKE2b-512 を全面採用 + +- 判定基準: 「**合意側の誰かが再計算する値か**」。該当: `output_commitment`、checkpoint/execution root、 + route root、state root、schedule commitment — 端から端まで同一関数でなければ再計算照合が成立しない。 +- Receipt v3 の projection/body/id/nullifier/output/pair は、node が定義する 6 個の byte-domain と + fixed-width・宣言順・LE preimage を使う。可変長 field の汎用 framing を Receipt v3 wire へ流用しない。 +- runtime-local root の内部 domain は別管理とし、その出力を node-owned Receipt v3 field へ格納する。 +- runtime 現行の Hash32 / SHA3 系列 domain-hash はすべて対象(receipt フィールド型は Hash32→**Hash64**)。 +- BLAKE3 は **S0 実測で BLAKE2b が受領税目標(<1%)を割ると判明した場合にのみ**、宣言行付きの第二ハッシュ + として検討。既定は一本。 + +## 3. 衝突 3(射影)— MatchProjectionV2 として作り直す + +- V1 の `runtime_class_id` 一致要件をコピーすると「**異機種が bit 一致しても永久に mint されない**」= + 単一プール化(canonical-integer トラックの成果)の否定になる。 +- 対応は原子性規則の runtime 側適用: **射影の再定義と diversity 述語の削除を同一統合で**。 +- V2 の中身 = §8 exact-match 集合: + `compute_set_id / job_challenge / output_commitment / schedule root / execution root / route root / + state root / CU / token count / stop reason` +- `runtime_class_id` は **ImplementationTelemetry**(非合意)へ移動。 +- **V1 の名前を再利用しない**(「名前だけ同じで中身をすり替えると未来の実装者が泣く」)。V2 として切る。 + +> 整合メモ: これは本リポジトリの canonical-integer 実証(実 attention/層が CPU×2 ISA + CUDA でバイト一致、 +> [4b-2b](evidence/canonical-int-attention-4b2b-2026-07-19.md) / [4b-3](evidence/canonical-int-fulllayer-4b3-2026-07-20.md)) +> と噛み合う: 異ベンダ bit 一致を mint 可能にするのは V2 射影があってこそ。 + +## 4. Receipt v3 — 一括破壊変更(破壊は 1 回) + +上 3 点は個別に入れず、予約済みの一括バンドルに吸収: +**ML-DSA envelope / Hash64 / leaf フィールド(compute_set_id・job_challenge・checkpoint root・state root・ +stop reason 等)/ challenge 束縛 / checkpoint root**。 + +**runtime 側の実施順(完了済み)**: +1. node Receipt v3 の署名 context と 6 byte-hash domain を原本として固定 +2. canonical body/projection を fixed-width・宣言順・LE、Hash64 keyed BLAKE2b-512 へ移行 +3. ML-DSA-87 化(Ed25519 を v3 経路で表現不能にする) +4. MatchProjectionV2(+ `runtime_class_id` → ImplementationTelemetry) +5. production encoder/worker/verifier を node golden と byte 一致させ、admission expectations を必須化 + +## 5. 経済側の付記(node 側で対応、runtime は観測のみ) + +- `pump_ev_negative`: esc(k) は credential 単位のため**回転 sybil カルテル**(min-bond credential 多数で + no-show を分散、全イベント k=0)が esc をほぼ回避 → コストは平坦モデルに退化し pump が黒字化し得る + (140k < 167k)。よって真の目標状態は「**平坦モデル単独で EV 負**」、esc は defense-in-depth に降格。 + P₀ 再較正(1.3–1.5×+余裕)で到達可能。薄市場は N_STAT hold が π ごと凍結するので pump 不能。 +- テスト pin は再較正完了と同時に `flat_cost < gain` → `flat_cost > gain` へ反転。 +- `NoShowPenaltyDestination` に `ToRequester` を**設けない**型表現は維持(悪い状態を表現不能に)。 + +## 帰結(runtime-palw への具体的影響リスト) + +- `receipt.rs` / `protocol_v2.rs`: ComputeReceipt **v3**(Hash64 フィールド、新 leaf フィールド)。 + Ed25519 envelope は feature `legacy-ed25519`(既定 OFF)。v1/v2 デコードは検証ツール専用に隔離。 +- 全 domain-hash site(schedule chain / trace roots / nullifiers / commitments): keyed BLAKE2b-512 へ。 + タグは node 台帳の行として登録し、runtime 内のタグ定数は台帳から生成 or 参照に一本化。 +- `matcher.rs`: `MatchProjectionV2`(§8 集合)。`runtime_class_id` は `ImplementationTelemetry` 構造体へ。 +- `mint.rs`: 署名方式ゲート(ML-DSA でないものは mint 経路で不適格)を追加。 +- receipts/final-v7 は破棄対象(アーカイブとしてのみ保持、照合対象にしない)。 +- [palw-node-integration-spec.md](palw-node-integration-spec.md) は本 ADR を上位とし、v1/Ed25519/V1 射影の + 記述は「現行(較正期)実装の記録」に降格(同文書冒頭に supersession banner)。 + +## 実装状況(2026-07-20 更新) + +- **node 原本実装済み**: `MISAKA.../mil/palw/src/receipt_v3.rs` が `ComputeReceiptV3` / projection / + output/nullifier/pair の fixed-width LE bytes と byte-domain、selected-chain expectations 付き検証を所有する。 + fixture は `mil/palw/test-data/receipt_v3_golden_v1.json`。 +- **runtime production 経路を整合済み**: `runtime-palw/src/receipt_v3.rs` と integer/testnet worker CLI は + node と同じ canonical bytes、Hash64、credential binding、execution nullifier、output commitment、 + ML-DSA-87 context `misaka-palw-v3/receipt/mldsa87`(vk 2592B / sig 4627B)を使用する。node fixture の + body/digest/pair-id golden と一致し、selected-chain admission API は network/compute-set/challenge/slot/ + epoch/registered credential の `ReceiptV3Expectations` を必須とする。self-consistency と low-level matcher は + admission verdict ではない。 +- **残り(ネットワーク統合)**: submission transport/scheduler、node genesis network identity と job challenge、 + provider credential registry 接続、global nullifier/lifecycle、receipt DA と監査・決済。暗号 backend の + 「次段配線」ではなく、これら外部ライフサイクルが公開・有価値 network の blocker である。 diff --git a/docs/canonical-integer-compute-design.md b/docs/canonical-integer-compute-design.md new file mode 100644 index 0000000000000000000000000000000000000000..411c81ad0903773dbaabcf7365d8db0e95c44541 --- /dev/null +++ b/docs/canonical-integer-compute-design.md @@ -0,0 +1,188 @@ +# Canonical Integer Compute — 決定性クラス統合の設計(RTX20/40 を 1 クラスに) + +**日付**: 2026-07-19 **目的**: 異なる GPU 世代(RTX 20 `sm_75` / RTX 40 `sm_89`)・異なるベンダを、**1 つの +determinism / pool クラス**に統合する。**手段**: 浮動小数点をやめ、**整数(固定小数点)で計算**し、`int32` +累積の結合律を使って hardware 非依存のビット同一性を「祈り」ではなく **設計上の性質**として保証する。 + +> 本設計は [spec-snapshot §6.1(QW36-int)](spec-snapshot-qwen36-2026-07-19.md) の「整数意味論」章を、 +> クラス統合という目的から具体化したもの。**現状は未実装**(設計 + 基礎 primitive の実機実証まで)。 + +--- + +## 0. 結論(先に) + +- fp のままでは RTX20/40 を同一クラスにできる保証はない。実測でも candle 0.9.2→0.11 の bump だけで + **Ada の token commitment が flip**(`f151cd16…`→`60a662e2…`)した — 低マージンの決定手が存在し、kernel + 差で argmax が動く([COMMITMENTS 参照](../README.md))。Metal は偶然保持しただけ。 +- **整数化すれば、順序・kernel・世代・ベンダに依らずビット同一**になる。`int32` 加算は 2³² を法とする + 可換環で、wrapping 込みで結合律・交換律が厳密に成立するため、reduction 順序を固定する必要すらない。 +- 本セッションで基礎 primitive を実装し、**Apple aarch64 と NVIDIA host x86_64 で canonical 整数出力が + バイト同一**(commitment `743652581b0b9725`)、fp は順序で割れることを実機確認した(§4)。 +- 追い風: **int8 tensor core は Turing(`sm_75`)/Ampere/Ada/Hopper に共通**(bf16 と違い Turing にもある) + ので、整数パスなら RTX20〜40+ を 1 クラスに束ねられる筋がある。 + +--- + +## 1. なぜ fp では世代を跨げないか + +- **根本原因**: IEEE-754 加算は**非結合** `(a+b)+c ≠ a+(b+c)`。GPU の matmul は tensor core の MMA、warp + reduction、タイル分割で**加算順序がカーネル/世代ごとに異なる** → 丸めが変わる → ビットが変わる。 +- **実測の裏付け(本リポジトリ)**: candle を 0.9.2→0.11 に上げただけで、**同じ Ada カード**上で CUDA の + token `vector_commitment` が変化した。これは job set 内に「top-1 マージンが極小の決定手」が存在する + 直接証拠。世代違い(Turing vs Ada、tensor core 世代が違う)ではさらに割れやすい。 +- argmax 自体は頑健(Apple M1 Max と M4 Pro は logits が違っても token 一致)。しかし**低マージン手が + 1 つでもあれば**、その手で世代差が argmax を反転させ、token 列全体のバイト一致が壊れる。 + +--- + +## 2. 核心定理 — 整数累積は hardware 非依存 + +**定理.** `int32` 上の加算は、二の補数 wrapping を含めて結合的かつ可換(2³² を法とする可換環)。 +したがって `int8 × int8 → int32` GEMM の各出力要素は、**k 方向の reduction 順序に依らず同一ビット**。 + +**系.** 同じ量子化入力を与えれば、逐次和・ペアワイズ木・SIMD レーン・GPU warp-shuffle・マルチタイル、 +どの実装でも、CPU / Apple GPU / NVIDIA(Turing/Ampere/Ada/Hopper) のどれでも**同一ビット**を出す。 +fp と違い、順序を仕様で固定する必要すらない(整数であること自体が不変性を与える)。 + +--- + +## 3. Canonical 整数計算パスの仕様(QW36-int) + +肝は「**どこにも fp を混入させない**」こと。各段を整数関数として一意に定義する。 + +### 3.1 量子化 +- **重み**: W4(or W8)対称量子化。スケールを `(multiplier: i32, shift: u32)` の固定小数点で pin。 + 丸めは round-half-to-even 固定。決定論的で 1 回だけ(GGUF に相当)。 +- **活性**: A8 対称。per-tensor(or per-token)スケールを固定小数点で。**量子化の除算も固定小数点** + (`x * inv_scale_M0 >> shift`)で行い fp を使わない。 + +### 3.2 GEMM +- `int8 × int8 → int32` 累積。tensor core 使用可(§5 caveat)。順序自由(結合律)。 + +### 3.3 再量子化 `int32 → int8`(**最重要・fp 再混入の最大リスク**) +- gemmlowp / TFLite 形式: `q = clamp((acc * M0 + (1<<(shift-1))) >> shift, -127, 127)`。 + `M0: i32`, `shift: u32` を層ごとに pin。**整数のみ・丸め点固定**。ここで fp を使うと台無しになる。 + +### 3.4 非線形 +- softmax / SiLU / sigmoid / exp / GELU は**固定小数点 or LUT**(量子化ドメイン上のテーブル)で定義し、 + 入力→出力を整数関数として一意化。 +- RMSNorm / LayerNorm は平方根を含む → 固定小数点 `rsqrt`(Newton 反復回数を固定)or LUT。 + +### 3.5 routing / argmax / sampler +- MoE router logits は整数 → **top-k は整数比較**、同点は低 index 優先で決定的。 +- 最終 argmax も整数比較 + 決定的 tie-break。greedy は完全決定的で、これが token `vector_commitment` + を backend 非依存にする。 + +### 3.6 DeltaNet(hybrid model のみ) +- 状態更新を**トークン単位で再量子化**([spec-snapshot §6.1(2)](spec-snapshot-qwen36-2026-07-19.md))。 + 丸め点を固定し、逐次デコードとチャンク並列 prefill が同一の丸め列を辿るようにする。 + +--- + +## 4. 実証(本セッション・実機) + +- prototype: `docs/evidence/canon_int.rs`(依存なし、`rustc -O canon_int.rs && ./canon_int`)。 +- `int8×int8→int32` GEMM を **逐次和と木和の 2 つの reduction 順序**で計算し比較、加えて同じソースを + 2 ISA で実行。結果([evidence md](evidence/canonical-int-cross-isa-2026-07-19.md)): + +| machine | ISA | INT 順序不変 (seq vs tree) | fp32 順序不変 | canonical commitment | +|---|---|---|---|---| +| Mac (Apple) | aarch64 | **MATCH** | DIFFER | `743652581b0b9725` | +| RTX box (NVIDIA host) | x86_64 | **MATCH** | DIFFER | `743652581b0b9725` ✅ 同一 | + +- 読み取り: **整数 → 順序にも ISA にも非依存**(同一ビット)。**fp → 順序で割れる**(`9019ee04…` vs + `e68f80da…`)= GPU の kernel 差がそのままビット差になる、という世代跨ぎ不能の正体。 + +**increment 2(GPU、完了)**: 同じ int8 GEMM を **custom GPU kernel** で実装し実 GPU 2 ベンダで実行 +([GPU evidence](evidence/canonical-int-gpu-3backend-2026-07-19.md)、`canon_metal.swift` / `canon_cuda.cu`): + +| backend | hardware | ISA / arch | INT fnv | +|---|---|---|---| +| CPU | Apple M1 Max | aarch64 | `4b138ed70d137781` | +| CPU | NVIDIA/WSL2 host | x86_64 | `4b138ed70d137781` | +| **Apple Metal GPU** | Apple M1 Max | Apple GPU | `4b138ed70d137781` | +| **NVIDIA CUDA GPU** | RTX 4060 Ti | Ada `sm_89` | `4b138ed70d137781` | + +→ **Apple GPU ≡ NVIDIA Ada ≡ x86 CPU ≡ ARM CPU** が同一ビット。canonical 整数 GEMM は vendor・microarch・ +ISA 非依存であることを実 GPU 上で実証。fp では candle bump だけで Ada の commitment が動いたのと対照的。 +残る Turing(`sm_75`)実機確認は増分 6(HW 未所持)。 + +--- + +## 5. これで RTX20/40 が 1 クラスになる理由 と caveat + +**なる理由** +- int8 tensor core は **Turing(`sm_75`)/Ampere/Ada/Hopper 共通**。bf16 が Turing に無い問題( + [docs/cuda-turing の背景](../README.md))を回避し、全世代で**同じ整数演算**を実行できる。 +- MMA int8 命令の内部累積は決定的。タイル間の `int32` 累積は結合律で順序不問 → **kernel/世代差を吸収**。 + +**caveat(正直に)** +1. 量子化・再量子化・非線形を**全 backend で同一定義**にすること。ここに fp が 1 箇所でも混ざると不変性は + 壊れる(§3.3 が最大リスク)。 +2. MMA の入力レイアウト・飽和/オーバーフロー挙動が仕様通りか、**各 backend で実測検証**が要る。 +3. 前提は「同じ量子化モデル」。量子化そのものが決定論的でなければ出発点で割れる。 +4. **Turing 実機は未所持** → 最終的な RTX20 でのビット一致確認はユーザ側 HW ステップ。 + +--- + +## 6. 実装ロードマップ(increment) + +- **increment 1(完了)**: 定理 + prototype + cross-ISA CPU 実証(§4)。 +- **increment 2(完了)**: GPU int8 canonical GEMM kernel(Metal shader + CUDA `.cu`)を実装し、CPU / Apple + Metal / NVIDIA CUDA(Ada) の 4 経路でビット一致(`4b138ed70d137781`)を実機実証(§4)。※ candle 0.11 の + 公開 Tensor API は int8×int8→int32 の GPU matmul を持たないため custom kernel を書いた。次の硬化は + tensor-core / split-K 版で GPU 並列 reduction を直接検証すること。 +- **increment 3(完了)**: 1 layer(int8 linear + int32 bias + 固定小数点 requant + 整数 LUT 非線形)を + canonical 化し、CPU aarch64 / CPU x86_64 / Apple Metal / NVIDIA CUDA(Ada) の 4 経路で LAYER fnv + `440bd1b93d991fe0` に一致([evidence](evidence/canonical-int-layer-4backend-2026-07-19.md)、 + `canon_layer.{rs,swift,cu}`)。requant は整数 mul+shift、非線形は pinned 整数 LUT lookup — **fp 再混入の + 2 大リスク(requant の fp 除算・on-device transcendental)を両方排除**。q は `[-127,127]` 全域に分布 + (非退化)。 +- **increment 4a(完了)**: 残る op 型を canonical 整数化 — softmax(固定小数点 exp、整数 Horner 生成の + pinned テーブル)/ RMSNorm(整数 rsqrt = 整数 Newton の isqrt)/ argmax — を CPU aarch64 / CPU x86_64 / + Apple Metal / NVIDIA CUDA(Ada) で COMBINED fnv `4bf581ead4d57139` に一致 + ([evidence](evidence/canonical-int-ops-4backend-2026-07-19.md)、`canon_ops.{rs,swift,cu}`)。 + **inc2(GEMM)+inc3(bias/requant/LUT)+inc4a(softmax/RMSNorm/argmax)= dense transformer forward に要る + 全 op 型が backend 非依存**。 +- **increment 4b**: 上記 primitive を dense Qwen2.5-0.5B の **実重み forward**(24 層 / RoPE / GQA / + SwiGLU / tied LM head)に組み上げ、verifier に `canonical` クラスを追加し、token `vector_commitment` が + CPU/Metal/CUDA で同一になることを実証。量子化の品質保持も含む大きな統合。 +- **increment 5**: 本番 hybrid Qwen3.6(MoE + DeltaNet)へ。[spec-snapshot §6.1](spec-snapshot-qwen36-2026-07-19.md) + の整数意味論を実装。 +- **increment 6**: 実機 Turing(RTX 2060 等)で increment 4 の値一致を確認 → **RTX20/16 系を Ada と同一 + クラスに登録**。 + +--- + +## 7. 正直な限界 + +- 本セッションは **CPU 整数の cross-ISA 実証**まで。GPU int8 kernel のビット一致は未実証(increment 2)。 +- Turing 実機未所持 → 最終確認はユーザ HW ステップ。 +- candle デモ verifier は fp で動くので、canonical class は**別実装(runtime 側)**になる。 +- 性能/精度トレードオフ: 整数パスは精度検証が要る。int8 tensor core は速いが、量子化誤差が argmax を + 変えないか(perplexity・token 一致率)を実測で確認する必要がある。 + +--- + +## 8. Metal 移植のための i128→i64 refactor 設計(2026-07-20 追記) + +MSL には 128-bit 整数が無い(`long`=i64 まで)。hifi 構成は Q60 スケールで i128 積(最大 ~2^102)を +使うため、そのままでは Metal に載らない。設計: + +1. **スケールを正規化 (m0, e) 表現に**: `scale = m0 × 2^e`、`m0 ∈ [2^30, 2^31)`(31-bit 正規化)。 + すべてのスケール演算後に renormalize(shift+指数調整)。 + - `mul(s1, s2)`: `m0 = (m0a×m0b) >> 31`(積は < 2^62 → **i64 に収まる**)、`e = ea+eb+31`、renorm。 + - `ratio = s1/s2`: `m0 = (m0a << 31) / m0b`(分子 < 2^62)、`e = ea−eb−31`、renorm。 +2. **apply(value, scale)** = `value × m0 × 2^e`: `value`(≤2^42)× `m0`(2^31)= ~2^73 — i64 超。 + → **mulhi ベースの 64×64→128 エミュレーション**: MSL の `metal::mulhi(long, long)` + `x*y`(下位)で + 128-bit 積を 2 word で構成し、`>> (−e)` を 2-word shift で実施(丸めは下位 word から桁上げ)。 + 固定パターンなので ~15 行のヘルパで済む。 +3. **残る 128/64 除算は 2 箇所だけ**: (a) A·V の `acc·2^SC / den` — `acc ≤ 2^42`、`den ≤ 2^20` 程度なので + 先に `acc / den`(i64 除算)と剰余補正で分解可能(丸め規約を参照実装と厳密一致させること)。 + (b) 残差の `(colsq << SC)/xsq` — 正規化表現では ratio 演算(上記 1)に置換され消える。 +4. **検証手順**: まず Rust で「正規化 (m0,e) + mulhi エミュ」版を書き、既存 i128 版と**同一 fingerprint** + (`906729150a5a926a` / full-model FNV `65d2e3adaf2dce9e`)を出すことを確認 → その Rust 版を MSL に + 直訳(意味論が i64+mulhi に閉じているので直訳可能)→ M1 Max で layer → full-model の順に一致確認。 + +状態: **設計確定・実装は未着手**(このセッションでは CPU×2 ISA + CUDA の 4 実装一致を優先)。 +Apple ハードは CPU(aarch64)で full-model 一致済みのため、残るのは Apple **GPU** 固有分のみ。 diff --git a/docs/canonical-integer-inc4b-fullmodel-design.md b/docs/canonical-integer-inc4b-fullmodel-design.md new file mode 100644 index 0000000000000000000000000000000000000000..5d3c468a762db5d4d806779e67ae308a37c87b0f --- /dev/null +++ b/docs/canonical-integer-inc4b-fullmodel-design.md @@ -0,0 +1,114 @@ +# inc4b 設計 — dense Qwen2.5-0.5B の canonical 整数 forward + +**日付**: 2026-07-19 **前提**: [canonical-integer-compute-design.md](canonical-integer-compute-design.md) +の increment 4b。inc2–4a で **全 op 型の backend 非依存を実機実証済み**(GEMM / bias / requant / LUT / +softmax / RMSNorm / argmax)。本書はそれらを **実重みの Qwen2.5-0.5B forward** に組み上げ、verifier で +token `vector_commitment` を CPU/Metal/CUDA 同一にするための設計。 + +## 0. 目的 + +- 実モデルの forward 全体を **fp を一切使わず整数**で通し、**backend 非依存の token commitment** を出す。 +- これが完成すると Apple/Metal・NVIDIA/CUDA・AMD/ROCm・CPU が **1 pool** で相互検証でき、「Mac の job を + CUDA で丸ごと検証」がモデル全体で成立する。 + +## 1. 対象アーキ(`/root/models/qwen0.5b-q4.gguf` GGUF 実測) + +| 項目 | 値 | +|---|---| +| architecture | `qwen2`(Qwen2.5-0.5B-Instruct) | +| block_count(層) | 24 | +| embedding_length(hidden d) | 896 | +| feed_forward_length(FFN) | 4864(SwiGLU: gate/up 896→4864, down 4864→896) | +| head_count / head_count_kv | 14 / 2(GQA, group=7), head_dim=64 | +| rope.freq_base | 1000000.0 | +| rms_norm_eps | ~1e-6 | +| file_type | 15(Q4_K_M) | +| vocab | 151936(tied: `output` = `token_embd`)※ token_embd tensor shape で確認 | +| bias | Qwen2 は q/k/v proj に **bias 有**、o_proj は bias 無 ※ tensor 名で確認 | + +## 2. 量子化方式の決定 — **W8A16**(4b-2a 実測後の改訂) + +> **改訂([4b-2a evidence](evidence/canonical-int-attention-4b2a-2026-07-19.md))**: 実測で W8A8(int8 活性)は +> attention の**活性外れ値**で誤差 ~21% と大きく、素の SmoothQuant(α=0.5)では改善せず(要 calibration 調整)。 +> **W8A16(int8 重み × int16 活性 → int64 累積)**なら外れ値を潰さず、RMSNorm を高精度化(下記)した上で +> **attention 誤差 2.6%** に収まる。int8×int16 は tensor core 非対応(通常 ALU、決定性は不変)なので、 +> tensor core を使う W8A8+調整済み SmoothQuant は後段の性能最適化に回す。**まず W8A16 で正確な canonical を作る。** +> また RMSNorm の整数実装は「正規化値を高精度固定小数点(例 Q20)で保持」すること必須(`>>` で潰すと精度崩壊。 +> これが 4b-2a で 36% 誤差の真因だった)。 + +### 旧 §2(W8A8 前提、tensor core 最優先版 — 参考) + +- **重み int8・活性 int8・累積 int32**。理由: (1) int8 tensor core は Turing〜Hopper 共通で世代統合に最適、 + (2) 0.5B が W8A8 で ~0.5GB と軽い、(3) 4bit より精度余裕があり最初の正しさ確認向き。W4A8(~0.25GB)は + サイズ最適化として後続。 +- **重み量子化はオフラインで 1 回**(bf16 原重み → per-output-channel 対称 int8、scale は + `(multiplier:i32, shift:u32)` の固定小数点で pin)。**fp はこのオフライン工程のみで、runtime には一切 + 入らない**。scale は成果物(`.canon` ファイル)に焼き込む。Q4_K_M の fp16 block scale は使わず、bf16 + 原重みから再量子化する(Q4_K_M は既に lossy かつ fp scale なので canonical には不適)。 +- **活性量子化**: per-token 動的対称 int8。scale 計算は「max(|x|) を整数で求め、`127/max` を固定小数点 + reciprocal」で行い fp を使わない(= inc3 の requant 機構の一般化)。または静的 calibration scale を pin。 +- **residual stream は int32(固定小数点)で保持**し、各 matmul の入口でだけ int8 に量子化する(層をまたぐ + ドリフトを防ぐ)。 + +## 3. 各 op の整数仕様(すべて実証済み primitive の合成) + +| stage | 整数仕様 | 依拠 primitive | +|---|---|---| +| token embedding | int8 埋め込み表を gather(tied) | 決定的 gather | +| RMSNorm | `Σx²`(int64)→ 整数 rsqrt(isqrt)→ γ(pin, 固定小数点) | **inc4a** | +| Q/K/V proj (+bias) | int8 GEMM → +int32 bias → 固定小数点 requant | **inc2+inc3** | +| RoPE | cos/sin を **位置×dim/2 の固定小数点 pinned テーブル**化(on-device で fp 三角関数を計算しない)→ 整数 mul-add 回転 | 新(LUT と同型、pin) | +| QK^T | int8 GEMM(head 毎)→ int32 scores | **inc2** | +| scale 1/√dₕ | 固定小数点定数の mul-shift | inc3 | +| causal mask | masked 位置に大負整数を加算 | 整数 | +| softmax | max 減算 → exp は pinned 固定小数点表 → int64 和 → 整数除算 | **inc4a** | +| A·V | probs(Q16) × V(int8) 累積(int64) → requant | inc2+inc3 変種 | +| o_proj + residual | int8 GEMM → int32 residual 加算 | inc2 | +| SwiGLU FFN | gate/up = int8GEMM+requant;act = SiLU LUT(gate);h = act·up → requant;down = int8GEMM+requant | **inc2+inc3** | +| final RMSNorm | 同上 | inc4a | +| LM head(tied) | int8 GEMM x × embedddingᵀ → int32 logits | inc2 | +| argmax(greedy) | 整数比較・決定的 tie-break | **inc4a** | + +**RoPE テーブル**: 32768 pos × 32 dim-pair × (cos,sin) を固定小数点(例 Q15)で pin ≈ 8MB。必要位置のみ +生成でも可だが、生成式は整数固定小数点で決定的にすること(fp `sin/cos` を on-device で使わない)。 + +## 4. 決定論の担保 + +全 stage が「inc2–4a で backend 一致を実証済みの整数 op」+「pin された固定小数点テーブル(RoPE/LUT/scale)」 ++「決定的 gather」の合成。**backend 一致な op の合成は backend 一致**(fp が 1 箇所も無い限り)。よって +full forward は **構成上** backend 非依存。残るのは「実装で fp を混入させないこと」の徹底のみ。 + +## 5. 精度検証計画(別軸: 決定論 ≠ 精度) + +- canonical 整数 forward と **fp 参照(candle)**を同一プロンプトで greedy 実行し、**token 一致率**と + **perplexity** を測る。W8A8 は多くの token で fp に一致するはずだが、低マージン token は分岐し得る + (それは fp とは別の、より低精度だが**決定的**なモデル、という理解)。 +- **重要**: canonical class の golden token commitment は fp class の値とは**別物**(新しいクラス)。価値は + 「全 backend で同一」であること。品質が不足なら per-group scale / W8A16 活性へ。 + +## 6. verifier 統合 + +- `misaka-verify` に `--canonical`(整数 forward)モードを追加。`published.json` に **`canonical` クラス** + を 1 つ追加し、その token `vector_commitment` は **backend 非依存の単一値**(Apple+NVIDIA+AMD+CPU 共通)。 +- fp の per-vendor クラス(metal/cuda)はそのまま残し、canonical を並置(移行期)。 + +## 7. 実装 sub-increments + +- **4b-1**: オフライン量子化器 — bf16 原重み → W8A8 `.canon` 成果物(pin: 重み int8 + per-channel 固定小数点 + scale + RoPE/LUT テーブル + config)。 +- **4b-2**: attention ブロック 1 個(RoPE + GQA QK^T + scale + causal + softmax + A·V + o_proj)を整数実装、 + 4-backend 一致(inc4a の延長)。 +- **4b-3**: full transformer layer 1 個(attn + SwiGLU + 2×RMSNorm + residual)を整数実装、4-backend 一致。 +- **4b-4**: 24 層 + LM head + argmax の full forward、実プロンプトで token 列を出し 4-backend 一致。 +- **4b-5**: fp 参照との token 一致率/perplexity 測定、verifier に `canonical` クラス統合・golden 記録。 + +## 8. リスクと正直な限界 + +- **精度**: W8A8 が 0.5B の出力品質を落とす可能性 → 4b-5 で定量化。不足なら per-group/W8A16 に段階調整。 +- **residual 精度**: int32 固定小数点の bit 幅設計を誤るとドリフト → 4b-3 で検証。 +- **工数**: これは実質「整数推論エンジン 1 本」。単発カーネルの数倍で、複数セッション想定。 +- **Turing 実機**: 依然未所持。4b-4 の値は定理上 Turing でも一致するが確認は HW ステップ(inc6)。 +- **本番 hybrid(Qwen3.6 MoE+DeltaNet, inc5)**は 4b の後。MoE 整数 top-k・DeltaNet per-token 再量子化は + [spec-snapshot §6.1](spec-snapshot-qwen36-2026-07-19.md) 参照。 +- 本書の scale/表は demonstrator。**本番は正確な scale・表・反復回数を pin し、モデル成果物と consensus に + 含める**必要がある。 diff --git a/docs/canonical-integer-qwen36-design.md b/docs/canonical-integer-qwen36-design.md new file mode 100644 index 0000000000000000000000000000000000000000..afcb0396db2c926bb72671fae582b49ffbb86fcf --- /dev/null +++ b/docs/canonical-integer-qwen36-design.md @@ -0,0 +1,181 @@ +# QI35 — Qwen3.6-35B-A3B canonical 整数 forward 設計(inc5 具体化) + +**日付**: 2026-07-20 **対象モデル**: `models/Qwen3.6-abliterated-35b-Claude-4.7/Qwen3.6-abliterated-35b-Claude-4.7-Q4_K_M.gguf` +(23,938,321,728 bytes — 受領/receipt トラックが commit している本番 blob そのもの) +**前提**: [canonical-integer-compute-design.md](canonical-integer-compute-design.md)(inc1–4b、0.5B hifi で +fp 12/12 追跡 + 4 実装ビット一致を実機実証済み)。本書はロードマップ **increment 5**(本番 hybrid)の設計。 +**方針転換**(ユーザ指示 2026-07-20): 以後 Qwen3.6-35B のみを対象とし、0.5B デモは基盤実証としてのみ参照。 + +--- + +## 0. 結論(先に) + +1. **重み源は Q4_K_M blob そのもの**。bf16 再量子化はしない。Q4_K/Q6_K の fp16 スケール + (d, dmin)は**ビット操作だけで exact な整数×2^e に分解できる**ため、GEMM を「4/6-bit 整数 × + A24 整数活性 × 整数スケール代数」として**一切の丸め追加なしに**定義できる(§4)。 + → 変換成果物ゼロ、receipt が commit する blob hash がそのまま canonical 整数モデルの重み hash。 +2. **整数化により spec-snapshot §6.3 の最難関(DeltaNet の per-token 丸め列を chunked prefill が + 再現する条件)は構造的に消滅**する。整数加算は結合的なので、GEMM をバッチしつつ状態再帰を + トークン順に回せば、逐次 decode と prefill は**自動的にビット一致**する(§7)。 +3. Mac Studio M1 Max 32GB は本モデルにとって**構造的に最良のローカル機**: 23.9GB blob が unified + memory に丸ごと載る(RTX 4060 Ti 16GB VRAM には載らない)。Metal offload(QI35-5)の速度目標は + fp Metal llama.cpp と同オーダーの tok/s。 + +--- + +## 1. アーキテクチャ ground truth(GGUF + vendored llama.cpp 実測) + +メタデータ(`docs/evidence/` の gguf_full_dump 実行結果、2026-07-20): + +| 項目 | 値 | +|---|---| +| arch / params | `qwen35moe` / 35,951,822,704 | +| 層 | 40(`full_attention_interval=4`: **(i+1)%4==0 の層 3,7,…,39 が full attention(10 層)、他 30 層が gated delta net**。`qwen35moe.cpp:39` `is_recr_impl[i] = (i+1)%4 != 0`) | +| d_model / vocab | 2048 / 248,320(eos 248046) | +| attention | 16 head × head_dim 256、GQA kv 2 head(`head_count_kv` 配列)、q/k per-head RMSNorm、**Q 射影は query+gate 融合出力 8192**(interleave: head 毎 [q256, gate256])、kq_scale = 1/√256 = **1/16(有理数)**、出力は sigmoid(gate) で elementwise gating | +| RoPE | IMROPE だが**テキスト入力では素の NEOX RoPE に退化**(3 位置成分が同値、`llama-graph.cpp:128-139`)。回転は先頭 64/256 次元のみ、ペア (i, i+32) i∈0..31、θ_k = pos·(10^7)^(−2k/64)。**YaRN 無効**(GGUF に rope.scaling 無し → ext_factor=0, freq_scale=1, attn_factor=1) | +| DeltaNet | conv kernel 4 / k-head 16×128 / v-head 32×128(d_inner 4096)/ state 128×128×32head / `ssm_dt`[32] bias / `ssm_a`[32](= −exp(A_log) 変換済み F32)/ β,α 射影 [2048,32] / 出力 gate z [2048,4096] / per-head RMSNorm(ssm_norm[128])× SiLU(z) / ssm_out [4096,2048] | +| MoE(全 40 層) | router `ffn_gate_inp` [2048,256] **F32**、**softmax→top-8**(argsort desc)、選択重みを和で正規化(fp は F16 min でクランプ)、expert_weights_scale **無し**(0.0 → スキップ確認)、expert FFN 512(gate/up Q4_K・down Q6_K、fused gate_up 無し)、shared expert 512 + sigmoid(gate_inp_shexp[2048,1] F16) | +| head | `output.weight` Q6_K [2048,248320](**非 tied**)、output_norm RMSNorm | +| 量子化型 | Q4_K 349 / Q6_K 94 / F32 544 / F16 207 tensor | +| 同梱・非実行 | vision tower `v.*` 441 tensor、MTP `mtp.*` 20 tensor(nextn KV 無し → n_layer_nextn=0、テキスト経路は 40 trunk 層のみ。ローダは個数整合のみ処理 `qwen35moe.cpp:184-192`) | +| その他 | `ssm.v_head_reordered=True` は**ランタイム未参照**(変換時に並べ替え済み、実測 grep 0 件) | + +fp 意味論の出典(全て vendored `llama.cpp`、行番号 2026-07-20 時点): +- 層構造 / attention / MoE 呼び出し: `src/models/qwen35moe.cpp:202-599` +- DeltaNet per-token 再帰(canonical の基準): fused op CPU 実装 `ggml/src/ggml-cpu/ops.cpp:10768-10809` + および非融合 autoregressive `src/models/delta-net-base.cpp:289-371` — 数式は同一: + `S ← S·exp(g)`(head 毎スカラー減衰)→ `Δ[j] = (v[j] − Σ_i S[i][j]k[i])·β` → `S[i][j] += k[i]Δ[j]` + → `o[j] = (Σ_i S[i][j]q[i])·(1/√128)` +- MoE: `src/llama-graph.cpp:1799-2148`(softmax 先行 → argsort top-8 → 和正規化 → silu(gate)·up → + down → ×weight → 8 和; tie-break は fp 実装では**未定義**(std::sort 不安定)) +- 単項 op: silu/sigmoid/softplus(x>20→x)/l2_norm(x/max(√Σx²,eps))/rms_norm(1/√(mean+eps))/ + ssm_conv(窓 4 前方 dot)/soft_max: `ggml-cpu/vec.h:936-1048, unary-ops.cpp:80, ops.cpp:3795-5502,9494,10671` +- RoPE: `ggml-cpu/ops.cpp:5766-6031`(theta_scale = freq_base^(−2/64)、NEOX ペア、64 以降素通し) + +--- + +## 2. Canonical class の定義(何を不変量とするか) + +- **重み**: blob のバイト列(既存 receipt の model hash と同一物)。 +- **整数意味論**: 本書 §4–§9 の固定小数点演算列。丸め点は全て pin(RNE または明示 floor)。 +- **決定性**: 全段整数 → reduction 順序・カーネル・ベンダ・ISA 非依存(inc1–4a で実証済みの性質)。 +- **出力**: greedy argmax token 列(低 index tie-break)と、その上の commitment。 +- fp 版とは**別クラス**(fp 品質をどこまで追跡するかは G3 で実測報告 — 捏造しない)。 + +## 3. スケール表現 + +- 0.5B hifi の **Q60 固定小数点(SC=60, HALF, RNE)+ i128 ホスト演算**をそのまま採用(実証済み)。 +- fp16 → exact 整数: 値 = mant·2^(e−25)(normal: mant = 1024+frac10, e=exp5; subnormal mant=frac10, + e=0 扱い)。fp32 → mant·2^(e−150)。**ビット操作のみ、fp 演算ゼロ**。 +- Metal オンデバイスでスケール代数が要る箇所は設計済みの正規化 (m0,e)+mulhi(親設計 §8)を使う。 + ただし QI35-5 の第一段は CUDA-mirror 構成(GEMM=GPU 整数、スケール代数=host i128)なので不要。 + +## 4. GEMM — Q4_K / Q6_K を**直接**整数消費する + +**Q4_K**(superblock 256 = 32×8 sub-block; 16B ヘッダ + 128B nibble): +`w[i] = d·sc_j·q[i] − dmin·m_j`(q∈[0,15], sc_j,m_j∈[0,63] 6-bit)。行方向(k 方向)dot は + +``` +dot_j = Σ_{i∈sub_j} x_i·q_i (x: A24 → |x|≤2^23; |dot_j| ≤ 2^32.3 → i64) +sum_j = Σ_{i∈sub_j} x_i (|sum_j| ≤ 2^28) +acc = Σ_j [ (d_mant·sc_j)·dot_j·2^{e_d} − (dmin_mant·m_j)·sum_j·2^{e_dm} ] +``` + +d_mant·sc_j < 2^17、積 ≤ 2^49.3 → **i64 に収まり、指数アライン(行内最小指数へシフト)後の総和は +CPU=i128 / GPU=2-word (hi,lo) i64 で exact 累積**。丸めは発生しない(有理数の exact 和)。 +**Q6_K**(256 = 16×16; sc int8 符号付き, w = d·sc_j·(q−32)): 同型で dmin 項なし。 +行スケールは Q60 に正規化した `colsq[o]`(= 2^{e_row}·…)として吐き、以降は 0.5B hifi の +`proj()`(colsq/cmax → ratio → apply → aq)と同一代数。**bias は本モデルに存在しない**(全 GEMM bias 無し)。 + +F32/F16 重み(router / norm γ / ssm_a / ssm_dt / conv1d / shexp gate)は**ロード時に exact 分解 → +per-tensor Q 固定小数点へ pinned RNE 変換**(決定的、blob バイトのみに依存)。 + +## 5. 活性量子化 — **per-token A24** + +- aq() は hifi と同じ(ALVL=2^23−1、符号付き丸め、scale=amax·cur/ALVL)だが、**必ず token 行単位**。 + 理由: prefill(複数 token 同時)と decode(1 token)で量子化が変わらないこと = KV cache / + 逐次一致の前提。0.5B 参照(テンソル一括 amax)からの意図的変更点として pin する。 +- residual stream は i64 値 + per-token スケール(hifi と同じ持ち方)。 + +## 6. 非線形 — 全て pinned 整数 LUT / 整数反復 + +| op | 実装(**G2 で較正・確定** — [evidence](evidence/canonical-int-qwen36-layer-g2-2026-07-20.md)) | 状態 | +|---|---|---| +| softmax exp | **[0,8) 1/2048 刻み 16384 entry Q20**(0.5B の [0,4)×1/512 から拡張 — 35B router 差は ~7.2)、範囲外→0 | 確定 | +| sigmoid / SiLU / β / attention gate / shexp gate | 16384×(1/512) [−16,16) Q16、SiLU = x·σ表[x]、域外 clamp | 確定 | +| softplus(α+dt) | 16384×(1/512) [−16,16) 出力 Q16、x>20→x(ggml 準拠) | 確定 | +| decay γ=exp(g), g≤0 | **[−16,0] 1/2048 刻み 32769 entry Q30**、g<−16→0(実測 g absmean 8.8) | 確定 | +| RMSNorm(attn_norm/post_norm/q,k_norm/ssm_norm) | Σx² → 整数 isqrt。**eps=1e-6 を整数で pin**: `EPS_Q40=1,099,512` / `EPS_Q48=281,474,977`、可変スケール入力は per-token `eps_int=EPS_Q40·U²>>40`(U=2^60/scale)。**+1 近似は不可** — 減衰 head は eps-floor 支配(G2 で 29× ずれる実測) | 確定 | +| L2 norm(DeltaNet q,k) | x·2^20 / max(isqrt(Σx²),1)(ggml の分母 floor 意味論)→ 単位 Q20 | 確定 | +| RoPE cos/sin | Q30、32 ペア × SMAX pos、テーブルのバイト列を成果物として pin | 確定 | +| 1/√128(DeltaNet 読み出し) | Q30 定数 RNE pin | 確定 | +| 1/16(attention scale) | 右シフト 4 の exact 有理数 | — | + +## 7. DeltaNet 整数意味論(per-token が定義、prefill は同ビット) + +状態 S: **Q24 固定小数点・i64 演算で閉じる**(G2 実測 |S|≤10 → 積和 ≤2^55。GPU に i128 不要) +128×128×32head×30 層 ≈ 63MB(i32 格納可)。conv 窓と v は**絶対 Q16**(実測 ±35)、 +q/k は L2 norm 後 **Q20 単位ベクトル**、β/softplus Q16、γ Q30。v-head h ↔ k-head **h%16**(タイル)。 +1 token の列(丸め点 = ★のみ、他は exact): + +1. qkv_mixed = W_qkv GEMM(§4)→ per-token aq ★ +2. conv 窓 [state(3), 今 token] の 4-tap 整数 dot(conv 重みは pinned 整数)→ SiLU LUT ★ → conv state 更新(整数値そのまま保存) +3. q,k 分割 → L2 norm ★(head 128 毎)、k-head 16 → v-head 32 へ複製 +4. β = σ表[β_proj] ★、g = ssm_a ⊙ softplus表[α_proj + dt] ★、γ = exp表[g] ★ +5. S ← (S·γ) ★RNE→QS; Δ[j] = (v[j] − Σ_i S[i][j]k[i]·2^{−…}) ★·β ★; S[i][j] += k[i]Δ[j] ★RNE→QS +6. o[j] = Σ_i S[i][j]q[i] ★ × (1/√128)Q30 ★ +7. o を per-head RMSNorm(ssm_norm)★ × SiLU(z)★ → aq ★ → W_out GEMM(§4) + +**prefill**: 射影 GEMM(1,2 の GEMM 部と 7)は全 token バッチで GPU 実行可(token 独立)。 +状態再帰(5,6)は token 順に実行(≈0.5M 積和/token で軽い)。整数なので**逐次 decode と bit 一致**。 +状態 checkpoint(spec §7 の N=64 token 毎 hash)は S+conv state+KV の整数バイト列に対して自然に定義できる。 + +## 8. Full attention 層整数意味論 + +1. QG = W_q GEMM → interleave 分割(q: head h の [512h, 512h+256)、gate: [512h+256, 512h+512)) +2. q,k per-head RMSNorm ★(γ=attn_q_norm/attn_k_norm)→ RoPE(Q30 表、64 次元、ペア (i,i+32))★ +3. KV cache: k,v を**絶対 Q16 固定小数点**で保存(per-token スケールだと A·V の異 token 混合で + 共通化が要る — 絶対 Q16 なら逐次 decode ≡ prefill が自明。範囲は q/k norm 後 ±10 で安全) +4. scores = q·k dot(i128)× 1/16 → hifi 方式 max 減算 + exp 表 ★ → A·V 整数除算 ★(fdiv+HALF pin) +5. × σ表[gate] ★ → aq ★ → W_o GEMM +GQA: kv head = h/8。mask: 因果のみ(pre-softmax の別 bias 無し、sinks 無し)。 + +## 9. MoE 整数意味論 + +1. logits = router GEMM(F32→pinned Q24 重み、i128 acc) +2. softmax: max 減算 → exp 表 ★ → Σ → 各 top-8 重み w_i = fdiv(e_i·2^30 + Σ/2, Σ) ★ +3. **top-8 = 整数 desc 比較、同点は低 index**(fp の std::sort 不安定性を canonical 側で確定) +4. 分母 Σw_top8 で再正規化 ★(fp の F16-min クランプは整数では Σ≥1 floor として pin) +5. 各 expert: gate/up GEMM(Q4_K §4)→ SiLU 表 ★ × up → aq ★ → down GEMM(Q6_K)→ ×w_i ★ → 8 本 exact 和 +6. shared expert: 同型 + σ表[shexp gate] ★、routed 和に exact 加算 +expert GEMM は「選ばれた 8 expert の行だけ」を実行(offset 計算は tensor 3 次元目 stride)。 + +## 10. メモリ予算(M1 Max 32GB) + +| 項目 | 量 | +|---|---| +| blob mmap(重み直読) | 23.9GB(クリーン evict 可) | +| DeltaNet S + conv state | ~66MB | +| KV cache(4K ctx, 10 層) | ~35MB(整数) | +| LUT/RoPE 表/スケール表 | <50MB | +| 活性ワーク | <200MB | +| 合計 | ~24.3GB — 32GB 内、余裕 ~7GB | + +## 11. 実装ゲート(hardware-validated、捏造なし) + +- **G1(op)**: Q4_K/Q6_K 直接整数 GEMM を Python bigint exact 参照と突合、Rust x86/aarch64 でビット一致 +- **G2(layer)【完了 2026-07-20】**: blk.0(DeltaNet)+ blk.3(attn)+ MoE を llama.cpp CPU fp32 の + 実活性と全 77 チェックポイント照合 — **layer 出力 ≤3%・状態 rmse ~1e-4・MoE top-8 全一致** + ([evidence](evidence/canonical-int-qwen36-layer-g2-2026-07-20.md)、`qi35_int_ref.py`) +- **G3(model)**: 40 層 greedy を fp(CPU/Metal)と複数 prompt で token 一致率測定(結果は実測値のみ記載) +- **G4(Metal)**: 整数 Metal ≡ 整数 CPU のビット一致(fingerprint) +- **G5(CUDA)**: RTX box 復帰時に同 fingerprint(モデルは VRAM 16GB に載らないため CPU offload 構成) + +## 12. 正直な限界 + +- fp 品質追跡(G3 の一致率)は**未測定**。0.5B hifi は 12/12 だったが、35B/Q4_K/長系列で同等とは限らない。 +- 本書の丸め点(★)のうち QS 等の定数は QI35-3 の実データ較正後に pin(現時点は仕様スロット)。 +- 長文(YaRN 域)・vision・MTP・非 greedy sampler は対象外(テキスト greedy のみ)。 +- fp 実装の MoE tie-break 未定義により、logit 完全同点の手では fp と本 class は原理的に分岐し得る。 +- 32GB 機で fp llama.cpp と整数版を**同時**駐在させると swap する(逐次実行で運用)。 diff --git a/docs/evidence/canon_attn.cu b/docs/evidence/canon_attn.cu new file mode 100644 index 0000000000000000000000000000000000000000..f1903c2dbf8ae78f159837e6f6175951d86f83cf --- /dev/null +++ b/docs/evidence/canon_attn.cu @@ -0,0 +1,136 @@ +// inc4b-2b: CUDA (NVIDIA GPU) port of the fully-integer attention block. Host parses layer0.bin and uploads +// clean device arrays; a single-thread kernel runs the block with __int128 scales — same integer semantics +// as the Rust/Python reference. Must reproduce fingerprint dea9a75ba85b1352. +// Build/run: nvcc -O3 -arch=sm_89 canon_attn.cu -o canon_attn_cu && ./canon_attn_cu layer0.bin +#include +#include +#include +#include +#include +typedef long long i64; typedef __int128 i128; +static const int SC=30, D=896, NH=14, NKV=2, HD=64, S=4, FF=4864; +static const unsigned long long TARGET=0x906729150a5a926aULL; // full-layer (4b-3) + +// ---------- host parse ---------- +struct Buf{ std::vector b; long p=0; + unsigned u32(){unsigned v;memcpy(&v,&b[p],4);p+=4;return v;} int i32(){int v;memcpy(&v,&b[p],4);p+=4;return v;} + i64 i64_(){i64 v;memcpy(&v,&b[p],8);p+=8;return v;} unsigned char u8(){return b[p++];} +}; +__host__ __device__ static i128 qm(i64 m0,i64 sh){ i64 e=sh-1; return e>=0?((i128)m0<>(-e)); } + +// device scratch (single thread, no races) +__device__ i64 dyg[S*D], dxnq[S*D], dqq[S*NH*HD], dkq[S*NKV*HD], dvq[S*NKV*HD]; +__device__ i64 dscores[NH*S*S], dpn[NH*S*S], dps[NH*S], dattn[S*D], daqa[S*D], doacc[S*D], dacc[S*FF]; +__device__ i128 dcolsq[FF]; +__device__ i64 dhout[S*D], dxn2[S*D], dgate[S*FF], dup[S*FF], dact[S*FF], dhmul[S*FF]; // 4b-3 MLP scratch + +__device__ static i128 HALFd=(i128)1<<29, ALVLd=32767; +__device__ i128 ap(i128 a,i128 s){return (a*s+HALFd)>>SC;} +__device__ i128 mul128(i128 a,i128 b){return (a*b+HALFd)>>SC;} +__device__ i128 sfrom(i128 amax,i128 cur){return (amax*cur)/ALVLd;} +__device__ i128 fdiv(i128 a,i128 b){i128 q=a/b,r=a%b; if(r!=0&&((r<0)!=(b<0)))return q-1; return q;} +__device__ i128 isqrtd(i128 n){if(n<=0)return 0; i128 x=n,y=(x+1)/2; while(yam)am=a;} i128 amc=am; + for(int i=0;i0?amc:(xi<0?-amc:0); v[i]=(i64)fdiv(xi*ALVLd*2+sg,2*amc);} return sfrom(amc,cur);} + +// proj: acc into dacc/out, returns scale via out param +__device__ i128 projd(const signed char* w,int oc,int inn,const i128* wsq,const i64* bv,i128 bsq, + const i64* xnq,i128 xnsq,i64* out){ + for(int s=0;scmax)cmax=dcolsq[o]; + for(int o=0;o>1),cs)>>SC):0; + i128 r=(cs<>15;} } + int gg=NH/NKV; + for(int h=0;hm)m=v;} + i64 sm=0; for(int t=0;t<=s;t++){ i128 dd=ap((i128)(m-dscores[(h*S+s)*S+t]),c8); if(dd<0)dd=0; if(dd>255)dd=255; + i64 e=expl[(int)dd]; dpn[(h*S+s)*S+t]=e; sm+=e;} dps[h*S+s]=sm; } + for(int h=0;h>1),den)>>SC);}}} + for(int i=0;i4095)idx=4095; dact[i]=dgate[i]*sig[(int)idx]; } + i128 spre=aqd(dact,S*FF,gsq); i128 asq=spre>>16; // dact now int16 + for(int i=0;i& w, std::vector& sq){ + b.u8(); int r=b.u32(), co=b.u32(); w.resize((long)r*co); for(long i=0;i<(long)r*co;i++) w[i]=(signed char)b.u8(); + int ns=b.u32(); sq.resize(ns); for(int i=0;i& v, i128& sq){ + b.u8(); int n=b.u32(); b.u32(); v.resize(n); for(int i=0;i1?argv[1]:"layer0.bin"; FILE* f=fopen(path,"rb"); fseek(f,0,SEEK_END); long n=ftell(f); fseek(f,0,SEEK_SET); + Buf b; b.b.resize(n); fread(b.b.data(),1,n,f); fclose(f); b.p=8; + std::vector gv,qb,kb,vb,pav; i128 gamsc,qbsc,kbsc,vbsc,pasc; + std::vector qw,kw,vw,ow,xv,gw_,uw_,dw_; std::vector qsc,ksc,vsc,osc,xsc,gsc,usc,dsc; + rd_v1d(b,gv,gamsc); rd_w2d(b,qw,qsc); rd_v1d(b,qb,qbsc); rd_w2d(b,kw,ksc); rd_v1d(b,kb,kbsc); + rd_w2d(b,vw,vsc); rd_v1d(b,vb,vbsc); rd_w2d(b,ow,osc); rd_w2d(b,xv,xsc); + rd_v1d(b,pav,pasc); rd_w2d(b,gw_,gsc); rd_w2d(b,uw_,usc); rd_w2d(b,dw_,dsc); + std::vector cosq(S*HD),sinq(S*HD),expl(256),sig(4096); + for(auto* a:{&cosq,&sinq}) for(auto& x:*a) x=b.i64_(); for(auto& x:expl) x=b.i64_(); for(auto& x:sig) x=b.i64_(); + // upload + auto up=[&](auto& v){ using T=typename std::decay::type::value_type; T* d; cudaMalloc(&d,v.size()*sizeof(T)); cudaMemcpy(d,v.data(),v.size()*sizeof(T),cudaMemcpyHostToDevice); return d; }; + auto *dqw=up(qw),*dkw=up(kw),*dvw=up(vw),*dow=up(ow),*dxv=up(xv); + auto *dqsc=up(qsc),*dksc=up(ksc),*dvsc=up(vsc),*dosc=up(osc),*dxsc=up(xsc); + auto *dgv=up(gv),*dqb=up(qb),*dkb=up(kb),*dvb=up(vb),*dcos=up(cosq),*dsin=up(sinq),*dexp=up(expl); + auto *dpav=up(pav),*dsig=up(sig); // i64* + auto *dgw=up(gw_),*duw=up(uw_),*ddw=up(dw_); // signed char* + auto *dgsc=up(gsc),*dusc=up(usc),*ddsc=up(dsc); // i128* + i64* dhout; cudaMalloc(&dhout,S*D*sizeof(i64)); + run<<<1,1>>>(dqw,dqsc,dqb,qbsc, dkw,dksc,dkb,kbsc, dvw,dvsc,dvb,vbsc, dow,dosc, + dxv,dxsc,dgv,gamsc, dpav,pasc,dgw,dgsc, duw,dusc,ddw,ddsc, dcos,dsin,dexp,dsig, dhout); + cudaError_t e=cudaDeviceSynchronize(); if(e){printf("cuda err %s\n",cudaGetErrorString(e));return 1;} + std::vector hout(S*D); cudaMemcpy(hout.data(),dhout,S*D*sizeof(i64),cudaMemcpyDeviceToHost); + unsigned long long h=0xcbf29ce484222325ULL; for(i64 v:hout){ unsigned char* p=(unsigned char*)&v; for(int k=0;k<8;k++){h^=p[k]; h*=0x100000001b3ULL;} } + cudaDeviceProp pr; cudaGetDeviceProperties(&pr,0); + printf("backend = NVIDIA CUDA GPU (%s sm_%d%d)\n",pr.name,pr.major,pr.minor); + printf("block-out FNV= %016llx\n",h); + printf("TARGET = %016llx\n",TARGET); + printf("MATCH = %s\n", h==TARGET?"YES (byte-identical to Rust/Python reference)":"NO"); + printf("hout[0..4] = %lld %lld %lld %lld\n",(long long)hout[0],(long long)hout[1],(long long)hout[2],(long long)hout[3]); + return 0; +} diff --git a/docs/evidence/canon_attn.py b/docs/evidence/canon_attn.py new file mode 100644 index 0000000000000000000000000000000000000000..4ec57a5532f59d3c2fded72bab0b9cb08726d29d --- /dev/null +++ b/docs/evidence/canon_attn.py @@ -0,0 +1,188 @@ +#!/usr/bin/env python3 +# inc4b-2a: integer attention block (Qwen2.5-0.5B layer 0) consuming the real `.canon` W8A8 artifact. +# Validates CORRECTNESS: the integer block (int8 matmuls + integer RMSNorm/RoPE/softmax) matches a full-fp +# reference computed from the SAME dequantized weights. Scalar scales are carried as fp here (they become +# fixed-point in the 4b-2b cross-backend port). Heavy tensor compute is all integer. +import sys, json, struct, numpy as np + +CANON = sys.argv[1] if len(sys.argv) > 1 else "/root/models/qwen05.canon" +D, NH, NKV, HD, THETA = 896, 14, 2, 64, 1000000.0 +S = 4 # prefill length + +def read_canon(path): + with open(path, "rb") as f: + assert f.read(8) == b"CANONv1\n" + n = struct.unpack("