T.TAO
返回博客
/8 min read/Graphics Engine

Metal #12 粒子系统

#ComputerGraphics#GraphicsEngine#Metal

这篇笔记主要写一下 GPU 粒子系统:把模拟和渲染都留在 GPU 上,CPU 每帧只需要发一条指令。

为什么放在 GPU 上

粒子是 GPU 的理想工作负载:成千上万个彼此独立、执行完全相同逻辑的实体。在 CPU 上模拟十万个粒子,光是每帧把位置上传到 GPU 就足以成为瓶颈(十万个 float3 是 1.2 MB,每帧传一次就是 72 MB/s)。

放在 GPU 上,粒子数据从头到尾不离开显存。CPU 每帧只 dispatch 一次 compute,再 draw 一次。

数据布局

Ctypedef struct {
    vector_float3 position;
    vector_float3 velocity;
    vector_float4 color;
    float         age;
    float         lifetime;    // age >= lifetime 表示死亡
    float         size;
} Particle;

agelifetime 分开存,而不是存一个"剩余寿命",是因为很多效果需要归一化的生命进度 age / lifetime 来驱动颜色、大小、透明度的曲线。

这个结构是 64 字节,正好是一条 cache line。这不是巧合——让结构体大小对齐到 cache line 边界,能避免一个粒子跨两条 cache line 的情况。

模拟:compute shader

MSLkernel void particle_update(device Particle *particles [[buffer(0)]],
                            constant SimParams &params [[buffer(1)]],
                            uint id [[thread_position_in_grid]])
{
    if (id >= params.count) return;      // 网格大小未必整除线程组大小

    device Particle &p = particles[id];

    p.age += params.deltaTime;
    if (p.age >= p.lifetime) {
        respawn(p, params, id);          // 回收:直接重置,不做增删
        return;
    }

    // 半隐式欧拉:先更新速度,再用新速度更新位置
    p.velocity += params.gravity * params.deltaTime;
    p.velocity *= (1.0 - params.drag * params.deltaTime);
    p.position += p.velocity * params.deltaTime;

    float t  = p.age / p.lifetime;
    p.color  = mix(params.startColor, params.endColor, t);
    p.size   = mix(params.startSize,  params.endSize,  t);
}

几个要点。

边界检查不能省。 dispatch 的线程数会被向上取整到线程组大小的整数倍,多出来的线程会越界写内存。

半隐式欧拉(先更新速度再用新速度更新位置)比显式欧拉稳定得多,而且代码量完全一样。做任何粒子或物理模拟时都应该用它。

回收而不是删除。 GPU 上增删元素代价很高,正确的做法是让粒子池大小固定,死掉的粒子直接重置为新生状态。缓冲区永远不会碎片化,也不需要压缩。

CPU 端的调度:

Swiftlet encoder = commandBuffer.makeComputeCommandEncoder()!
encoder.setComputePipelineState(updatePipeline)
encoder.setBuffer(particleBuffer, offset: 0, index: 0)
encoder.setBytes(&params, length: MemoryLayout<SimParams>.stride, index: 1)

let threadsPerGroup = min(updatePipeline.maxTotalThreadsPerThreadgroup, 256)
encoder.dispatchThreads(MTLSize(width: particleCount, height: 1, depth: 1),
                        threadsPerThreadgroup: MTLSize(width: threadsPerGroup, height: 1, depth: 1))
encoder.endEncoding()

dispatchThreads(非均匀线程组)只在支持的设备上可用,但它省掉了手动向上取整。老设备要用 dispatchThreadgroups 并自己算组数。

maxTotalThreadsPerThreadgroup 是这个 pipeline 在当前设备上的上限,由寄存器占用决定。硬编码 1024 在寄存器压力大的 kernel 上会失败。

发射

发射新粒子的关键是不要在 CPU 上维护"下一个空闲索引"——那需要读回 GPU 的状态。两种常见做法:

环形缓冲。 维护一个原子递增的发射游标:

MSLkernel void particle_emit(device Particle *particles [[buffer(0)]],
                          device atomic_uint *cursor [[buffer(1)]],
                          constant EmitParams &params [[buffer(2)]],
                          uint id [[thread_position_in_grid]])
{
    if (id >= params.emitCount) return;
    uint slot = atomic_fetch_add_explicit(cursor, 1, memory_order_relaxed) % params.poolSize;
    particles[slot] = makeParticle(params, id);
}

简单可靠,代价是可能覆盖还活着的粒子。只要池子够大就不会有观感问题。

自由列表。 死掉的粒子把自己的索引 push 进一个栈,发射时 pop。更精确,但需要额外的原子操作和一个索引缓冲区。

生成随机数用基于索引的哈希,而不是需要状态的随机数发生器:

MSLfloat hash(uint n) {
    n = (n << 13U) ^ n;
    n = n * (n * n * 15731U + 789221U) + 1376312589U;
    return float(n & 0x7fffffffU) / float(0x7fffffff);
}

每个线程用 hash(id + frameSeed) 得到自己的随机数,无需共享状态,完全可并行。

渲染:实例化

粒子通常渲染成永远面向相机的公告板(billboard)。不需要为每个粒子准备四个顶点——用实例化,顶点函数从 instance_id 取粒子数据,从 vertex_id 算四边形的角:

MSLvertex ParticleOut particle_vertex(uint vid [[vertex_id]],
                                   uint iid [[instance_id]],
                                   device const Particle *particles [[buffer(0)]],
                                   constant Uniforms &u [[buffer(11)]])
{
    device const Particle &p = particles[iid];

    // 由 vertex_id 生成 (-1,-1) (1,-1) (-1,1) (1,1)
    float2 corner = float2((vid & 1) * 2.0 - 1.0, (vid >> 1) * 2.0 - 1.0);

    // 用摄像机的右向量和上向量展开,公告板就永远面向相机
    float3 right = float3(u.viewMatrix[0][0], u.viewMatrix[1][0], u.viewMatrix[2][0]);
    float3 up    = float3(u.viewMatrix[0][1], u.viewMatrix[1][1], u.viewMatrix[2][1]);
    float3 posWS = p.position + (right * corner.x + up * corner.y) * p.size;

    ParticleOut out;
    out.position = u.viewProjection * float4(posWS, 1);
    out.uv       = corner * 0.5 + 0.5;
    out.color    = p.color;
    return out;
}

rightup 是从视图矩阵的转置里取的——视图矩阵的旋转部分是世界到视图,转置就是视图到世界,它的列正好是摄像机在世界空间中的基向量。

绘制:

Swiftencoder.drawPrimitives(type: .triangleStrip, vertexStart: 0,
                       vertexCount: 4, instanceCount: particleCount)

排序问题

半透明粒子需要从远到近绘制,否则混合结果是错的。但在 GPU 上给几万个粒子排序不便宜。

实践中的三条路:

用加法混合,绕开排序。 火焰、魔法、发光轨迹都适合,因为加法是可交换的,顺序不影响结果。这是最常用的答案。

GPU 上排序。 双调排序(bitonic sort)可以在 compute shader 里实现,复杂度 O(n log²n)。对于几万个粒子这是可行的,但它会吃掉相当一部分帧时间。

软粒子与深度写入的取舍。 无论怎么排序,粒子与场景几何相交时都会出现硬边。软粒子通过比较粒子深度和场景深度,在接近时淡出:

MSLfloat sceneDepth = depthTexture.sample(s, screenUV).r;
float fade = saturate((linearize(sceneDepth) - linearize(in.position.z)) / softness);
color.a *= fade;

这个技巧比正确排序重要得多——它解决的是观感上最刺眼的那个问题。

下一篇讲曲面细分。

本系列文章

Metal
  1. 01Metal #0 Swift 回顾
  2. 02Metal #1 初始化
  3. 03Metal #2 渲染管线
  4. 04Metal #3 顶点函数
  5. 05Metal #4 片元函数
  6. 06Metal #5 纹理
  7. 07Metal #6 摄像机与交互
  8. 08Metal #7 光照
  9. 09Metal #8 材质
  10. 10Metal #9 渲染通道
  11. 11Metal #10 阴影
  12. 12Metal #11 延迟渲染
  13. 13Metal #12 粒子系统
  14. 14Metal #13 曲面细分
  15. 15Metal #14 后处理
  16. 16Metal #15 反射与折射
  17. 17Metal #16 动画
  18. 18Metal #17 光线追踪(一)渲染算法
  19. 19Metal #18 光线追踪(二)阴影与光照
  20. 20Metal #19 光线追踪(三)性能优化
  21. 21Metal #21 [附录] 计算着色器
  22. 22Metal #22 [附录] SwiftUI 中的 Metal