WEBKT

突破性能瓶颈：基于 WebGPU Compute Shader 的 GPU 显存级空间哈希粒子碰撞算法实现

4 0 0 0

在传统的 Web 浏览器端，当粒子系统数量达到数万级别时，CPU 与 GPU 之间的带宽传输（WebMap/BufferSubData）以及 CPU 端的物理碰撞计算就会成为致命的性能瓶颈。

WebGPU 的到来彻底改变了这一现状。借助于 Compute Shader（计算着色器），我们不仅能把渲染留在 GPU，还能将复杂的物理模拟、空间划分和碰撞检测完全在 GPU 显存内（In-Place）完成。

本文将深入探讨如何利用 WebGPU 编写 WGSL（WebGPU Shading Language）计算着色器，通过**空间哈希网格（Spatial Hash Grid）**算法，实现 $O(N)$ 时间复杂度的超大规模粒子显存级碰撞检测。

一、 为什么放弃 $O(N^2)$ 与 CPU 方案？

在粒子物理模拟中，碰撞检测最直观的写法是双重循环：
$$\text{Complexity} = O(N^2)$$

当 $N = 100,000$ 时，需要进行 $10,000,000,000$（100亿）次碰撞判定。即使在 GPU 里并行展开，庞大的线程交织也会导致硬件计算单元（ALU）瞬间过载，引发 GPU 掉帧甚至驱动重置（TDR）。

GPU 空间哈希（Spatial Hashing）的救赎

空间哈希的核心思想是**“分而治之”**：

  1. 将 2D/3D 空间划分为规整的网格（Grid），网格单元（Cell）的尺寸等于或略大于粒子的最大直径。
  2. 将每个粒子映射到其所处的 Cell 中。
  3. 碰撞检测时，每个粒子只需要与自身所在 Cell 以及相邻的周边 Cell（2D 空间下为 9 个，3D 空间下为 27 个）中的粒子进行比对。

通过这种方式，算法复杂度被陡降至接近 $O(N)$。

二、 显存数据结构设计与 WGSL 对齐规范

在 WebGPU 中，Compute Shader 直接读写存储型缓冲区（Storage Buffer）。为了保证 GPU 读取效率，数据结构的设计必须遵循 WGSL 的内存对齐规则（如 std430 布局）。

1. 粒子数据结构 (Particle)

我们定义单个粒子属性，并进行手动的 16 字节对齐：

struct Particle {
    position: vec2<f32>,
    velocity: vec2<f32>,
    radius: f32,
    mass: f32,
    padding: vec2<f32>, // 保持 16 字节对齐
}

2. 空间网格数据结构 (Grid Cell)

在 GPU 上构建动态链表非常低效。一种高并发友好的方案是：固定容量桶（Fixed-capacity Bucket）。每个网格单元拥有一个原子计数器（Atomic Counter）和一个固定长度的索引数组。

struct GridCell {
    count: atomic<u32>,
    // 假设每个网格单元最多容纳 8 个粒子（取决于粒子密度和网格大小）
    particle_indices: array<u32, 8>, 
}

注意：atomic<u32> 只能存在于存储缓冲区（Storage Buffer）中，且读写时必须使用 atomicAdd 等原子操作函数。

三、 GPU 空间哈希的三阶段管线设计

高效的 GPU 碰撞检测需要分阶段调度不同的 Compute Pipeline：

[ 阶段 1: 清空网格 ] -> [ 阶段 2: 填充网格 (哈希映射) ] -> [ 阶段 3: 碰撞检测与解算 ]

Phase 1: 清空网格 (Clear Grid Shader)

在每一帧物理步进前，必须清零所有网格单元的计数器。

@group(0) @binding(0) var<storage, read_write> grid: array<GridCell>;

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    let cell_index = global_id.x;
    if (cell_index >= arrayLength(&grid)) {
        return;
    }
    // 将网格计数器安全置零
    atomicStore(&grid[cell_index].count, 0u);
}

Phase 2: 填充网格 (Grid Insertion Shader)

每个粒子计算自己的空间坐标，找到对应的网格索引，并使用 atomicAdd 将自己的粒子 ID 存入该网格的 particle_indices 中。

struct Uniforms {
    grid_size: vec2<u32>,
    cell_size: f32,
    particle_count: u32,
}

@group(0) @binding(0) var<storage, read> particles: array<Particle>;
@group(0) @binding(1) var<storage, read_write> grid: array<GridCell>;
@group(0) @binding(2) var<uniform> params: Uniforms;

// 将世界坐标转换为网格一维索引
fn get_cell_index(pos: vec2<f32>) -> u32 {
    let grid_x = u32(clamp(pos.x / params.cell_size, 0.0, f32(params.grid_size.x - 1u)));
    let grid_y = u32(clamp(pos.y / params.cell_size, 0.0, f32(params.grid_size.y - 1u)));
    return grid_x + grid_y * params.grid_size.x;
}

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    let p_idx = global_id.x;
    if (p_idx >= params.particle_count) {
        return;
    }

    let p = particles[p_idx];
    let cell_idx = get_cell_index(p.position);

    // 原子增加该 Cell 的计数器，获取当前写入位置
    let slot = atomicAdd(&grid[cell_idx].count, 1u);

    // 防止超出桶容量上限（例如 8）
    if (slot < 8u) {
        grid[cell_idx].particle_indices[slot] = p_idx;
    }
}

Phase 3: 邻域搜索与碰撞解算 (Collision Resolution Shader)

在这个阶段，每个工作线程负责一个粒子的碰撞判定。它会计算出当前粒子所在的网格，遍历周围 9 个网格，提取粒子 ID 进行窄相（Narrow-phase）碰撞检测，并计算冲量。

@group(0) @binding(0) var<storage, read_write> particles: array<Particle>;
@group(0) @binding(1) var<storage, read> grid: array<GridCell>;
@group(0) @binding(2) var<uniform> params: Uniforms;

fn resolve_collision(p1_idx: u32, p2_idx: u32, pos1: vec2<f32>, radius1: f32, mass1: f32, pos2: ptr<function, vec2<f32>>, vel2: ptr<function, vec2<f32>>, radius2: f32, mass2: f32) -> vec2<f32> {
    if (p1_idx == p2_idx) { return vec2<f32>(0.0); }
    
    let delta = *pos2 - pos1;
    let dist = length(delta);
    let min_dist = radius1 + radius2;
    
    if (dist < min_dist && dist > 0.0) {
        let normal = delta / dist;
        // 简化的弹性碰撞冲量计算
        let overlap = min_dist - dist;
        // 穿透修正（防止重叠黏连）
        *pos2 = *pos2 + normal * overlap * 0.5;
        
        // 返回施加在当前粒子上的斥力位移/速度变化量
        return -normal * overlap * 0.5;
    }
    return vec2<f32>(0.0);
}

@compute @workgroup_size(64)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    let p_idx = global_id.x;
    if (p_idx >= params.particle_count) { return; }

    var p = particles[p_idx];
    
    let center_grid_x = i32(p.position.x / params.cell_size);
    let center_grid_y = i32(p.position.y / params.cell_size);
    
    var pos_offset = vec2<f32>(0.0);

    // 遍历周边 3x3 网格
    for (var dy = -1; dy <= 1; dy++) {
        for (var dx = -1; dx <= 1; dx++) {
            let gx = clamp(center_grid_x + dx, 0, i32(params.grid_size.x - 1u));
            let gy = clamp(center_grid_y + dy, 0, i32(params.grid_size.y - 1u));
            let cell_idx = u32(gx) + u32(gy) * params.grid_size.x;
            
            let cell_count = min(atomicLoad(&grid[cell_idx].count), 8u);
            
            for (var i = 0u; i < cell_count; i = i + 1u) {
                let other_idx = grid[cell_idx].particle_indices[i];
                var other_p = particles[other_idx];
                
                pos_offset = pos_offset + resolve_collision(
                    p_idx, other_idx, 
                    p.position, p.radius, p.mass,
                    &other_p.position, &other_p.velocity, other_p.radius, other_p.mass
                );
            }
        }
    }
    
    // 应用物理更新并写回显存
    p.position = p.position + pos_offset;
    // 边界碰撞反弹
    if(p.position.x < 0.0 || p.position.x > f32(params.grid_size.x) * params.cell_size) { p.velocity.x *= -1.0; }
    if(p.position.y < 0.0 || p.position.y > f32(params.grid_size.y) * params.cell_size) { p.velocity.y *= -1.0; }
    
    particles[p_idx] = p;
}

四、 WebGPU 主程序调度（JavaScript 侧）

在 JS 侧，我们需要创建 3 个 pipeline，并在每一帧中依次调用它们。关键点在于利用 Pipeline Barrier（管线屏障） 确保显存读写顺序。

// 渲染循环中的物理更新步骤
function updatePhysics(device, commandEncoder, resources) {
  const computePass = commandEncoder.beginComputePass();
  
  // 1. 清空网格
  computePass.setPipeline(clearGridPipeline);
  computePass.setBindGroup(0, gridBindGroup);
  const gridDispatchSize = Math.ceil(GRID_CELLS_COUNT / 64);
  computePass.dispatchWorkgroups(gridDispatchSize);
  
  // 写入栅栏：确保清空后再写入（WebGPU 的 Pass 内部由显卡自动优化，但由于我们使用 storage buffer 写入，需依赖隐式屏障或显式分 Pass）
  // 在 WebGPU 中，同一个 ComputePass 内不同 pipeline 间若有读写依赖，推荐拆分为不同的 Pass
  computePass.end();
  
  // 2. 插入网格
  const insertPass = commandEncoder.beginComputePass();
  insertPass.setPipeline(insertGridPipeline);
  insertPass.setBindGroup(0, mainBindGroup);
  const particleDispatchSize = Math.ceil(PARTICLE_COUNT / 64);
  insertPass.dispatchWorkgroups(particleDispatchSize);
  insertPass.end();
  
  // 3. 碰撞检测与位置解算
  const collisionPass = commandEncoder.beginComputePass();
  collisionPass.setPipeline(collisionPipeline);
  collisionPass.setBindGroup(0, mainBindGroup);
  collisionPass.dispatchWorkgroups(particleDispatchSize);
  collisionPass.end();
}

注意：WebGPU 的 GPUEnergy（由 beginComputePass 生成）在结束时，会自动向 GPU 驱动提交依赖屏障。因此，通过拆分 Pass 的形式，我们可以非常安全地保证第一步的 Grid 清空在第二步开始前已彻底完成。

五、 性能压榨与底层调优

要想在浏览器里跑通 10 万+ 的粒子实时碰撞，光写出算法还不够，还要针对 WebGPU/GPU 硬件架构进行深度微调。

1. 规避网格原子溢出

固定桶容量（例如上面的 array<u32, 8>）会导致在粒子极度密集的地方产生溢出，即部分粒子无法被塞入网格，从而穿透。

  • 解决方案：动态调整 cell_size，使其略微大于粒子最大直径的 1.2 倍。
  • 双重保障：在 WGSL 插入代码中，使用 clamp 限制槽位，防止写爆显存边界：
    let slot = atomicAdd(&grid[cell_idx].count, 1u);
if (slot < 8u) {
    grid[cell_idx].particle_indices[slot] = p_idx;
}

2. 利用 Workgroup Shared Memory（工作组共享内存）

在碰撞阶段，频繁从 Global Memory（Storage Buffer）读取相邻网格的粒子数据会造成巨大的内存带宽消耗。

  • 进阶优化：可以利用 var<workgroup> 声明一块共享内存，让同一个工作组（Workgroup）的线程合作把临近的粒子数据一次性批量载入到片上 L1 缓存（Shared Memory），然后再从 L1 缓存中检索碰撞。这样可以减少近 70% 的全局显存带宽。

3. 使用一维紧凑数组（Prefix Sum）代替固定容量桶（非必选）

如果你追求极致的稳定性，不想承担固定容量 8 个粒子溢出的风险，可以采用更为复杂的**前缀和（Prefix Sum / GPU Scan）**算法：

  1. 统计每个 Cell 中的粒子数量。
  2. 在 GPU 上运行 Parallel Prefix Sum 计算每个网格在全局数组里的起始偏移量（Offset）。
  3. 根据偏移量将粒子重新排序，存入紧凑的一维连续显存中。
    该方案无容量上限，但增加了 GPU 的计算管线开销，适用于粒子分布极度不均（如大面积团聚）的场景。

结语

基于 WebGPU Compute Shader 实现的空间哈希碰撞算法，完全切断了 CPU 与 GPU 在物理循环中的频繁交互。粒子状态直接保存在显存的 VRAM 中，渲染管线（Vertex Shader）直接读取相同的 Storage Buffer 进行位置绘制。

这种**“计算与渲染零数据拷贝”**的架构，正是高性能 Web 3D 引擎、流体模拟（SPH）以及群落仿真（Boids）走向实用化的技术基石。

极客显卡说 WebGPU碰撞检测

评论点评