S4b — read_mean_ssim を gpu に移す ssim_reduce flag、採用確定 (2026-08-30 gate)

`p3-s4-host-shadow` §7 は「`read_mean_ssim` (1.92M 要素の CPU 総和、`loss = "l1_ssim"` の host 仕事 1.871ms/iter で最大の未着手候補) を動かせば、S4 の実測見積り (-1.5〜-3%) を元の -3〜-5% レンジに近づけられる可能性が高い」と申し送っていた。本 finding はその follow-up 実装 — `[backend] ssim_reduce` flag に `host` (default) / `host_lagged` / `gpu` の 3 variant を追加し、消費箇所の分類・2 案の比較・test・gate config 準備までを行う。

1. `read_mean_ssim` の呼び出し経路と消費箇所の分類

`loss_history` (`TrainOutcome.loss_history`)

`train_loop` が毎 iteration 1 個 push する `Vec<f32>`。**唯一の実質的な消費者はさらに 2 段先**: (1) `outcome.loss_history.last()` → CLI の `final_loss_value` (stderr 表示 + `result.toml` `metrics.final_loss`)、(2) `snapshots` (`log_every` ヒット時の `TrainSnapshot { iter, loss, splats }`) → `result.toml` `[curve.loss]` (web の run summary グラフ)。**refine tick / 早期停止判定 / eval は loss を一切参照しない** (`splat-train-v1/src/refine.rs` に `loss` という識別子は 0 件、grep で確認済み) — refine は `it % every` のみで駆動する純粋な iteration counter 駆動

1 iteration 遅延しても意味論が変わらないもの

`loss_history` の**中間**エントリ (`i` 番目の値が eventually 正しい iteration `i` の値になっていれば、host が『いつ』計算するかは result.toml/stderr の見た目に一切影響しない) — 唯一の制約は『配列の長さと各要素の帰属を保つ』こと

1 iteration 遅延すると意味論が変わるもの

(1) **最終 iteration (`it == max_steps`) の loss** — `loss_history.last()` が `result.toml` の `final_loss` として報告される実質的な成果指標なので、1-step stale になると『学習が実際にどこで終わったか』とズレた値を報告してしまう。(2) **`snapshots` (`log_every` ヒット時)** — `curve.loss` は『iteration `it` 時点の loss』として web に描画されるグラフなので、defer されたまま古い値が入ると per-iteration 帰属が壊れる (advisor 指摘)。(3) **refine tick 直前の値** — 本実装には該当する消費者が無い (refine は loss を見ない) ので理論的な懸念に留まる、念のため §3 の設計はこの経路も汚さないよう `is_snapshot_iter` 相当のガードを refine tick には不要と判断した (refine tick 自身の iteration は `param_mutated_after_step` により `prefetched=false` になるため、defer 条件が構造的に成立せず自動的に安全)

2. 案 (a) GPU reduction と (b) 1 iteration lag の比較 — 選定は「両方実装」

(a) GPU reduction (`ssim_sum_reduce`)

実装してみると想定より工数が低かった — `loss.metal` の `loss_l1_only` (`splat-metal/src/kernels/loss.rs` の `LossEvaluator`) が**すでに同型の kernel** (`atomic<float>` + `simd_sum`、`loss_sum_buf` の zero-clear→dispatch→4 byte readback) を実装・本番稼働させていたため、そのパターンを丸ごと再利用するだけで済んだ。**利点**: pipelining/host_schedule に一切依存せず、`prefetched == false` の iteration (refine tick・grad RMW・MCMC noise・最終 iteration) でも効く — host 仕事を『どかす』のではなく『消す』。**欠点**: bit-exact ではない (atomic 加算順が dispatch ごとに非決定的)。ただし `loss = "l1_ssim"` の L1 項はすでに `loss_l1_combine_ssim` の atomic 加算で同じ性質の ULP 揺れを持っており (`loss.metal` 冒頭コメントで既存文書化済み)、**新しい種類の非決定性を持ち込むわけではない**。gradient (`ssim_backward*`) には一切触れないので学習結果 (PSNR) への影響は理論上ゼロ

(b) 1 iteration lag (`host_lagged`)

`host_schedule = "backward_shadow"` が既に実装した『grad readback を 1 iteration 遅らせて次の backward CB shadow で flush する』パターンをそのまま SSIM 総和に適用。`work_ssim` は `Rasterize::work_grad` と同じ `GenScratch` 2 世代 ping-pong に乗っているため (`ssim.rs` の既存コメントがこれを明記していた)、grad buffer と全く同じ安全性論証が成り立つ。**利点**: アルゴリズム自体は `Host` と完全同一 (CPU `f64` 逐次総和) なので、同一 buffer に対しては値が bit-exact (`read_mean_ssim_from_buf_matches_synchronous_read` test で pin)。**欠点**: `host_schedule = "backward_shadow"` かつ『この iteration が次を prefetch した』ときにしか効かない — refine tick 直前 / RMW / 最終 iteration では構造的に no-op。実装コストは (a) より高い: `loss_history` の bit-exact 性を保ったまま `history.push` 自体を 1 slot 遅らせる設計 (`StepOutcome.loss: Option<f32>` + `StepOutcome.deferred_loss: Option<f32>` の 2 本立て、`StepCtx.is_snapshot_iter` で snapshot/log 消費点だけ defer を止める) が必要だった

選定

**advisor レビューでの結論に従い、両方を `ssim_reduce` の 3 variant (`host`/`host_lagged`/`gpu`) として実装した**。「工数と bit-exact 性で有利な方」を単独で選ぶなら bit-exact 性は (b)、実測工数は (a) の方が僅かに低かった (loss.metal の precedent を丸ごと再利用できたため) — 実装時点では甲乙つけがたく、かつ (a) は (b) が構造的にカバーできない iteration (refine tick 等) もカバーするため両方を実装した。**その後判明した追加情報 (`p3-s4-host-shadow` 本体の 12.5k gate 結果) が選定を確定させた**: `host_schedule = "backward_shadow"` が gate で棄却された (+1〜5% 遅い、期待符号と逆) ため、それに依存する (b) `host_lagged` は現状の production スタックでは**発火する場面が無い** (`backward_shadow` を有効にしないなら `host_lagged` は常に no-op)。よって **(a) gpu を唯一の gate 対象 (on arm)** とし、(b) host_lagged は実装・test のみ完了させた代替 flag として残す (追加工数はかけない、`p3-s4-host-shadow` の verdict が将来覆った場合の再利用候補)

2b. gate 前の Δ 予測と PSNR guard の位置づけ (advisor 指摘、gate 実施前に明記)

Δ 予測 (wallclock、`ts_ssim_host_sum` 実測後の確定版)

計装 (§4) の結果、1000 iter smoke (800×800、splats ~6.8k) で `ts_ssim_host_sum` (host 経路の同期 CPU 総和) = **1.093ms/iter** (TOTAL[sum] の 31.0%) — `p3-c2-bubble-attribution` §4 の事前見積り (grad copy 150µs + accumulate 407µs を引いた残り ≈1.3ms) とほぼ整合。SSIM 総和は **`W*H*3` の image-size-bound コストで splat 数に依存しない** (forward/backward/binning は splat 数に比例) ため、この 1.093ms/iter という絶対値は 12.5k 本番 regime (400k splats) でもほぼ同じ値で持ち越せると見積る。base の forward shadow (2.00ms) に対し host 仕事 1.871ms/iter (fill 93.6%、`p3-c2-bubble-attribution` §4) のうち SSIM 総和が 1.093/1.871=**58%** を占める計算になる。`gpu` はこれを ~0 (実測 `ts_ssim_host_sum`=29.6µs + `ts_ssim_sum_reduce` encode 326ns、合計 <0.04ms) まで削るので、fill 率は 93.6% → 93.6%×(1-0.58)=**39%程度**まで下がる。残差 7.2〜8.1s (wall比 5.8〜6.5%、`p3-c2-bubble-attribution` §5) が host-late にほぼ 1:1 対応するという機構仮説のもとでは、回収量は概ね『動かした割合』(58%) に比例: 残差の 58% ≈ **4.2〜4.7s (wall 124.75s 比 3.4〜3.8%)**。ceiling ではなく比例見積りであり、fill 率が 100% を割った時点で理論上は頭打ちになるはずなので実際の回収はこれより小さく出る可能性がある一方、fill 率が 39% まで下がれば host-late 自体がほぼ消える可能性もある。**レンジとして -2〜-4% と見積る**

smoke 規模での実測 (参考、production 規模への外挿には注意)

1000 iter smoke (splats ~6.8k) では **wall 3.427s (host) → 2.927s (gpu)、busy 76.0% → 91.4%** という大きな差が観測できた (Δ -14.6%)。これは `p3-s4-host-shadow` の smoke が『grad copy/accumulate は splat 数に比例するコストなので splats ~6.9k では小さすぎて run-to-run noise に埋もれる』と結論づけたのと対照的 — **SSIM 総和は image-size-bound なので splat 数が少ない smoke 規模でも絶対値 (1.093ms/iter) がほぼ変わらず、むしろ per-iteration budget が小さい smoke ではその『相対的な重さ』が大きくなる**。この smoke の -14.6% という数字自体は 12.5k 本番の Δ (-2〜-4% と見積り) には直接外挿できない (400k splats では forward/backward/binning の絶対コストが増え、SSIM 総和 1ms が相対的に薄まる) が、**instrumentation が正しく効いていること・機構仮説 (host 仕事の削減が busy% を直接押し上げる) 自体は smoke でも定性的に確認できた** という点で価値がある

PSNR guard の位置づけ (advisor 指摘 — 誤読すると誤った verdict を出す)

`work_ssim` は `read_mean_ssim`/`ssim_sum_reduce` からしか読まれず、そのまま `mean_ssim` → `total_loss` → `history` に流れるだけで、**gradient には一切合流しない** (gradient は別経路: `encode_forward_and_grad` が書く `work_grad` → `encode_l1_combine_ssim` の `dldr` → backward)。refine/LR schedule/checkpoint も loss を参照しない (§1 で確認済み)。**したがって `ssim_reduce` は構造的に学習された param を一切変えられない** — off/on で PSNR に系統的な差が出ることは原理的にありえない。gate で観測される off-vs-on の PSNR 差は **すべて backward atomic の run-to-run noise** (`per_kernel_single_texel_run_to_run_deterministic_1iter` と同型) であり、`p3-s4-host-shadow` の smoke が使った比較基準 (off-vs-off の splats 数ばらつき 6886 vs 6896 等) と同じ土俵で見るべきもの。**`±0.1 dB` の PSNR guard を『抵触したら NO-GO』の基準として使うのは誤り** — off-vs-off の repeat 内ばらつきと同程度の差なら『guard 抵触』ではなく『ノイズ』と正しく判定すること。次段が gate 結果を読むときにこの一段落を必ず参照すること

3. flag 配線と依存関係

• `splat/crates/splat-train-v1/src/config.rs`: `SsimReduceKind { Host (#[default]), HostLagged, Gpu }`、`BackendConfig::ssim_reduce` field。`validate()` に `ssim_reduce = "host_lagged"` かつ `host_schedule != "backward_shadow"` の **warning** (host_schedule 自身の hard error とは非対称 — defer 条件が `self.backward_shadow && self.ssim_lagged && ...` の AND で構造的に false に落ちて `Host` と完全に同じ挙動になる『効かないだけで安全』なケースなので、config を reject するのではなく警告に留めた、advisor 指摘どおり fallback を doc に明記)
• `splat/crates/splat-metal/src/kernels/ssim.rs`: kernel-level mirror `SsimReduceKind { Host, Gpu }` (2 variant のみ — `HostLagged` は算法的に `Host` と同一なので mirror が要らない)。新規: `pso_sum_reduce` PSO、`sum_buf: GenValue` (`loss.metal` の `loss_sum_buf` と同型)、`encode_sum_reduce()`、`read_mean_ssim(img_size, reduce)` (旧 API に reduce 引数を追加)、`read_mean_ssim_from_buf()` (deferred flush 専用、生 Buffer から直接総和)、`ssim_buf_handle()` (deferred flush 用に現在の `work_ssim` handle を clone)
• `splat/shaders/loss/ssim.metal`: `ssim_sum_reduce` kernel を追加 (`loss.metal` の `loss_l1_only` と同一パターン、warp `simd_sum` → lane0 `atomic_fetch_add`)
• `splat/crates/splat-train-v1/src/pipeline.rs`: `Depth2Pipeline` に `ssim_lagged: bool` / `deferred_ssim: Option` を追加。`DeferredSsim { ssim_buf: Buffer, l1: f32, lambda: f32, img_size }`。`flush_deferred_ssim()` を新設し `flush_deferred_accumulate()` と同じ flush point (手順 4 の直後) で常時呼ぶ。`StepOutcome` を `{ loss: Option, deferred_loss: Option, prefetched_forward: bool }` に変更 — defer した iteration では `loss = None`、次の呼び出しの flush で確定した値が `deferred_loss` に載る。`StepCtx` に `is_snapshot_iter: bool` を追加 (snapshot/log 消費点では defer しない)。`encode_loss_backward_chain` に `trainer.ssim_reduce.get() == Gpu` のときだけ `encode_sum_reduce` を追加 encoder として積む処理を追加
• `splat/crates/splat-train-v1/src/train_loop.rs`: `is_snapshot_iter` を `log_every` の条件式から計算し `StepCtx` に渡す。`let loss: Option = match {...}` に変更、`Some(p)` 枝で `out.deferred_loss` を先に `history.push`、その後 `out.loss` があれば push (`loss: Option` → `if let Some(l) = loss { history.push(l); }`)。snapshot/log ブロックは `is_snapshot_iter` をガードにし、`loss.expect(...)` で defer されていないことを pin する。`ssim_lagged` bool を `backend.ssim_reduce == HostLagged` から計算し `Depth2Pipeline::new` の第 4 引数に渡す (`host_schedule` と同じ配線パターン — `Trainer::Cell` を経由せず `backend: &BackendConfig` から直接)
• `splat/crates/splat-train-v1/src/trainer.rs` / `splat/crates/splat-cli/src/util.rs`: `Trainer::ssim_reduce: Cell` を追加、`apply_backend_config` の exhaustive destructure に `ssim_reduce` を追加し `Host|HostLagged → Host`、`Gpu → Gpu` で配線 (host_lagged の defer タイミングは `train_loop` が直接読む別経路)

4. 検証結果

5. 未解決点 / 次段への申し送り

• ~~**12.5k A/B gate (n=3) 未実施**~~ → **2026-08-30 実施済み、採用確定**。§6/§7 参照
• **`host_lagged` は gate 対象外のまま (実装のみ、追加工数なし)** — `p3-s4-host-shadow` (backward_shadow) が棄却されたため現状の production スタックでは発火しない。`p3-s4-host-shadow` の verdict が将来覆るか、backward_shadow に依存しない別の defer 機構が生まれた場合の再利用候補として残す (verdict は §5 のまま変更なし、コードの enum doc にも同旨を追記済み)
• **GPU reduction の regime 依存性が未検証のまま採用** — 128×128 (49,152 要素) の unit test では `Gpu` vs `Host` の相対誤差 <1e-4 だったが、12.5k 本番 regime (800×800×3 = 1.92M 要素、K=7 window) では atomic 競合の規模が 2 桁近く大きい。gate 自体は wallclock/PSNR の両方で pass しており (§6)、PSNR guard は「gradient 非変更のため理論上ノイズ」という設計上の主張どおりの挙動 (§7) だったため、この不確実性は採否判断をブロックしなかった
• ~~gate run 後は本 finding に verdict を追記~~ → **完了 (§6/§7)**。§2b の Δ 予測 (-2〜-4%) は実測 (-1.4347%) より過大だったため、§7 で機構解釈を追記した
• **`ssim_reduce_host_lagged_matches_host_loss_history_bit_exact` は ~1/30 の確率で `assert_params_close` の rel tolerance をわずかに超えて fail しうる** (§4 発見 2 参照)。既存 sibling test も同水準の flake rate を持つ pre-existing な fixture 性質であり S4b 固有のバグではないため対処不要と判断したが、CI で偶発的に赤くなった場合はこの finding を参照して再実行で切り分けること (real regression なら閾値を大きく超える・再現性がある、のいずれかで区別できるはず)

6. 12.5k n=3 gate (2026-08-30、s4b-12500 series) — 採用根拠

前提

build: `cargo build --release`、HEAD `1c54a0b6a064` 上でビルド (binary sha256 `a4317ea72e8b35e259075859c3e312fac2930a28f7d92db7849894087ccfef4e`)。起動直前に `pgrep -fl "target/release/splat"` を確認したところ直前の re200 gate (別 finding `p3-refine-every-200`) の r2 pair retry がまだ稼働中だったため、その完了 (pgrep 空) を待ってから series を開始した。config 6 本は事前に `splat check` で ok を確認済み。AC 接続を series 通じて維持 (`pmset -g batt` で `AC Power` を都度確認、driver 自身も check_ac で監視)

driver

`splat/scripts/gate-ab-n.sh configs/2026-08-30-s4b-12500-off configs/2026-08-30-s4b-12500-on s4b-12500 3` (nohup+disown で detach、`s4b-12500-gate.DONE` sentinel で完了検出)。**1 回目起動は prefix に `configs/` を付け忘れて off-r1 が `config IO error` で即時 rc=1 終了 (GPU work 無し、0.018s) — ログは `s4b-12500-*-BADPREFIX-attempt1.*` として保存し、prefix を修正して再起動した (2 回目が本 gate の実測)**

binary/HEAD provenance

series 開始時 binary sha256 `a4317ea7…` / HEAD `1c54a0b`。**series 完了時、driver が記録した binary sha256 は開始時と同一 (`a4317ea7…`、不変)**、一方 `git rev (post-series)` は `b88e4dcb86c3` に変わっていた — 別 finding (`p3-refine-every-200`、re200 gate の結果) の docs-only commit が series 実行中に main へ入ったため。code 変更を含まないコミット (`git show --stat` で確認、変更は `docs/findings/*` と `splat/configs/2026-08-30-re200-*` と `splat/runs-logs/re200-*` のみ) であり、かつ binary は series 中一度も再ビルドしていないため、**6 run はすべて単一 binary (`a4317ea7…`) 上で実行されており gate 結果への影響はない**

flag 実測確認 (6/6)

全 6 run の `[backend] BackendConfig { … }` から他の全 flag が off/on 間で完全に同一であることを確認 (`backward=Simd, loss_path=Gpu, binning=Resident, pipelining=Depth2, host_schedule=ForwardShadow, refine_scratch_persist=true, refine_direct_write=true` 等)、差分は `ssim_reduce: Host` (off 3/3) / `ssim_reduce: Gpu` (on 3/3) のみ

汚染 / retry 判定

事前登録条件 (`p3-tier1-bit-exact.toml` §5 の教訓どおり (i) load 証拠 AND (ii) 自 arm 内 3σ 超 の **AND**) を両 arm 全 6 run に適用。off (125.094, 124.605, 123.966, mean 124.555, sample σ 0.566s=0.454%) — 最大偏差は off-r1 の 0.95σ、off-r1 の run 窓内 load1 (mean 3.99, max 5.12) は他の run より高いが (i) のみで (ii) 不成立。on (123.987, 122.062, 122.254, mean 122.768, sample σ 1.060s=0.864%) — 最大偏差は on-r1 の 1.15σ。**いずれの run も条件 (ii) の 3σ を大きく下回るため retry 対象なし、6 run すべてそのまま採用**

7. Δ 実測 (-1.4347%) が §2b 予測 (-2〜-4%) より小さかったことの機構解釈

8. default flip の実装詳細 — `p3-s4-host-shadow.toml` §10 の罠の回避

• `p3-s4-host-shadow.toml` §10 が指摘した罠 (field-level `#[serde(default)]` を付けた enum field は、TOML 省略時に `BackendConfig::default()` の literal ではなく enum 自身の `#[default]` 属性 [`SsimReduceKind::default()`] から解決される — bool field の `bool::default()`=false 罠と同型) を踏まえ、**4 箇所を一括で変更した**: (1) `SsimReduceKind` enum の `#[default]` 属性を `Host` から `Gpu` に移動、(2) `BackendConfig::default()` の `ssim_reduce: SsimReduceKind::Host` を `SsimReduceKind::Gpu` に変更、(3) `BackendConfig::ssim_reduce` field の doc comment (旧: `default = Host (従来動作)。gate 待ち。`) を gate 結果込みで更新、(4) `SsimReduceKind` の 3 variant doc (`Host` から「現行」を外す、`HostLagged` に「gate 対象外のまま」の明示 verdict を追加、`Gpu` に `#[default]` と gate 実測を追記)
• **実装後に実機で検証**: `ssim_reduce` を省略した config (12.5k gate の off-r1 config から `ssim_reduce = "host"` 行を削除、`max_steps = 1` にした 1-iter smoke) を実行し、`splat check`/`splat train` が出力する `[backend] BackendConfig { …, ssim_reduce: Gpu, … }` を目視確認した — TOML 省略時の解決先が期待どおり `Gpu` になっていることを Rust の型システムの推論だけでなく実行時出力で pin した (A7 の default flip 前検証と同型の手順)
• kernel-level mirror (`splat_metal::kernels::ssim::SsimReduceKind`、2 variant のみ) 側の `#[default] Host` 属性は**変更していない** — この mirror の default 値は `Trainer::new` が `Cell::new(SsimReduceKind::default())` で初期化する一時値に過ぎず、`apply_backend_config` (`splat-cli/src/util.rs`、`train.rs`/`bench.rs` の両方の config load 経路から無条件で呼ばれることを確認済み) が config 側の値で必ず上書きするため、mirror 側 default の値自体は実行結果に影響しない。config 側 (`splat-train-v1::config::SsimReduceKind`) と mirror 側の 2 つの enum が別物であることに注意 (mirror 側は変更不要、と結論)
• `config_paths` にある gate config family (`2026-08-30-final-12500-r{1,2,3}.toml` 等、ヘッダに『全 [backend] flag は default に任せる』と明記されたもの) は、本 flip 後は `ssim_reduce` を明示していない限り**黙って `Gpu` 経路になる** — 意図した挙動 (default 追従) だが、これらの config が過去に生成した『baseline』wallclock 数値は host 経路のものであり、本 flip 後に同じ config を再実行すると異なる (gpu 経路の) 数値になる点に注意。既存の gate config はすべて `ssim_reduce` を明示済み (off/on とも) なので、本 flip によって過去の gate 結果自体が無効になることはない
• `cargo test --workspace --release` を計 4 回連続実行し、うち 2 回で `train_loop::tests::ssim_reduce_host_lagged_matches_host_loss_history_bit_exact` / `ssim_reduce_host_lagged_without_backward_shadow_matches_host` が低頻度で fail した (`rot_quats[2]` rel≈1.0e-4、§4 発見 2 が既に文書化した ~1/30 pre-existing flake と同一 signature)。両 test は `ssim_reduce` を fixture 内で明示的に `Host`/`HostLagged` に設定しており default 値には依存しないため、**本 flip が原因ではない** (flip 前から存在する rasterize backward atomic 非決定性、既知の flake)。最終確認の連続実行では全 pass (splat-metal 95、splat-train-v1 90、他クレート含め failure 0) を得た

• mixedC2 実測 — 残差 bubble は host-late 支配 (7.2〜8.1s)、refine tick drain は A7 で半減済み (3.0〜3.8s)。handoff は定常帯で headroom ほぼ無し2026-08-30
• positiveM7 最終スタック確定 — 12.5k 124.39s (独立 n=3)、msplat 等価 gap 0.927×→0.755×、throughput 1.135×負け→0.926× (msplat 側 avg splats 借用値 ±8% 未満なら逆転)2026-08-30
• negativeS4 — post-backward accumulate を backward CB shadow に配置する host_schedule flag — 棄却確定 (2026-08-30 gate)2026-08-30
• positiveS3 iteration pipelining depth 2 — 12.5k -18.1% / busy 78 → 91% (default flip)2026-08-17