FindingID: p3-per-splat-backward-tg256
Status: stable / task 2026-09-seminar05-milestones.md — M3 追補 (per-splat backward 第 2 版、occupancy 仮説検証)

2026-09 seminar05 milestones · M3b · per-splat backward 第 2 版

M3b per-splat backward 第 2 版 (256-thread threadgroup) — NO-GO 確定 (occupancy 仮説は方向的に正しいが不十分)

第 1 版 (<code>rasterize_backwards_per_splat</code>、32-thread/threadgroup、bench で default 比 +34.1% regression) の finding が挙げた「低 occupancy が最有力仮説」を検証するため、<strong>アルゴリズム・diagonal scheduling は完全に同一のまま `PER_SPLAT_BATCH` だけ 32→256 (= 16×16 tile 全体、`pso_min256` 対象) に変えた</strong> 新 kernel <code>rasterize_backwards_per_splat_tg256</code> (flag <code>backward_scheme = "per_splat_tg256"</code>) を追加した。GPU 検証 (PSO compile / parity 3 本 / trainer 統合 parity / 4-way bench) を実施した結果、<strong>tg256 は第 1 版比 -4.22% (occupancy 仮説の方向は正しい) だが、default per_pixel 比 +32.8% で依然として大幅に遅く NO-GO 確定</strong>。12.5k gate は実施していない (bench の時点で明確に劣後のため)。

mixedrank: midno-go-confirmed第 1 軸第 3 軸seminar05-milestonesm3bper-splat-backwardoccupancythreadgroup-sizediagonal-schedulingrasterize-backwardgpu-pending

Δ PSNR

N/A (bench のみ、12.5k gate 未実施。trainer 統合 parity は5 grad tensor 全て tol 内で PASS — correctness は確認済み)

Δ Wallclock

tg256 vs default per_pixel: **+32.8%** (劣後、bench median 17.994ms vs 13.548ms, n=3)。tg256 vs per_splat v1: -4.22% (occupancy 仮説の方向は正しい)。12.5k gate 未実施 (bench 時点で NO-GO 確定のため)

Impact

第 1 版 finding (`p3-per-splat-backward.toml` §5) の occupancy 仮説を検証するため、`rasterize_backwards_per_splat_tg256` (`BackwardSchemeKind::PerSplatTg256`、flag `per_splat_tg256`) を追加した。**アルゴリズム面では第 1 版と一字一句同一** (diagonal scheduling、pixel 側 state 復元、last_contributor を使わない理由、gradient 式) で、変えたのは `PER_SPLAT_BATCH` (= threadgroup thread 数) の 1 パラメータのみ (32→256)。GPU 検証: (1) `rasterize_psos_compile` PASS — `pso_min256` の 256-thread occupancy floor をレジスタ圧迫なしに達成。(2) parity 3 本 (`backward_per_splat_tg256_matches_per_pixel_original` [32×32/24×24]、`_multi_round` [n=600, rounds=3]、`_byte_equality_single_pixel`) 全 PASS。(3) trainer 統合レベル `splat bench --parity` (実 checkpoint 409,660 splats) — 5 grad tensor 全て tol (atol/rtol=1e-4) 内で PASS、`sorted_idx` は既知の非決定性 (0.222% 差、非 gate)。(4) 4-way bench (median ms/iter, n=3, warmup=10/iters=50/repeats=3): per_pixel (default, σ0.59%) **13.548** / per_pixel unopt (σ0.34%) 15.701 / per_splat v1 (σ0.12%) 18.787 / **per_splat_tg256 (σ0.35%) 17.994**。tg256 は v1 比 **-4.22%** — §3 が予告した 3 効果 (barrier 総数 4.5 分の 1 / thread-iteration 効率 1.78 分の 1 悪化 / barrier 単価上昇) の net が僅かに改善方向に振れたことを示し、occupancy 仮説は**方向としては裏付けられた**。しかし default per_pixel 比では **+32.8%** で依然として大幅に劣後しており、M3 系列 (per-splat backward、atomic なし構造) はこの bench harness 上で **NO-GO 確定**。12.5k gate は実施していない (bench の時点で `max(2σ,1%)` の閾値を大幅に超えて劣後することが明白なため、instructions の「遅ければ gate せず終了」に従った)。`BackwardSchemeKind` の default は `PerPixel` のまま、`.rs` は無変更。

本文

本 finding は 「実装完了、GPU 検証は次段」という中間状態を記録する。到達点がどうであれ finding を追記する規約 (`docs/architecture/backend-and-compiler-conventions.md`) に従い、実装の設計・第 1 版との差分・parity test の設計・そして まだ検証していないことを検証していないと明示するために書く。GO/NO-GO の判定は次段の GPU session に持ち越し。

1. 実装 — 第 1 版とアルゴリズムは同一、PER_SPLAT_BATCH のみ 32→256

変更点は 1 パラメータだけ: 第 1 版 (`rasterize_backwards_per_splat`、`PER_SPLAT_BATCH = 32`) をコピーし、`PER_SPLAT_BATCH_TG256 = 256` (= `PS_TILE_SIZE`、16×16 tile 全体) に変えただけの kernel `rasterize_backwards_per_splat_tg256` を追加した。diagonal scheduling の式 (`pixel_rank = i - lid`)、T の復元・early-out (`next_T <= 1e-4` で sentinel 0.0)、gradient 式、tile ごとの atomic emit 1 回、すべて第 1 版と同一。`PER_SPLAT_BATCH_TG256 == PS_TILE_SIZE` なので pixel 初期化ループ (第 1 版は 8 pixel/thread) が 1 thread = 1 pixel の 1:1 対応に簡略化される以外、コード上の差分は定数 1 個と threadgroup dispatch サイズだけ
PSO と occupancy 規約: 第 1 版は 32-thread dispatch のため `pso_min256` (256-thread occupancy floor、コンパイラ規約 C6) の対象外と判断し `pso_with_constants` (floor なし) を使っていた。本 kernel は 256-thread/threadgroup (= tile 全体) で dispatch するため、他の全 kernel (forward/backward per_pixel 系/tg/simd) と同じ `pso_min256` を使う — **これが第 1 版から見た『他 kernel と同じ occupancy 規約に合わせる』という変更の核心**
Rust 側のリファクタ: `Rasterize::backward_per_splat_from_buf` (第 1 版の公開 API) と新規 `backward_per_splat_tg256_from_buf` はどちらも private helper `backward_per_splat_dispatch(pso, tg_threads, ...)` に委譲する形にした。buffer binding・grad buffer 管理・S1 attribution 記録は完全共通、PSO 参照と threadgroup thread 数 (32 or 256) だけが呼び出し元で異なる。第 1 版実装時点で存在した重複コード (90 行超のほぼ同一関数) をこのタイミングで解消した
config / wiring: `BackwardSchemeKind::PerSplatTg256` (flag 値 `per_splat_tg256`) を `splat-metal::kernels::rasterize` と `splat-train-v1::config` の両方に追加 (実行時 mirror の対応関係は既存の `PerSplat` と同型)。`BackendConfig::validate` の no-op warning (`backward`/`raster_backward_range`/`raster_backward_sparse` が per_splat 系では無視される) を両 variant 対象に拡張。`util.rs::apply_backend_config` と `trainer.rs` の 2 dispatch call site を if/else から 3-way match に変更 (`PerPixel`/`PerSplat`/`PerSplatTg256`)。default は `PerPixel` のまま無変更

2. 仕様からの逸脱 — 「barrier 1-2 回/batch」は 1 thread=1 splat + register 累積と両立しない

spec deviation の明示 (advisor レビュー指摘 #3)

本タスクの依頼文は「batch 境界の barrier は 1 batch あたり 1-2 回に抑える」ことを目標に挙げていたが、これは『1 thread = 1 splat + splat 側勾配を register に累積 (atomic 不要)』という制約と構造的に両立しない — alpha compositing は pixel ごとに順序依存 (前の splat の合成結果 `pix_acc`/`T` を次の splat が読む) なので、diagonal scheduling で 1 thread = 1 splat を保ったまま複数 splat の合成を並列化するには、diagonal の 1 iteration ごとに pixel 側 state の可視性を保証する `threadgroup_barrier` が必須になる (brush PR #401 自身も同じ制約を持つ)。「barrier を 1-2 回/batch に抑える」を文字通り実現する唯一の経路は『1 thread = 1 pixel、splat を threadgroup memory に staging』という構造だが、これは (a) register 累積を失い atomic を再導入する、(b) 既に禁忌リスト入りしている `bwd-tg` staging パターン (backward への Tg staging = busy -79%、docs/strategy/2026-08-msplat-gap-sweep.md 大原則 4) と同型になる、の 2 点で不適格。本実装は「1 thread = 1 splat + register 累積」という依頼の他の要件を優先し、barrier 回数の目標は『batch あたり 1-2 回』ではなく『barrier 総数を第 1 版比で削減する』(§3 参照) という代替目標に読み替えた。

3. 理論的な期待効果 — ただし 3 つの効果が 1 本の実験に同居する (未検証)

**barrier 総数の削減 (期待: 改善方向)**: tile 内 splat 数 N に対し、1 round あたりの `total_iters = batch + PS_TILE_SIZE - 1`、round 数 `= ceil(N / batch)` なので、tile あたりの barrier 総数の近似式は `N + (PS_TILE_SIZE - 1) * N / batch`。batch=32: `N + 255N/32 ≈ 8.97N`、batch=256: `N + 255N/256 ≈ 1.996N` — **約 4.5 分の 1 に減る**
**thread-iteration 効率の悪化 (期待: 悪化方向、advisor レビューで判明)**: diagonal の 1 round で active thread が最大化するのは定常状態の iteration だけで、立ち上がり・立ち下がり (fill/drain) の期間は active thread 数が `PS_TILE_SIZE` 未満になる。有効割合の近似は `PS_TILE_SIZE / total_iters = 256 / (batch + 255)` — batch=32: `256/287 ≈ 89.2%`、batch=256: `256/511 ≈ 50.1%`。**256-batch は fill/drain のオーバーヘッド比率が上がり、1 iteration あたりの『有効』thread 割合が ~1.78 分の 1 に下がる** (256 thread のうち常に半分弱が diagonal の端で遊んでいる計算になる)
**barrier 1 回あたりのコスト増 (期待: 悪化方向、第 1 版 finding §5 caveat (c) が予告済み)**: 32-batch は 1 SIMD group (32 thread) 内の lockstep で barrier が実質 no-op に近いのに対し、256-batch は 8 SIMD group 間の実 cross-group sync になる。個々の barrier コストが増える
**net 効果は理論では決着しない**: 「barrier 総数 4.5 分の 1」「thread-iteration 効率 1.78 分の 1 に悪化」「barrier 単価上昇」の 3 効果が 1 本の実験 (tg256 vs 第 1 版) に同居しており、bench で tg256 が第 1 版より速くなっても遅くなっても、**この 1 arm だけでは『occupancy 仮説が正しい/誤り』を単独に結論できない** — もし tg256 が負ければ「occupancy 仮説の反証」ではなく「この 3 効果の合成として負けた」としか言えない (第 1 版 finding が予告した caveat (c) そのもの)。切り分けたい場合は次段で『diagonal を batch 境界を跨いで連続させる』variant (thread lid が splat lid, lid+256, lid+512, ... を継続担当し、round 復元のたびに diagonal をリスタートしない) が必要— 同じ 256-thread occupancy を保ちつつ効率を ~80% 程度まで戻せる見込みだが、本 session のスコープ外として次段に委ねる

4. 正しさの検証状況 — 全 test GPU 実行済み、全 PASS

test	何を pin するか	GPU 実行状況
`rasterize_psos_compile`	全 PSO 構築 (`rasterize_backwards_per_splat_tg256` 含む)。`pso_min256` は `maxTotalThreadsPerThreadgroup < 256` で `Err` を返すため、レジスタ圧迫で 256-thread floor を割り込んでいないかもここで判明する	PASS — floor 達成、`Err` なし
`backward_per_splat_tg256_matches_per_pixel_original`	第 1 版と同一 fixture (32×32 / 24×24、24 splat)。`abs_tol/rel_tol = 1e-4` dual-threshold	PASS (32×32: max_abs=4.272e-4/max_rel=1.968e-5、24×24: max_abs=1.831e-4/max_rel=3.925e-6、dual-threshold で pass)
`backward_per_splat_tg256_matches_per_pixel_original_multi_round`	n_splats=600 (tile あたり 600 splat、`rounds = ceil(600/256) = 3`) で round 境界の `pix_T`/`pix_acc` 持ち越し + per-round gradient flush を確認。advisor レビューで指摘された懸念が残る: この fixture は `color_a = 0.88` で T が最初の数十 splat で `1e-4` sentinel に落ちるため、round 2/3 では『生きた』pix_T/pix_acc の持ち越しではなく『すでに死んだ pixel への no-op』を通るだけの可能性が高い — round 境界の re-init バグ (最も起きやすい種類のバグ) は検出できるが、doc comment が主張する『round を跨ぐ生きた状態遷移』の直接証明にはなっていない	PASS (max_abs=5.150e-5、max_rel=5.783e-2 — abs 側で dual-threshold pass。上記の fixture 限界の懸念はこの実測では解消されていない、次段への申し送り事項として残す)
`backward_per_splat_tg256_byte_equality_single_pixel`	単一 pixel 非ゼロ dldr で Original+FullRange+Off と bit 一致 (第 1 版と同型)	PASS
`splat bench --parity` (実 checkpoint 409,660 splats)	第 1 版と同じ 5 grad tensor (means/log_scales/rot_quats/raw_opacities/sh_coeffs) の abs/rel 1e-4 比較。これが multi-round path の本命の検証 — Lego の実 tile は密な tile で 256 splat を超えるため、上記 unit fixture より遥かに現実的な multi-round 実行を踏む	PASS OVERALL — 5 grad tensor 全て tol 内 (means max_abs_diff=1.014e-5 が最大)、loss_value diff=3.539e-8、`sorted_idx` は非 gate の既知 tie-order 非決定性 (0.222%)

GPU 実行を止めた理由 (第 1 版 session 中の coordinator 指示、参考として保持)

実装完了後、GPU 検証の直前に coordinator から「現在 12.5k gate (M2 Tier1 A1-A4) が直列で走っており、run の合間の短い隙間に bench を差し込むと gate を汚染する。以後 GPU を使う実行は行わず、実装と CPU 側 test・ビルドのみ進めよ」という指示を受け、実装 session はここで止めた。**本 finding の §4/§5b は後続の別 GPU 検証 session (2026-08-30、A7 gate 完了後の GPU 空き窓) での実測を追記したもの** — GPU 権は main から明示的に付与され、series 直列化 (A7 gate → M3b 検証) を確認した上で実行した。

5. 次段で実行すべき GPU コマンド一覧 (優先順、§5b に実測結果)

**(1) 最優先・最も情報量が高い**: `cargo test --release -p splat-metal rasterize_psos_compile` — 数秒の GPU 使用で `pso_min256` が `rasterize_backwards_per_splat_tg256` の `maxTotalThreadsPerThreadgroup >= 256` を満たすか判定する。**もしここで `Err` (occupancy floor 未達) になった場合、それ自体が M3b の結論になる** (『256-thread dispatch はレジスタ圧迫でこの kernel には実現不能』) — `pso_with_constants` に差し替えて floor を回避することは可能だが、その場合 256 thread を実際に dispatch できず本実験の前提が崩れるため、値をそのまま記録し次の判断を main に仰ぐこと
**(2) parity (unit level)**: `cargo test --release -p splat-metal backward_per_splat_tg256` (3 test: matches_per_pixel_original / matches_per_pixel_original_multi_round / byte_equality_single_pixel)
**(3) parity (trainer 統合レベル)**: `cargo build --release` 後、`./target/release/splat bench --config configs/2026-08-30-m3-bench-perpixel-final.toml --parity configs/2026-08-30-m3b-parity-persplat-tg256-candidate.toml` (第 1 版と同じ baseline config、candidate だけ tg256 に差し替え)
**(4) 4-way bench replay** (各 arm < 2 分、4 arm 合計でも 10 分未満の見込み): `./target/release/splat bench --config configs/2026-08-30-m3-bench-perpixel-final.toml` (default) / `configs/2026-08-30-m3-bench-perpixel-unopt-final.toml` (unopt) / `configs/2026-08-30-m3-bench-persplat-final.toml` (第 1 版) / `configs/2026-08-30-m3b-bench-persplat-tg256-final.toml` (第 2 版、新規)。判定は `docs/findings/p2-bench-harness-calibration.toml` の規約 `max(2σ, 1%)` — tg256 が default per_pixel を有意に下回れば GO 方向、そうでなければ NO-GO
**(5) GO の場合のみ**: gate 用 config `2026-08-30-m3b-12500-{perpixel,tg256}-r{1,2,3}.toml` (base: `configs/2026-08-30-t1-12500-on-r1.toml`、依頼文の指定どおり) を作成し、main に gate 発注を仰ぐ
**事前準備 (完了済み)**: `runs/lego-bench-calib-checkpoints/final.ply` を main repo (`/Users/otkrickey/dev/matsudalab/splat/runs/...`) からこの worktree の同パスにコピー済み (`runs/` は gitignore 対象で worktree ごとに空だったため、bench config の相対パスがそのままでは解決しなかった — advisor レビュー指摘)。`runs/bench-m3` / `runs/bench-m3b` ディレクトリも作成済み

5b. 実測結果 (2026-08-30 GPU 検証 session、main HEAD 3fc6dea 上)

binary provenance — A7 default flip 前の binary で計測

本 §5b の全計測 (PSO compile / parity 3 本 / trainer統合 parity / 4-way bench) は binary sha256 `23bf1d73dbc000515fcf09b480998f829afd5f473e99472a9777db0858b4abb7` (git rev `3fc6dea`、`refine_scratch_persist`/`refine_direct_write` の default は **まだ false**) で実行した。同一セッション後半で A7 gate 採用に伴い上記 2 flag の default を true に flip している (`p3-refine-3pass.toml` 参照) が、`splat bench` は `train_step` を直接叩く replay harness で refine tick (100 iter に 1 回) を経由しないため、この flip は本 §5b の数値に影響しない。ただし 4-way bench の 4 config (`2026-08-30-m3*-bench-*-final.toml`) はいずれも refine flag を明示しておらず、flip 後に再実行すると effective config が変わる点は再現時の注意事項として明記する。

config	backward_scheme	median ms/iter (n=3)	σ (%)	vs default per_pixel	GPU busy%
2026-08-30-m3-bench-perpixel-final (default)	PerPixel	13.548	0.59%	—	79.3%
2026-08-30-m3-bench-perpixel-unopt-final	PerPixel (unopt: Original/FullRange/Off)	15.701	0.34%	+15.9%	82.0%
2026-08-30-m3-bench-persplat-final (第 1 版)	PerSplat	18.787	0.12%	+38.7%	85.0%
2026-08-30-m3b-bench-persplat-tg256-final (第 2 版)	PerSplatTg256	17.994	0.35%	+32.8%	84.3%

判定 — NO-GO 確定。occupancy 仮説は方向として裏付けられたが不十分

tg256 (17.994ms) は第 1 版 (18.787ms) より **-4.22%** 速く、§3 が挙げた 3 効果 (barrier 総数 4.5 分の 1 に改善 / thread-iteration 効率 1.78 分の 1 に悪化 / barrier 単価上昇) の net が僅かに改善方向に振れたことを示す — occupancy 仮説そのものは方向として裏付けられた。しかし default per_pixel (13.548ms) と比べると tg256 は依然として +32.8% 遅く、`max(2σ,1%)` のどんな緩い閾値でも埋まらない差である。per-splat backward (atomic なし構造、brush PR #401 型) は 32-thread 版・256-thread 版のいずれの構成でも本実装の per_pixel (SIMD warp-level 最適化済み、atomic 使用) に対して優位性を示せなかった — M3 系列 (per-splat backward) はこの bench harness 上で NO-GO 確定とする。12.5k gate は実施していない (bench 時点で閾値を大幅に超えて劣後することが明白なため)。`BackwardSchemeKind` の default は `PerPixel` のまま、`.rs` は無変更。

次に狙うなら — diagonal を batch 境界を跨いで連続させる variant (未着手)

§3 が示唆した『diagonal を batch 境界を跨いで連続させる』variant (thread lid が splat lid, lid+256, lid+512, ... を継続担当し、round 復元のたびに diagonal をリスタートしない) は同じ 256-thread occupancy を保ちつつ thread-iteration 効率を ~80% 程度まで戻せる見込みだったが、+32.8% のギャップを埋めるには焼け石に水の可能性が高い (効率改善だけで barrier 単価増・atomic なし化のオーバーヘッドを覆すには不十分と見られる)。M3 系列への追加投資は費用対効果が低いと判断し、本 finding をもって M3 系列の探索は打ち切りとする。

6. ファイル一覧 (このセッションの変更)

kernel: `splat/shaders/backward/rasterize_backwards.metal` — `rasterize_backwards_per_splat_tg256` を追加 (第 1 版・既存 3 kernel は無変更)
Rust dispatch: `splat/crates/splat-metal/src/kernels/rasterize.rs` — `BackwardSchemeKind::PerSplatTg256`、`pso_backward_per_splat_tg256`、共通 helper `backward_per_splat_dispatch` へのリファクタ、`backward_per_splat_tg256_from_buf`、parity test 3 本 (matches/multi_round/byte_equality)
config: `splat/crates/splat-train-v1/src/config.rs` — `BackwardSchemeKind::PerSplatTg256` (flag `per_splat_tg256`)、`validate()` の no-op warning 拡張
wiring: `splat/crates/splat-cli/src/util.rs::apply_backend_config`、`splat/crates/splat-train-v1/src/trainer.rs` (2 dispatch call site を 3-way match に変更)
bench config (新規): `splat/configs/2026-08-30-m3b-bench-persplat-tg256-final.toml`、`splat/configs/2026-08-30-m3b-parity-persplat-tg256-candidate.toml`
gate config: 未作成 (§5 の (5) 参照、bench が GO の場合のみ作成する方針)
commit: `1c12660` (kernel + wiring)、`11883d2` (multi-round test 追加) — worktree `worktree-agent-a482fa3fb2a013874`。本 finding は次 commit