SSブログ

OS XのOpenCL - その25 [OS XのOpenCL]

OS X独自のOpenCL拡張についてずっと見てきた。難しいことを考えずに簡単に使えるように、という配慮がみられる。それはAppleのOS X戦略の一環で、シェアの低いOS Xがシェアを増やすためにするべきことをしているにすぎない。しかしたとえAppleがいろいろ配慮しても、OpenCL、というかGPUのパフォーマンスを十分に引き出して、CPUより速くかつCPUの負荷を下げるには結局、かなりのチューニングが必要で、そのためにはそれなりの知識と、最適化の実験が必要だということがわかってきた。

今日のところはAppleのガイドにも触れられていない、OpenCLのコードがあったときXcodeが自動生成するふたつのファイルが何なのか中身を見てみる。

16  OS XでのOpenCLの拡張はなにをしているのか

OpenCL関数のうちcl_で始まる関数はOpenCLオリジナルで、gcl_で始まるものはAppleの独自拡張の関数である。gcl_にはcl_と似たような機能の関数が並んでるけど、GCDを使うようになっているのと、メモリ確保などオリジナルより簡単にできるようになっている。

それはいいとして、Apple独自拡張ではそれだけではなくて.clのカーネルソースをXcodeでコンパイルするとふたつのソースが自動生成される。例えばCopyKernel.clをコンパイルすると
CopyKernel.cl.h
CopyKernel.cl.c
ができる。これは何なのか。具体的に見てみる。

CopyKernel.clとしてこんなの
kernel void copy(global const float * in,
                 global float * out,
                 int w,int h)
{
    int x = get_global_id(0); // ワークアイテム内の
    int y = get_global_id(1); // ピクセルの位置 = (x,y)
    out[x+y*w] = in[x+y*w]; // 読み込んで書き出す
}
を書いたとしよう。このカーネルはただ2次元のfloatをinからoutにコピーするだけである。これをXcodeでコンパイルしてできたCopyKernel.cl.hが
/***** GCL Generated File *********************/
/* Automatically generated file, do not edit! */
/**********************************************/

#include <OpenCL/opencl.h>
extern void (^copy_kernel)(const cl_ndrange *ndrange, cl_float* in, cl_float* out, cl_int w, cl_int h);
こんなものである。カーネルの名前copyをcopy_kernelとして、GCDのキューの中で呼び出すことができるブロックの外部宣言になっている。

もう一方のCopyKernel.cl.cはというとちょっと長いけど全部引用すると
/***** GCL Generated File *********************/
/* Automatically generated file, do not edit! */
/**********************************************/

#include <stdio.h>
#include <stdlib.h>
#include <assert.h>
#include <dispatch/dispatch.h>
#include <OpenCL/opencl.h>
#include <OpenCL/gcl_priv.h>
#include "CopyKernel.cl.h"

static void initBlocks(void);

// Initialize static data structures
static block_kernel_pair pair_map[1] = {
      { NULL, NULL }
};

static block_kernel_map bmap = { 0, 1, initBlocks, pair_map };

// Block function
void (^copy_kernel)(const cl_ndrange *ndrange, cl_float* in, cl_float* out, cl_int w, cl_int h) =
^(const cl_ndrange *ndrange, cl_float* in, cl_float* out, cl_int w, cl_int h) {
  int err = 0;
  cl_kernel k = bmap.map[0].kernel;
  if (!k) {
    initBlocks();
    k = bmap.map[0].kernel;
  }
  if (!k)
    gcl_log_fatal("kernel copy does not exist for device");
  kargs_struct kargs;
  gclCreateArgsAPPLE(k, &kargs);
  err |= gclSetKernelArgMemAPPLE(k, 0, in, &kargs);
  err |= gclSetKernelArgMemAPPLE(k, 1, out, &kargs);
  err |= gclSetKernelArgAPPLE(k, 2, sizeof(w), &w, &kargs);
  err |= gclSetKernelArgAPPLE(k, 3, sizeof(h), &h, &kargs);
  gcl_log_cl_fatal(err, "setting argument for copy failed");
  err = gclExecKernelAPPLE(k, ndrange, &kargs);
  gcl_log_cl_fatal(err, "Executing copy failed");
  gclDeleteArgsAPPLE(k, &kargs);
};

// Initialization functions
static void initBlocks(void) {
  const char* build_opts = "";
  static dispatch_once_t once;
  dispatch_once(&once,
    ^{ int err = gclBuildProgramBinaryAPPLE("OpenCL/CopyKernel.cl", "", &bmap, build_opts);
       if (!err) {
          assert(bmap.map[0].block_ptr == copy_kernel && "mismatch block");
          bmap.map[0].kernel = clCreateKernel(bmap.program, "copy", &err);
       }
     });
}

__attribute__((constructor))
static void RegisterMap(void) {
  gclRegisterBlockKernelMap(&bmap);
  bmap.map[0].block_ptr = copy_kernel;
}
というようなCのコードになっている。

この中身は
  1. initBlocks()というstatic宣言された関数
  2. pair_mapというやはりstatic宣言された構造体データ
  3. 同様にbmapという構造体データ
  4. copy_kernelブロック本体
  5. initBlocks()の定義
  6. RegisterMap()というstatic関数
になっている。

GCDのキューの中でカーネル(つまりこの場合copy_kernel)を呼び出すと、この中のcopy_kernelブロックが呼ばれることになる。このブロックでは
  1. 初期化
  2. 引数の設定
  3. OpenCL標準の形でのカーネルの呼び出し
  4. 引数の解放
をしていることがなんとなくわかる。

この初期化というのは本体がinitBlocks()関数で、その中身は
  1. カーネルのコンパイル
  2. bmap構造体へのコンパイル済みカーネルコードの設定
をやってることがわかる。pair_map、あるいはbmapというデータはgclExecKernelAPPLE()という関数に渡して実行させるためのブロック定義とカーネルコードを結びつけるためのものであることが想像できる。

ようするにそれほど難しいことはしていない。自動生成って人が読み書きするのが大変だったり、より上位の記述からブレークダウンされたコードを作るために使われることが多いので、なんとなく肩透かしみたいな感じのコードではある。

なぜいったんCのコードを自動生成してそれをコンパイルするような形にしたのか、というとカーネルへの引数の型と数によらず.cl.hに定義されたブロックを呼べばいいだけにするためだろう。OS XでOpenCLを使いたいユーザにとってあまり何も考えずに使えるように、というAppleらしいやりかたではある。

Appleのたくさんあるframeworkのうち、かなりのものがこういう思想でできてる。ようするに素人さんでもとりあえず難しいことを考えずに使えますよ、という感じである。Linuxカーネルのプログラマみたいなガチガチのテクニシャンから見ると素人臭くてやりたいことができない、という感じだろうけど、僕みたいな素人の、しかも記憶力理解力思考力が低下したジジイにはありがたい。

でもそうやってスポイルされてOS Xしか使えないヘタれプログラマになってしまうんだなあ、と思う。もうすでに十分遅いけど。

nice!(0)  コメント(0)  トラックバック(0) 

nice! 0

コメント 0

コメントを書く

お名前:
URL:
コメント:
画像認証:
下の画像に表示されている文字を入力してください。

トラックバック 0

面白い記事連休二日目だが... ブログトップ

この広告は前回の更新から一定期間経過したブログに表示されています。更新すると自動で解除されます。