在 3D 高斯泼溅（3D Gaussian Splatting）的实时渲染系统中，CUDA 内核的性能直接决定了渲染帧率与训练速度。原始 3DGS 实现虽然实现了突破性的实时渲染，但其 CUDA 内核存在显著的性能瓶颈：全局内存争用、空间局部性差、寄存器压力过大。本文深入分析这些微观层面的优化技术，提供可落地的工程参数与监控要点。
1. 原始 3DGS CUDA 内核的性能瓶颈分析
1.1 后向光栅化的全局内存争用
根据 LiteGS 框架的分析，原始 3DGS 实现中后向光栅化（backward rasterization）占主导执行时间。主要瓶颈在于梯度累加操作：原始实现使用AtomicAdd指令直接在全局内存（global memory）上操作，导致：
- 高争用：多个线程同时访问相同内存地址，造成存储 / 加载队列饱和
- Warp 停滞：由于原子操作延迟，warp 被迫等待，平均每条指令需要 30 + 周期
- MIO/LG 节流：内存输入 / 输出（MIO）和本地 / 全局（LG）节流显著
在训练初期，单个高斯核覆盖区域较大，导致大量像素的梯度需要累加到相同内存地址，进一步放大争用问题。
1.2 空间局部性问题
原始实现的高斯参数管理策略破坏了空间局部性：
- 缓存线利用率低：NVIDIA GPU 缓存线为 128 字节，但相关数据分散存储
- L2 缓存命中率下降：随着训练进行，L2 缓存命中率从约 80% 下降至 60% 以下
- Warp 级发散：可见点与剔除点在内存中交错存储，即使单个点可见，整个 warp 仍需执行
2. 共享内存访问模式优化
2.1 多批次归约算法
LiteGS 框架引入多批次归约（multi-batch reduction）算法，优化梯度累加模式：
// 算法核心思想：在共享内存中累加梯度，减少全局内存原子操作
__shared__ float gradient[9][tilesize*tilesize];
// 每个像素涉及9个浮点梯度，分组处理减少循环开销
threads_per_property = threads_per_block / property_num;
property_id = threadIdx / threads_per_property;
ele_offset = threadIdx % threads_per_property;
// 在共享内存中执行归约
if (property_id < property_num) {
float sum = 0;
for (int i = ele_offset; i < tilesize*tilesize; i += threads_per_property) {
sum += gradient_buffer[property_id][i];
}
gradient_buffer[property_id][ele_offset] = sum;
__syncthreads();
// 最终结果写入全局内存
if (threadIdx < property_num) {
float final_sum = 0;
for (int i = 0; i < threads_per_property; i++) {
final_sum += gradient_buffer[threadIdx][i];
}
atomicAdd(&global_gradient[threadIdx], final_sum);
}
}
2.2 有效梯度压缩
仅将有效片段（非空片段）的梯度压缩到共享内存：
__shared__ int validPixNum;
__shared__ float gradient[9][tilesize*tilesize];
foreach point in VisiblePoints {
// 计算当前片段的α和透射率
float alpha = calculateAlpha(point);
float transmittance = calculateTransmittance(point);
if (transmittance * alpha > 1.0/255.0) {
int idx = atomicAdd(&validPixNum, 1);
GradientPack grad = calcGradient(point);
// 仅存储有效梯度
gradient[0][idx] = grad.ndc[0];
// ... 存储其他8个梯度分量
gradient[8][idx] = grad.alpha;
}
}
2.3 共享内存分配策略
| 参数 | 推荐值 | 说明 |
|---|---|---|
| 梯度数组维度 | [9][tilesize*tilesize] | 每个像素 9 个梯度分量 |
| 瓦片大小 | 8×8 或 16×16 | 小场景用 8×8，大场景用 16×16 |
| 共享内存使用 | ≤ 48KB / 块 | 避免超过硬件限制 |
| 线程块大小 | 256-512 线程 | 平衡占用率与寄存器压力 |
3. 线程块调度策略优化
3.1 Morton 码排序与块划分
LiteGS 采用 Morton 码（Z-order 曲线）对高斯点进行空间排序：
- 块大小：128 个高斯点作为一个处理块
- 空间一致性：Morton 码确保空间邻近的点在内存中连续存储
- 动态重排：仅在密度控制（克隆 / 分裂 / 修剪）时触发重新分块
# 简化版Morton码排序实现
def morton_sort(gaussian_points):
# 计算每个点的Morton码
morton_codes = []
for point in gaussian_points:
# 将3D坐标量化到固定网格
x_idx = quantize(point.x, grid_resolution)
y_idx = quantize(point.y, grid_resolution)
z_idx = quantize(point.z, grid_resolution)
# 交错比特生成Morton码
code = interleave_bits(x_idx, y_idx, z_idx)
morton_codes.append((code, point))
# 按Morton码排序
morton_codes.sort(key=lambda x: x[0])
# 分组为128点的块
blocks = []
for i in range(0, len(morton_codes), 128):
block = [point for _, point in morton_codes[i:i+128]]
blocks.append(block)
return blocks
3.2 集群级剔除与压缩
相比逐点剔除，LiteGS 采用集群级（cluster-level）剔除：
- AABB 包围盒：为每个 128 点块计算轴对齐包围盒
- 视锥体剔除：在集群级别执行，减少剔除操作数量
- 内存压缩：将可见点压缩到连续内存区域，提高访问局部性
3.3 线程块调度参数
| 调度参数 | 优化值 | 性能影响 |
|---|---|---|
| 块大小 | 128 点 | 平衡并行度与局部性 |
| 线程数 / 块 | 256 | 50% 占用率（RTX 3090） |
| 共享内存 / 块 | 32-40KB | 留出寄存器溢出空间 |
| 网格维度 | 动态计算 | 基于可见块数量 |
4. 寄存器压力平衡技术
4.1 CUDA 13.0 共享内存寄存器溢出
NVIDIA 在 CUDA 13.0 中引入了共享内存寄存器溢出功能：
__global__ void gaussian_splatting_kernel(...) {
// 启用共享内存寄存器溢出
asm(".pragma .enable_smem_spilling");
// 内核代码...
// 编译器会优先将溢出的寄存器存储到共享内存
// 而非本地内存（全局内存）
}
4.2 性能收益分析
启用共享内存寄存器溢出带来以下优势：
- 延迟降低：共享内存延迟约 20-30 周期，本地内存（全局内存）延迟 200 + 周期
- L2 压力减少：减少对 L2 缓存的访问压力
- 吞吐量提升：在测试案例中实现 7.76% 的内核持续时间改善
4.3 寄存器使用优化策略
| 优化技术 | 实施方法 | 预期收益 |
|---|---|---|
| 循环展开控制 | #pragma unroll 4 |
平衡寄存器使用与 ILP |
| 局部变量复用 | 重用临时变量 | 减少寄存器占用 |
| 数据打包 | 使用 float4/float2 | 提高内存吞吐量 |
| 计算融合 | 合并相关计算 | 减少中间变量 |
5. 工程实践参数与监控要点
5.1 性能监控指标
# 使用NVIDIA Nsight Systems监控的关键指标
monitoring_metrics = {
"warp_cycles_per_instruction": "目标<30周期",
"l2_cache_hit_rate": "目标>70%",
"shared_memory_utilization": "目标40-80%",
"register_pressure": "目标<64寄存器/线程",
"atomic_operation_count": "尽量减少",
"memory_throughput": "接近理论峰值"
}
5.2 参数调优清单
共享内存配置：
- 梯度缓冲区：9 × tilesize² × 4 字节
- 有效像素计数：4 字节
- 临时变量：预留 20-30% 空间
- 总使用量：控制在 48KB 以内
线程块配置：
- 块大小：128 点（空间局部性优化）
- 线程数：256（平衡占用率）
- 共享内存：32KB（留出溢出空间）
- 寄存器限制：使用
--maxrregcount 64编译选项
内存访问模式：
- 合并访问：确保线程访问连续内存地址
- 银行冲突：避免共享内存银行冲突（32 银行）
- 预取策略：对全局内存使用预取指令
5.3 调试与验证流程
- 性能分析：使用 Nsight Compute 分析内核瓶颈
- 正确性验证：与原始实现逐像素比较输出
- 回归测试：确保优化不影响渲染质量
- 压力测试：大场景、高分辨率下的稳定性
6. 实际性能数据与对比
根据 LiteGS 框架的实验结果，上述优化技术带来显著性能提升：
- 训练速度：相比原始 3DGS 实现，实现 3.4 倍加速
- 内存使用：GPU 内存使用减少约 30%
- 渲染质量：PSNR、SSIM、LPIPS 指标保持或略有提升
- 可扩展性：在 RTX 3090 和 A100 上均表现良好
结论
高斯泼溅的 CUDA 内核优化需要从微观层面解决三个核心问题：共享内存访问模式、线程块调度策略、寄存器压力平衡。通过多批次归约算法优化梯度累加、Morton 码排序维持空间局部性、共享内存寄存器溢出降低延迟，可以显著提升实时渲染性能。
工程实践中需要注意共享内存容量限制、银行冲突避免、寄存器使用平衡等技术细节。持续的性能监控与参数调优是保持优化效果的关键。随着 CUDA 技术的不断发展，特别是共享内存寄存器溢出等新特性的应用，高斯泼溅的实时渲染性能仍有进一步提升空间。
资料来源：
- LiteGS: A High-Performance Modular Framework for Gaussian Splatting Training (arXiv:2503.01199)
- How to Improve CUDA Kernel Performance with Shared Memory Register Spilling (NVIDIA Developer Blog, 2025)