[iOS]MetalでGPUコンピューティング(13) スレッドについて

Posted at 2016-12-25

この記事は、Metal Advent Calendar2016の24日目です。Metal Advent Calendar2016への投稿は、今回が最後となります。

扱うサンプルコードは、前回と同じライフゲームのアプリ、MetalGameOfLifeです。
 MetalGameOfLife

（実行画面）

今回は、MetalによるGPU並列演算の肝となるのですが今ひとつ分かり難い、スレッドについて解説を行います。
スレッドの数は、GPUで並列に演算される演算の数です。
スレッドは、スレッドグループに含まれます。従いまして、スレッドの設定時はスレッドグループ数とそこに含まれるスレッドの数を指定する必要があります。
この設定は、Metalのパフォーマンスに影響を及ぼします。（参考）

このサンプルコードでは、二通りのスレッドグループの設定を行なっていますので、そのあたりの違いについて解説していきたいと思います。

予め、スレッド関連の用語の解説を行いたいと思います。

Thread:
並列で実行される多くのカーネル関数の実行の一つです。一つのスレッドは、スレッドグループの一部として複数の処理要素を実行します。カーネル関数の中で、あるスレッドは他のスレッドとグリッド中のスレッド位置により識別されます。

Thread group:
単一のコンピューティング用のユニット上で実行されるスレッドの集合体です。同一スレッドグループ内のスレッドは、同じカーネル関数を実行し同じスレッドグループ用のメモリを共有します。なお、スレッドグループ内のスレッド数の上限は、MTLComputePipelineStateのmaxTotalThreadsPerThreadgroupで取得することができます。

SIMD group:
同時にSIMDで実行されるスレッドの集合体です。SIMDとは、命令は1つですがそれを複数のデータに適用する、コンピュータの並列化の形態のことです。スレッドグループの中のスレッドはSIMDグループの集合体として実行されます。

Thread Execution Width:
SIMDグループのサイズです。すなわち、SIMDグループ中のスレッド数です。

Grid:
スレッドグループが形作るグリッドです。1D、2D、3Dの配列のデータになります。2Dの画像処理においては、グリッドは処理する画像のピクセルに対応します。

Threadgroups-per-grid:
グリッドに含まれるスレッドグループの総数です。

MTLComputePipelineStateのthreadExecutionWidthで、一つの命令で同時に実行可能なスレッド数を取得することができます。
例えば、スレッドグループ内のスレッド数の上限が512で、threadExecutionWidthが32の場合、以下のようにスレッドグループ内のスレッド数を設定するのが適切なようです。

let threadsPerThreadgroup = MTLSizeMake(32, 16, 1)

それでは、サンプルコード内のスレッドの設定箇所を見ていきます。
繰り返しになるのですが、このサンプルコードは、主に以下のファイルで構成されています。

AAPLRender.h
AAPLRender.m
AAPLViewController.h
AAPLViewController.m
Sharder.metal

このうち、AAPLRender.mには並列コンピューティング及び描画のCPU側のロジックが、Shader.metalには頂点シェーダー、フラグメントシェーダー、GPUコンピューティング用のシェーダーが書かれています。

ここからは、AAPLRender.m内におけるスレッドの設定箇所を解説していきます。
以下は、ライフゲームのロジックのために画像データのピクセル全てに処理を行うための箇所に記述されたスレッドの設定です。

AAPLRender.m

MTLSize threadsPerThreadgroup = MTLSizeMake(16, 16, 1);
MTLSize threadgroupCount = MTLSizeMake(ceil((float)self.gridSize.width / threadsPerThreadgroup.width),
                                           ceil((float)self.gridSize.height / threadsPerThreadgroup.height),
                                           1);

各スレッドグループ内には16x16x1のスレッドが含まれ、全てのグリッドはこのスレッドグループでカバーされることになります。二次元の画像データに対する処理を行うので、スレッドグループの設定は16x16の二次元的になっています。

また、以下の箇所はタップした周囲のグリッドにしか処理を行わない箇所に記述されたスレッドの設定です。

AAPLRender.m

MTLSize threadsPerThreadgroup = MTLSizeMake(self.activationPoints.count, 1, 1);
MTLSize threadgroupCount = MTLSizeMake(1, 1, 1);

必要最小画面のスレッドしか設定されておらず、そもそもスレッドグループも一つしか設定されていません。スレッドの数は並列で行われる演算数なので、必要数あれば十分ということなのでしょう。

そして、MTLComputeCommandEncoderのdispatchThreadgroups:threadsPerThreadgroup:メソッドにより、スレッドグループにおけるコンピューティング用の関数の実行がエンコードされます。

AAPLRender.m

[commandEncoder dispatchThreadgroups:threadgroupCount threadsPerThreadgroup:threadsPerThreadgroup];

スレッドグループのサイズとパフォーマンスの関係ですが、以前に簡易な方法で実測してみたところ、スレッドグループの最適化によりかなりパフォーマンスが改善するようでした。この辺りはいずれきっちりと実験を行なってみたいと思います。

今回はスレッドの解説を行いましたが、このようにGPU並列演算のためにはスレッドを適切に設定することが一つの要になります。

一連の投稿で、MetalのGPUコンピューティング関連技術の解説を行いましたが全て完全に理解する必要はないですし、私自身も分からないことだらけです。
しかしながら、他の技術と同様、頭で理解しようとするよりも実際に手を動かしてみたほうが早いですので、興味のある方はこちらなどを参考にぜひコードを書いてみてくださいね。

なお、こちらもご参考にどうぞ。
MetalでiOSアプリに宿る生命

それでは皆さま、よいお年を！

参考: MTLComputeCommandEncoder

You get articles that match your needs
You can efficiently read back useful information
You can use dark theme

What you can do with signing up