这篇笔记主要写一下 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;
把 age 和 lifetime 分开存,而不是存一个"剩余寿命",是因为很多效果需要归一化的生命进度 age / lifetime 来驱动颜色、大小、透明度的曲线。
这个结构是 64 字节,正好是一条 cache line。这不是巧合——让结构体大小对齐到 cache line 边界,能避免一个粒子跨两条 cache line 的情况。
模拟:compute shader
MSLkernel void particle_update(device Particle *particles [[buffer(0)]],
constant SimParams ¶ms [[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(¶ms, 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 ¶ms [[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;
}
right 和 up 是从视图矩阵的转置里取的——视图矩阵的旋转部分是世界到视图,转置就是视图到世界,它的列正好是摄像机在世界空间中的基向量。
绘制:
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;
这个技巧比正确排序重要得多——它解决的是观感上最刺眼的那个问题。
下一篇讲曲面细分。