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
