rasterize forward + adam kernel 監査 — forward は帯域到達率 ~9% (ALU-bound)、adam は ~90〜93% (帯域飽和済み)、bit-exact 候補はいずれも harness 分解能未満で段 2 省略

backward_raster 監査 (p3-backward-audit、M8 atomic emit lane 分散を実装するも NO-GO) に続き、per-kernel share v2 (p3-kernel-share-post-s3 §6: backward_raster 40.07% > adam 15.28% > rasterize 14.13% > ssim 13.12% > radix 12.01%) で 2・3 位の splat/shaders/forward/rasterize.metal::render_splats_f32 と splat/shaders/opt/adam.metal::adam_step を監査した。task 指定の観点 (a)〜(h) を検討し、特に (f) 「adam 12.7% が帯域床なら visible-only 以外の削減余地は理論上ない」を実際に検算することを主眼に置いた。

1. 監査 — rasterize forward (a)〜(e)

(a) tile 単位 simd_any による loop 早期打ち切り (backward W1a と同型か)

**該当なし — 構造が異なるため W1a は forward には移植できない**。backward の W1a `simd_any(contrib)` は「全 lane が寄与ゼロの iteration で 9 回の `simd_sum`+atomic emit をまとめて省略する」ためのもの — <em>省略できる高コストな操作 (reduction/atomic)</em> が前提にある。forward の per-splat 本体は reduction も atomic emit も持たない (`pix_out`/`T` は thread-local register のみ)。forward が既に持つ早期終了は 2 層: (1) per-thread `!done` 条件による loop 自体の早期脱出 (`for (t=0; !done && t<remaining; t++)`) — GPU の execution mask 収束により、simd group 内の全 lane が done になれば hardware がその simd group の残り iteration を自然にスキップする (`simd_all(done)` を明示的に追加しても原理的に redundant)。(2) `num_done_atomic` によるタイル全体 (256 thread) 単位の outer batch loop 打ち切り。**候補として成立しない (N/A)**

(b) threadgroup staging 幅と barrier 数

**bit-exact な 1 barrier 削減を発見・安全性確認済み、ただし期待 Δ は harness 分解能未満と判断**。現行は tile 開始時に barrier 2 本 (①`num_done_atomic` の 0 初期化を全 lane に可視化、②全 lane の `done` 由来 `atomic_fetch_add` を loop 先頭の `atomic_load` (rasterize.metal:117) より前に可視化) + batch ごとに barrier 2 本 (staging read の前後)。`grep -rn num_done splat/shaders/ splat/crates/` で `num_done_atomic` の読み手が rasterize.metal:117 (loop 先頭の early-break 判定) の 1 箇所のみであることを確認した — この読み出しは<strong>単調増加のカウンタに対する近似判定</strong> (early-break の可否だけに影響、`pix_out`/`T`/`stop_idx` などの出力値には一切影響しない) なので、②の barrier を削除しても最悪ケースは「その回の early-break を 1 batch 分だけ逃す」だけで <strong>正しさは不変 (bit-exact)</strong>。ただし forward は (f)/(d) と同じ byte-accounting の帰結として ALU/latency-bound (帯域到達率 ~9%、§2 参照) であり、barrier 自体が bottleneck である証拠はない — 削除できる barrier は 1 tile あたり典型 6 本中 1 本 (avg pairs/tile ≈306 → batch 数 2、2(init)+2×2(batch)=6 → 5、-16.7%) で、これが wallclock に翻訳される保証はなく harness 分解能 (~1%) を超える可能性は低いと判断した

(c) 出力 (rendered/final_T/last_contributor) の書き込み帯域

**候補として成立しない (禁忌+精度制約の二重ブロック)**。task 冒頭の禁忌リストが f16 storage/forward を明示的に禁止しており、`rendered` の quantize は不可。`final_T`/`last_contributor` は backward の勾配式 (`rasterize_backwards.metal` 冒頭の論証) が f32 精度をそのまま要求するため精度を落とせない。3 buffer 個別書き込み (24B/pixel、640,000px = 15.36MB/iter、image size 不変なので毎 iter 固定) を 1 struct にまとめても総 byte 数は不変、$2 の帰結 (forward は帯域 1 割未満) からもこの経路の削減は wallclock にほぼ効かない

(d) conic/color の load 幅 (float4 化)

**低 payoff と判断、候補から実質除外**。§2 の byte-accounting で forward の実効帯域到達率が公称 410GB/s の <strong>~8.5〜8.7%</strong> と判明した — backward audit (d) の broadcast-read 論証 (L1 が既に broadcast 済み) と同型の帰結だが、forward の staging read は per-tile-visit の gather read (broadcast ではない) という点で backward とは性質が違う。それでも「そもそも帯域の 1 割も使っていない」以上、36B→float4 aligned 化で削れる byte 数がどうであれ wallclock 改善には直結しない (律速は ALU/latency 側)。producer (`project.rs`) と consumer (forward kernel) 双方に触る高工数でもあり、期待 Δ/工数比が悪い

(e) occupancy (register 数)

**ツール限界、追加候補立てられず — backward 監査と同じ結論**。この (headless) 環境には Xcode GPU Frame Capture 相当のツールがなく、on-device register 割り当て (driver JIT) を直接観測できない (p3-backward-audit §1(e) と同じ限界)。静的に確認できたのは TG memory 使用量のみ: `local_batch[256]` (Splat2D 36B×256=9,216B) + `atomic_uint` (4B) = 9,220B、Apple GPU の TG memory 上限 (通常 32KB) の 28% — TG memory 律速の可能性は低いと推測できるが、register spill は別問題で観測手段がない。副次的に AIR level (`xcrun -sdk macosx metal -O3 -std=metal3.1 -ffast-math -S`) で `exp()` が既に `air.fast_exp.f32` (近似 intrinsic) に lowering されていることを確認した — これは `-ffast-math` を明示した本コマンドの結果だが、`splat-metal::Device::compile` は `MTLCompileOptions.fastMathEnabled` を `shader_fast_math` (device.rs、**default = true**、Metal API 自体の default でもある) から設定しており、production 既定のビルドで<strong>既に</strong> fast exp が使われている。つまり「exp を fast path に切り替える」という誰もが次に思いつく候補は<strong>既に default で実施済み</strong>であり、追加の余地はない (これも記録する価値のある closure)

2. adam の帯域床 — 自前 byte-accounting による再計算

(f) 帯域床の再計算

**再計算完了、~90〜93% 到達 (ほぼ飽和) — 旧「65.6%」は出典不明、今後は本値を正とする**。1 splat あたりの adam element 数: means 3 + log_scales 3 + rot_quats 4 + raw_opacities 1 + sh_coeffs (`max_coeffs`=16、`sh_degree=3`、`num_sh_coeffs`=`(d+1)^2`) × 3 = 48 → <strong>計 59 element/splat</strong> (`sh_progressive` (init=0/max=3/unlock_interval=1000) は active SH degree を段階解放するだけで、`adam.rs::encode_step`/`step_buf` は常に `param.max_coeffs` 全体を dispatch するため sh_coeffs の element 数はこれで不変であることを adam.rs 読解で確認済み)。1 element あたりの byte 数: read (grad 4B + m 4B + v 4B + param 4B = 16B) + write (m 4B + v 4B + param 4B = 12B) = <strong>28B/element</strong> → <strong>1,652 B/splat</strong>。<strong>平均 splat 数 321,637</strong> (p3-final-stack、M7 final-stack r2-r4 の trapezoid 積分による自前算出値、旧 322k 推定と 0.1% 差で一致) を使うと <strong>531.3 MB/iter</strong>。per-kernel share v2 (`p3-kernel-share-post-s3` §6) の `ts_adam` gpu_ms/iter (v1 1.438ms、M4/A7 前 / v2 1.400ms、M4/A7 後 — いずれも adam 自体は M4 (radix)/A7 (refine) に触れられていないので近い値のはず、という予想どおり) と組み合わせると実効帯域 <strong>369.5〜379.5 GB/s</strong>。この Mac の実機仕様 (`system_profiler`: Apple M4 Max、14-core CPU/32-core GPU、36GB — WebSearch で Apple 公式値 410GB/s と確認、上位 16-core/40-core 構成の 546GB/s とは別) を分母にすると <strong>到達率 90.1〜92.6%</strong>。**この計算には「gpu_ms/iter が 12,500 iter 全体の平均値である以上、splat 数も同じ平均 (321,637) を使わねばならず、最終 splat 数 (407,028、同じく p3-final-stack) を使ってはいけない」という罠がある** — 実際に最終 N で計算すると 672.4MB/iter ÷ 1.400ms = 480.3GB/s = <strong>117%</strong> という物理的に不可能な値が出る (これ自体が「最終 N を使うと過大評価になる」ことの直接証拠)。90〜93% という到達率は「pow hoist/f16 のような ALU 側最適化に意味のある headroom はない」という既存の closure 判断を裏づける — ただし §で述べる通り、その closure が引用してきた具体的な数値「65.6%」自体の出典は本監査では確認できなかった

旧「65.6%」の出典調査

`docs/strategy/2026-08-next-phase-tasklist.md:46` (「adam の pow hoist/f16 (帯域床 65.6% + grad 1e-14 の FTZ で数値的に閉鎖)」) と `docs/strategy/2026-09-seminar05-milestones.md:188` (「adam 13.8% は帯域床 65.6% で閉鎖済み」) が同じ 65.6% を引用しているが、`grep -rn adam docs/strategy/raw/*.json | grep -iE '帯域|bandwidth|GB/s|床'` (4 軸再監査の生ログ) では adam 固有の帯域算出は 1 件もヒットしなかった。一方 `docs/findings/p2-counter-session.toml:264` の per-kernel 表に <code>["iter5000 / per_kernel", "17.869", "65.6%", ...]</code> という行があり、この列は adam 個別ではなく<strong>その run 全体の GPU busy%</strong> (非 production な `per_kernel` submission regime) である。`2026-09-seminar05-milestones.md:19` 自身も「c1v2 run の 65.6% は per_kernel/per_stage の非 production 測定専用 regime の値であり使えない」と明記しており、これは <strong>adam の帯域到達率とは無関係な、別の run の GPU busy% を指している可能性が高い数値</strong>である。同一の数値が 2 つの異なる意味 (「ある run 全体の GPU busy%」と「adam kernel 固有の帯域到達率」) で引用されている状態であり、<strong>誤用と断定できる証拠はない (出典ログが既に失われている可能性がある) が、今後は §(f) で自前導出した 90〜93% を adam の帯域到達率として引用すべきで、65.6% を再利用しないことを推奨する</strong>

(g) buffer 数 / dispatch 数

**既に対応済み — 追加候補なし**。`adam.rs::encode_step` (W2 Stage A) が 5 param group を 1 command buffer / 1 encoder / 5 dispatch に統合済み (task 記載の「5 CB→1 済み」を `chained_step_matches_per_kernel_byte_exact` test の存在で再確認)。5 dispatch を 1 dispatch に melt する (SoA flatten + offset table) 案も検討したが、各 group の grad tensor が別 buffer object であり同一 dispatch に混在させるには indirect buffer / offset table 経由の間接参照が要る (= ArgBuffer 系、task 禁忌 !31 に抵触するリスクが高い構造) — 却下。5 dispatch 自体の launch overhead は n が数十万〜百万要素あるため相対的に無視できる規模と推測され、削っても bandwidth-bound kernel の wallclock には効かない (§(f) の帰結と整合)

(h) grad=0 splat の visible-only skip (非 bit-exact、候補のみ)

task 指定どおり **非 bit-exact のため段 2 実装対象外、候補としてのみ記録**。`2026-09-seminar05-milestones.md` M6 (J2 ユーザー判断待ち) が同じ着想を「visible-only Adam (gsplat `sparse_grad` 相当)」として既に spec 化しており、本監査はこれを追認するに留まる。§(f) の byte-accounting から言えること: adam が到達率 90〜93% で帯域飽和している以上、<strong>bytes を減らす</strong> (h) 以外に理論的な削減手段はない (pow hoist 等の ALU 側最適化は帯域の影で完全に隠れる)。期待 Δ の上限は「不可視 splat の割合 × adam share (15.28%)」— 具体的な不可視割合は本 session では計測していない (visibility mask 自体の生成コストとのトレードオフも未評価)。J2 の判断が下りるまで実装しない

3. rasterize forward の帯域到達率 — 詳細計算

staging read

`splats[sorted_idx[batch_start+lid]]` は tile ごとに visible (splat, tile) pair を 1 回ずつ 36B (Splat2D = 9 float) 読む gather read。407,267-splat checkpoint での実測 pair 数は 969,838 (`splat bench --parity` 実測、p3-radix-ballot §0 の <code>sorted_idx</code> 記述より)。avg splat 数 (321,637) / final splat 数 (407,267) ≈ 0.7898 を線形スケールした概算 avg pairs ≈ <strong>765,900</strong> (pair 数は splat の footprint に比例し splat 数に厳密比例しないため、この概算は<strong>桁のオーダーとしてのみ</strong>信頼できる — ただし後述のとおり結論の頑健性はこの誤差に依存しない)。staging read ≈ 765,900×36B = 27.57MB、加えて `sorted_idx` 自体の読み出し 765,900×4B = 3.06MB

output write

`out` (float4, 16B) + `final_T_out` (float, 4B) + `last_contrib_out` (u32, 4B) = 24B/pixel、800×800 = 640,000 pixel は image size 不変で毎 iter 固定 → 15.36MB/iter

合計と到達率

46.0MB/iter (概算) ÷ `ts_fwd_rasterize` gpu_ms/iter (v1 1.320ms / v2 1.295ms、p3-kernel-share-post-s3 §6) = <strong>34.9〜35.5 GB/s</strong> → 410GB/s 比 <strong>8.5〜8.7%</strong>。avg pairs の概算誤差がたとえ数倍あっても (仮に 10 倍過小評価だったとしても 85〜87% にしかならず帯域飽和の adam と同水準に達する可能性はゼロではないが)、969,838 という実測 pair 数 (final checkpoint) をそのまま使っても staging read は 34.9MB (÷46.0MB 全体の 76%) にしかならず、10 倍の過小評価は物理的にありえない (avg splat 数が avg pairs の下限を規定する — avg splats 321,637 で avg pairs が final の 969,838 を<strong>超える</strong>ことはまず無い) ため、<strong>「forward は ALU/latency-bound であり帯域律速ではない」という結論は概算の誤差に対して頑健</strong>と判断した

4. 候補表 (上位 5 件 + 既存の敗退/禁忌の参考枠)

5. 段 2 (実装) を省略した判断根拠

• task 指定条件: 「候補が全て <1% なら段 2 は省略し、その旨を finding に書く」— 本監査は forward 5 候補・adam 3 候補いずれについても、bit-exact-class かつ期待 Δ が harness 分解能 (~1%) を明確に超えるものを 1 件も見つけられなかった
• forward の唯一の bit-exact 候補 (barrier 1 本削除、§1(b)) は安全性を確認済みだが、forward 自体が ALU/latency-bound (帯域到達率 8.5〜8.7%、§3) という自己整合的な理由により、barrier 数を削っても wallclock 改善が harness 分解能を超える可能性は低いと判断した — 実装コストは低い (shader 2 箇所の行削除) ため、将来 harness 分解能が向上した場合や他の ALU 側施策との合算 gate で再検討する価値は残る
• adam は §2(f) の再計算で帯域到達率 90〜93% (ほぼ飽和) と確認でき、bit-exact な候補は理論上存在しない (bytes を減らす (h) は非 bit-exact で J2 判断待ち) ことを実測的に裏づけた — これは『pow hoist/f16 は閉鎖済み』という既存判断を、出典不明だった「65.6%」の代わりに自前の数値で裏づけ直したことに相当する
• GPU run は本監査を通じて 1 度も起動していない (byte-accounting は既存の per-kernel share v2 / p3-final-stack の実測値の再利用のみ、AIR 確認は `xcrun metal` のオフラインコンパイルのみ) — `pgrep -fl "target/release/splat"` の空確認は当初想定していたが、GPU run を起動しなかったため不要だった

6. 未解決点・次段への引き継ぎ

• **forward の avg pairs 概算 (765,900) は 1 点の final checkpoint (407,267 splats → 969,838 pairs) を平均 splat 数比でスケールした概算であり、pair 数は splat footprint (densification で split されるほど個々の footprint は縮む) に比例するため splat 数への線形スケールは正確ではない** — ただし §3 で示した通り、結論 (ALU-bound、帯域到達率 1 割未満) はこの誤差に対して頑健 (10 倍近い過小評価がない限り帯域飽和には届かない)。正確な avg pairs を得るには 12.5k run 全体を通した `ts_fwd_emit_pairs` の出力サイズ trajectory を記録する計装が要る (本 session では未実施)
• **forward の register spill / occupancy は本環境では直接観測できない** (backward audit §1(e) と同じ限界、Xcode GPU Frame Capture 相当のツールが無い) — GUI 環境がある session であれば `render_splats_f32` の実際の register 使用数を見て (e) を再監査する価値がある
• **adam の byte-accounting は 2 つの異なる run から数値を合成している** (avg splat 数 321,637 は M7 final-stack r2-r4、gpu_ms/iter 1.400/1.438ms は c1v2/c1 の per_kernel/per_stage 非 production run) — 両 run は同じ `[refine]` hyperparameter で走らせているため densification trajectory は近いはずだが、同一 run で両方を計測する追試 (12.5k 1 run で `ts_adam` gpu_ms trajectory と splat 数 trajectory を両方ログする) を行えば 90〜93% という到達率レンジをより狭い区間に絞れる
• **「65.6%」の出典は本監査で追跡しきれなかった** — `docs/strategy/raw/*.json` (4 軸再監査の生ログ) に adam 固有の帯域算出が見当たらず、`p2-counter-session.toml` の非 production run GPU busy% と同一数値であることのみ確認できた。誤用と断定する証拠 (例えば当該 tasklist 行を書いた session の生ログや commit メッセージ) は本 session の探索範囲では見つからず、断定は避けた — 次段で `git log -p -- docs/strategy/2026-08-next-phase-tasklist.md` などで当該行の追加 commit を辿れば出典が確定できる可能性がある
• **adam (h) visible-only skip の期待 Δ 上限が未計測** — 不可視 splat の割合 (view frustum + tile 非交差) を実際に計測すれば M6 着手時の優先度判断に使える具体的な数字が出せる (GPU run 不要、既存 checkpoint + visibility mask 計算のみで可能)

• negativebackward_raster kernel 監査 — per-pixel 構造保持の残余候補 6 件、atomic emit lane 分散 (M8) を実装するも 3/3 pair で小さい regression (NO-GO)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
• mixedM1/v2 — post-S2/S3 per-kernel share 再計測: backward_raster 33.9%→40.1% (M4/A7 後)、radix 26.4%→12.0% (M4 の絶対削減を反映) → M3 GO 判定は不変。C2 残差 4.54s は driver latency / CB gap に分解不能 (計装欠如、file:line 提案のみ)2026-08-30
• positiveM4/A5 radix_scatter simd_ballot 2段化 — 12.5k gate 採用、-14.9%、default flip2026-08-30
• negativeM5a A6 SSIM float4 化 — stage1 (float4) は byte-exact 確認済み・PSNR riskゼロ、stage2 (float4_rcp) は 12.5k gate で PSNR guard 抵触・不採用2026-08-30
• negativeM2 Tier1 (A1-A4 bit-exact) 12.5k gate v1+v2 — 棄却維持 (v1: Δ-5.0%だが2σ未達+PSNR guard抵触 / v2 n=5 radix ballot上: PSNR分離は反証、wallclockは効果量≤1%で判定不能)2026-08-30
• mixedP2 counter session — default スタックの GPU busy% は 68〜69% で確定、H1 (bubble) 単独で W 比 0.70 を説明。H3 (DVFS) は棄却、卒論の「エネルギー 0.7-0.8× で勝つ」は不成立 (GPU energy 0.90×・package energy 1.06×)2026-08-16