OS XのOpenCL - その23 [OS XのOpenCL]
前回の続きである。もう2ヶ月近くまえのことだけど、ちゃんと覚えているのである。僕にしてはすごいのである。ということでOpenCLのプログラミングガイドに出てくる言葉を僕なりに整理するのである。
残りのメンバはあとで。
2次元では
例えばRGBA画像を渡してそのピクセルごとになんらかの処理をするカーネルの場合、画像上のxとyのインデクスを得て、そのピクセル位置を処理するように書けばいい。
これがそれぞれのワークアイテムに割り振られるが、global_work_sizeで指定された個数とおなじだけ計算エレメントがあるわけではない。global_work_sizeで指定された個数よりも計算エレメントの数が多いとき、残りの計算エレメントは別のことができるか、というとそうではなくてただ遊ぶだけになる。またglobal_work_sizeよりも計算エレメントの数が少ないとき、全部が終わるまで並列可能な演算をループさせて実行されるようである。
そういった計算エレメントへの割り振り方はプログラマからは隠蔽されている。
しかし、GPUのハードウェアで決まっている場合もあるので、効率を追求するならハードウェアとアルゴリズムに最適な値を選ばなければならないようである。
難しいことを考えなければNULLを渡せばOpenCLがよきに計らってくれるらしいし、ハードウェアに最適な値をOpenCLに問い合わせることもできるが、効率はカーネルのアルゴリズムにも依存するはずである。難しいのでよくわからない。
またグローバルIDを何次元にするか、それぞれの次元の大きさをいくつにするかで、最適なlocal_work_sizeの値は違ってくるように思えるんだけど、その辺のところもよくわからない。
必然性が理解できないので、どうすればいいのかもよくわからない。まあとりあえずNULLを渡しておくことにする。
もちろん、文字通りの意味ではないだろう。
15.2.10 ND-Range NDレンジ
ホストがワークアイテムごとに処理する内容を指定するための構造体。定義は次のようなものである。typedef struct _cl_ndrange { size_t work_dim; size_t global_work_offset[3]; size_t global_work_size[3]; size_t local_work_size[3]; } cl_ndrange;まずワークアイテムを何次元に構成するか、という構造体のメンバwork_dimがある。1、2、3次元の3通りが使える。もちろん計算エレメントが3次元に配置されているわけはないので、プログラミングモデル上の話である。ということは扱うデータが2次元の画像であってもホスト側でピクセルを1列に並べれば1次元としてNDレンジを定義できるし、その逆もあり得る。プログラミング上の使い勝手で選べばよい。
残りのメンバはあとで。
15.2.11 get_global_id()関数
カーネルコードの中でワークアイテムが自分のインデクスを得る関数。NDレンジ構造体で指定された形で返される。つまり1次元の場合size_t index = get_global_id(0);でそのカーネルがどのワークアイテムで実行されているかを知ることができる。
2次元では
size_t xIndex = get_global_id(0); size_t yIndex = get_global_id(1);のように、3次元では同様になる。
例えばRGBA画像を渡してそのピクセルごとになんらかの処理をするカーネルの場合、画像上のxとyのインデクスを得て、そのピクセル位置を処理するように書けばいい。
15.2.12 global_work_sizeメンバ
NDレンジ構造体のみっつめのメンバ。全部でいくつのワークアイテムが処理すべきかを指定する。次元ごとに指定できて、2次元の場合は合計で(global_work_size[0] × global_work_size[1])個になる。3次元でも同様である。これがそれぞれのワークアイテムに割り振られるが、global_work_sizeで指定された個数とおなじだけ計算エレメントがあるわけではない。global_work_sizeで指定された個数よりも計算エレメントの数が多いとき、残りの計算エレメントは別のことができるか、というとそうではなくてただ遊ぶだけになる。またglobal_work_sizeよりも計算エレメントの数が少ないとき、全部が終わるまで並列可能な演算をループさせて実行されるようである。
そういった計算エレメントへの割り振り方はプログラマからは隠蔽されている。
15.2.13 global_work_offsetメンバ
NDレンジ構造体のふたつめのメンバ。どこから開始するかを指定する。たとえばndRange.global_work_offset[0] = 173;と指定することはカーネルコードの中で
Index = get_global_id(0) + 173;とするのとまったくおなじのようである。わざわざこのような屋上屋を作るのは、定数演算はなるべくカーネルの外で行ってカーネルの効率を上げよう、ということなのかもしれない。このへんよくわからない。
15.2.14 local_work_sizeメンバ
NDレンジ構造体の最後のメンバ。これが一番難しい。メモリ階層にローカルメモリというのがあった。これはひとつのワークグループに属する計算エレメントが共有するメモリだった。このlocal_work_sizeはワークグループに計算エレメントがいくつ属するかを決める値のようである。しかし、GPUのハードウェアで決まっている場合もあるので、効率を追求するならハードウェアとアルゴリズムに最適な値を選ばなければならないようである。
難しいことを考えなければNULLを渡せばOpenCLがよきに計らってくれるらしいし、ハードウェアに最適な値をOpenCLに問い合わせることもできるが、効率はカーネルのアルゴリズムにも依存するはずである。難しいのでよくわからない。
またグローバルIDを何次元にするか、それぞれの次元の大きさをいくつにするかで、最適なlocal_work_sizeの値は違ってくるように思えるんだけど、その辺のところもよくわからない。
必然性が理解できないので、どうすればいいのかもよくわからない。まあとりあえずNULLを渡しておくことにする。
15.2.15 ウォームアップ
GPUのチューニングのところは面白かったけど、そこにパフォーマンス測定のためにはウォームアップが必要だ、ということが書いてある。これはどういうことだ?もちろん、文字通りの意味ではないだろう。
2015-10-11 15:34
nice!(0)
コメント(0)
トラックバック(0)
コメント 0