WebGPU超大规模物理仿真：突破Storage Buffer绑定上限的架构设计与实战

在利用 WebGPU 进行大规模物理仿真（如百万级粒子群、SPH 流体、高分辨率软体物理等）时，开发者几乎不可避免地会撞上一堵墙：Storage Buffer 的绑定尺寸上限。

即便你在请求设备（requestDevice）时申请了最高的 requiredLimits，WebGPU 在不同平台和硬件上对单个 Storage Buffer 的最大绑定尺寸（maxStorageBufferBindingSize）依然有着严格的限制。在很多移动端设备或集显上，这个上限可能只有 128MB，即便在高端独显上，也通常被限制在 1GB 到 2GB 左右。

对于动辄需要存储粒子位置、速度、密度、受力、邻域网格以及历史状态的物理仿真系统来说，单一扁平化的 Buffer 结构很快就会撑爆这个限制。本文将深入探讨如何在 WebGPU 中优雅地绕过这一硬件红线，实现更大规模的物理仿真。

一、 核心痛点解析：我们在和什么限制打交道？

在 WebGPU 中，内存限制主要由以下三个指标决定：

1. maxBufferSize：设备允许创建的单个 GPUBuffer 的绝对最大尺寸。
2. maxStorageBufferBindingSize：单个绑定点（Binding）能够绑定的 Storage Buffer 的最大尺寸。通常这个值会远小于 maxBufferSize。
3. minStorageBufferOffsetAlignment：Storage Buffer 绑定的偏移量对齐要求（通常为 256 字节），在做动态偏移绑定时必须遵守。

// 查询当前设备的极限值
const limits = adapter.limits;
console.log("最大Buffer大小:", limits.maxBufferSize);
console.log("单次绑定Storage Buffer最大大小:", limits.maxStorageBufferBindingSize);

当我们的物理粒子数量达到千万级，或者网格体素（Voxel）分辨率达到 $256^3$ 以上时，单个大数组的尺寸会轻松突破 maxStorageBufferBindingSize。如果强行绑定，WebGPU 会在创建 BindGroup 时直接抛出 Validation Error。

二、 方案一：分片缓冲区（Buffer Chunking）与多 BindGroup 路由

既然单个绑定点有尺寸限制，最直接的想法就是化整为零：将数据拆分到多个物理 GPUBuffer 中，并将它们绑定到不同的 Binding 点上，或者在运行时通过动态偏移（Dynamic Offset）进行分片处理。

1. WGSL 中的多 Buffer 路由设计

由于 WGSL 目前不支持在着色器中声明动态大小的绑定数组（例如 binding(0) var<storage> buffers: array<binding_type> 是非法的），我们必须在 Shader 中静态声明多个绑定，并在代码逻辑中通过索引进行路由。

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

// 拆分成 4 个独立的 Storage Buffer 段
@group(0) @binding(0) var<storage, read_write> particles_chunk0: array<Particle>;
@group(0) @binding(1) var<storage, read_write> particles_chunk1: array<Particle>;
@group(0) @binding(2) var<storage, read_write> particles_chunk2: array<Particle>;
@group(0) @binding(3) var<storage, read_write> particles_chunk3: array<Particle>;

struct SimulationParams {
    chunk_size: u32,
};
@group(0) @binding(4) var<uniform> params: SimulationParams;

// 辅助函数：根据全局粒子索引路由到对应的 Buffer 读写
fn get_particle(global_idx: u32) -> Particle {
    let chunk_idx = global_idx / params.chunk_size;
    let local_idx = global_idx % params.chunk_size;

    switch (chunk_idx) {
        case 0u: { return particles_chunk0[local_idx]; }
        case 1u: { return particles_chunk1[local_idx]; }
        case 2u: { return particles_chunk2[local_idx]; }
        case 3u: { return particles_chunk3[local_idx]; }
        default: { return Particle(vec3<f32>(0.0), vec3<f32>(0.0)); } // 越界安全保护
    }
}

fn set_particle(global_idx: u32, p: Particle) {
    let chunk_idx = global_idx / params.chunk_size;
    let local_idx = global_idx % params.chunk_size;

    switch (chunk_idx) {
        case 0u: { particles_chunk0[local_idx] = p; }
        case 1u: { particles_chunk1[local_idx] = p; }
        case 2u: { particles_chunk2[local_idx] = p; }
        case 3u: { particles_chunk3[local_idx] = p; }
        default: { } // 忽略越界
    }
}

2. 方案评估

• 优点：逻辑直观，可以直接利用多核 GPU 的并发能力。
• 缺点：分支开销。虽然现代 GPU 对统一分支（Coherent Branching）优化得很好，但如果同一个 Workgroup 内的线程访问了不同的 Chunk，会导致线程束发散（Warp Divergence），带来一定的性能惩罚。
• 适用场景：粒子索引高度连续、计算时局限在自身分片内的物理仿真（如简单的粒子生命周期模拟）。

三、 方案二：高密度数据压缩与量化（Quantization）

很多时候，我们并不需要用 f32（32位浮点数）去精确存储每一个物理量。通过空间换时间、精度换空间的思路，我们可以把数据极限压缩，从而在不改变 Buffer 绑定的情况下，存储 2 到 4 倍的数据量。

1. 浮点数转定点数（Fixed-Point Packing）

假设物理仿真的世界边界是已知的（例如在一个 $[-100, 100]$ 的立方体边界内）。我们可以将 f32 的坐标量化为 u16 或 u8。

例如，将 3D 坐标 vec3<f32> 压缩进两个 u32 中：

// 将 [-100.0, 100.0] 的 vec3 压缩到两个 u32 字段中
// 占用 8 字节，相比 vec3<f32> 占用 16 字节（带对齐），内存直接减半
struct PackedParticle {
    packed_pos_xy: u32, // x 占 16 位, y 占 16 位
    packed_pos_z_vel: u32, // z 占 16 位, 其他空间可存其他低精度物理量
};

fn pack_f32_to_u16(val: f32, min_val: f32, max_val: f32) -> u32 {
    let normalized = clamp((val - min_val) / (max_val - min_val), 0.0, 1.0);
    return u32(normalized * 65535.0);
}

fn unpack_u16_to_f32(val: u32, min_val: f32, max_val: f32) -> f32 {
    let normalized = f32(val) / 65535.0;
    return min_val + normalized * (max_val - min_val);
}

fn unpack_position(packed: PackedParticle) -> vec3<f32> {
    let x_raw = packed.packed_pos_xy & 0xFFFFu;
    let y_raw = (packed.packed_pos_xy >> 16u) & 0xFFFFu;
    let z_raw = packed.packed_pos_z_vel & 0xFFFFu;

    return vec3<f32>(
        unpack_u16_to_f32(x_raw, -100.0, 100.0),
        unpack_u16_to_f32(y_raw, -100.0, 100.0),
        unpack_u16_to_f32(z_raw, -100.0, 100.0)
    );
}

2. 启用 shader-f16 扩展

如果你的运行设备支持 f16（半精度浮点数），可以在 WebGPU 中开启此扩展。这不仅能减少一半的内存占用，还能提升显存带宽的吞吐效率。

// 在请求设备时请求 f16 支持
const device = await adapter.requestDevice({
    requiredFeatures: ["shader-f16"]
});

在 WGSL 中声明：

enable f16;

struct Particle {
    position: vec3<f16>, // 仅占用 6 字节（由于内存对齐，实际表现视 Struct 布局而定）
    velocity: vec3<f16>,
};

四、 方案三：多通道（Multi-pass）分片计算与虚拟化

对于极其复杂的仿真（例如流体与固体耦合），将所有数据一次性加载到显存中进行计算是不切实际的。我们可以借鉴地形系统中的 Virtual Geometry 思想，引入 Multi-pass 空间分片仿真。

1. 空间哈希网格与动态页表

我们将仿真空间划分为 3D 网格（Grid），不再使用一维扁平粒子数组，而是维护一个激活网格页表（Page Table）。

1. Active List（激活列表）：只将当前帧中发生了物理运动、碰撞、有粒子存在的空间网格送入 GPU 运算。
2. 虚拟 Buffer 换入换出（Host-to-Device Streaming）：对于处于静止状态、或者远离视角的物理区域，将其数据保留在 CPU 端（或较低性能的普通 Buffer 中），在需要时利用 device.queue.writeBuffer 动态更新活动插槽。

[ 物理世界 (无限大) ] -> 空间哈希离散化 -> [ 仅保留激活的 16 个 Chunk ]
                                                  |
                                                  v
                                     [ WebGPU Storage Buffer ]
                                     (Size <= maxStorageBufferBindingSize)

2. Multi-Pass 帧内循环

在单一 Frame 内，通过重用同一个 BindGroup 和同一个 Storage Buffer，配合不同的 setBindGroup(0, bindGroup, dynamicOffsets) 指令，分批次调度 Compute Shader 执行。

const commandEncoder = device.createCommandEncoder();
const passEncoder = commandEncoder.beginComputePass();
passEncoder.setPipeline(computePipeline);

// 假设我们有 4 个空间分区，它们共享同一个超大 GPUBuffer，但通过 offset 区分
const chunkSize = 64 * 1024 * 1024; // 64MB 一分区
const alignment = device.limits.minStorageBufferOffsetAlignment; 

for (let i = 0; i < 4; i++) {
    const dynamicOffset = i * chunkSize;
    // 绑定时传入动态偏移量
    passEncoder.setBindGroup(0, bindGroup, [dynamicOffset]);
    passEncoder.dispatchWorkgroups(workgroupCount);
}

passEncoder.end();
device.queue.submit([commandEncoder.finish()]);

对应的 WGSL 定义需要声明为 var<storage, read_write> 且支持偏移访问：

// 这里的 array 没有指定大小，会根据 BindGroup 的 size 和 dynamic offset 动态切片
@group(0) @binding(0) var<storage, read_write> active_particles: array<Particle>;

五、 实战演练：构建一个超限粒子更新系统

下面提供一套完整的、兼顾了 分片缓冲 (Chunking) 与 定点压缩 思想的 WebGPU 完整流程搭建示例。

1. 运行端 JS 数据初始化与拆分

const PARTICLE_COUNT = 10_000_000; // 一千万粒子
const STRUCT_SIZE = 16; // 压缩后每个粒子占用 16 字节 (8字节位置 + 8字节速度)
const TOTAL_SIZE = PARTICLE_COUNT * STRUCT_SIZE; // 160 MB

// 假设 maxStorageBufferBindingSize 限制为 128MB
// 我们必须将其拆分为两个 Buffer，每个 80MB
const chunkSize = 5_000_000 * STRUCT_SIZE; 

const bufferA = device.createBuffer({
    size: chunkSize,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
});

const bufferB = device.createBuffer({
    size: chunkSize,
    usage: GPUBufferUsage.STORAGE | GPUBufferUsage.COPY_DST,
});

// 创建统一的 BindGroup
const bindGroup = device.createBindGroup({
    layout: pipeline.getBindGroupLayout(0),
    entries: [
        { binding: 0, resource: { buffer: bufferA } },
        { binding: 1, resource: { buffer: bufferB } },
        // ... 其他参数
    ]
});

2. 计算着色器实现 (WGSL)

struct CompressedParticle {
    pos_xy: u32,
    pos_z_vel_x: u32,
    vel_yz: u32,
};

@group(0) @binding(0) var<storage, read_write> chunk_a: array<CompressedParticle>;
@group(0) @binding(1) var<storage, read_write> chunk_b: array<CompressedParticle>;

// 解析压缩数据为标准物理量
fn unpack_particle(p: CompressedParticle) -> vec3<f32> {
    let x = f32(p.pos_xy & 0xFFFFu) / 655.35 - 50.0;
    let y = f32((p.pos_xy >> 16u) & 0xFFFFu) / 655.35 - 50.0;
    let z = f32(p.pos_z_vel_x & 0xFFFFu) / 655.35 - 50.0;
    return vec3<f32>(x, y, z);
}

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    let idx = global_id.x;
    let chunk_limit = 5000000u;

    if (idx >= chunk_limit * 2u) {
        return;
    }

    var particle: CompressedParticle;
    
    // 显式分支路由，规避单Buffer绑定上限
    if (idx < chunk_limit) {
        particle = chunk_a[idx];
    } else {
        particle = chunk_b[idx - chunk_limit];
    }

    // 执行物理仿真计算
    var pos = unpack_particle(particle);
    
    // 极其简化的重力更新示例
    pos.y = pos.y - 0.01; 
    if (pos.y < -50.0) {
        pos.y = 50.0; // 边界重置
    }

    // 重新打包写入
    let p_x = u32((pos.x + 50.0) * 655.35) & 0xFFFFu;
    let p_y = u32((pos.y + 50.0) * 655.35) & 0xFFFFu;
    
    var updated: CompressedParticle;
    updated.pos_xy = p_x | (p_y << 16u);
    updated.pos_z_vel_x = particle.pos_z_vel_x; // 保持其余轴和速度不变
    updated.vel_yz = particle.vel_yz;

    if (idx < chunk_limit) {
        chunk_a[idx] = updated;
    } else {
        chunk_b[idx - chunk_limit] = updated;
    }
}

六、 总结与最佳性能实践

在处理超大型 WebGPU 物理仿真时，不要一味奢望通过申请更高的 requiredLimits 来解决问题。遵循以下开发准则，才能保证你的物理引擎在各种级别的显卡上平稳运行：

1. 优先使用“定点数/半精度”压缩：这是投入产出比最高的手段。将 f32 压缩至 u16 或 u8，不仅能突破绑定上限，还能由于减少了显存带宽占用，带来 30% 以上的计算性能红利。
2. 合理规划 GPU 对齐规则：设计 Struct 时时刻注意 align。如果不小心制造了大量的内存空隙，会无端消耗宝贵的绑定额度。
3. 结合 Compute 与 Render 通道复用：尽量在 Compute Pass 中完成所有物理模拟，并通过 GPUBufferUsage.INDIRECT 间接绘制直接递交给 Render Pass，避免数据从显存回流（Readback）到 CPU。