Metal で算術演算(MTLComputeCommandEncoder) を行った際のデバッグ方法

Metal といえばテクスチャなど GPU を使ったレンダリングの処理をおこなうものと思う方が大半かと思います。そのような利用の時は「Metalのデバッグまとめ（随時更新）」に書かれているようなデバッグ方法でことたります。

今回は算術演算で GPU を使いたくて MTLComputeCommandEncoder をつかったのですが、Instruments で表示しても上記のような情報がなくデバッグの方法が全くわからなかったので自分なりに試して、ある程度目処がたちましたので、その方法のメモとなります。

また、このエントリーでは特に Metal についてはふれませんので、さっと学びたい方は「iOSDC2017で「飛び道具ではないMetal」という話をしました #iOSDC 」を読んでもらえれば大枠は理解できるかと思います。

サンプルコードと問題

それでは早速 noppoMan/Shaders.metal のコードで試していきましょう。 Shaders.metal の kernel は Metal で実装されたシグモイド関数です。何か計算しているんだなレベルで問題ありません。

これをプロジェクトに追加して試したときにデバッグできなくて困りました。簡単に紹介するとkernel 内で breakpoint がはれない、print しようとしてもそんな関数がない、stdio を import しようとするとそんなものはない、となりどのようにデバッグすればいいんだと発狂しそうになりました。

これについてはある程度、挙動が確認できた状態のものを kernel に移植すれば、まだある程度動作は担保されます。しかし実行時のスレッドの情報(thread_position_in_gridなど)は確認することができないのは困ります。

見かけたデバッグ

「【Swift Metal】thread_position_in_grid等の属性について解説」のエントリーを見て out buffer につめれば、実行時のスレッドの情報(thread_position_in_threadgroup, thread_position_in_gridなど)を取得できるのでと思い試してしたが、取得できはするのですがデバック時に毎回これをやるのは正直億劫です。

MTLCaptureManager を使用

そんな中「Capturing GPU Command Data Programmatically」というドキュメントにたどりつきましたが、どうせ MTLComputeCommandEncoder だとまともにキャプチャされないんだろうなと思いつつ試してみましたが、なんとこれで確認することができました！！

使い方は簡単で MTLCaptureManager.shared().startCapture(with: ) を実行するだけです。 sigmoid_on_gpu.swift に反映すると下記のようになります。

func sigmoid_on_gpu(_ input: [Float]) throws -> [Float] {
    var input = input

+    let captureDescriptor = MTLCaptureDescriptor()
+    captureDescriptor.captureObject = device
+    try! MTLCaptureManager.shared().startCapture(with: captureDescriptor)

...（省略）

+    MTLCaptureManager.shared().stopCapture()
    return output
}

実行すると gputrace が開かれ情報を見ることができます。

f:id:dealforest:20200602193333p:plain

上記のような画面が表示され Buffer をダブルクリックすると kernel への in, out を確認することができます。便利ですね。あとは左側のナビゲーションエリアにある ComputeCommandEncoder を選択するとプレビューが表示されサマリーが表示されます。また、このプレビューからも out を表示することができます。

f:id:dealforest:20200602194158p:plain