How 3D Gaussian Splatting got fast. 3D Gaussian Splatting 是怎么变快的

§1 The atom: a single 3D Gaussian, dissected 原子：拆开一个 3D Gaussian 看看

Forget neural networks. The fundamental unit of a 3DGS scene is a fuzzy ellipsoid:

把神经网络放一边。3DGS 场景里的基本单位是一个模糊的椭球：

Two parameters describe its shape and place: a center $\boldsymbol{\mu} \in \mathbb{R}^3$ and a $3\times 3$ covariance $\boldsymbol{\Sigma}$ that says how stretched and oriented the blob is. Plus an opacity $o \in (0,1)$ and a color $\mathbf{c}$ that depends on viewing direction (so reflections work) — more on that in a moment.

两个参数刻画它的形状与位置：一个中心 $\boldsymbol{\mu} \in \mathbb{R}^3$， 和一个$3\times 3$ 协方差 $\boldsymbol{\Sigma}$ 表示这个 blob 怎么拉伸、怎么旋。 再加一个 opacity $o \in (0,1)$ 和一个视角相关的颜色 $\mathbf{c}$（这样反射才有得玩）—— view-dependent color 那块等下再说。

Why factor $\Sigma$ into $\text{rotation} \times \text{scale}$? 为什么要把 $\Sigma$ 拆成 $\text{rotation} \times \text{scale}$？

Covariance matrices have to be symmetric positive semi-definite. If you optimize the 6 independent entries of $\boldsymbol{\Sigma}$ directly, gradient descent will happily push them into an invalid region (negative eigenvalues = imaginary ellipsoid). Kerbl et al. dodge this by storing $\Sigma$ as

Covariance matrix 必须对称正定。如果你直接优化 $\Sigma$ 的 6 个独立项，gradient descent 完全有可能把它推进 非法区域（负特征值就意味着虚数椭球，几何上不存在）。Kerbl 等人的对策是把 $\Sigma$ 写成：

where $S = \mathrm{diag}(s_x, s_y, s_z)$ is a diagonal scale and $R$ is a rotation matrix built from a unit quaternion $\mathbf{q}$. Any 3D ellipsoid can be written this way, and the parameterization is automatically valid no matter what gradient lands on $(\mathbf{q}, \mathbf{s})$. This trick — pick coordinates so the constraints are free — appears again and again in the field.

其中 $S = \mathrm{diag}(s_x, s_y, s_z)$ 是对角的 scale，$R$ 由单位四元数 $\mathbf{q}$ 构成。 任何 3D 椭球都能这么写，且不管梯度怎么砸到 $(\mathbf{q}, \mathbf{s})$ 上，参数化都结构性合法。 这种「换坐标让约束自动满足」的招数，在 3DGS 这一脉里反复出现。

The splatting projection 把椭球「拍扁」到屏幕上

To draw the ellipsoid into a 2D image, project the center $\boldsymbol{\mu}$ through the camera (standard pinhole) and the covariance too — that's the EWA splatting trick from Zwicker et al. (2001). If $W$ is the world-to-camera matrix and $J$ is the Jacobian of the perspective projection at $\boldsymbol{\mu}$, then the 2D screen-space covariance is

要把这个椭球画到 2D 图像上，先把中心 $\boldsymbol{\mu}$ 通过相机投到屏幕（标准 pinhole）， covariance 也得跟着投 —— 这就是 Zwicker 等人 2001 年 EWA splatting 的做法。 设 $W$ 是 world-to-camera，$J$ 是透视投影在 $\boldsymbol{\mu}$ 处的 Jacobian， 那 2D screen-space covariance 就是：

Drop the bottom row and right column and you get a $2\times 2$ covariance — the ellipse the 3D ellipsoid casts onto the image. Every pixel within that ellipse gets a contribution $\alpha\cdot G_{\text{2D}}(\mathbf{p})$ from this Gaussian. That's it. That's "splatting."

扔掉最下一行和最右一列就得到 $2\times 2$ 的 covariance —— 这就是 3D 椭球投影到屏幕上的那个椭圆。 椭圆覆盖到的每个像素都从这个 Gaussian 拿到 $\alpha\cdot G_{\text{2D}}(\mathbf{p})$ 的贡献。 没了。这就是「splatting」全部含义。

View-dependent color via spherical harmonics 用 SH 系数表达 view-dependent color

Same problem NeRFs solved with a view-direction input to the MLP — gloss and specularity demand $c$ depend on $\mathbf{d}$. 3DGS handles it without an MLP: store color as spherical harmonic coefficients up to degree 3 (16 per channel, 48 numbers). Evaluate at the per-Gaussian viewing direction, get an RGB. Critically this is evaluated once per Gaussian per frame in the preprocess kernel — not per pixel — so the inner compositor loop sees only a precomputed color and pays nothing extra for view-dependence.

这是 NeRF 用「把 view direction 喂进 MLP」解决的同一个问题 —— 高光、镜面反射要求 $c$ 依赖 $\mathbf{d}$。 3DGS 用 SH（spherical harmonics）绕开 MLP：每个 Gaussian 存 degree 3 以内的 SH 系数（每通道 16 个，共 48 个数）。 在 Gaussian 的视角方向上 evaluate 一下，得到 RGB。 关键是：这一步在 preprocess kernel 里每帧每 Gaussian 算一次，不是每像素一次 —— 所以内层 compositor 看到的就是已经算好的 RGB，view-dependence 没花额外算力。

§2 Stacking blobs: how a pixel gets its color 叠合 blob：一个像素的颜色是怎么算出来的

Same volume-rendering equation you've used for NeRFs, but evaluated over a sorted list of Gaussians instead of MLP samples along a ray. A pixel typically sits "under" hundreds of Gaussians; the renderer sorts them front-to-back by depth and composites:

跟 NeRF 用的是同一个 volume-rendering 方程，只是 evaluation 的对象从「沿射线的 MLP 采样」 换成了「一个排好序的 Gaussian 列表」。一个像素下面通常压着几百个 Gaussian， renderer 按 depth front-to-back 排好，再合成：

$T_i$ is the transmittance — how much light from Gaussian $i$ survives the Gaussians in front of it. Notice $T_1 = 1$ (nothing in front), and $T_i$ shrinks fast once you hit opaque Gaussians. As soon as $T$ falls below a threshold ($10^{-4}$ in the paper), no later Gaussian can contribute meaningfully — the pixel terminates early. This single observation is the backbone of the entire speed story.

$T_i$ 是 transmittance —— 第 $i$ 个 Gaussian 的光能透过前面 $i-1$ 个的比例。 $T_1 = 1$（前面没东西挡），一旦遇到 opaque 的 Gaussian，$T_i$ 就掉得飞快。 $T$ 一旦小于阈值（论文里取 $10^{-4}$），后续任何 Gaussian 的贡献都没意义 —— 这个像素就提前结束。整个 3DGS 的速度故事，骨架就是这一个观察。

§3 One frame, six CUDA kernels 一帧六个 CUDA kernel，全貌

Now we have the pieces. A trained scene has roughly 1–6 million Gaussians. To render one frame, the original diff-gaussian-rasterization CUDA pipeline does six things:

零件齐了。一个训练好的场景大约有 1–6M 个 Gaussian。要渲染一帧， 原版 diff-gaussian-rasterization 的 CUDA pipeline 做这六件事：

1. Preprocess. For every Gaussian, project $\mu$ to screen, project $\Sigma$ to 2D, decide which $16\times 16$ tiles it overlaps.
2. Duplicate-with-keys. Every Gaussian writes one (tile_id, depth) key per tile it touches into a flat array.
3. Sort. One global radix sort on those 64-bit keys.
4. Find tile ranges. Locate each tile's run in the sorted array.
5. Render. One CUDA block per tile, threads cooperate to walk front-to-back.
6. Backward. The same in reverse, with atomic gradient adds.

1. Preprocess（预处理）。 每个 Gaussian：把 $\mu$ 投到屏幕、把 $\Sigma$ 投到 2D、确定它会盖哪些 $16\times 16$ 的 tile。
2. Duplicate-with-keys。 每个 Gaussian 把它覆盖到的每个 tile 都展开成一条 (tile_id, depth) 的 key，写进一个 flat 数组。
3. Sort。 对这些 64-bit key 跑一次全局 radix sort。
4. Find tile ranges。 在排好序的数组里定位每个 tile 占的那段 [start, end)。
5. Render。 每个 tile 一个 CUDA block，里面 256 个 thread 协同 walk front-to-back。
6. Backward。 反着来一遍，梯度用 atomic add 累加。

The single most important architectural decision is the tile. Instead of asking "for each pixel, which Gaussians touch me?" (a many-to-many nightmare), the renderer asks "for each $16\times 16$ tile, which Gaussians touch me?" and a CUDA block of 256 threads handles that tile in lockstep. $16\times 16 = 256$ threads = exactly 8 warps = one perfectly-sized CUDA block.

整个架构里最关键的一个选择是 tile。不是问「每个像素，哪些 Gaussian 碰到我？」（一个多对多的灾难）， 而是问「每个 $16\times 16$ tile，哪些 Gaussian 碰到我？」 一个 256 thread 的 CUDA block 同步搞定一个 tile。 $16\times 16 = 256$ thread = 正好 8 个 warp = 一个尺寸完美的 CUDA block。

The radix-sort trick Radix-sort 的妙招

Here is the cleverest line in the paper. The duplicate-list keys are 64 bits — high 32 bits are the tile id, low 32 bits are the depth (as a float bit-pattern). One cub::DeviceRadixSort call sorts the whole thing. The result? Every tile's Gaussians are now contiguous in memory AND sorted by depth. No per-tile sort. No two-stage anything. A global sort of ~10–30 M keys runs in a few hundred microseconds on a modern GPU.

这是原论文里最聪明的一手。duplicate list 的 key 是 64 bit —— 高 32 位是 tile id， 低 32 位是 depth（浮点的 bit pattern）。一次 cub::DeviceRadixSort 把整个列表排完。 结果：每个 tile 的 Gaussian 既在内存里连续，又按 depth 排好了。 不用 per-tile sort、不用两段式什么的。 现代 GPU 上几千万 key 的全局排序也就几百微秒。

// Pack tile id into the high bits, depth into the low bits.
uint64_t key = ((uint64_t)tile_id << 32) | __float_as_uint(depth);
keys_unsorted[idx] = key;
values_unsorted[idx] = gaussian_id;

// One global sort. CUB picks bit ranges for us.
cub::DeviceRadixSort::SortPairs(
    workspace, workspace_bytes,
    keys_unsorted, keys_sorted,
    values_unsorted, values_sorted,
    n_entries);

Rendering a tile Render 单个 tile

The render kernel launches one block per tile. Inside the block, 256 threads share the work of loading Gaussians from global memory into shared memory in batches of 256; every thread (one per pixel) walks through that batch front-to-back, accumulating its own pixel's color. When every thread in the block hits $T \lt 10^{-4}$, the block exits.

render kernel 每个 tile 起一个 block。块内 256 个 thread 分摊从 global memory 把 Gaussian 加载到 shared memory 的活，每次 256 个一批； 然后每个 thread（一个 pixel 一个）front-to-back 遍历这批 Gaussian，累加自己 pixel 的颜色。 当 block 里每个 thread 都 $T \lt 10^{-4}$ 时，整个 block 退出。

// One block per tile, 256 threads (16x16 pixels)
__shared__ Gaussian batch[BATCH];          // cooperative load
float T = 1.0f, C[3] = {0,0,0};
for (int b = start; b < end; b += BATCH) {
    // cooperative load: each thread fetches one Gaussian
    batch[threadIdx] = gaussians[ sorted_ids[b + threadIdx] ];
    __syncthreads();

    for (int k = 0; k < BATCH; ++k) {
        float g = eval_2d_gaussian(batch[k], pixel);     // exp(-1/2 x^T Σ⁻¹ x)
        float a = batch[k].opacity * g;
        C[0] += batch[k].color[0] * a * T;
        C[1] += batch[k].color[1] * a * T;
        C[2] += batch[k].color[2] * a * T;
        T *= (1.0f - a);
        if (T < 1e-4f) { done = true; break; }
    }
    if (__syncthreads_count(done) == blockDim.x) break;   // whole tile finished
    __syncthreads();
}

Three details to notice. (1) Each Gaussian is fetched from slow global memory once per tile, not once per pixel — shared memory turns 256 redundant loads into one. (2) The pixel loop has zero branches except the early-out, so the warp stays coherent. (3) When most of a tile finishes early, the whole block exits with a single __syncthreads_count; CUDA's block-wide ballot is essentially free.

三个细节值得拎出来。(1) 每个 Gaussian 从慢的 global memory 取出来每个 tile 只取一次， 不是每个 pixel 一次 —— shared memory 把 256 次冗余 load 变成 1 次。 (2) pixel 循环除了 early-out 没有分支，warp 内部完全 coherent。 (3) 一个 tile 大多数 pixel 提前结束时，整个 block 通过一次 __syncthreads_count 集体退出 —— CUDA 的 block-wide ballot 基本上 0 开销。

§4 The backward pass: chain rule, by hand 反向传播：chain rule 全部手写

What makes 3DGS a learning system is that the rasterizer is differentiable. Every step above — the SH evaluation, the EWA projection, the alpha composition — has a hand-written backward kernel. PyTorch never sees the gradients; CUDA computes them directly.

3DGS 之所以是个可学习系统，关键在于 rasterizer 可微。 上面每一步 —— SH evaluation、EWA projection、alpha composition —— 都有一份手写的 backward kernel。 PyTorch 完全看不到梯度从哪儿来，CUDA 自己一路算到底。

The trick: instead of recording every intermediate value (which would cost gigabytes), the backward pass replays the forward composition in reverse, using two stored pieces of state per pixel — final $T$ and the index of the last Gaussian that contributed. From those two numbers it can reconstruct every $T_i$ on the fly via $T_i = T_{i+1} / (1 - \alpha_i)$.

招数：与其把每个中间值都存下来（要占几个 GB），backward 选择把 forward 的合成倒着重放一遍， 每个 pixel 只存两份状态 —— 最后的 $T$ 和最后一个起作用的 Gaussian 的 index。 靠这两个数，可以即时反推每一个 $T_i$：$T_i = T_{i+1} / (1 - \alpha_i)$。

The unpleasant part: many Gaussians contribute to many pixels, so the gradient updates $\partial L / \partial \mu_i$ for one Gaussian arrive from many threads at once. The original implementation uses atomicAdd on global memory — correct, but a contention nightmare. DISTWAR and gsplat's fused backward later attack exactly this (see §7.3).

讨厌的地方在于：一个 Gaussian 会被很多 pixel 用到，于是它的梯度 $\partial L / \partial \mu_i$ 会从很多 thread 同时砸过来。原版直接 atomicAdd 到 global memory，正确， 但争用得一塌糊涂。DISTWAR 和 gsplat 的 fused backward 后来攻的就是这一块（见 §7.3）。

§5 What "fast" meant in 2023 2023 年的「快」是什么概念

NeRF at the time: ~1 FPS, days of training, a few hundred MB of network weights. 3DGS was three orders of magnitude faster to render and $\sim 50\times$ faster to train. And the file you shipped was not a black-box neural network — it was a transparent point cloud you could open in a debugger.

同时代的 NeRF：~1 FPS、几天训练、几百 MB 的网络权重。3DGS 的渲染快了 三个数量级，训练快了 $\sim 50\times$。 而且交付出去的文件不是一个黑盒神经网络，而是一个可以直接用 debugger 打开看的点云。

§6 Three years of speedups, on one rail 三年间的加速史，一条 rail 串起来

The 2023 paper opened a floodgate. Below is a curated timeline of the works that pushed rendering or training efficiency the most. Click any entry for the one-line CUDA idea.

2023 那篇论文开了个闸。下面是对渲染或训练效率推动最大的工作的精选时间线。 点任何一条看「一句话讲清的 CUDA idea」。

§7 The CUDA tricks, one by one 六个 CUDA 招数，逐个拆开看

Picking from the timeline, six ideas have done the most to make 3DGS render faster without changing what 3DGS fundamentally is. Each is its own little CUDA lesson.

从时间线里挑出来，这六个 idea 在不动 3DGS 本质的前提下，对渲染加速贡献最大。 每一个都是一节小小的 CUDA 课。

7.1 · gsplat — the open rewrite开源重写版

7.2 · StopThePop — sort, per pixel按 pixel 重新排序

// Per-pixel insertion buffer of size K=4
Gaussian queue[4];
int  qlen = 0;
for (each gaussian g in tile order) {
    float d_pix = depth_along_ray(g, pixel);
    int pos = qlen;
    while (pos > 0 && queue[pos-1].d > d_pix) {
        queue[pos] = queue[pos-1]; --pos;
    }
    queue[pos] = {g, d_pix};
    if (++qlen > 4) { composite(queue[0]); shift_left(queue); --qlen; }
}

7.3 · DISTWAR — fixing the backward atomic stormtreat 反向 atomic 雪崩

// Warp-level reduction before atomicAdd
float v = local_grad_mu_x;
#pragma unroll
for (int off = 16; off > 0; off >>= 1)
    v += __shfl_xor_sync(0xffffffff, v, off);
if ((threadIdx.x & 31) == 0)            // lane 0 of each warp
    atomicAdd(&grad_mu_x[g_id], v);

7.4 · Speedy-Splat — count tiles correctly把 tile 数对

7.5 · 3DGS-MCMC — kill the densification hacks把 densification 启发式干掉

7.6 · LightGaussian & RadSplat — fewer Gaussians, free FPSGaussian 越少，FPS 越白嫖

§8 2026 and onwards 2026 之后

The pattern is consistent: every year someone notices a constant factor and removes it. Tiles got tighter, sorts got finer, atomics got aggregated, bookkeeping got smarter. The remaining big targets:

套路很稳：每年都有人盯出一个常数因子并干掉它。tile 越收越紧、sort 越做越细、atomic 越收越拢、 bookkeeping 越来越聪明。剩下还能啃的几个大方向：

• Memory bandwidth. Even with all the above, the render kernel is memory-bound. SoA layouts (struct-of-arrays) and FP16 attributes are already in flight (EAGLES, Compact-3DGS, Reduced-3DGS).
• Hardware rasterization. Several 2025 papers (Petrov et al.) show that you can emit Gaussians as tessellated triangles to the standard hardware rasterizer and let the GPU's fixed-function pipeline do the heavy lifting. Faster on consumer cards, friendlier to VR and AR.
• Dynamic and editable scenes. Spacetime Gaussians, Deformable-GS, and 4D-GS extend the primitive to time. The CUDA story is mostly "the same pipeline, indexed by $(\text{gaussian}, t)$" with new tricks for temporal coherence.
• Mobile and web. Several open-source WebGL / WebGPU implementations now hit 60 FPS on a phone for small scenes (1M Gaussians). The bottleneck shifts from arithmetic to texture-cache hit rate.

• Memory bandwidth。 即便把上面所有招数都用了，render kernel 还是 memory-bound。 SoA 布局（struct-of-arrays）和 FP16 属性已经在路上了（EAGLES, Compact-3DGS, Reduced-3DGS）。
• Hardware rasterization。 2025 年几篇论文（如 Petrov 等）表明可以把 Gaussian tessellate 成 triangle 喂给 GPU 自带的硬件 rasterizer，让 fixed-function pipeline 接手最重的活。 消费级显卡更快，对 VR/AR 也更友好。
• Dynamic / editable scene。 Spacetime Gaussians, Deformable-GS, 4D-GS 都把 primitive 拉到时间维。 CUDA 故事大体上是「同一套 pipeline，用 $(\text{gaussian}, t)$ 索引」加几个时间一致性的新招数。
• 移动端和 web。 一些开源 WebGL / WebGPU 实现已经能在手机上跑小场景（1M Gaussian）到 60 FPS。 这里的瓶颈从算术换成了 texture-cache 命中率。

What's striking is how durable the original 2023 architecture has proven. Three years and dozens of papers later, every fast renderer still has a tile loop, a sorted Gaussian list, a front-to-back walk, and an early-out. The constants change. The shape doesn't.

最让人惊讶的是 2023 年那套架构的耐用程度。三年、几十篇论文之后， 每一个跑得快的 renderer 都还是：tile loop、排好序的 Gaussian list、front-to-back walk、early-out。 常数变了，形状没变。

Reading list 扩展阅读

• Kerbl et al. — 3D Gaussian Splatting (SIGGRAPH 2023) — the foundation paper.
• Zwicker et al. — EWA Splatting (IEEE TVCG 2002). Where the 2D covariance projection comes from.
• diff-gaussian-rasterization — the original CUDA code, ~1500 lines worth reading in full.
• gsplat — the well-engineered modern base.
• awesome-3D-gaussian-splatting — community-maintained paper list.

• Kerbl 等 — 3D Gaussian Splatting (SIGGRAPH 2023) — 奠基论文。
• Zwicker 等 — EWA Splatting（IEEE TVCG 2002）。2D covariance projection 的源头。
• diff-gaussian-rasterization —— 原版 CUDA，~1500 行，值得通读。
• gsplat —— 工程化的现代 baseline。
• awesome-3D-gaussian-splatting —— 社区维护的论文列表。