突破WebGPU性能瓶颈：海量粒子碰撞中的Compute Shader极限与内存对齐陷阱

在 WebGL 时代，在网页端实现数十万级别的粒子物理碰撞几乎是天方夜谭。WebGL 缺乏通用的 GPU 计算能力（GPGPU），开发者不得不利用顶点着色器和纹理读取这种“曲线救国”的方式来模拟物理。

随着 WebGPU 标准的落地，Compute Shader（计算着色器）赋予了前端开发者直接操控 GPU 通用计算的能力。然而，真正将 10 万、甚至 50 万以上的粒子碰撞检测搬上 Web 浏览器时，开发者往往会遭遇严重的性能雪崩和数据错位。这两大核心瓶颈分别对应了：GPGPU 的算法极限与硬件制约，以及 WGSL（WebGPU Shading Language）中极其严苛且隐蔽的内存对齐陷阱。

一、 内存对齐陷阱：WGSL的“隐形杀手”

在 WebGPU 中，CPU（JavaScript/TypeScript）与 GPU（WGSL）之间的数据交互通过 GPUBuffer 进行。很多开发者在初涉 WebGPU 时，会遇到渲染画面撕裂、粒子行为异常等诡异 Bug。这绝大多数是因为忽视了 WGSL 的**内存对齐（Memory Alignment）**规则。

1. 经典惨案：消失的 4 个字节

在 TypeScript 中，我们习惯将粒子属性紧凑地排布在一个 Float32Array 中。例如，一个粒子包含位置（vec3）和半径（f32），以及速度（vec3）和质量（f32）：

// JS / TS 中的紧凑排布（企图）
const particleData = new Float32Array([
    // px,  py,  pz,  radius,  vx,  vy,  vz,  mass
       0.0, 1.0, 2.0,  0.5,    10.0,11.0,12.0, 1.0
]);

在 WGSL 中，你可能会本能地定义如下结构体：

struct Particle {
    position: vec3<f32>,
    radius: f32,
    velocity: vec3<f32>,
    mass: f32,
}

当你把 Float32Array 绑定到 Storage Buffer 并启动 Compute Shader 后，你会发现 velocity 读出来的数据完全是错的。

2. 为什么会这样？WGSL 的对齐规则

WGSL 遵循统一的内存布局标准（类似 std430 布局）。在这个标准中，数据类型的对齐要求（Alignment）和大小（Size）有着严格的物理界定：

注意 vec3<f32>：它的尺寸是 12 字节，但它的对齐基数是 16 字节！

这意味着，在结构体 Particle 中：

• position 占用第 0~11 字节。
• radius（f32，对齐为4）可以直接排在后面，占用第 12~15 字节。此时累计占用 16 字节。
• 紧接着是 velocity（vec3<f32>，对齐为 16）。由于当前的偏移量是 16，恰好是 16 的倍数，所以 velocity 占用第 16~27 字节。
• mass（f32，对齐为 4）排在后面，占用第 28~31 字节。
• 致命痛点：因为 Particle 结构体中包含对齐最大为 16 的数组成员，整个结构体的大小（Size）必须是 16 字节的倍数。因此，WGSL 会在末尾自动补齐（Padding）4 字节，使结构体总长度变为 32 字节。

到目前为止，TS 排布和 WGSL 排布看似一致（都是 32 字节，即 8 个 float）。

然而，如果你把 radius 移到最前面：

struct IncorrectParticle {
    radius: f32,          // 偏移 0，占用 0~3
    position: vec3<f32>,  // 偏移必须是16的倍数！因此 4~15 被强制填充为空白 Padding！16~27 才是 position
    velocity: vec3<f32>,  // 偏移必须是16的倍数！前面到了27，所以 28~31 填充空白！32~43 是 velocity
    mass: f32,            // 占用 44~47
    // 整体对齐至16的倍数，总大小 48 字节！
}

此时，TS 的紧凑数组（32字节）与 WGSL 的实际内存结构（48字节）发生彻底错位，导致 GPU 读到的全是脏数据。

3. 最佳实践：手动填充（Padding）与 vec4 降维打击

为了规避极其繁琐的对齐偏移计算，在海量粒子系统的开发中，推荐采用显式对齐。最直接且高效的方法是将所有的 vec3 全部升级为 vec4，并利用 w 分量存储单精度标量。

修改后的结构体设计：

struct Particle {
    // x, y, z 表示位置；w 表示半径
    pos_and_radius: vec4<f32>,
    // x, y, z 表示速度；w 表示质量
    vel_and_mass: vec4<f32>,
}

对应的 TS 端数据构建极其直观，没有任何对齐陷阱，且读写效率极高：

// 每个粒子正好占用 8 个 Float32 (32字节)，无需任何补零计算
const particleData = new Float32Array([
    px, py, pz, radius,  // pos_and_radius (vec4)
    vx, vy, vz, mass     // vel_and_mass (vec4)
]);

二、 海量粒子碰撞检测的算法瓶颈

在解决了内存对齐问题后，真正的性能挑战来到了物理算法本身。如果我们要对 10 万个粒子进行碰撞检测，最暴力的 $O(N^2)$ 双重循环算法需要执行 $10^{10}$（100亿）次计算。在 60 FPS 的帧率要求下（每帧 16.6ms），即便是顶级的独立显卡也会瞬间卡死导致 GPU 驱动重置（TDR 异常）。

要在 WebGPU 中攻克这一难关，必须建立一套适合 GPU 架构的**空间分区（Spatial Partitioning）**管线。

1. 空间哈希网格（Spatial Hash Grid）的 GPU 适配

在 CPU 中，我们常用动态链表或 B 树来实现空间划分。但在 GPU 中，显存指针不连续且无法动态申请内存（没有 malloc），因此我们必须使用固定大小的数组和并行排序来实现网格。

一个完整的 WebGPU 粒子碰撞计算管线通常划分为以下 4 个 Compute Pass：

+------------------+     +-----------------------+     +-------------------+     +---------------------+
| Pass 1: 计算 Hash | --> | Pass 2: 并行排序 (Sort) | --> | Pass 3: 建立索引表 | --> | Pass 4: 局部碰撞检测 |
+------------------+     +-----------------------+     +-------------------+     +---------------------+

1. Pass 1：计算粒子所属网格并生成 Key-Value 对
   • 每个工作线程对应一个粒子。
   • 根据粒子位置算出其所在的网格索引（Cell ID），将 (Cell ID, Particle Index) 写入 Storage Buffer 中。
2. Pass 2：并行排序
   • 对生成的 Cell ID 数组进行排序。由于是在 GPU 上运行，首选 双调排序（Bitonic Merge Sort）。
   • 排序后，属于同一个网格的粒子会被连续排布在一起。
3. Pass 3：建立网格起止索引（Cell Starts / Ends）
   • 启动一个一维线程组，比较排序后数组中相邻两个元素的 Cell ID。
   • 如果第 i 和第 i-1 个元素的 Cell ID 不同，说明这里是新网格的起点，记录 CellStarts[current_id] = i，同时记录前一个网格的终点。
4. Pass 4：窄相碰撞检测（Narrow-Phase Collision）
   • 每个粒子线程根据自己当前的位置，计算周边相邻 27 个网格的 Cell ID。
   • 通过 CellStarts 和 CellEnds 快速获取这 27 个网格内所有邻近粒子的实际索引，进行碰撞解算。这样，碰撞复杂度从 $O(N^2)$ 直接暴降到接近 $O(N)$。

三、 探索WebGPU的性能极限：突破硬件约束

即便使用了空间哈希，在海量粒子场景下（如 50 万个粒子物理碰撞），我们依然会触及 WebGPU 和底层硬件的物理限制。

1. Workgroup Size（工作组尺寸）的甜点区（Sweet Spot）

在 WGSL 中，声明 Compute Shader 入口时必须指定工作组大小：

@compute @workgroup_size(64, 1, 1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) { ... }

这是一个直接关系到硬件调度（SIMD 单元执行）的关键参数。

• 硬件背景：NVIDIA GPU 的 Warp 尺寸为 32，AMD GPU 的 Wavefront 尺寸为 64。
• 极限调优：工作组大小应当始终设置为 32 或 64 的倍数。在海量粒子一维并行计算中，@workgroup_size(64, 1, 1) 或 @workgroup_size(256, 1, 1) 通常能榨干硬件的执行效率。如果将其设为 13 或 57，会导致大量的硬件执行单元空转，吞吐量骤降。

2. 避免局部存储器（Workgroup Memory）的银行冲突（Bank Conflict）

为了加速 Pass 4 中的碰撞检测，我们可以利用工作组共享变量：

var<workgroup> local_particles: array<Particle, 256>;

将频繁访问的网格内粒子数据一次性拷贝到比 Storage Buffer 快得多的共享内存（Workgroup Memory）中。但必须注意银行冲突（Bank Conflict）。

工作组内存被划分为 32 个可以并行访问的内存条（Banks）。如果工作组内的多个线程同时请求访问同一个 Bank 里的不同数据，访问就会串行化。

• 避坑手段：在多线程协作载入数据时，确保线程 local_id 与数组索引是一一对应且步长为 1 的线性映射：

  local_particles[local_id.x] = global_particles[global_id.x];
  

3. 原子操作（Atomic Operations）的锁争用

在构建网格计数器时，开发者常常会滥用原子操作：

// 极低效的设计：所有线程直接向全局计数器发起原子自增
var<storage, read_write> global_cell_count: array<atomic<u32>, total_cells>;
...
atomicAdd(&global_cell_count[cell_id], 1u);

这会导致严重的硬件冲突与等待，数十万个线程同时抢占极少数显存地址的读写锁，导致 GPU 吞吐量呈现指数级下滑。

• 优化策略：千万不要使用全局 atomic 来做多通道排序。应当通过前文提到的前缀和（Prefix Sum / Scan）算法或双调排序来实现无锁网格重构。

四、 实战：极简极速的 Compute Shader 碰撞物理模拟

下面提供一段经过内存对齐调优、基于 WebGPU 标准的粒子碰撞物理模拟 Compute Shader 核心代码。

本示例采用了简化版的网格分区/边界弹跳机制，重点演示如何安全、无对齐风险地实现粒子间的相互碰撞解算：

// 完美的 16 字节整数倍对齐粒子结构体设计
struct Particle {
    pos_and_radius: vec4<f32>, // [x, y, z] 为位置, [w] 为半径
    vel_and_mass: vec4<f32>,   // [x, y, z] 为速度, [w] 为质量
};

@group(0) @binding(0) var<storage, read> input_particles: array<Particle>;
@group(0) @binding(1) var<storage, read_write> output_particles: array<Particle>;

struct Config {
    particle_count: u32,
    dt: f32,
    damping: f32,
    gravity: f32,
};
@group(0) @binding(2) var<uniform> config: Config;

@compute @workgroup_size(256, 1, 1)
fn main(@builtin(global_invocation_id) id: vec3<u32>) {
    let index = id.x;
    if (index >= config.particle_count) {
        return;
    }

    var self_particle = input_particles[index];
    var pos = self_particle.pos_and_radius.xyz;
    let radius = self_particle.pos_and_radius.w;
    var vel = self_particle.vel_and_mass.xyz;
    let mass = self_particle.vel_and_mass.w;

    // 1. 物理环境计算（重力及边界碰撞）
    vel.y += config.gravity * config.dt;
    pos += vel * config.dt;

    // 简单边界碰撞限制 (假设包围盒为 [-100, 100])
    let bound = 100.0;
    if (pos.x - radius < -bound) { pos.x = -bound + radius; vel.x = -vel.x * config.damping; }
    if (pos.x + radius > bound)  { pos.x = bound - radius;  vel.x = -vel.x * config.damping; }
    if (pos.y - radius < -bound) { pos.y = -bound + radius; vel.y = -vel.y * config.damping; }
    if (pos.y + radius > bound)  { pos.y = bound - radius;  vel.y = -vel.y * config.damping; }
    if (pos.z - radius < -bound) { pos.z = -bound + radius; vel.z = -vel.z * config.damping; }
    if (pos.z + radius > bound)  { pos.z = bound - radius;  vel.z = -vel.z * config.damping; }

    // 2. 粒子间碰撞检测（此示例采用窄域优化查找或示范性遍历）
    for (var i = 0u; i < config.particle_count; i = i + 1u) {
        if (i == index) { continue; }
        
        let other = input_particles[i];
        let other_pos = other.pos_and_radius.xyz;
        let other_radius = other.pos_and_radius.w;
        
        let dir = pos - other_pos;
        let dist = length(dir);
        let minDist = radius + other_radius;
        
        if (dist < minDist && dist > 0.0) {
            let normal = normalize(dir);
            // 解决重叠渗透
            pos += normal * (minDist - dist) * 0.5;
            
            // 弹性碰撞冲量计算
            let relative_vel = vel - other.vel_and_mass.xyz;
            let vel_along_normal = dot(relative_vel, normal);
            
            if (vel_along_normal < 0.0) {
                let restitution = 0.8; // 弹力系数
                var impulse_scalar = -(1.0 + restitution) * vel_along_normal;
                impulse_scalar /= (1.0 / mass + 1.0 / other.vel_and_mass.w);
                
                let impulse = impulse_scalar * normal;
                vel += (1.0 / mass) * impulse;
            }
        }
    }

    // 3. 结果写回
    output_particles[index].pos_and_radius = vec4<f32>(pos, radius);
    output_particles[index].vel_and_mass = vec4<f32>(vel, mass);
}

五、 总结性能调优清单

如果你正在开发一个海量粒子的 WebGPU 项目，以下清单能帮你避开 90% 的性能暗礁：

1. 结构体检查：在 WGSL 结构体中，尽量消除 vec3，一律使用 vec4 代替。如果一定要用 vec3，请在其后手动补充一个 f32 作为 Padding。
2. 避免大循环：禁止在 Compute Shader 中直接进行 $O(N^2)$ 的暴力全遍历。使用并行双调排序和空间网格哈希将复杂度降至 $O(N)$。
3. 工作组对齐：根据目标用户的硬件基准，将 @workgroup_size 设为 64 或 256。
4. 单通道内存写回：减少 Compute Pass 之间不必要的 GPU-to-CPU 数据拷贝。将所有的物理计算、网格构建和碰撞响应完全保持在显存中（VRAM），直至渲染管线直接读取物理 Buffer。

通过彻底征服内存对齐和精心设计的 GPU 并行算法，你可以在现代浏览器中轻松流畅运行数十万个相互碰撞的沙粒、流体或物理刚体，真正释放 WebGPU 这头性能野兽的真实威力。