Genta Hirauchi

公開日:2020/03/09
更新日:2020/03/23

【Swift Metal】MetalでGPGPUを実装する方法を解説

  • MetalでGPGPUを実装する方法が知りたい。

Metalとは、GPU(Graphics Processing Unit)へのローレベルなアクセスにより、グラフィックスや演算処理を実装する際に、非常に高いパフォーマンスを発揮してくれるフレームワークです。

本記事では、そのうちのMetalによる演算処理(GPGPU)の実装方法を解説致します。

目次

MetalでGPGPUを実装する方法

基本クラスの紹介

GPGPUの実装の前に、Metalフレームワークに用意されている基本的なクラスの簡単な紹介を致します。

MTLDevice
グラフィックスの描画や並列処理の実行の際に使用される、るGPUへのMetalインターフェイスです。
MTLCreateSystemDefaultDevice()で取得することができます。

MTLLibrary
MTL(Metal Shading Language)で実装された関数のコレクションです。
MTLDeviceのmakeDefaultLibrary()等で取得することができます。

MTLFunction
MTL(Metal Shading Language)で実装された関数オブジェクトです。
MTLLibraryのmakeFunction()で取得することができます。

MTLCommandQueue
GPUによって実行されるコマンドバッファーをの実行順を管理するキューです。
MTLDeviceのmakeCommandQueue()で取得することができます。

MTLComputePipelineState
コンパイルされた計算プログラムを参照するために使用されるオブジェクトです。
MTLDeviceのmakeComputePipelineState()で取得することができます。

MTLCommandBuffer
GPUで実行するためにエンコードされたコマンドを、格納するコンテナーです。
MTLCommandQueueのmakeCommandBuffer()で取得することができます。

MTLComputeCommandEncoder
並列演算処理のコマンドをエンコードするために使用されるオブジェクト。
MTLCommandBufferのmakeComputeCommandEncoder()で取得することができます。

GPGPUの実装

以下は、Floatの配列を、別のFloatの配列に移し変えるという処理を、CPUとGPUで実装したサンプルコードです。

  • ViewController
  • Shader.metal
import UIKit

// Metalをインポートする
import Metal

class ViewController: UIViewController {

    // MTLDeviceの生成
    private let device = MTLCreateSystemDefaultDevice()!
    private var library: MTLLibrary!
    private var commandQueue: MTLCommandQueue!
    private var computePipelineState: MTLComputePipelineState!

    private let yCount = 10
    private let xCount = 10

    private var inputData: [Float] = []

    override func viewDidLoad() {
        super.viewDidLoad()

        for _ in 0..<yCount*xCount {
            inputData.append(Float.random(in: 0..<10))
        }

        initMetal()
        startCPU()
        startGPU()
    }

    // Metalの初期化処理
    private func initMetal() {
        // (1) MTLLibraryの生成
        guard let libUrl = Bundle.main.url(forResource: "default.metallib", withExtension: nil) else { return }
        library = try! device.makeLibrary(URL: libUrl)

        // (2) MTLFunctionの生成
        let function = library.makeFunction(name: "testFunction")!

        // (3) MTLComputePipelineStateの生成
        computePipelineState = try! device.makeComputePipelineState(function: function)

        // (4) MTLCommandQueueの生成
        commandQueue = device.makeCommandQueue()!
    }

    // CPU処理
    private func startCPU() {
        var outputData = [Float](repeating: 0, count: inputData.count)

        let start = Date().timeIntervalSince1970

        for y in 0..<yCount {
            for x in 0..<xCount {
                outputData[y * xCount + x] = inputData[y * xCount + x]
            }
        }

        let end = Date().timeIntervalSince1970
        print("CPU || outputData.first: \(outputData.first ?? 0), outputData.last: \(outputData.last ?? 0), time: " + String(format: "%.5f ms", (end - start) * 1000))
    }

    // GPU処理
    private func startGPU() {
        var outputData = [Float](repeating: 0, count: inputData.count)

        let start = Date().timeIntervalSince1970

        // (5) 入力バッファと出力バッファの生成
        let inputBuffer = device.makeBuffer(bytes: inputData, length: MemoryLayout<Float>.stride * inputData.count, options: [])
        let outputBuffer = device.makeBuffer(bytes: outputData, length: MemoryLayout<Float>.stride * outputData.count, options: [])

        // (6) MTLCommandBufferの生成
        let commandBuffer = commandQueue.makeCommandBuffer()!

        // (7) MTLComputeCommandEncoderの生成
        let computeCommandEncoder = commandBuffer.makeComputeCommandEncoder()!
        computeCommandEncoder.setComputePipelineState(computePipelineState)
        computeCommandEncoder.setBuffer(inputBuffer, offset: 0, index: 0)
        computeCommandEncoder.setBuffer(outputBuffer, offset: 0, index: 1)

        // (8) スレッドグループ数、スレッドの数の設定
        let width = computePipelineState.threadExecutionWidth
        let threadgroupsPerGrid = MTLSize(width: (outputData.count + width - 1) / width, height: 1, depth: 1)
        let threadsPerThreadgroup = MTLSize(width: width, height: 1, depth: 1)
        computeCommandEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)

        // (9) エンコードの終了
        computeCommandEncoder.endEncoding()

        // (10) コマンドバッファを実行
        commandBuffer.commit()
        commandBuffer.waitUntilCompleted()

        // (11) 結果の取得
        let resultData = Data(bytesNoCopy: outputBuffer!.contents(), count: MemoryLayout<Float>.stride * outputData.count, deallocator: .none)
        outputData = resultData.withUnsafeBytes { Array(UnsafeBufferPointer(start: $0.baseAddress!.assumingMemoryBound(to: Float.self ), count: $0.count / MemoryLayout<Float>.size)) }

        let end = Date().timeIntervalSince1970
        print("GPU || outputData.first: \(outputData.first ?? 0), outputData.last: \(outputData.last ?? 0), time: " + String(format: "%.5f ms", (end - start) * 1000))
    }
}
#include 
using namespace metal;

kernel void testFunction(const device float* inputData [[ buffer(0) ]],
                         device float* outputData      [[ buffer(1) ]],
                         uint thread_position_in_grid [[thread_position_in_grid]] // Gridにおけるthreadの位置
) {
    outputData[thread_position_in_grid] = inputData[thread_position_in_grid];
}

(1) では、MTLLibraryの生成を行っております。
default.metallibというファイルのURLを参照しております。このファイルは、metalファイルを作成しビルドを行うと、アプリが自動的に作成するファイルです。
このURLを引数にし、MTLDeviceのmakeLibrary()でMTLLibraryを生成しております。

(2) では、MTLFunctionの生成を行っております。
makeFunction()の引数には、Shader.metalファイルに実装した関数名を指定しております。

(3) では、MTLComputePipelineStateの生成を行っております。
makeComputePipelineState()の引数には、先ほど生成したMTLFunctionを指定します。

(4) では、MTLCommandQueueの生成を行っております。

(5) では、Floatの配列を元に、入力バッファと出力バッファの生成を行っております。

(6) では、MTLCommandBufferの生成を行っております。

(7) では、MTLComputeCommandEncoderの生成を行っております。
そして、setComputePipelineState()でMTLComputePipelineStateの設定と、
setBuffer()でバッファの設定を行っております。

(8) では、スレッドグループ数、スレッドの数の設定を行っております。
スレッド数の設定に関しましては、公式ドキュメントのCalculating Threadgroup and Grid Sizesという記事にてわかりやすく解説されておりますので、そちらをご覧頂けたらと思います。

(9) では、エンコードを終了メソッドを呼んでおります。

(10) では、コマンドバッファを実行しております。

(11) では、outputBufferから、実行結果を取得しております。

Shader.metalでは、thread_position_in_gridをindexにし、inputDataをoutputDataに移す処理を実行しております。

GPUの処理では、CPUの処理のようなfor文が実装されておりませんね。これは、GPUでは複数スレッドによる並列処理が行われるからです。CPUでは10*10で100回の移し変え処理を行いますが、GPUでは、100のスレッドが平行して1回の処理を行うイメージです。

ここまでが、サンプルコードの大まかな内容です。

以下は、yCountとxCountの値を変更して実行した結果です。

10 * 10
CPU || outputData.first: 3.0272727, outputData.last: 0.8784741, time: 0.27895 ms
GPU || outputData.first: 3.0272727, outputData.last: 0.8784741, time: 4.32491 ms

50 * 50
CPU || outputData.first: 9.532412, outputData.last: 7.8694916, time: 4.61698 ms
GPU || outputData.first: 9.532412, outputData.last: 7.8694916, time: 2.49100 ms

100 * 100
CPU || outputData.first: 2.0698965, outputData.last: 1.419571, time: 8.14319 ms
GPU || outputData.first: 2.0698965, outputData.last: 1.419571, time: 2.00367 ms

500 * 500
CPU || outputData.first: 2.014732, outputData.last: 5.279184, time: 96.63892 ms
GPU || outputData.first: 2.014732, outputData.last: 5.279184, time: 2.34222 ms

1000 * 1000
CPU || outputData.first: 3.9278536, outputData.last: 0.89728, time: 380.40805 ms
GPU || outputData.first: 3.9278536, outputData.last: 0.89728, time: 4.67777 ms

5000 * 5000
CPU || outputData.first: 3.1419396, outputData.last: 7.6604886, time: 9489.20012 ms
GPU || outputData.first: 3.1419396, outputData.last: 7.6604886, time: 62.89768 ms

10*10の場合はCPUの方がパフォーマンスがいい結果となりました。このように、並列処理が少ないと、CPUで処理させる方がいい場合もあります。

それ以外では、GPUの方がパフォーマンスが良かったです。リピート数が増えれば増えるほど、その差が顕著に現れる結果となりました。

まとめ

  • MTLDeviceは、MTLCreateSystemDefaultDevice()で取得できる
  • MTLLibraryは、MTLDevice.makeDefaultLibrary()等で取得できる
  • MTLFunctionは、MTLLibrary.makeFunction()で取得できる
  • MTLCommandQueueは、MTLDevice.makeCommandQueue()で取得できる
  • MTLComputePipelineStateは、MTLCommandQueue.makeComputePipelineState()で取得できる
  • MTLCommandBufferは、MTLDevice.makeCommandBuffer()で取得できる
  • MTLComputeCommandEncoderは、MTLCommandBuffer.makeComputeCommandEncoder()で取得できる
  • metalファイルを作成すると、自動でdefault.metallibファイルが作成される
  • thread_position_in_gridには、Gridにおけるthreadの位置が格納されている
  • 並列処理が少ないと、CPUの方がパフォーマンスがいい場合がある