This note covers a GPU particle system: keeping both simulation and rendering on the GPU, so the CPU issues a couple of commands per frame and nothing more.
Why it belongs on the GPU
Particles are an ideal GPU workload: thousands of independent entities running identical logic. Simulating a hundred thousand particles on the CPU makes the per-frame upload alone a bottleneck — a hundred thousand float3s is 1.2 MB, and once per frame that is 72 MB/s of traffic for positions only.
On the GPU the particle data never leaves device memory. The CPU dispatches one compute pass and issues one draw.
Data layout
Ctypedef struct {
vector_float3 position;
vector_float3 velocity;
vector_float4 color;
float age;
float lifetime; // age >= lifetime means dead
float size;
} Particle;
Storing age and lifetime separately, rather than a single "time remaining", is deliberate: most effects want the normalized life progress age / lifetime to drive curves for colour, size and opacity.
This struct is 64 bytes, exactly one cache line. That is not a coincidence — aligning the struct size to a cache line boundary avoids a particle straddling two lines.
Simulation: the 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; // the grid may not divide evenly by the threadgroup
device Particle &p = particles[id];
p.age += params.deltaTime;
if (p.age >= p.lifetime) {
respawn(p, params, id); // recycle: reset in place, never add or remove
return;
}
// semi-implicit Euler: update velocity first, then position with the new velocity
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);
}
Several points.
The bounds check is not optional. The dispatched thread count is rounded up to a multiple of the threadgroup size, and the surplus threads will write out of bounds.
Semi-implicit Euler — update velocity, then use the new velocity for position — is dramatically more stable than explicit Euler and costs exactly the same number of lines. Use it for any particle or physics integration.
Recycle instead of deleting. Adding and removing elements on the GPU is expensive; the right structure is a fixed-size pool where dead particles are reset in place. The buffer never fragments and never needs compacting.
Dispatching from the 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 (non-uniform threadgroups) is only available on supporting devices, but it removes the manual rounding. Older devices use dispatchThreadgroups and compute the group count themselves.
maxTotalThreadsPerThreadgroup is this pipeline's limit on this device, determined by register occupancy. Hard-coding 1024 fails on a kernel with heavy register pressure.
Emission
The key to emitting new particles is not maintaining a "next free index" on the CPU — that would require reading GPU state back. Two common approaches:
A ring buffer. Keep an atomically incremented emission cursor:
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);
}
Simple and robust; the cost is that it may overwrite a particle that is still alive. With a large enough pool this is never visible.
A free list. Dead particles push their index onto a stack and emission pops from it. More precise, at the cost of extra atomics and an index buffer.
Generate random numbers with an index-based hash rather than a stateful generator:
MSLfloat hash(uint n) {
n = (n << 13U) ^ n;
n = n * (n * n * 15731U + 789221U) + 1376312589U;
return float(n & 0x7fffffffU) / float(0x7fffffff);
}
Each thread derives its randomness from hash(id + frameSeed) — no shared state, perfectly parallel.
Rendering with instancing
Particles are usually rendered as billboards that always face the camera. There is no need to build four vertices per particle: with instancing, the vertex function reads the particle from instance_id and derives the quad corner from 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 generates (-1,-1) (1,-1) (-1,1) (1,1)
float2 corner = float2((vid & 1) * 2.0 - 1.0, (vid >> 1) * 2.0 - 1.0);
// expand along the camera's right and up vectors, so the billboard always faces it
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 and up come from the transpose of the view matrix: the rotation part of the view matrix maps world to view, so its transpose maps view to world, and its columns are exactly the camera's basis vectors in world space.
The draw:
Swiftencoder.drawPrimitives(type: .triangleStrip, vertexStart: 0,
vertexCount: 4, instanceCount: particleCount)
The sorting problem
Transparent particles must be drawn back to front or the blend is wrong. But sorting tens of thousands of particles on the GPU is not cheap.
Three routes in practice:
Use additive blending and sidestep sorting. Fire, magic and glowing trails all suit it, because addition commutes and order does not matter. This is the usual answer.
Sort on the GPU. A bitonic sort can be implemented in a compute shader at O(n log²n). For tens of thousands of particles it is feasible, but it consumes a real slice of the frame budget.
Soft particles, and what actually matters. No amount of sorting fixes the hard edge where a particle intersects scene geometry. Soft particles compare particle depth against scene depth and fade out as they approach:
MSLfloat sceneDepth = depthTexture.sample(s, screenUV).r;
float fade = saturate((linearize(sceneDepth) - linearize(in.position.z)) / softness);
color.a *= fade;
This trick matters far more than correct sorting — it removes the artefact the eye actually notices.
The next note covers tessellation.