Metalは、iOS8で導入された画像処理や並列演算などを効率的に行うためのフレームワークです。本記事では、そのうち画像処理の方ではなくCPU/GPUの連携による並列演算の方を扱います。
GPUコンピューティング(GPGPU)は、一般に機械学習やリアルタイム画像解析、オーディオなどパフォーマンスがクリティカルであり、なおかつ並列演算が有効である場面で多く使われてきましたが、Metalの登場によりiOSでも手軽にGPUコンピューティングが扱えるようになりました。
今後どんなアプリに応用できるのかとてもワクワクしているのですが、Metalを用いた開発は少々敷居が高く、また通常のCPUによるシーケンシャルな演算とは特性が異なります。
Metalをもっと身近なものにするために、今回はまず可能な限り簡単な実装を用いてMetalによるGPU演算を行い、その後CPUによる演算とパフォーマンスと特性の比較を行ってみたいと思います。
まず、Metal演算の登場人物紹介です。
MTLDevice:
単一のGPUを抽象化したプロトコルです。GPUへ送るデータのバッファを作ったり、コマンドのキューを作ったりします。
MTLCommandBuffer:
特定のGPUで実行されるコマンドを格納します。
MTLCommandQueue:
MTLCommandBufferのキューです。コマンドの実行順を管理します。
MTLFunction:
Metalシェーディング言語で書かれた演算処理です。
MTLLibrary:
MTLFunctionのインターフェイスです。
MTLComputePipelineState:
MTLFunctionをコンパイルされたコードに変換します。
MTLComputeCommandEncoder:
データ並列演算処理用のコマンドエンコーダーです。
次にコードです。まずViewControllerの方から。
import UIKit
import Metal
class ViewController: UIViewController {
let inputDataCount = 100
var device: MTLDevice!
var defaultLibrary: MTLLibrary!
var commandQueue: MTLCommandQueue!
var computePipelineState: MTLComputePipelineState!
override func viewDidLoad()
{
super.viewDidLoad()
//初期化
device = MTLCreateSystemDefaultDevice()
defaultLibrary = device.newDefaultLibrary()
commandQueue = device.makeCommandQueue()
let ml2Func = defaultLibrary.makeFunction(name: "addOne")!
computePipelineState = try! device.makeComputePipelineState(function: ml2Func)
}
@IBAction func calculate(sender: AnyObject)
{
//入力データ
var inputData:[Float] = []
for _ in 0...inputDataCount-1 {
inputData.append(Float(arc4random_uniform(UInt32(inputDataCount))))
}
//コマンドバッファとエンコーダの作成と設定
let commandBuffer = commandQueue.makeCommandBuffer()
let computeCommandEncoder = commandBuffer.makeComputeCommandEncoder()
computeCommandEncoder.setComputePipelineState(computePipelineState)
//入力バッファの作成と設定
let inputDataBuffer = device.makeBuffer(bytes: inputData, length: inputData.byteLength, options: [])
computeCommandEncoder.setBuffer(inputDataBuffer, offset: 0, at: 0)
//出力バッファの作成と設定
let outputData = [Float](repeating: 0, count: inputData.count)
let outputDataBuffer = device.makeBuffer(bytes: outputData, length: outputData.byteLength, options: [])
computeCommandEncoder.setBuffer(outputDataBuffer, offset: 0, at: 1)
//スレッドグループの数、スレッドグループ内のスレッドの数を設定。これにより並列で実行される演算数が決定される
let width = 64
let threadsPerGroup = MTLSize(width: width, height: 1, depth: 1)
let numThreadgroups = MTLSize(width: (inputData.count + width - 1) / width, height: 1, depth: 1)
computeCommandEncoder.dispatchThreadgroups(numThreadgroups, threadsPerThreadgroup: threadsPerGroup)
//エンコーダーからのコマンドは終了
computeCommandEncoder.endEncoding()
//コマンドバッファを実行し、完了するまで待機
commandBuffer.commit()
commandBuffer.waitUntilCompleted()
//結果をresultDataに格納
let data = Data(bytesNoCopy: outputDataBuffer.contents(), count: outputData.byteLength, deallocator: .none)
var resultData = [Float](repeating: 1, count: outputData.count)
resultData = data.withUnsafeBytes {
Array(UnsafeBufferPointer<Float>(start: $0, count: data.count/MemoryLayout<Float>.size))
}
//結果の表示
print("[Input data]: \(inputData)")
print("[Result data]: \(resultData)")
}
}
//配列要素のバイト数を取得
private extension Array {
var byteLength: Int {
return self.count * MemoryLayout.size(ofValue: self[0])
}
}
ランダムな数が100個格納されたの配列のデータをGPU側に渡し、addOneという名前の関数で演算を行なった結果をコンソールに出力します。
なお、Metalベストプラクティスガイドに従い、アプリ起動時に宣言して保持すべき値と、実行ごとに宣言する値を区別しています。
次はシェーダーです。シェーダーはGPU側で実行されるコードなのですが、C++ベースのMetal Shading Languageで記述します。
#include <metal_stdlib>
using namespace metal;
kernel void addOne(const device float *inputData [[ buffer(0) ]],
device float *outputData [[ buffer(1) ]],
uint id [[ thread_position_in_grid ]])
{
float result = inputData[id];
result += 1.0f;
outputData[id] = result;
}
ここでは単に、入力に1を足して出力としています。idはスレッドの位置です。
なお、Metal用のシェーダーファイルは、
New -> File -> iOS -> Meta File
で作成することができます。
出力結果はこのようになります。
[Input data]: [26.0, 18.0, 19.0, 38.0, 44.0, 72.0, 51.0, 72.0, 56.0, 15.0, 16.0, 61.0, 46.0, 18.0, 54.0, 68.0, 18.0, 73.0, 44.0, 25.0, 29.0, 81.0, 63.0, 21.0, 77.0, 19.0, 59.0, 28.0, 73.0, 79.0, 94.0, 43.0, 82.0, 88.0, 10.0, 0.0, 88.0, 60.0, 17.0, 1.0, 77.0, 8.0, 44.0, 97.0, 51.0, 91.0, 5.0, 33.0, 24.0, 89.0, 63.0, 90.0, 60.0, 93.0, 77.0, 97.0, 36.0, 27.0, 84.0, 12.0, 35.0, 33.0, 38.0, 13.0, 34.0, 99.0, 58.0, 90.0, 3.0, 35.0, 40.0, 6.0, 66.0, 27.0, 29.0, 44.0, 26.0, 72.0, 7.0, 51.0, 30.0, 25.0, 75.0, 37.0, 27.0, 33.0, 52.0, 18.0, 3.0, 90.0, 55.0, 13.0, 64.0, 69.0, 95.0, 18.0, 16.0, 76.0, 56.0, 38.0]
[Result data]: [27.0, 19.0, 20.0, 39.0, 45.0, 73.0, 52.0, 73.0, 57.0, 16.0, 17.0, 62.0, 47.0, 19.0, 55.0, 69.0, 19.0, 74.0, 45.0, 26.0, 30.0, 82.0, 64.0, 22.0, 78.0, 20.0, 60.0, 29.0, 74.0, 80.0, 95.0, 44.0, 83.0, 89.0, 11.0, 1.0, 89.0, 61.0, 18.0, 2.0, 78.0, 9.0, 45.0, 98.0, 52.0, 92.0, 6.0, 34.0, 25.0, 90.0, 64.0, 91.0, 61.0, 94.0, 78.0, 98.0, 37.0, 28.0, 85.0, 13.0, 36.0, 34.0, 39.0, 14.0, 35.0, 100.0, 59.0, 91.0, 4.0, 36.0, 41.0, 7.0, 67.0, 28.0, 30.0, 45.0, 27.0, 73.0, 8.0, 52.0, 31.0, 26.0, 76.0, 38.0, 28.0, 34.0, 53.0, 19.0, 4.0, 91.0, 56.0, 14.0, 65.0, 70.0, 96.0, 19.0, 17.0, 77.0, 57.0, 39.0]
結果は入力に1を加えたものとなっており、正しく演算されていることが分かります。
それでは次に、CPUとパフォーマンスの比較を行ってみます。
Metal側、CPU側でそれぞれ同様の処理を行ない、所要時間を比較します。
まず、Metal側のコードです。
import UIKit
import Metal
class ViewController: UIViewController {
let inputDataCount = 10000
var device: MTLDevice!
var defaultLibrary: MTLLibrary!
var commandQueue: MTLCommandQueue!
var computePipelineState: MTLComputePipelineState!
override func viewDidLoad()
{
super.viewDidLoad()
//初期化
device = MTLCreateSystemDefaultDevice()
defaultLibrary = device.newDefaultLibrary()
commandQueue = device.makeCommandQueue()
let ml2Func = defaultLibrary.makeFunction(name: "addAndSubtract")!
computePipelineState = try! device.makeComputePipelineState(function: ml2Func)
}
@IBAction func calculate(sender: AnyObject)
{
//入力データ
var inputData:[Float] = []
for _ in 0...inputDataCount-1 {
inputData.append(Float(arc4random_uniform(UInt32(inputDataCount))))
}
//所要時間の測定開始
let startingDate = Date()
//コマンドバッファとエンコーダの作成と設定
let commandBuffer = commandQueue.makeCommandBuffer()
let computeCommandEncoder = commandBuffer.makeComputeCommandEncoder()
computeCommandEncoder.setComputePipelineState(computePipelineState)
//入力バッファの作成と設定
let inputDataBuffer = device.makeBuffer(bytes: inputData, length: inputData.byteLength, options: [])
computeCommandEncoder.setBuffer(inputDataBuffer, offset: 0, at: 0)
//出力バッファの作成と設定
let outputData = [Float](repeating: 0, count: inputData.count)
let outputDataBuffer = device.makeBuffer(bytes: outputData, length: outputData.byteLength, options: [])
computeCommandEncoder.setBuffer(outputDataBuffer, offset: 0, at: 1)
//スレッドグループの数、スレッドグループ内のスレッドの数を設定。これにより並列で実行される演算数が決定される
let width = 64
let threadsPerGroup = MTLSize(width: width, height: 1, depth: 1)
let numThreadgroups = MTLSize(width: (inputData.count + width - 1) / width, height: 1, depth: 1)
computeCommandEncoder.dispatchThreadgroups(numThreadgroups, threadsPerThreadgroup: threadsPerGroup)
//エンコーダーからのコマンドは終了
computeCommandEncoder.endEncoding()
//コマンドバッファを実行し、完了するまで待機
commandBuffer.commit()
commandBuffer.waitUntilCompleted()
//結果をresultDataに格納
let data = Data(bytesNoCopy: outputDataBuffer.contents(), count: outputData.byteLength, deallocator: .none)
var resultData = [Float](repeating: 1, count: outputData.count)
resultData = data.withUnsafeBytes {
Array(UnsafeBufferPointer<Float>(start: $0, count: data.count/MemoryLayout<Float>.size))
}
//結果の表示
print("[Time] \(Date().timeIntervalSince(startingDate))")
print("[Input data] Count: \(inputData.count), First value: \(inputData.first!), Last value: \(inputData.last!)")
print("[Result data] Count: \(resultData.count), First value: \(resultData.first!), Last value: \(resultData.last!)")
}
}
//配列要素のバイト数を取得
private extension Array {
var byteLength: Int {
return self.count * MemoryLayout.size(ofValue: self[0])
}
}
先程のコードとの違いは、配列の要素数が10000に増えた点と、所要時間測定のコードが書かれた点です。
所要時間の測定にはXcode標準のテスト機能を用いてもよかったのですが、極力シンプルなコードを書きたかったのでDate()を用いました。
シェーダーの方はこちらです。
#include <metal_stdlib>
using namespace metal;
constant int repeatCount = 100000;
kernel void addAndSubtract(const device float *inputData [[ buffer(0) ]],
device float *outputData [[ buffer(1) ]],
uint id [[ thread_position_in_grid ]])
{
float result = inputData[id];
for (uint i=0; i<repeatCount; i++){
result += 1.0f;
result -= 1.0f;
}
outputData[id] = result;
}
1を足して引くという意味のない処理をrepeatCountの数だけ繰り返します。
このrepeatCountを変えて、所要時間の変遷を記録します。
次に、CPU側のコードです。
import UIKit
class ViewController: UIViewController {
let dataCount = 10000
let repeatCount = 1000000
@IBAction func calculateWithFor()
{
var inputData:[Float] = []
for _ in 0..<dataCount {
inputData.append(Float(arc4random_uniform(UInt32(dataCount))))
}
let startingDate = Date()
var resultData = [Float](repeating: 0, count: inputData.count)
for i in 0..<dataCount
{
var result = inputData[i]
for _ in 0..<repeatCount
{
result += 1
result -= 1
}
resultData[i] = result
}
print("[Time] \(Date().timeIntervalSince(startingDate))")
print("[Input data] Count: \(inputData.count), First value: \(inputData.first!), Last value: \(inputData.last!)")
print("[Result data] Count: \(resultData.count), First value: \(resultData.first!), Last value: \(resultData.last!)")
}
}
ループを用いてMetal側のコードと同様の処理を行なっています。こちらも同様に、repeatCountを変えて所要時間の変遷を記録します。
なお測定する際にはEdit Schemeで以下の設定を行い、Metal側、CPU側ともに本来のパフォーマンスが発揮できるようにしておきます。
Build configurationをReleaseに。
GPU Frame CaptureとMetal API ValidationをDisabledに。
検証機にははiPhone 6 Plus、iOS10を用いました。内蔵されているチップはApple A8です。
測定結果はこちらです。
このままだと分かりにくいので、グラフにしてみます。
リピート数が大きいと、パフォーマンスに50-60倍もの極端に大きな差が出ることがわかります。
同じデータを対数表示にすると、このようになります。
リピート数が小さい場合はCPUの方がMetalよりも所要時間が短いことが分かります。
これは、処理のサイズが小さい場合は、処理そのものよりもGPU側へのデータの転送や設定が所要時間を決定するためだと考えれれます。
一般的に、GPU演算ではGPUへのドローコールを最低限にするのが望ましいとされていますが、スレッド内の処理のサイズが小さいと全処理中に占めるドローコールの割合が大きくなってしまいます。
Meatlの性能を十分に発揮するには、CPU/GPU間のやりとりを最小限にし、kernel関数である程度大きな処理を行うことが大事なようです。
GPUメーカー、NVIDIAのジェンスン・ファンCEOがGTC Japanで「脳はGPU」と述べていましたが、ひょっとしたらMetalの登場により個々のアプリに知性が宿る時代が近づいているのかもしれません。
Metalの能力をフルに引き出すためには、その動作原理を理解し、特性を把握することが大事ななので、今回は極力シンプルなMetalのコードを書いて実際に演算を行い、その特性を検証してみました。
次回は、Metalの応用例を紹介したいと考えています。
参考:
Metal for Developers
Metalベストプラクティスガイド
Metal Shading Language Guide
ゼロからはじめるGPUコンピューティング
歴代iOSデバイスのGPUまとめ
iOS Metalで流体シミュレーション！
Wikipedia PowerVR
Data-Parallel Programming with Metal and Swift for iPhone/iPad GPU
MetalKitでGPUを使いこなす
waifu2xをMetalで書いてみた
かつてSFで描かれた世界はもはやフィクションではない