iPhone12シリーズに搭載されている Apple A14 Bionic では GPU処理において効率的に合計値、最大値、最小値を計算する関数が利用できるので、その動作確認をしてみました。
また、GPUとCPU各々で100万個の整数の合計にかかる時間も計測してみます。
※「SIMD group」「Thread Group」といった用語が出てきます。これらの意味やGPUでの処理の単位については こちらの記事 で解説しています。
１．SIMDグループ関数
1-1．A14 Bionic で追加された命令の種類
A14 Bionic で追加された命令（SIMDグループ関数)には次のものがあります。
| 関数名 | 機能 | 補足 |
|---|---|---|
| simd_sum | 和 | 整数型と浮動小数点型で利用可能 |
| simd_product | 積 | 整数型と浮動小数点型で利用可能 |
| simd_min | 最小値 | 整数型と浮動小数点型で利用可能 |
| simd_max | 最大値 | 整数型と浮動小数点型で利用可能 |
| simd_and | 論理積 | |
| simd_or | 論理和 | |
| simd_xor | 排他的論理和 |
1-2．追加された命令（SIMD scope reduction instructions）の特徴
A14 Bionicで追加されたSIMDグループ関数の特徴は次の通りです。
- SIMDグループ内の全てのスレッドに渡って値の合計等を算出し、その結果を全てのスレッドにブロードキャストする。
- オーバーヘッドが小さい。
- 命令１つでスレッド間の値の計算が可能。
1-3．用途
GPUでの並列処理がより効率的となるため、次のような用途での活用が期待できます。
- 配列のすべての要素を合計する → 平均を計算するのに使える
- 配列の最小値、最大値を取得する → トーンマッピングアルゴリズムで使える
２．配列の合計を計算する
本記事では次のステップで合計を算出します。
※データ数は 1,048,576(=1,024 x 1,024)
- SIMD group毎にsimd_sum命令で合計値を算出（→32スレッド分の合計を得る）
- 1の値を threadgroup のメモリの配列に格納（配列のサイズは32。1Thread Groupあたり32 x 32のスレッドを前提）
- Thread Group内の全てのスレッドが2を実行するのを待つ。これは『バリア』1と呼ばれる機能を利用する。
- 2 の値を合計する。この合計の際も simd_sum 命令を用いて合計をとる。全てのSIMD Groupで計算する必要はないので、
simdgroup_index_in_threadgroupが0番目のSIMD Groupに処理をさせる。 - ここまででThread Group内の合計が算出できたので、その値をバッファに格納。
- Swift側で 5 の値を合計する（1,024個の合計データを合計）
6について、これもGPUで処理させるのも一案ですが端折りました（実際にサンプルプログラムを実行してみると、CPUでの1024個の計算時間は、サンプルプログラムの実行時間全体を俯瞰する上で、無視できるくらい小さな値であったため）。
上記1〜5のMSL側のコードは次のようになります。
typedef float DataType;
kernel void group_max(const device DataType* input_array [[ buffer(0) ]],
device DataType* output_array [[ buffer(1) ]],
uint position [[thread_position_in_grid]],
uint group_pos [[threadgroup_position_in_grid]],
uint simd_group_index [[simdgroup_index_in_threadgroup]],
uint thread_index [[thread_index_in_simdgroup]])
{
// 1Thread Groupに32のSIMD groupがあるのでその各々の計算結果を格納するメモリを確保
threadgroup DataType simd_sum_array[32];
// 各スレッドのinput値の合計を求める
DataType simd_group_max = simd_sum(input_array[position]);
// 合計を一時保存
simd_sum_array[simd_group_index] = simd_group_max;
// Thread Group内のすべてのスレッドの合計の計算＆一時保存を待つ
threadgroup_barrier(mem_flags::mem_threadgroup);
// 一時保存した値をすべてを足す。
// 1つのSIMD Groupにで32スレッドあるので、１つのSIMD Groupのみで処理。
if (simd_group_index == 0) {
output_array[group_pos] = simd_sum(simd_sum_array[thread_index]);
}
}
３．動作確認
100万個のランダムな整数値をCPU(Swift)とGPUの両方で計算して、実行結果と実行時間を確認します。
【検証環境】iPhone12Pro/iOS14.6
3-1) 計測結果
３回計測した結果は次の通りです（単位はs）。
| 計測回 | GPU(*1) [s] |
GPU(*2) [s] |
CPUのみ [s] |
GPU(*1)とCPUのみの実行時間差 [倍] |
|---|---|---|---|---|
| 1 | 0.00267 | 0.00285 | 0.02430 | 9.1 |
| 2 | 0.00280 | 0.00284 | 0.02341 | 8.4 |
| 3 | 0.00284 | 0.00287 | 0.02430 | 8.5 |
*1 threadGroupSizeIsMultipleOfThreadExecutionWidth = true
*2 threadGroupSizeIsMultipleOfThreadExecutionWidth = false
3-2) 実機計測でわかったこと
- 今回の計算方法の場合、A14 Bionicでの実行を前提として、GPUで計算した方がCPUのみで計算するより8〜9倍速い。
- threadGroupSizeIsMultipleOfThreadExecutionWidthの明らかな効果は確認できなかった。Appleのドキュメントには、Thread Groupのサイズが常に
threadExecutionWidthの整数倍なら true を設定することでパフォーマンスが向上するとある。今回のサンプルは該当している。2次元や3次元配列なら効果があるのかも。
４．注意するところ
・浮動小数点値の合計が結構ずれる ＆ double型が使えない
これは致し方ないのですがfloat型で100万個の浮動小数点値を足し込むと、かなり誤差が生じます。そこでdouble型で動作確認しようとしたらMSLのコンパイル通りませんでした。MSLの仕様(2.1 Scalar Data Types)を確認したところ「Metal does not support the double, long long, unsigned long long, and long double data types.」の記載があります。doubleは使えないのですね。。（CPUの処理をGPUに置き換える際には注意が必要です。Metalで倍精度の計算をするときはどうするのでしょう。。。）
・浮動小数点値の合計がCPUとGPUでずれる
今回のサンプルアプリで型をfloatで実行するとCPUとGPUで合計結果が異なります。
どのくらいズレるかがわかりやすいように、ランダムな値ではなく百万個(1,048,576)の1.1を合計した結果が次の通りです。
| デバイス | 合計値 | 補足 |
|---|---|---|
| GPU | 1,153,421.8 | |
| CPU | 1,165,568.5 | |
| 電卓 | 1,153,433.6 | これが正しい |
この差はfloat型の精度が低いことと合計の仕方が異なることに起因しています。
GPUでは、SIMD Group毎に合計 → Thread Group毎に合計 → 全てのThread Groupを合計、のように段階を踏みます。CPU(Swift)でも次のようにGPUと同様に段階を踏んで計算すると、GPUの計算結果と一致します（CPUとGPUでどちらの計算結果がより期待値に近いのか、はおいて）。
let val: Float = 1.1
var gridSum: Float = 0
for _ in 0..<1024 {
var threadGroupSum: Float = 0
for _ in 0..<32 {
var simdSum: Float = 0
for _ in 0..<32 {
simdSum += val
}
threadGroupSum += simdSum
}
gridSum += threadGroupSum
}
print("gridSum [\(gridSum)]") // 1153421.8
５．最後に
MSLのコードもSwiftのコードもシンプルなので「大量データの単純処理で軽量化を図りたい」というようなケースで、GPUへの移行の敷居が低くなったのではないかと思います（A14でなければ、参考記事の[Compute sum of array values in parallel with metal swift]の方法も使えると思います）。
なお、色々と初めてな内容を調べて記事を書いたので内容に誤りがあるかもしれません。ご指摘いただけると助かります。
参考にした記事
-
MSL仕様 6.9.1 Threadgroup and SIMD-group Synchronization Functions
「A barrier function (threadgroup_barrier or simdgroup_barrier) acts as an execution and memory barrier. All threads in a threadgroup (or SIMD-group) executing the kernel must encounter the threadgroup_barrier (or simdgroup_barrier) function. If threadgroup_barrier (or simdgroup_barrier) is inside a conditional statement and if any thread enters the conditional statement and executes the barrier function, then all threads in the threadgroup (or SIMD-group) must enter the conditional and execute the barrier function.」 ↩