Metal #12 パーティクルシステム

このノートでは GPU パーティクルシステムを扱います。シミュレーションも描画も GPU に置き、CPU は毎フレーム数個のコマンドを出すだけにします。

なぜ GPU に置くのか

パーティクルは GPU にとって理想的な負荷です。同じロジックを実行する、互いに独立した数千・数万の実体だからです。10 万個を CPU でシミュレーションすると、毎フレームのアップロードだけでボトルネックになります。float3 を 10 万個で 1.2 MB、毎フレームなら位置だけで 72 MB/s のトラフィックです。

GPU 上ならパーティクルデータはデバイスメモリから出ません。CPU はコンピュートを 1 回ディスパッチし、描画を 1 回発行するだけです。

データ配置

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 バイト、ちょうど 1 キャッシュラインです。偶然ではありません。構造体サイズをキャッシュライン境界に合わせることで、1 パーティクルが 2 ラインにまたがる事態を避けられます。

シミュレーション：コンピュートシェーダー

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);
}

要点をいくつか。

境界チェックは省略できません。 ディスパッチされるスレッド数はスレッドグループサイズの倍数に切り上げられ、余分なスレッドは範囲外に書き込みます。

半陰的オイラー法（速度を更新し、その新しい速度で位置を更新する）は陽的オイラー法よりはるかに安定しており、行数はまったく同じです。パーティクルや物理の積分には常にこちらを使ってください。

削除ではなく再利用。 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 はこのパイプラインの、このデバイスにおける上限で、レジスタ占有量で決まります。1024 をハードコードすると、レジスタ圧の高いカーネルでは失敗します。

発生

新しいパーティクルを発生させる鍵は、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) から乱数を得ます。共有状態なしで完全に並列です。

インスタンシングによる描画

パーティクルは通常、常にカメラを向くビルボードとして描かれます。パーティクルごとに 4 頂点を用意する必要はありません。インスタンシングを使い、頂点関数が 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);

    // カメラの right と up に沿って展開すれば、ビルボードは常にカメラを向く
    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 上でソートする。 バイトニックソートはコンピュートシェーダーで O(n log²n) で実装できます。数万個なら現実的ですが、フレーム予算の相応の割合を消費します。

ソフトパーティクル、そして本当に効くもの。 どれだけソートしても、パーティクルがシーンのジオメトリと交差する箇所の硬い縁は直りません。ソフトパーティクルはパーティクルの深度とシーンの深度を比較し、近づくにつれてフェードさせます。

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

この手法は正しいソートよりはるかに重要です。目が実際に気づく破綻を取り除いてくれます。

次回はテッセレーションを扱います。