Apple Silicon native MSL、kernel は最小単位で読みやすさ重視、refine / MCMC を CPU 寄りに置き unified memory zero-copy を活用、Phase H/I で算法面 (MCMC paper §3.2 full、scene-adaptive iter budget) を厚くする
本プロジェクト splat (Rust + native Metal) と brush (Rust + Burn + CubeCL + wgpu) と msplat (rayanht v1.1.3、Apple Silicon Metal-only) の 7 パイプライン軸 + メモリレイアウト + 学習ループ哲学を Pareto 直交性の観点で整理。
3 実装は Pareto 平面上で直交する 3 つのベクトルを持つ。これは 2026-06-26 の msplat M4 Max 実測で確定した。
Apple Silicon native MSL、kernel は最小単位で読みやすさ重視、refine / MCMC を CPU 寄りに置き unified memory zero-copy を活用、Phase H/I で算法面 (MCMC paper §3.2 full、scene-adaptive iter budget) を厚くする
Burn + CubeCL + WGSL、wgpu 経由で WebGPU portability を保つ、Burn の symbolic differentiation に乗っかり kernel 数を最小化、StopThePop で per-tile pruning、storage に f16 packed RGB
Apple Silicon Metal-only、42 個の巨大 fused kernel に振り切り、kernel fusion (project+SH 1-dispatch、SH backward+Adam 1-kernel、separable SSIM 1-kernel)、pre-allocated per-tile bins + bitonic sort で global sort を捨てる
ホスト言語 / GPU カーネル言語 / GPU API / Autodiff / カーネル数 / 混在精度 / MCMC paper coverage の 10 軸で並べる。
| 軸 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| ホスト言語 | Rust | Rust | C++ + Objective-C++ |
| GPU カーネル言語 | Metal Shading Language (MSL) | WGSL 手書き + CubeCL | Metal Shading Language (MSL) |
| GPU API | Metal (Apple Silicon native) | wgpu → Metal/Vulkan/DX12 (naga 変換) | Metal (Apple Silicon native) |
| Autodiff | 手書き backward kernel + CPU chain | Burn の symbolic autodiff + custom backward | 手書き backward kernel (CUDA 系統踏襲) |
| Python binding | 無し (Rust CLI) | 無し (Rust binary) | nanobind 経由 _core.so |
| Tensor 抽象 | 自前 Param / AdamState (SoA Metal buffer) | Burn Tensor<B, N> | MTensor (unified memory shared MTLBuffer wrapper) |
| カーネル数 (rasterize 系) | ~12 (project / tile_bin / fwd / bwd×3 variant / ssim / loss / adam / refine_opacity) | ~25 (FidelityFX radix sort 5 段 + prefix sum 4 段 + render 7 段 + bwd 2 段 + fused_ssim 2) | 42 (fused-heavy、CHANGELOG v1.1.3 で kernel 統合進行中) |
| クロスプラットフォーム | 不可 (Metal only) | 可 (WebGPU portability、Apple Silicon / NVIDIA / Linux) | 不可 (Apple Silicon + macOS 14+ only) |
| 混在精度 | f32 base、render_splats_f16 variant あり (wallclock 効果 ~1%) | storage f16 packed (ProjectedSplat.color_{r,g,b}: f16)、compositing は f32 | 完全 f32 only (DType に Float16 無し、fusion で速度を稼ぐ哲学) |
| MCMC paper §3.2-§3.4 | full 実装 (relocate Eq.9 + SGLD noise + scale L1 + opacity L1) | 簡約版 (clone/split 統一 + o_new=1-√(1-o) + (1-o)^150·visible noise gate) | 無し (pure Kerbl classic + 内部 fusion 最適化のみ) |
cargo workspace = 6 crates + 9 MSL kernel files (path は /Users/otkrickey/dev/matsudalab/splat/ ルート)。
| Crate | Path | 役割 | LoC |
|---|---|---|---|
splat-core | crates/splat-core/ | Pure Rust 型 + 数学 (Metal 非依存)、CameraGpu、sh.rs、scene | ~419 |
splat-metal | crates/splat-metal/ | Metal device wrapper + 7 kernel dispatcher (project / tile_bin / rasterize / backward / ssim / loss / adam) | ~431 |
splat-io | crates/splat-io/ | PLY load + NeRF synthetic dataset reader | ~634 |
splat-train-v1 | crates/splat-train-v1/ | Trainer / train_loop / refine / adam / loss / regularize / mcmc / relocate | ~3,957 |
splat-cli | crates/splat-cli/ | clap subcommand (train / check / render / hello / convert / eval)、TOML root config | ~407 |
splat-summary | crates/splat-summary/ | Post-training metrics / logging stub (Phase D 用) | ~179 |
ルート: /Users/otkrickey/dev/3dgs-workspace/brush/。WGSL kernel は #[wgsl_kernel] proc macro で Rust 型と整合性チェック、CubeCL の cube! マクロは Rust 関数を直接 GPU IR にコンパイル。
| Crate | 役割 | LoC |
|---|---|---|
brush-render | Forward: projection / culling / depth sort / SH eval / rasterization | Rust ~2,200 + WGSL ~950 |
brush-render-bwd | Backward: per-pixel grad accumulation + chain rule to splat params | Rust ~590 + WGSL ~590 |
brush-sort | FidelityFX Radix Sort 6 stage (depth sort + tile ID sort、subgroup prefix sum) | Rust ~530 + WGSL ~618 |
brush-prefix-sum | Cumulative sum 3-pass (scan + scan_sums + add_scanned) | Rust ~200 + WGSL ~85 |
brush-fused-ssim | Fused SSIM (CubeCL DSL) + 11-tap Gaussian blur + 全 SSIM 式 | Rust ~988 (cube! マクロ、shared memory 26×26 halo) |
brush-train | Learning loop: loss / backward / optimizer / refinement (prune + split + clone) | Rust ~1,941 |
brush-kernel | Common GPU utility: workgroup size calc、dynamic dispatch buffer | Rust ~120 + WGSL 39 |
brush-wgsl | proc macro: #[wgsl_kernel] codegen、naga_oil shader composition | Rust ~700 |
| Module | Path | 役割 | LoC |
|---|---|---|---|
msplat_metal.metal | core/metal/ | 42 fused Metal compute kernels (projection+SH fused / sort / rasterize / SSIM / backward / GPU-resident densify) | 3,741 MSL |
model.hpp | core/include/ | Model クラス: Gaussian params + Adam state (6 group) + densify flag | 102 C++ |
metal_tensor.hpp | core/include/ | MTensor (lightweight GPU tensor wrapper、unified memory) | 158 C++ |
model.cpp | core/src/ | Training loop orchestration / optimizer / densify / scheduler | 597 C++ |
msplat_api.mm | core/src/ | C++ API bridge (PIMPL) / Swift / Python binding / step() orchestration | 393 ObjC++ |
msplat_metal.mm | core/metal/ | Metal command encoder setup / kernel dispatch / timestamp profiling | 1,592 ObjC++ |
| loaders | core/src/loaders/ | COLMAP / Nerfstudio / Polycam / PLY | ~800 C++ |
radix_sort_histogram / radix_sort_scan / radix_sort_scatter / map_gaussian_to_intersects_kernel は metal source に残るが .mm の dispatch chain には登場せず legacy として保持。v1.1.3 の active path は scatter_to_prealloc_bins_kernel → bitonic_sort_per_tile_kernel。Adversarial verify #3 で確認済。 forward projection から backward refinement まで 7 軸に分けて、各実装の実装位置・数式・shared memory 使用量・fusion 度を表で並列展開する。
3D Gaussian (位置 + 共分散 + SH) を 2D screen-space の (xy, conic, color, alpha) に変換する forward pass。cov3d / cov2d 数式、Jacobian の frustum clamp、SH 評価との fusion、出力レイアウト、backward の cov3d 復元経路がここで分かれる。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Pass 構成 | 1-pass: SH + geometry + culling + SH evaluation を完結 | 2-pass split: project_forward (visibility / atomic compaction 専任) + project_visible (SH + packed 出力) | 1-pass + SH fused (project_and_sh_forward_kernel、cov3d は register、means3d 二度読み排除) |
| 入力レイアウト | SoA 5 buffer (means / log_scales / rot_quats / raw_opacities / sh_coeffs) | packed transforms[N*10] 1 buffer (mean / quat / log_scale concat) + 別 SH + raw_opac | SoA 別 buffer (means3d / scales / quats / features_dc / features_rest) + viewmat / projmat / intrins |
| cov3d の materialize | しない: V = J·R·S·diag(scale)、raw_cov = V·Vᵀ を直接 | しない (calc_cov2d in helpers.wgsl) | upper-triangle 6 floats を register IR 化 (forward は register、backward は SoA 拡張に開かれる) |
| Cov dilation | cov_xx += 0.3, cov_yy += 0.3 固定 | compensate_cov2d、MIP_SPLATTING 時 0.1 / 通常 0.3、filter_comp を opacity に乗ずる | cov2d.x += 0.3, cov2d.z += 0.3 を return 文に埋め込み |
| Jacobian frustum clamp | 無し (-fx·mx·rz² 素のまま、暴走は depth / finite cull 任せ) | calc_cam_J で 1.15/-0.15·img_size 非対称 lim、clamp(uv, lims_neg, lims_pos) | lim = 1.3·tan_fov の classic 3DGS 値 |
| Numerical guard | isfinite(xy) + det≤0 + alpha<1/255 の 3 段 | 3 つの中で最強: is_finite_f32 bit-level、max_abs > 1e18 → PSD-preserving scale-down、det2_strict | det == 0 cull のみ、disc = sqrt(max(0.1, b²-det)) eigenvalue safeguard |
| Tile bbox | 無し (後段 tile_bin.metal::emit_pairs で 3σ bbox emit-time 計算) | opacity-aware: compute_bbox_extent + StopThePop will_primitive_contribute per-tile contribution test | axis-aligned 3σ ceil(3·sqrt(cov2d.x/z)) (opacity 非依存) |
| 出力 | AoS Splat2D[9 float] + depth[N] 別 buffer | AoS ProjectedSplat[32B] (xy/conic/alpha f32、RGB のみ f16 packed)、atomic compaction で num_visible 個に圧縮 | 完全 SoA 多 buffer (xys / depths / radii / conics / colors / aabb / num_tiles_hit) |
| Depth gradient 受け口 | 無し | 無し (project_backwards.wgsl に "Disabled as there is no depth supervision currently" 明示) | 有り (v_depth[idx] を受け、v_mean3d += viewmat_row3 · v_z、depth supervision を将来サポート可能) |
| SH → viewdir → means chain | 有り (normalize_vjp_3d で v_pos_sh を v_pos に加算、Phase G.3 で degree≥1 完備設計) | 無し (sh_coeffs_to_color_vjp は coeffs gradient のみ) | 無し (SH gradient のみ) |
| Backward の cov3d 復元 | V 経由のみ (project_backwards.metal:402-430) | register に materialize: covar = M·Mᵀ、v_M = (v_covar + v_covarᵀ)·M | upper-triangle 6 floats を IR に書き戻す (v_cov3d[0..5]) |
3D Gaussian を tile に振り分け、depth でソートして rasterize 前段を組む。bin allocation の戦略、sort algorithm、global vs per-tile の選択がここで分かれる。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Tile size | 16×16 固定 | 16×16 固定 | 16×16 (binning) + 8×8 (rasterize sub-tile) |
| tiles_touched 計算 | per-emission atomic_fetch_add(pair_count, 1)、3σ AABB | 2 度評価 (project_forward で count、map_gaussian_to_intersects で同じ判定再走) + StopThePop will_primitive_contribute | per-tile atomic scatter_to_prealloc_bins_kernel で直接 bin に書く (count + scatter 融合) |
| Bin allocation 戦略 | フラット 1 本 (pair_capacity flat buffer、overflow で host retry) | フラット 1 本 (tile_id_from_isect[num_intersections] + compact_gid_from_isect[]) | per-tile pre-alloc (prealloc_bins[num_tiles * MAX_TILE_ELEMS=2048] uint64、overflow flag) |
| Prefix sum | 自前 MSL radix_prefix_scan (16 thread × 1 TG, simd_prefix_exclusive_sum) | 専用 crate brush-prefix-sum (3 pass, 512 threads/group, Hillis-Steele) | per-tile counter scan のみ (num_tiles 規模) |
| Sort key | u64 = (tile_id<<32) | f32_to_bits(depth) single global sort | 2 段 sort: depth (u32 32-bit) + tile (ceil(log2(num_tiles)) bit 動的) | uint64 = (depth_bits<<32) | gaussian_id per-tile に閉じた sort、global sort 無し |
| Sort algorithm | 4-bit LSD radix sort × 16 pass、自前 MSL | FidelityFX 6-stage radix sort (4-bit、subgroup prefix scan、dual-LDS intra-WG shuffle) | per-tile bitonic sort (threadgroup uint64_t data[2048] 16 KB shared memory) + sort/pack 融合 |
| Tile offsets 抽出 | per-tile binary search (tile_bin.metal:348-380) | post-sort get_tile_offsets | sort kernel が tile_bins[tile_id] = (start, end) を直接書く |
tile ごとに splat を front-to-back で α-合成。shared memory への cooperative load、early-out、backward 用 state の保存方法がここで分かれる。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Threads per tile | 16×16 = 256 thread / 1 pixel / thread | 16×16 = 256 thread / 1 pixel / thread + Morton (Z-order) pixel layout | 8×8 = 64 thread / 1 pixel / thread (RAST_BLOCK_X/Y=8) — sort tile より細かい sub-tile |
| Shared memory | ~9 KB (local_batch[256] Splat2D AoS 36B × 256) + atomic counter、f32 only on shared | 8 KB (local_batch[256] ProjectedSplat 32B × 256、RGB f16 packed) + BWD_INFO 時 +1 KB load_gid[256] | 2.25 KB (xy_opacity_batch[64] + conic_batch[64] + rgbs_batch[64]、全 f32) — 最小 |
| Pixel layout | row-major | Morton Z-order (分岐 divergence 削減) | row-major (sub-tile 8×8 = 1 workgroup) |
| Cooperative load | lid 番 thread が 256 splat 単位で詰める | 同上、workgroupUniformLoad で WGSL portability | sort kernel が既に packed buffer に展開済、64 splat batch を直接 read |
| Workgroup-wide early exit | if atomic_load(num_done_atomic) >= TILE_SIZE break | 同上、緩い同期 (stale read OK) | 無し (代わりに if (sigma >= 5.55f) continue で exp 評価を skip、94% hit empirical) |
| Per-pixel early-out | next_T <= 1e-4f → done | 同上 | 同上 + last_contributor 記録 |
| α 上限 | alpha = min(0.999, color.a * exp(-sigma)) | 同上 | 同上 (fma 多用、packed_rgb は raw SH なので max(raw + 0.5, 0.0) で DC offset+clamp) |
| Background composite | 無し (host 側 or 後段で加算) | pix_out + T*background.rgb を最終加算 | per-pixel saturate(fma(bg, T, pix_out)) で fused |
| Backward 用 state | final_T_out[pix_id] のみ | BWD_INFO mode で full out_img: vec4f + per-global-splat visible[gid] = 1.0 mask (last_contributor なし) | final_Ts[pix_id] + final_index[pix_id] = last_contributor の両方 → backward が [range.x, last_contributor] だけ replay でき ~半分削減 |
| Precision | f32 / f16 mixed の 2 variant、accumulator のみ half | storage-only f16 (ProjectedSplat.color_{r,g,b}: f16)、compositing は f32 | f32 only |
per-pixel に atomic add で勾配を書き戻すか、warp / threadgroup reduction で集約してから書くか。Inria 系の per-tile shared memory load 踏襲度、replay 方向、Refine 信号 fusion の有無が論点。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Thread parallelism | 1 thread / 1 pixel (Original / Simd / Tg 全 variant 共通) | 64 thread × 4 pixel/thread (register pressure 抑制) | 256 thread × 1 pixel + 64-splat threadgroup batch |
| Shared memory | Original / Simd は使わず、Tg variant は partial-sum reduction 用のみ | local_batch[64] + local_compact_gid[64] で splat input を cache + barrier | id / xy_opac / conic / rgb 4 種 cooperative load + simd_broadcast で lane 0 read を 32 lane に配布 |
| SIMD reduction | simd_sum で 256 atomic → 8 atomic / splat (32× 削減、Simd variant が production) | subgroupAdd で 10 scalar (xy(2)+conic(3)+rgb(3)+alpha(1)+refine(1))、subgroup_invocation_id == 0 のみ atomic | warpSum / warpSum2 / warpSum3 で 32× 削減 + warp_reduce_all_or(valid) で warp 全員 invalid なら loop 反復 skip + warp_max_finals で leading batch 全 skip |
| Atomic backend | atomic<float> native (M4 Max 前提、CAS fallback 無し) | HARD_FLOAT ifdef で native atomicAdd<f32> か atomicCompareExchangeWeak CAS fallback (WebGPU portability) | atomic_float native |
| Replay 方向 | forward 方向 replay (pix_acc 逐次更新) | forward 方向 replay (rgb_pixel_finals 経由) | 逆方向 traversal (T *= 1/(1-alpha) で reverse 復元、Inria 原典) |
| 出力レイアウト | splats_grad[N*9] (xy / conic / rgb / color_a) dense | v_splats[num_visible*10] (xy(2)+conic(3)+rgb(3)+alpha(1)+refine_weight(1)) sparse compact_gid 添字 | 4 別 buffer (v_xy[N*2] / v_conic[N*3] / v_rgb[N*C] / v_opacity[N]) SoA |
| Sigmoid backward | 別 kernel (project_backwards.metal:22a) | autodiff graph 経由 | fused (v_opacity = -v_sigma * (1 - opac) を kernel 内) |
| clamp_min RGB | 別経路 | autodiff の select(0, vis * v_outs.rgb, color.rgb >= 0) で fused | b_rgb + 0.5 < 0 で atomic skip (fused) |
| Refine 信号 fusion | CPU refine が後段で別途集計 | v_refine = length(v_xy_local * img_size) / final_a を backward kernel 内で累積 (fusion) | 別 accumulate_grad_stats_kernel で per-gaussian grad norm |
| Chunked rasterization | 無し | 無し | 有り (rasterize_backward_chunked_kernel、tile splat >512 で発火) |
| StopThePop | 不採用 | forward でのみ (per-tile bbox 段階)、backward 本体は通常 path | 不採用、しかし warp_max_finals + tile_max_bin_final の subtile-level early exit が機能的に等価 |
球面調和関数の評価。最大 degree、progressive band schedule、DC / higher band の構造分離、SH backward → Adam fusion の有無が論点。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Max SH degree | 0..=4 (25 coeffs) | 0..=4 (25 coeffs) | 0..=4 (25 bases) |
| Progressive band schedule | 有り: ShProgressiveConfig{init, max, unlock_interval}、active = init + (iter-1)/unlock_interval、init オフセット指定可 | 無し (full degree 固定、代わりに lr_coeffs_sh_scale=10 の AdamScaled per-element scaling で吸収) | 有り (s.degreesToUse = step / shDegreeInterval、init オフセット指定不可、常に 0 開始) |
| Buffer / Active 分離 | CameraGpu.{sh_degree, active_sh_degree} 2 フィールド (buffer stride vs eval loop bound) | 単一 sh_degree (buffer = active) | s.degree (buffer = featuresRest.size(-2) + 1 から逆算) + s.degreesToUse (active) |
| DC / higher band の構造分離 | 単一 sh_coeffs tensor (flat、stride = max_coeffs*3) | 単一 sh_coeffs[N,C,3] tensor、scaling tensor で per-band LR | 2 buffer 分離 (featuresDc[N,3] + featuresRest[N,C-3,3]、別 Adam state group) |
| Forward fusion | project と rasterize 別 dispatch、SH は project 内で評価 + Splat2D に packed color として書く | project_forward / project_visible / rasterize 3 段、SH は project_visible | project + SH を 1 kernel に fuse (project_and_sh_forward_kernel) |
| SH backward → Adam 経路 | SH 勾配を sh_coeffs_grad device buffer に書き → 別 dispatch adam_step | Burn autodiff graph で v_coeffs[N,C,3] dense scatter → AdamScaled optimizer の 3 step | project_and_sh_backward_kernel の register 内で SH 勾配計算 → 同 thread 内 adam_update_element を直接呼び features_dc / features_rest を in-place 更新 (~600 MB/iter at 1.5M splats 削減) |
| Adam dispatch 数 | 5 group × 1 dispatch | 3 step (transforms / sh_coeffs / raw_opacities) | 4 group dispatch (means / scales / quats / opacities)、SH 2 group は backward kernel に fused |
| viewdir → means chain | 有り (normalize_vjp_3d で v_pos_sh) | 無し ("Camera gradient is not done yet" TODO) | 無し (SH gradient のみ) |
| Numerical precision | f32 (SH eval / grad) | ProjectedSplat の color を f16 pack、SH grad は f32 | f32 only |
L1 + SSIM を組み合わせた loss と、その勾配。SSIM の separable convolution の有無、forward / backward の fusion 粒度、kernel 間の memory round-trip 排除の哲学が分かれる。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Loss config | LossKind::{L1Only, L1Ssim} + lambda、SsimConfig で window K=5/7/11 可変 | ssim_weight / lpips_loss_weight (VGG perceptual, optional) / match_alpha_weight | ssim_weight のみ、kernel ハードコード式 |
| L2 | 無し | 無し | 無し (3DGS 慣例) |
| LPIPS | 無し | 有り (VGG perceptual, optional) | 無し |
| SSIM convolution 構造 | direct 2D 11×11 (per-pixel 121 global load、separable 無し、shared memory 無し) — 純朴 baseline | separable 11+11 1D × 2 pass + 16×16 shared memory タイル化 (26×26 halo)、CubeCL cube! マクロで Rust→GPU 自動コンパイル | 2 系統共存: (a) v1.0 fused_loss_*** = direct 2D、(b) v1.1.3 separable + tiled (3 ch まとめて threadgroup load 9.75 KB) |
| Forward / backward fusion | SSIM forward と backward 別 kernel、L1 と SSIM combine も別 kernel (3 kernel) | forward が SSIM map + 3 偏微分を 1 kernel で保存、backward が tiled separable conv で完結 (2 kernel)。L1 は Burn tensor op | 最も aggressive fusion: ssim_fused_v_fwd_h_bwd_kernel 1 つで「vertical forward 再計算 + horizontal backward + L1 + threadgroup reduce + loss atomic」を完遂、中間 loss_intermediates (130 MB/iter at 1.5M splats) を排除 |
| Loss combine 数式 | dldr = (1-λ)·sign(r-t)/N - λ·ssim_grad、loss は (1-λ)·L1_mean + λ·(1-mean_ssim) | tensor op l1*(1-w) - ssim_err*w (autodiff で chain) | ssim_w·(1-ssim/3) + (1-w)·L1/3 を kernel 内ハードコード、backward も kernel 内 |
| Gaussian tap | device buffer window_1d[K] 経由 (K 可変) | const GAUSS: [f32; 11] comptime 埋め込み (σ=1.5 固定) | constant float GAUSS_1D[11] MSL ソースハードコード (σ=1.5 固定) |
| Reduction | warp 単位 simd_sum → 1 atomic / warp | 直接 backward image 書き込み、loss scalar は Burn mean() tensor op | threadgroup tg_sum[256] → 1 atomic / workgroup |
| 数値精度 | f32 only | f32 only (OOB zero-pad、SSIM raw を [-1, 1] clamp) | f32 only (dy_min/dx_min clamp で boundary 帯まで window 縮小) |
Adam tensor layout、LR schedule、densify / prune の閾値、MCMC paper coverage、opacity reset の挙動。算法面の差異が最も大きく出る軸。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Adam tensor layout | 5 component SoA (means / log_scales / rot_quats / raw_opacities / sh_coeffs の m, v) | AdamScaled: 3 group (transforms[N,10] / sh_coeffs[N,C,3] / raw_opacities[N])、per-component LR は scaling tensor の broadcast 乗算 | 6 group SoA (means / scales / quats / featuresDc / featuresRest / opacities)、ring-buffered with capacity |
| Adam dispatch 数 | 1 kernel × 5 回 dispatch (per-component LR 差し替え) | 3 step | 4 group dispatch (SH 2 group は backward kernel 内で fused) |
| Bias correction | kernel 内 pow(beta1, t) 毎 step | autodiff graph で symbolic | host precompute step_size = adam_lr[g] / bc1、bc2_sqrt = √(1-β2^t) を kernel に注入 (kernel 内 pow 不要) |
| LR schedule | ExpDecay { init, final_val, max_steps }、means のみ exp decay | ExponentialLrScheduler、lr_mean は毎 step * median_scale で scene scale 適応 | means のみ exponential interpolation、scene-invariant |
| Densify 実装場所 | CPU: read_vec/write_vec で snapshot → compact → writeback (unified memory zero-copy) | Burn tensor ops 経由 (symbolic 構築)、ただし multinomial_sample や stats reduction は host (into_data_async() readback) — ハイブリッド | 完全 GPU resident (densify_classify / append_split / append_dup / cull / compact_scatter / compact_copy_back の 6 kernel + GPU prefix sum)、CPU readback ゼロ |
| Densify schedule | 100 iter / stop 15000、capacity overflow で 2x grow | 200 iter / stop 15000 / max 10M | refineEvery && step > warmupLength && step < stopSplitAt (=maxSteps/2) |
| Split 数式 | log_scale -= ln(1.6) + 軸ごと jitter (子 2 つ Adam m/v 0、Clone は m/v 保持) | cur_log_scale + log(1/√2) + Mirror split、parent opacity o_new = 1 - √(1-o) で補正 (3DGS-MCMC Eq.9 N=2) | parent も in-place -log(1.6)、child r = randn * exp(parent_scale)、child Adam m/v は全 0 |
| Prune | sigmoid(raw_opacity) < 0.005、max_scale > cap | sigmoid(raw_opac) < 1/255、any scale < 1e-10 / > 100*bounds.extent、bounds 中心から max 超え | σ(opacity) < 0.1、split parent は always 削除、check_huge で max_scale > 0.5 |
| Relocate / MCMC | paper §3.2 Algorithm 1 full 実装 (relocate.rs:63-228): dead/live threshold 分割、opacity-weighted sample、Eq.9 o_new = 1-(1-o_t)^(1/N_t) + log_scale += -0.5*ln(N_t)、Adam state は意図的に触らない | 専用 MCMC path 無し (clone/split 統一で Eq.9 を常時適用)、noise gate は inv_opac.powi_scalar(150).clamp(0,1) * visible (paper σ(-k(o-t)) と非互換) | 無し (pure Kerbl 2023 classic) |
| MCMC noise | regularize.rs:73-142 paper gate σ(-100·(o-0.995))、eff_lr_mean = lr_mean * median_scale、visible splat の means に N(0,1)*coeff 加算 | inv_opac^150 * visible gate、noise_weight = gate * (lr_mean * mean_noise_weight)、Normal(0,1)、max_noise=median_scale clamp | 無し |
| Scale L1 / Opacity L1 | apply_mcmc_scale_l1 / apply_mcmc_opacity_l1 | 無し | 無し |
| Opacity reset | 全 raw_opacity → logit(0.01) 下方リセット (default cadence 0=無効)、Adam m/v 保持 | 明示 reset 無く、毎 refine の sigmoid-space decay (new_opac = σ(raw) - decay*(1-train_t)) で代替、Adam m/v 保持 | op[i] > logit(0.25) の上方を クランプ + Adam group[5] m/v を全 zero 破棄 |
| Opacity decay | brush 互換 sigmoid-space (refine.rs:383-402) | refine 内で常時、opac_decay * (1-train_t)、scale_decay も同様 | 無し |
| Scene-adaptive iter budget | 無し (max_steps 固定、SH progressive は degree のみ段階解放) | 無し (LOD lod_levels で post-training 段階追加、base iter 固定 30000) | 無し (resolutionSchedule で rendering 解像度 downscale、iter 数固定) |
Gaussian 表現の SoA vs AoS、Adam state group 数、tile bin storage、backward saved state、SSIM intermediate、precision policy を 6 つの表に分解。
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Layout 哲学 | SoA per-component | partial AoS (transforms に concat) | SoA per-component, DC / rest 分離 |
| means | [capacity * 3] f32 | transforms[N, 10][0..3] | means3d[N, 3] f32 |
| log_scales | [capacity * 3] f32 | transforms[N, 10][7..10] | scales[N, 3] f32 (log-space) |
| rot_quats | [capacity * 4] f32 (brush 規約 w,x,y,z) | transforms[N, 10][3..7] | quats[N, 4] f32 |
| raw_opacities | [capacity] f32 | Tensor<B, 1>[N] | opacities[N, 1] f32 |
| SH coeffs | sh_coeffs[capacity * max_coeffs * 3] 単一 tensor | sh_coeffs[N, C, 3] 単一 tensor | featuresDc[N, 3] + featuresRest[N, C-3, 3] 2 buffer 分離 |
| Adam state | 5 component pair (m, v) | 3 group + scaling tensor | 6 group |
| Total per splat (sh_deg=3) | 252 B value + 252 B grad + Adam 504 B = 1008 B | ~236 B value | ~236 B value |
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Type | Splat2D AoS 9 floats = 36 B | ProjectedSplat AoS 32 B (xy/conic/alpha f32、RGB f16 packed) | SoA 多 buffer (xys / depths / radii / conics / colors / aabb / num_tiles_hit) |
| Culled marker | conic=(0,0,0), color_a=0, depth=+∞ | atomic compaction (num_visible 個のみ書き込み) | radii=0, num_tiles_hit=0 |
| Compaction | しない (dense N) | する (global_from_compact_gid[num_visible] で indirection) | しない (dense N、num_tiles_hit=0 で skip) |
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| 構造 | flat (tile_id, splat_idx) pair buffer | flat tile_id_from_isect[num_intersections] + compact_gid_from_isect | prealloc_bins[num_tiles * 2048] uint64 (per-tile pre-alloc) |
| Sort key | u64 packed (tile_id<<32 | depth_bits) | 2 段 (depth u32 + tile_id 動的 bit) | uint64 (depth_bits<<32 | gaussian_id)、per-tile に閉じる |
| tile_offsets | per-tile binary search で抽出 | post-sort get_tile_offsets | sort kernel が tile_bins[tile_id] = (start, end) 直接出力 |
| Packed sort output | 無し (rasterize が splats[sorted_idx[i]] で間接 read) | 無し (compact_gid → projected[] で間接 read) | 有り (sort kernel が xy / opacity / conic / rgb の packed buffer を fused 出力) |
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Forward saved | final_T_out[W*H] のみ | out_img[W*H*4] + visible[N] + BWD_INFO 時 load_gid[256] | final_Ts[W*H] + final_index[W*H] = last_contributor 両方 |
| Backward replay 範囲 | range 全部 (forward 方向) | range 全部 (forward 方向) | [range.x, last_contributor] だけ (~半分削減) |
| Chunked buffer | 無し | 無し | prefix_T_buf / chunk_T_buf / after_C_buf (>512 splat tile 用) |
| 観点 | splat | brush | msplat (v1.1.3) |
|---|---|---|---|
| Forward saved | 5-tuple stats [H*W*3*5] (μx, μy, vx, vy, vxy) — non-separable | SSIM map + 3 partials | v1.0: 5 stats loss_intermediates[H*W*15] / v1.1.3: ssim_h_buf[H*W*3*5] のみ、vertical 再計算 on-the-fly |
| 削減量 | — | tiled separable で global traffic 削減 | 130 MB/iter @ 1.5M splats 削減 |
| Precision | splat | brush | msplat |
|---|---|---|---|
| Parameter | f32 | f32 | f32 |
| Adam m/v | f32 | f32 | f32 |
| SH coeffs | f32 | f32 | f32 |
| Forward intermediate | f32 (f16 variant あり) | f16 packed (ProjectedSplat.RGB) | f32 |
| SSIM | f32 | f32 | f32 |
| Backward | f32 | f32 | f32 |
ProjectedSplat) と splat の experimental render_splats_f16 variant のみ。msplat は f16 を完全に避けて fusion 一本で速度を稼ぐ。 3 実装は "学習ループのどこに最適化を効かせるか" で哲学が分かれる。
read_vec/write_vec で snapshot → compact → writeback。algorithm 反復が速い (Rust で書ける)。HARD_FLOAT ifdef で CAS fallback を持つ、subgroup intrinsics を活かすが Apple Silicon 固有最適化はしない。ProjectedSplat)、lr_mean *= median_scale の毎 step scene scale 適応、refinement bounds の percentile 0.8 推定。simd_* intrinsics、unified memory、atomic_float native add に振り切る。| 哲学軸 | splat | brush | msplat |
|---|---|---|---|
| 最適化の主戦場 | 算法 (MCMC, SH progressive, scene-adaptive) | autodiff graph + cross-platform | kernel fusion + memory layout |
| 開発速度 | 速い (Rust + 小 kernel) | 中 (Burn graph 構築) | 遅い (巨大 fused kernel の保守) |
| scene robustness | 高 (Phase I narrative) | 中 (bounds-based prune) | 低 (materials 破綻実例) |
| Apple Silicon 適合 | 高 (native + unified memory CPU 活用) | 中 (wgpu 経由) | 最高 (Metal only + native intrinsics) |
| 卒研 narrative 適合 | A 軸 (Reframing) / B 軸 (axis 2 算法) の本体 | reference 実装、parity target | direct baseline (§5 実験)、materials 破綻 PR ネタ (E 軸) |
3 つの差別化軸を確定し、卒論 §5 (実験) と §4 (関連実装) にリンクする。
「splat は kernel-level 速さでは msplat に負けるが、scene-adaptive な品質追求と全 scene robustness で勝つ」を Pareto plot で killer figure として固定。
relocate.rs:63-228 + regularize.rs:73-142)CameraGpu で表現normalize_vjp_3d 経由で閉じている (degree≥1 完備設計)「kernel-level の fusion 競争 (msplat) には参戦しない、代わりに paper coverage と scene 適応で品質を取る」。
refine.rs:105-361) — これは離散 GPU では成立しない設計into_data_async() readback)「splat は Apple Silicon の unified memory を fast CPU densify に転用、これは Metal-only でも cross-platform でも独立した第 3 のポジション」— 卒論 §4 (関連実装) の独自軸として書ける。
Phase 3 で 13 の主張を adversarial verify したうち、4 つが訂正対象 (#2, #7, #8, #11)。残りは supported。
| # | 主張 | 判定 | 修正内容 |
|---|---|---|---|
| #1 | Pass 数 / fusion 哲学 (splat 1-pass、brush 2-pass、msplat 1-pass + SH fused) | supported | msplat v1.1.3 CHANGELOG は SH backward fusion を新機能として記載、forward fusion は pre-existing。"600 MB/iter 削減" は backward SH→Adam fusion @1.5M splats の数値。 |
| #2 | cov3d materialize の三分類 (msplat=IR、splat/brush=skip) | partial | brush backward は covar = M·Mᵀ を register に materialize する (3×3 行列、global IR ではないが "skip" でもない)。正しい三分類: (a) msplat = forward/backward 両方で upper-triangle 6-float global IR、(b) brush = forward は skip / backward は register materialize、(c) splat = forward/backward 両方で真に skip (V 行列経由の陰関数 chain rule)。 |
| #3 | Tile binning & sort axis (3 実装の本質的差) | supported | msplat の radix_sort_* kernel は metal source に残るが .mm で dispatch されない dead code であることを確認。 |
| #4 | Tile size 3 実装すべて 16×16 固定 | supported | inria-3dgs 由来は kernel コメント上は証明できないが妥当。 |
| #5 | splat forward rasterize 256 thread / ~9 KB shared / f32 only on shared | supported | Splat2D 9 floats = 36 B、36×256=9 KiB ぴったり。 |
| #6 | brush 256 thread + Morton + 8 KB shared + +1 KB load_gid | supported | ProjectedSplat 32 B × 256 = 8192 B ぴったり、BWD_INFO 時のみ +1 KB。 |
| #7 | Tile parallelism / shared memory の哲学 | partial | splat の Tg variant は threadgroup memory を gradient reduction 用のみ使う (splat input data は global 直 read)。msplat の threadgroup batch は RAST_BLOCK_SIZE = 64 (256 ではない)。simd_broadcast の役割は threadgroup-memory read を 1/32、global read 圧縮の主因は cooperative load (~4×)。 |
| #8 | 3 者とも 32-lane SIMD reduction、msplat だけが warp sparse skip | partial | brush も subgroupAny(hasGrad) + if anyGrad で warp-level sparse skip を持つ。warp-level sparse skip を 持たないのは splat の Simd kernel のみ。msplat 固有は (a) leading batch 全 skip (back-to-front の利点)、(b) warp 内 broadcast による batch ロード冗長削減。 |
| #9 | Progressive band schedule (splat/msplat 有、brush 無) | supported | brush は schedule 不在を lr_coeffs_sh_scale=10 の AdamScaled per-element scaling で代替。 |
| #10 | Buffer/Active layer 分離 (splat CameraGpu 2 field、msplat state struct) | supported | msplat の "state struct" は per-iter な CamSetup s (transient)、Model 本体は shDegree のみ永続。本実装の CameraGpu も per-call 組立。構造的並列が正確に成立。 |
| #11 | Loss/SSIM の 3 軸差異 (kernel-level + precision + fusion 哲学) | partial | precision は差なし — 3 実装すべて f32 throughout。「precision の本質的差異」と書くと refuted。正しい 2 軸は (A) アルゴリズム的差異 (brush/msplat が separable + tiled、splat だけが non-separable で 11× 重い)、(B) fusion 粒度の哲学差 (splat = kernel 間 fusion、brush = kernel 内 fusion + autodiff、msplat = loss-level fusion)。 |
| #12 | splat Fused SSIM は direct 2D 11×11 per pixel + 121 global load + separable 無し + shared memory 無し純朴 baseline | supported | K=11 default、shaders/loss/ssim.metal 全体に threadgroup/shared/simd の grep 0 hit、wy_g * window_1d[dx] の outer product 形だが loop は 2D のまま。 |
| #13 | (各種小 claim) brush の AdamScaled、msplat の <code>s.degreesToUse</code>、splat の <code>ShProgressiveConfig</code> など | supported | コード一致を確認済。 |