頭と尻尾はくれてやる!

パソコンおやじのiPhoneアプリ・サイト作成・運営日記


MetalでGPUの処理結果をCPUで受け取る

Metalで画像処理っぽいことをしようといろいろと試しているところなんだけどさ。

Metalだと頂点シェーダとフラグメントシェーダ以外にもう一つコンピュートシェーダってのが登場するのよ。シェーダファイル内ではkernelで始めてる部分ね。
このコンピュートシェーダでテクスチャ画像をごにょごにょ好きに画像処理するって流れがApple謹製サンプルコードにあって、ああなるほどこういう風に使うのねってのはまあいいんだわ。

血迷った俺はコンピュートシェーダの処理結果をCPU側で取得しようと思ったのよ。
コンピュートシェーダの入出力はMTLTextureオブジェクトなんだけどサイズが異なる、というケース。入出力のサイズが同じだと選択肢は広がるはず(シェーダの出力をreturn なんちゃらでもいけるので)。

出力するデータの型について次の三つを合わせないといけないのよ。

1)MTLTextureDescriptorオブジェクトで指定するフォーマット
2)コンピュートシェーダの出力
3)CPU側で受け取るデータ

この三つのデータの型が合わないとビルドが通らなかったり、値が取れなかったりする。
例えば8bit符号無し整数とかだとシェーダでの出力ができずビルドが通らない。
16bit浮動小数点数だとCPU側(Objective-CでもSwiftでも)は32bitか64bitでないと受け取ることができない。そもそもコンピュートシェーダは32bitも64bitの出力はできない。
そういうのをリファレンスみながらどうやったらできんのよー!と探した結果、結局16bit符号無し整数しかできなかったんだけど、、、ホントかな。

処理の流れはこんな感じ。

1)MTLTextureDescriptorオブジェクトで指定するフォーマット
let textureDescriptor = MTLTextureDescriptor.texture2DDescriptorWithPixelFormat(MTLPixelFormat.R16Uint, width: width, height: height, mipmapped: false)
↑一番最初の引数でコンピュートシェーダが出力するデータのフォーマットを指定。R16Uintは16bitの符号無し整数を一つってこと。ここで指定したwidthとheightのサイズ分コンピュートシェーダがコールされる。
texture = device.newTextureWithDescriptor(textureDescriptor)
でMTLTextureオブジェクト(って言っていいのかな?)を作成ね。

2)コンピュートシェーダの出力
kernel void hogeComputeShader(texture2d<half, access::read>  inTexture   [[ texture(0) ]],
                      texture2d<unsigned short, access::write> outTexture  [[ texture(1) ]],
                      uint2                          gid         [[ thread_position_in_grid ]]) {
コンピュートシェーダの名称と引数記述の部分なんだけど、一つ目が入力画像、二つ目が出力画像を指定してる。
前にも言ったように入出力のテクスチャサイズが変わるので、頂点シェーダみたいにreturnで値を返すようだとダメなので、ここは”write”としてる。この時のデータ型が”unsigned short”という符号無しの16bit整数。

3)CPU側で受け取るデータ
GPU側での処理が終わったらそれをCPUで受け取ってみようと。MTLTexture Protocol のリファレンス見ると

getBytes:bytesPerRow:bytesPerImage:fromRegion:mipmapLevel:slice:
getBytes:bytesPerRow:fromRegion:mipmapLevel:

といったメソッドがあって簡単そうな(?)下の方でやってみた。
{
    typealias UInt16Pointer = UnsafeMutablePointer
    let pointer = UInt16Pointer.alloc(w*h)
    let region = MTLRegionMake2D(0,0 , w,h)
    texture.getBytes(pointer, bytesPerRow: w*sizeof(UInt16), fromRegion: region, mipmapLevel: 0)
    println("\(pointer[0])”)
}
こんな感じでCPU側で値として受け取ることはできた。めでたしめでたし。

そうそう、MTLCommandBufferのaddCompleteHandler:メソッド使ってGPUの処理がちゃんと終わってからアクセスしないとCPUで受け取れなかったので最初、はまったよ。


memo:iOS 8.1.3

<< Swiftの配列の初期化  TopPage  iPad用スタンドを買ってみた >>

コメント


管理者にだけ表示を許可する
 

トラックバック

トラックバックURL
http://ringsbell.blog117.fc2.com/tb.php/954-a4f20b15




Copyright ©頭と尻尾はくれてやる!. Powered by FC2 Blog. Template by eriraha.

FC2Ad