Metal を使って10万個のパーティクルを描画しよう #Swift

はじめに

Metal を使うとたくさんの計算を並列で行うことができます. 60fps で画面の更新をする場合, 1フレームあたりの処理は約 16ms で収めなければいけません. UIKit を使って10万個の UIView の frame を更新しながら 60fps を保つのは難しいでしょう. Metal を使うとどんな計算が可能なのか, 簡単なサンプルアプリを実装したら勉強になった箇所がたくさんあったので, 共有と備忘録を兼ねてまとめたいと思います.

環境

Xcode 11.5
iPhone 11 Pro, iPhone XS, iPad Pro 11inch 第2世代 ¹

つくったもの

10万個のパーティクル	10万個のパーティクル(白黒)	設定画面

端末の画面いっぱいに最大で10万個のパーティクルがアニメーションします. ある程度可変なパラメータがあった方が理解に繋がると思ったので, パーティクルの色と背景色, パーティクルの個数は設定画面から指定できるようになっています.

ソースコードは naru-jpn/100000-particles で公開しています.

大まかな処理の流れ

大まかに処理の流れを図示すると上の図のようになります.

それぞれのパーティクルの位置情報などを更新 ²
1. で更新された情報をもとにして, それぞれのパーティクルを画面上に描画

という流れです. シンプルな構成です. Metal を使用する場合はパイプラインという処理の流れを記述する必要があり, 大まかな流れをイメージすることは大事なことだと思います. 60fps を維持したいので, これらの一連の処理は 16ms 以内で行う必要があります. 記事の後半にパフォーマンスについても記載しています.

実装の解説

具体的な実装について順を追って解説します. 記載しているコードは要点が分かるように部分的に簡略化しています.

パーティクルを表す構造体

ShaderTypes.h

#include <simd/simd.h>

typedef struct {
  vector_float4 color;
  vector_float2 position;
  vector_float2 velocity;
  float phase;
} particle_t;

上の構造体がパーティクルの実体です. RGBA の色情報, 2次元の座標上の位置と速度, 横方向の揺れを制御する為の変数から構成されています. この構造体は Swift のプログラムだけからではなくシェーダプログラムからも使いたいので, Swift の Struct ではなくこのように定義をする必要があります. 今回は定義を共有するものはこれだけなので, 直接このファイルを Bridging Header に指定します.

パーティクルを格納するバッファ

Renderer.swift

let length: Int = MemoryLayout<particle_t>.size * Renderer.maxNumberOfParticles
let buffer: MTLBuffer = device.makeBuffer(length: length, options: .storageModeShared)

必要な領域のサイズを指定して, バッファを作成しています. options に .storageModeShared を指定していますが, これは CPU と GPU の両方からこのバッファの内容を編集したいからです. ³

パーティクルの初期化

Renderer.swift

let particleBuffer = particleBuffers[0].contents().bindMemory(to: particle_t.self, capacity: numberOfParticles)
  for index in 0..<numberOfParticles {
    particleBuffer[index] = particle_t.create(with: setting, viewportSize: viewportSize)
  }

バッファの内容を編集するために bindMemory で型を指定して UnsafeMutablePointer<particle_t> に変換しています. 位置情報やそれぞれのパーティクルの色情報などをここで初期化しています.

描画のサイクル

Renderer.swift

func draw(in view: MTKView) { // 定期的に呼ばれるデリゲートメソッド
  let semaphore = inFlightSemaphore // トリプルバッファリングの制御
  _ = semaphore.wait(timeout: DispatchTime.distantFuture)

  do {
    let simulateSemaphore = simulationInFlightSemaphore // 
    _ = simulateSemaphore.wait(timeout: DispatchTime.distantFuture)

    guard let commandBuffer = commandQueue.makeCommandBuffer() else {
      fatalError("Failed to make command buffer.")
    }
    // 1. それぞれのパーティクルの位置情報などを更新
    simulate(in: view, commandBuffer: commandBuffer)
    commandBuffer.addCompletedHandler { _ in
      simulateSemaphore.signal()
    }
    commandBuffer.commit()
  }

  do {
    guard let commandBuffer = commandQueue.makeCommandBuffer() else {
      fatalError("Failed to make command buffer.")
    }
    // 2. 1. で更新された情報をもとにして, それぞれのパーティクルを画面上に描画
    render(in: view, commandBuffer: commandBuffer)
    commandBuffer.addCompletedHandler { _ in
      semaphore.signal()
    }
    commandBuffer.commit()
  }

  currentBufferIndex = (currentBufferIndex + 1) % Renderer.maxInFlightRenderingBuffers
}

上の関数は MTKViewDelegate に定義されている関数で, MTKView の再描画が必要なタイミングでこの関数が呼び出されます. セマフォの操作等 ⁴ をしていますが, ここで大事なのは simulate と render の2行です. この2つの処理が, 「大まかな処理の流れ」で説明した2つのステップに対応しています.

1. それぞれのパーティクルの位置情報などを更新

Swift側の処理

Renderer.swift

  guard let function = library.makeFunction(name: "simulate") else {
    fatalError("Failed to make function simulate.")
  }
  do {
    // 関数 'simulate' を使う ComputePipeline を定義
    simulatePipelineState = try device.makeComputePipelineState(function: function)
  }
  // ...


private func simulate(in view: MTKView, commandBuffer: MTLCommandBuffer) {
  // ...

  // 上で定義したパイプライン(simulatePipelineState)に従って処理をする
  computeEncoder.setComputePipelineState(simulatePipelineState)
  computeEncoder.setBuffer(particleBuffers[currentBufferIndex], offset: 0, index: 0) // 入力
  computeEncoder.setBuffer(particleBuffers[simulatedBufferIndex], offset: 0, index: 1) // 出力
  computeEncoder.setBytes(&viewportSize, length: MemoryLayout<vector_float2>.size, index: 2) // 画面サイズ
  computeEncoder.setThreadgroupMemoryLength(simulatePipelineState.threadExecutionWidth * MemoryLayout<particle_t>.size, index: 0)
  computeEncoder.dispatchThreads(dispatchThreads, threadsPerThreadgroup: threadsPerThreadgroup)
  computeEncoder.endEncoding()
}

library.makeFunction(name: "simulate") で simulate という関数をとってきて, device.makeComputePipelineState(function: function) でその関数を使うパイプラインを定義しています. 指定した情報は computeEncoder によって GPU に渡す為の命令にエンコードされます. 関数 simulate の定義は Shaders.metal ファイル内にあり, そのような関数はシェーダ関数と呼ばれます.

シェーダ側の処理

Shaders.metal

kernel void
simulate(device particle_t* currentParticles [[ buffer(0) ]], // 入力
         device particle_t* newParticles [[ buffer(1) ]], // 出力
         constant vector_uint2 *viewportSize [[ buffer(2) ]], // 画面サイズ
         const uint gid [[ thread_position_in_grid ]])
{
  // 更新前のパーティクル情報
  float2 position = currentParticles[gid].position;
  float2 velocity = currentParticles[gid].velocity;
  float4 color = currentParticles[gid].color;
  float phase = currentParticles[gid].phase;
  float end = (vector_float2(*viewportSize) / 2.0).y;

  position.x += sin(phase); // 横方向の移動
  if (position.y < -end) {　
    position.y = end; // 一番下まで到達したら一番上に戻る
  }

  // 更新後のパーティクル情報
  newParticles[gid].color = color;
  newParticles[gid].position = position + velocity; // velocityは下向きのy成分のみ.
  newParticles[gid].velocity = velocity;
  newParticles[gid].phase = phase + PHASE_INTERVAL;
}

上の2つのコードを見比べて, setBuffer や setBytes の index とシェーダ関数の引数にある数字とを比較すると対応関係が分かります. gid というのはパーティクルのインデックスで, それぞれの gid が並列で計算されていると思ってもらえばよいです. シェーダ関数 simulate に必要な情報を渡して, 関数の中で計算をしているという流れが分かります.

2. 1. で更新された情報をもとにして, それぞれのパーティクルを画面上に描画

Swift側の処理

Renderer.swift

  guard let vertexFunction = library.makeFunction(name: "particle_vertex") else {
    fatalError("Failed to make function particle_vertex.")
  }
  guard let fragmentFunction = library.makeFunction(name: "particle_fragment") else {
    fatalError("Failed to make function particle_fragment.")
  }

  // 関数 'particle_vertex', 'particle_fragment' を使う RenderPipeline を定義
  let renderPipelineStateDescriptor = MTLRenderPipelineDescriptor()
  renderPipelineStateDescriptor.vertexFunction = vertexFunction
  renderPipelineStateDescriptor.fragmentFunction = fragmentFunction
  renderPipelineStateDescriptor.colorAttachments[0].pixelFormat = ...
  // ...
  renderPipelineState = try device.makeRenderPipelineState(descriptor: renderPipelineStateDescriptor)
  // ...


private func render(in view: MTKView, commandBuffer: MTLCommandBuffer) {
  let renderPassDescriptor = MTLRenderPassDescriptor()
  renderPassDescriptor.colorAttachments[0].texture = view.currentDrawable?.texture
  renderPassDescriptor.colorAttachments[0].loadAction = .clear // 描画の前にテクスチャをクリア
  renderPassDescriptor.colorAttachments[0].clearColor = viewClearColor // 背景色の指定
  renderPassDescriptor.colorAttachments[0].storeAction = .store // 描画結果を保存

  guard let renderEncoder = commandBuffer.makeRenderCommandEncoder(descriptor: renderPassDescriptor) else {
    return
  }

  // 上で定義したパイプライン(renderPipelineState)に従って処理をする
  renderEncoder.setRenderPipelineState(renderPipelineState)
  renderEncoder.setVertexBuffer(particleBuffers[simulatedBufferIndex], offset: 0, index: 0) // パーティクル情報
  renderEncoder.setVertexBytes(&viewportSize, length: MemoryLayout<vector_float2>.size, index: 1)
  renderEncoder.drawPrimitives(type: .point, vertexStart: 0, vertexCount: numberOfParticles)
  renderEncoder.endEncoding()

  if let drawable = view.currentDrawable {
    commandBuffer.present(drawable)
  }
}

先ほどは ComputePipiline でしたが, ここでは RenderPipeline というものが出てきました. レンダリングについての説明をするのはとても大変なので割愛するのですが, ここでは2つのシェーダ関数を指定してパイプラインを定義しています. 上では renderEncoder が GPU に渡す為の命令へのエンコードを行います. setVertexBuffer と setVertexBytes で指定した情報を使って, drawable に .point のプリミティブを描画しています.

view.currentDrawable というのはこれから画面上に描画されるコンテンツです. view.currentDrawable?.texture を RenderPipeline の描画先に指定し, commandBuffer.present(drawable) によってコンテンツが描画された drawable を画面上に表示しています.

シェーダ側の処理

Shaders.metal

constant float PARTICLE_SIZE = 5.0f;

struct Point {
  float4 position [[position]]; // パーティクルの位置
  float size [[point_size]]; // パーティクルの大きさ
  float4 color;
};

vertex Point
particle_vertex(const device particle_t* particles [[ buffer(0) ]], // パーティクル情報
                constant vector_uint2 *viewportSizePointer [[ buffer(1) ]], // 画面サイズ
                unsigned int vid [[ vertex_id ]])
{
  Point out;
  out.position = vector_float4(0.0f, 0.0f, 0.0f, 1.0f);
  out.position.xy = particles[vid].position / (vector_float2(*viewportSizePointer) / 2.0f); 
  out.size = PARTICLE_SIZE;
  out.color = particles[vid].color;
  return out;
}

fragment float4
particle_fragment(Point in [[stage_in]])
{
  return in.color; // 領域内の色を指定
};

関数 particle_vertex, particle_fragment はそれぞれバーテックスシェーダ, フラグメントシェーダと呼ばれます. 先ほど .point のプリミティブを描画すると書きましたが, ここではポイントの位置や大きさ, 領域内の色などを決めています. [[position]] や [[point_size]] は Metal Shading Language ⁵ の中で定義されている attribute と呼ばれるもので, それぞれポイントの位置, ポイントの大きさに対応しています.

ここまでが大まかな処理の流れの実装を追ったものです. 細かい処理の内容は調べればキリがないですが, 全体の処理を俯瞰してみるととても単純なものだと分かります.

パフォーマンス

パーティクルを 100,000 個描画した場合の iPhone 11 Pro 上でのパフォーマンスを Xcode 上で確認しました.

右端のグラフが GPU が1フレームあたりにかけている処理時間を表していますが, まだ制限時間である 16.7ms の半分程度の余力を残しています. すごい.

展望

ここで紹介した処理は計算内容もパーティクルの描画処理も最小限のものでした. もともとはパーティクル同士の相互作用があったりなどの複雑な計算を行いたかったのですが, まずは Metal の操作に慣れるために今回の実装内容にまとめてみました. 物理演算的な計算をコンピュートシェーダ上で行って, リアルタイムシミュレーションのようなものができたら楽しいなと思っています.

もし記事の内容に不備がありましたら, 指摘していただけると幸いです😌

今回のプログラムは, GPU Family 5 以上の GPU が搭載されたデバイス上でのみ動作します. プログラム中で使用している MTLComputeCommandEncoder の関数 dispatchThreads(_:threadsPerThreadgroup:) が動作する必要があるからです. この並列処理に関する関数は GPU Family 4 でも動作するのですが、手元に端末がなくて動作が確認できなかったので対象外としています. dispatchThreadgroups(_:threadsPerThreadgroup:) で最適化を行うことによってもある程度は速度が出せると思うのですが, 最適化できるだけの知識がまだないので検証ができていません. ( Metal Feature Set Tables: https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf ) ↩
パーティクルの情報を更新する流れは必ずしも今回のようである必要はありません. 例えば Apple のサンプルコード MetalShaderShowcase では, birthOffset という時間経過に相当する変数から, 方程式を使ってパーティクルの位置を計算してそのまま描画を行います. 今回, パーティクルの情報を更新して一度バッファに保存するという流れにした理由は, 将来的にパーティクルの描画処理を物理シミュレーションなどに応用したいと思っているからです. ↩
MTLResourceOptions ↩
Apple によると, 描画をする際にはトリプルバッファリングを使用することが推奨されています. この部分では, トリプルバッファリングで全体の描画を制御し, パーティクル情報の更新の際には同時に1つの処理しか走らないような制御を行っています. 今回のパーティクル情報の更新では一つ前のフレームで更新した情報を計算の入力として用いるので, ここで同時に2つ以上の処理が走ってしまうと正常に計算が行えないからです. Metal Best Practices Guide: Triple Buffering ↩
Metal Shading Language Specification ↩