从 8.6 KB 到 4.2 KBSpirula Studio 3D高斯泼溅反向光栅化提速 19%【免费下载链接】spirula-studioCross-vendor 3D Gaussian Splatting trainer - video to splat to mesh, Vulkan or CUDA.项目地址: https://gitcode.com/GitHub_Trending/sp/spirula-studioSpirula Studio 是一款跨厂商的 3D 高斯泼溅3D Gaussian Splatting训练器支持视频 → Splat → 网格完整工作流Vulkan 与 CUDA 双后端。本文拆解它 2D 反向光栅化内核rasterize_bwd_2d的一次共享内存优化用 Slang 特化常量SpecializationConstant让可选功能数组随开关伸缩共享内存从 8.6 KB 降到 4.2 KB每个 SM 可驻留的 workgroup 从 11 个升到硬件上限 24 个内核耗时直降 19%。先看数字这一步优化拿到了什么 指标rasterize_bwd_2d2D 入口优化前优化后每个 workgroup 的共享内存8.6 KB4.2 KB单个 SM 可驻留 workgroup 数1124硬件上限内核耗时基线再降19%一句话讲占用率翻倍后反向光栅化在 Vulkan 后端上不再拖后腿。下面看它是怎么做到的。机制拆解groupshared 按 pipeline 分配数组必须跟着开关走 ✂️这个内核值得抠。rasterize_bwd_2d负责反向传播里 2D 屏幕坐标路径每个训练迭代都要执行一次。基准工具 raster_bench.cpp 的注释给出权重正向 反向光栅化合计占训练步 GPU 时间的~40%cap_max splats 场景省在这里每一迭代都兑现。开销藏在哪里。内核挂着一排可选功能失真校正distortion、中值深度median、累积权重accum weight等。初始 Slang 版本里这些功能的groupshared数组无论功能是否开启都按 64 个像素全量分配。CUDA 那边写的是__shared__ DStore pix_colors[dist_any(dist_type) ? BLOCK_SIZE : 1]RasterizationEval3DBwd_kernel.cuh功能关闭代价为零而 Slang 的groupshared是按 pipeline 分配不按分支分配——数组只要出现在 SPIR-V 里每条 pipeline 都照单全收功能没开也占满。怎么改。用 Slang 的[SpecializationConstant]同时充当分支条件和数组长度编译器会为每种功能组合生成专属 pipeline只分配该变体真正读到的共享内存。默认训练配置无失真、无中值、Avg 累积权重落进最小的那个变体。维度CUDANVIDIASlang / Vulkan 初始版本Slang / Vulkan 特化常量版可选功能共享数组? BLOCK_SIZE : 1关闭即零开销一律按 64 像素全量分配长度由特化常量决定关闭缩为 1分配粒度随编译期常量按 pipeline无条件全量按 pipeline但各带各的变体数组2D 入口共享内存—8.6 KB4.2 KB每 SM 驻留 workgroup—11 个24 个打满硬件上限最小改动一行特化常量约束数组长度 ✂️核心改动只有一句话把可选数组的长度写成跟随开关的特化常量而不是固定 64。// 功能关闭时数组长度为 1pipeline 只为该变体真正读到的数组分配共享内存 static const int RB_DIST_N kDistType ! 0 ? (int)RB_BLOCK : 1; static const int RB_MED_N kOutputMedian ! 0 ? (int)RB_BLOCK : 1; static const int RB_AW_N kAccumWeight ! 0 ? (int)RB_BLOCK : 1; groupshared float4 _rb_C[RB_DIST_N]; // distortion totals C groupshared float _rb_awmap[RB_AW_N]; // densify accum weights编译器为每条 pipeline 展开这些常量后关闭功能的数组缩到长度 12D 入口的共享内存随之减半。完整实现与注释在 rasterize_bwd.slang。效果归因为什么占用率翻倍就能快 19% 因果链很短特化常量把数组长度钉成每 pipeline 的编译期值 → 默认配置变体的 SPIR-V 不再携带 64 长度的失真/中值数组共享内存从 8.6 KB 掉到 4.2 KB → 共享内存是一个 SM 能塞几个 workgroup的预算预算翻倍后驻留数 11 → 24 打满硬件上限访存与原子竞争被更多 workgroup 摊薄内核整体省 19%。文档原话groupsharedis allocated per pipeline, not per branch …rasterize_bwd_2dwent 8.6 KB - 4.2 KB that way, which is 11 - 24 workgroups per SM and 19% off the kernel. —— docs/backends.md占用率起来后还叠了两笔收益把groupshared 树状归约、每层一次 barrier换成 wave 级统计 单次 barrierppl_fwd快 42%反超 CUDA 版幸存者缓冲SURV_CAP实测 2 / 4 / 8 个 chunk 分别是2.20/ 2.29 / 2.26 ms取 2 个——更深的缓冲只是推迟了同样的扫描rasterize_bwd.slang。亲手验证raster_bench 怎么跑、该看到什么 构建后直接跑基准工具固定 200 万个 splat、20 次迭代、1237×822 场景./build_develop.bash -DSS_BACKENDvulkan -DSS_BUILD_GUIOFF然后./build_vulkan/raster_bench跑完会看到四行结果raster fwd/raster bwd/ max wt/tile isect各自的 ms/call。想看这次优化的前后差距可把rasterize_bwd.slang里RB_DIST_N等三行临时改回固定 64 重新编译对比raster bwd一行——当前版本应明显更快。训练中则挂SS_PROFILE1拿每个内核的分段计时两个后端通用用法见 docs/testing.md。避坑要点 可选功能不要留一个分支就全量分配Slang 里数组长度必须由守护该功能的特化常量约束否则功能关闭时它仍是一条隐性共享内存税压住所有 pipeline 的占用率。不要在训练 run 里做 A/B训练轨迹不可复现rasterize_fwd两次运行能差 70%内核级对比用raster_bench的固定负载先用SS_PROFILE1拿内核级计时再动手。别盲目加深缓冲SURV_CAP取 2 个 chunk 就是实测最优更深的缓冲只是推迟了同样的扫描。资料清单 src/backend/vulkan/shaders/rasterize_bwd.slang —— 特化常量数组边界与 19% 优化注释所在处src/kernels/raster/RasterizationEval3DBwd_kernel.cuh —— CUDA 侧? BLOCK_SIZE : 1的对照写法docs/backends.md —— 双后端移植中四条免费午餐经验的完整总结按上面的路径构建一遍、跑一次raster_bench这些数字就都变成你自己看到的了。【免费下载链接】spirula-studioCross-vendor 3D Gaussian Splatting trainer - video to splat to mesh, Vulkan or CUDA.项目地址: https://gitcode.com/GitHub_Trending/sp/spirula-studio创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考