Phase 5 #5.31.x — GPU loss kernel + no-readback pipeline (wallclock -26.9% + variance 解明)

host pump pipeline (= wgpu 抽象典型) を排除し forward → loss → backward を GPU buffer で連結。2k iter で wallclock -26.9% (42.5 → 31.07 ms/iter)、30k では 4 sample variance test で wallclock -14〜-19% が再現。第 3 軸 (Apple unified memory) narrative の核となる成果。

実施

2026-04-30

Hardware

M4 Max (36 GB unified memory)

設定

F config + L1+SSIM (λ=0.20, K=7) 2000 iter、iter 50 reset、iter 50-2000 集計

profile script

<code>3dgs-rs/scripts/phase5/run_step31_x_gpu_loss_profile.sh</code>

30k validation script

<code>3dgs-rs/scripts/phase5/run_step31_x_gpu_loss_30k.sh</code>

variance test script

<code>3dgs-rs/scripts/phase5/run_step31_x_future_variance.sh</code>

Logs

<code>3dgs-rs/runs/phase5-step31-x*/</code> (gitignored)

TL;DR (4 sample variance 解明後)

1. wallclock -26.9% (42.5 → 31.07 ms/iter @ 2k iter)、30k で wallclock -14〜-19% (4 sample 計測)
2. 第 3 軸 (Apple unified memory) narrative の核 — host pump pipeline (= wgpu 抽象典型) を排除し、forward → loss → backward を GPU buffer で連結
3. 当初の "PSNR -0.65 dB regression" 主張は撤回 — #5.31.x.future 4 sample variance test で n=1 の sample noise と判明、systematic regression ではない

設計 (advisor 5 点全反映)

構造

forward_with_state ──→ ForwardState {
                         rendered_buf: Some(Buffer),  // 新フィールド
                         rendered: Vec<f32>,           // legacy 互換、L2 経路で使用
                         ...
                       }

iter 入口 (train_step):
  target_buf = upload(target)              // 1 回 / iter
  rendered_buf ─┐
                ├─→ ssim.forward_and_grad_buf(rendered_buf, target_buf)
                │       ├─→ work_grad: Buffer (Ref で返却)
                │       └─→ mean_ssim: f32 (CPU sum、1.28 ms 残置)
  target_buf ───┘
                            │
  rendered_buf ─┐           ▼
                ├─→ loss_eval.compute_l1_combine_ssim(
  target_buf ───┤              rendered_buf, target_buf, ssim_grad_buf, λ
  ssim_grad_buf─┘            ) → (l1_loss: f32, dldr_buf: Ref<Buffer>)
                                              │
  total_loss = (1-λ)·l1 + λ·(1-mean_ssim)    │
                                              ▼
  splats_grad: Vec = backward_rasterize_buf(fs, dldr_buf)
                                              │
                       (refine 用に readback、ここは将来 GPU 化)
                                              ▼
                                     project_backwards / Adam

2 つの新 kernel (advisor 推奨で branch flag 回避)

// loss.metal
kernel void loss_l1_only(
    const device float* rendered, target,
    constant LossUniforms& u,
    device atomic<float>* loss_sum,
    device float* dldr,
    ...) {
    // dldr = sign(r-t)/N
    // L1 reduction: simd_sum + warp atomic
}

kernel void loss_l1_combine_ssim(
    const device float* rendered, target, ssim_grad,
    constant LossUniforms& u,  // λ included
    device atomic<float>* loss_sum,
    device float* dldr,
    ...) {
    // dldr = (1-λ)·sign(r-t)/N - λ·ssim_grad
    // L1 reduction 同様
}

advisor 5 点の対応

計測結果 (2000 iter F config @ Lego)

Wallclock 比較

kernel 別 timing (新 path、抜粋)

削減源の内訳

Gate 判定 (pre-commit)

• pre-commit gate: wallclock ≥ 3% 改善 (= ≤41.2 ms/iter)
• 結果: 31.07 ms/iter、-26.9% wallclock = gate 9 倍
• PASS、本実装を採用

30k validation 結果 (初回 n=1)

初期 30k validation (2026-04-30、commit 22411d2 binary):

当初は「PSNR -0.65 dB regression」と判定、speed/quality tradeoff として記録した。しかしこれは n=1 sample で短絡した結論。#5.31.x.future で 4 sample variance test を行い、この regression 主張は撤回された (次節)。

30k variance test (#5.31.x.future で regression 主張を撤回)

動機

advisor 推奨で deep numerical analysis を実施。#5.31.x.future の最初の hypothesis は「f32 reduction order が training trajectory を狂わせている」だったが、diagnostic で:

1. iter 1 dldr per-element: GPU vs HOST で bit-exact match (2,560,000 elements 全て一致)
2. iter 1 ssim_grad per-element: 同じく bit-exact match
3. logged loss scalar: GPU 3.4371e-1 vs HOST 3.4327e-1 (+0.128% systematic offset) — simd_sum + warp atomic_fetch_add vs sequential left-fold の f32 reduction order 差、ただし表示用で訓練には影響なし
4. 同 GPU path 2 回実行: iter 1-700 完全一致、iter 800 から refine 結果 divergence — rasterize_backwards_simd の atomic_fetch_add は warp 順序非決定、refine state accumulator が iter 500-1500 の間に微小差を蓄積、iter 800 の refine 判定で異なる split/clone を選択 → 軌道分岐

→ dldr に GPU/HOST 差はない、後段の atomic 非決定性が refine 累積を介して divergence を生む。これは GPU/HOST 両 path 共通の挙動、#5.31.x 固有の問題ではない仮説。

4 sample variance 計測

--force-host-loss flag を追加 (commit 330c48a)、同 binary で GPU/HOST 切替 A/B が可能に。4 つの 30k validation を実施 (全て同 seed 42、同 F config):

• range: 26.27 - 24.86 = 1.41 dB
• mean: 25.43 dB
• stddev: ~0.61 dB

重要な観察

観察 1: 同 binary 330c48a 内で GPU > HOST (sample #4 vs #3) — HOST: 24.855 dB / GPU: 25.467 dB (+0.612 dB) → 当初の「GPU が host より劣る」という命題は反証された。

観察 2: 同 path (host) 別 binary で 1.41 dB 差 (sample #1 vs #3) — e2611df host: 26.27 dB / 330c48a host: 24.86 dB → 同じコード経路でも binary 違いで 1.41 dB のばらつき、atomic 非決定性が compile/scheduling を介して影響。

観察 3: GPU path の wallclock は consistently 速い — queue-reuse host (#1): 1592s (baseline) / GPU r1 (#2): 1297s (-19%) / HOST r1 (#3): 1548s (-3%、queue reuse 改善のみ) / GPU r2 (#4): 1371s (-14%) → GPU path は HOST path より 14-19% wallclock 短縮 (4 sample で再現)。

結論 (撤回 + 修正)

M-3 採用判断

#5.31.0b queue reuse は M-3 確定: 30k で 26.27 dB を出した sample (#1) は variance band の上端だが、同 path 別 binary でも 24.86 dB (#3) なので、期待 PSNR は ~25.4 dB ± 0.6 dB。queue reuse 自体の wallclock -13.7% は本質的、コード変更は健在。

#5.31.x GPU loss path は M-3 候補に再昇格 (mixed 撤回): PSNR は HOST path と同等 (variance 内、むしろ +0.61 dB 高い sample もある)、wallclock -14〜-19% (4 sample で再現)。ただし single sample の絶対 PSNR を milestone と呼ぶのは fragile、mean ± std での記述に切替。→ M-3.x として記録、stat-based wallclock 改善を主張。

卒論記述案 (honest version)

narrative 価値 (卒論章)

第 3 軸 (Apple unified memory) — 直撃の貢献

• brush (wgpu 経由): forward output → host Vec → upload to ssim → host Vec → upload to backward (=4 回 host pump)
• 本実装 (Metal 直接 + StorageModeShared): forward → rendered_buf → ssim → ssim_grad_buf + loss → dldr_buf → backward (= 0 回 host pump)
• 定量: -11.4 ms/iter、wallclock -26.9%、Lego F config 2000 iter
• narrative: 「unified memory が概念ではなく、実装の具体的な書き方で wallclock 性能に反映される」

第 2 軸 (wgpu 抽象 vs 直接 Metal) — 補強

• wgpu の制約 (host vec 経由前提) なしで、Metal の MTLBuffer を直接設計に組み込んだ
• ArgBuffer 単独 (#5.31 original) では 0.5% 上限だったところ、no-readback で 26.9%
• → 「抽象コスト」は API 命令の overhead より、API が前提とするデータフローパターンにある

第 1 軸 (native 実装) — 完成度補強

• M-1 (Phase 4 後続 25.59 dB) と並ぶ性能成果
• kernel level (M-2 SIMD reduction 2.43×) と pipeline level (M-3 GPU loss -26.9%) の両方を分離して評価できる

卒論への活用

本 step は mixed result で、卒論には 「速度の big win + PSNR regression という二面性」 の事例として収録 (※ 後の 4 sample variance test で PSNR 主張は撤回):

• 2k iter wallclock -26.9% (gate 9 倍): pipeline-level 最適化の有効性証明
• 30k PSNR -0.65 dB: f32 numerical drift が optimization trajectory に compound する事例 (※撤回済み)
• 第 3 軸 (unified memory) narrative の 設計的価値は維持 (host pump 排除は実装上正しい)
• 「速度の改善が必ずしも PSNR を保つわけではない、特に学習系では数値再現性が重要」という卒論主張

F (scale_reg 失敗) / G (MCMC 失敗) / #5.31 (ArgBuffer reject) の negative findings と、#5.31.0b (queue reuse +0.48 dB の純 win) の対比で「測定して進める研究の typology」が立つ:

#5.31 ArgBuffer の reject (0.5% 上限) と、#5.31.x の partial win (-22.8% wallclock、PSNR -0.65) を併せて「ベンチマーク前提の自己検証 → bottleneck 同定 → 想定外の上振れ → 30k validation で PSNR 副作用検出」 という方法論が立つ:

1. step 30 推定: ~30% が CPU encoding
2. step 31.0a 計測: encoder 部分は 1.4%、ArgBuffer reject
3. step 31.0b queue reuse: 13.7% wallclock (encoding ではなく queue 生成 + autoreleasepool)
4. step 31.0c 内訳計測: 真の bottleneck = host loss + readback (~10 ms)
5. step 31.x GPU loss + no-readback: -26.9% wallclock

= 「測ってから最適化し、何度も前提を更新した」5 段階の実証

未着手 (#5.31.x.後続 / future work)

合計潜在: 残り ~5 ms/iter、追加 -10〜15% wallclock 余地あり。本研究では #5.31.x の単独 win を採用、上記は future work 化。

この finding が観測された / 言及している実験 run。

• lego-sh3-30kA.8 SH degree ablation — sh_degree=3 (DC only, no view-dependence)PSNR 24.88 dBWallclock 23m 13sverdict partial

• mixedPhase C migration gate — splat workspace で M-3.x を再現する 30k 単発比較2026-04-30
• mixedPhase 5 step 30 — Instruments / Metal System Trace 分析結果2026-04-30
• positivePhase 5 step 30b — kernel-by-kernel timing 計測結果2026-04-30
• mixedPhase 5 #5.31 — CPU side profile + queue reuse 効果 + 真の bottleneck 同定2026-04-30