OS XのOpenCL - その7 [OS XのOpenCL]
OS XでのOpenCLの使い方を解説したガイドの抄訳の続き。今日はOS XのOpenCLの一番特徴的な部分であるGCDのところ。もとのOpenCL規格では専用のキュー（cl_command_queue）を使うことになってるけどOS Xでは並列処理の基本であるGCD（Grand Central Dispatch）を使う。でもGCDはしょせん、現代的に便利になったスレッドプールにすぎなくて（ブロックを使うことでさらにめちゃ簡単になるけど）、OpenCLのキューと実態は全く別物のはずなんだけどどういう実装になってるのかよくわからない。単なるラッパでGCDを呼ぶようになってるAPIはシンタクスシュガーなだけなのかもしれない。ソースを見ろよ、と言われるかもしれないけど．．．．
SIMDのアルゴリズムによってはCPUの方が速かったりGPUの方が速かったりするし、あるGPUによっても違ってくる。OS X10.7にあるツールを使えばデータを処理できるデバイスのタイプを探すことができる。
OpenCLカーネルが走る環境を知るために、デフォルトのグローバルコンテクストを扱わなければならない。コンテクストはデバイスや、それらのデバイスのアクセス可能なメモリやカーネルの十個王をスケジュールする複数のキューに関する情報を知ることができる。
アプリはシステムにあるデバイスのタイプや、カーネルを実行するための最適な構成に関する推奨をコンテクストから得ることができる。コンテクストを得ると、アプリは特定のタイプのキューを生成したり、あるデバイスのキューを生成したりできる。
デバイスを見つけるためには
注意：gcl_create_dispatch_queue(...)で作ったディスパッチキューを使い終わったらdispatch_release(...)関数を呼べ。すべてのサンプルコードはこの呼び出しを含んでいる。
いったんキューを作ったら、必要な数のカーネルを投入してよい。あるいは違う特性を持った追加のキューを作ってもよい。
gcl_get_kernel_block_workgroup_info関数が返した推奨ワークグループサイズは、特定のデバイス上の特定のカーネルのcl_ndrange.local_work_sizeに使うことができる。
注意：gcl_get_kernel_block_workgroup_info関数はgcl_create_dispatch_queue関数で作ったGCDディスパッチキューで実行されるブロックの中で呼ばなければならない。
5 OpenCLとGCDを使う
OS X10.7ではOpenCLの開発者はOpenCLカーネルとして書いた作業を、OpenCLデバイスに裏付けられたGCD（Grand Central Dispatch）で待ち行列投入できる。OpenCLとGCDを以下のようにすることで使うことができる。- OpenCLアプリが走る環境を調べる。特にシステムにあるどのデバイスが特定のOpenCL計算にとってパフォーマンスが高いか調べる。
- システムにあるOpenCLが動作するデバイスの計算パワーや技術的な特徴を調べる
- GCDは特定のカーネルにとってどのOpenCLデバイスが最良か提案できる
- カーネルをどう構成すればいいか推薦を受けることができる。例えば特定のデバイスでのあるカーネルのワークグループの最適なサイズの提案を受けることが出る
- カーネルをキューに投入する
- ホストとOpenCLデバイス、そしてデバイス間の同期を取る。
ホストは全部のキューの作業が終わるのを待つか一つのキューが終わってから他のキューを走らせるか、などができる。
5.1 使用可能な計算デバイスを見つける
OpenCLのカーネルは一つの命令で複数のデータを処理する（SIMD）並列計算モデルを仮定している。SIMDでは大きなデータをかたまりに分けて、それぞれのかたまりに同じ計算をすることになる。SIMDのアルゴリズムによってはCPUの方が速かったりGPUの方が速かったりするし、あるGPUによっても違ってくる。OS X10.7にあるツールを使えばデータを処理できるデバイスのタイプを探すことができる。
OpenCLカーネルが走る環境を知るために、デフォルトのグローバルコンテクストを扱わなければならない。コンテクストはデバイスや、それらのデバイスのアクセス可能なメモリやカーネルの十個王をスケジュールする複数のキューに関する情報を知ることができる。
アプリはシステムにあるデバイスのタイプや、カーネルを実行するための最適な構成に関する推奨をコンテクストから得ることができる。コンテクストを得ると、アプリは特定のタイプのキューを生成したり、あるデバイスのキューを生成したりできる。
デバイスを見つけるためには
- gcl_get_context関数で「グローバルな」OpenCLコンテクストを得る。
注意：このコンテクストはOpnCLが作るが、retain/releaseしてはいけない（あからさまに作ったコンテクストは必要） - clGetDeviceIds(...)関数（これはOpenCLの標準API）でコンテクストパラメータとして得られたコンテクストを特定する。この呼び出しは接続されたOpenCLデバイスのIDのリストを返す。
- デバイスの特徴や能力によって違うタイプの作業をデバイスに送るように選択することができる。コンテクストからデバイスのIDがわかれば、clGetDeviceInfo()関数を呼んでそれぞれのデバイスの情報が得られる。リスト6-1のサンプルコードはベンダ名とデバイス名を得ている。またclGetDeviceInfo()で計算コアの数やキャッシュラインサイズなどの技術的な情報を得ることもできる。詳細はOpenCLの規格で。
- コンテクストはデバイス間でメモリオブジェクトを共有する必要がある。もしOS Xのgcl_なんとかというAPIを使えば解放する必要があるが、デフォルトのグローバルコンテクストを使えばコテンクストを生成する必要はない。
- OpenCL規格のAPIを使うなら、専用のコンテクストを作る必要はない。
- OpenCLコンテクストはOpenGLシェアグループとよく似ている。シェアグループとはCPUGとGPUの両方からアクセスできるメモリのブロックを使えるようにするためのツール一式である。
5.2 カーネルをディスパッチキューに投入する
OpenCLの作業のためにはOpenCLコンパチのディスパッチキューを使わなければならない。システムにある特定のデバイスのキューを作ることができるし、特定のタイプのデバイスのキューを作ることもできる。好きなだけ違うキューを作ることができる。- 特定のタイプのどんなデバイスでも動作するディスパッチキューを作るには、CL_DEVICE_TYPE_CPU,かCL_DEVICE_TYPE_GPUはCL_DEVICE_TYPE_ACCELERATORを最初の引数にしてgcl_create_dispatch_queue関数を呼び出す。
注意：作ったディスパッチキューは特定のデバイスタイプに接続されなければならない。CL_DEVICE_TYPE_DEFAULTというデフォルトデバイスタイプのOpenCLコンパチのキューを作ることはできない。
OS XのOpenCLは指定されたデバイスタイプに従ってGPUかあるいはCPUを使うディスパッチキューを作る。ひとつ以上のGPUが使えるなら、OS XのOpenCLは指定されたタイプで最大の計算コアを持つデバイスにカーネルを投入する。
注意：CL_DEVICE_TYPE_GPUを指定してディスパッチキューを一旦作ってしまったら、どのGPUを使うことになるか知らなくていい。ディスパッチキューに接続されたデバイスを知るにはgcl_get_device_id_with_dispatch_queue関数を使う。 - clGetDeviceIds関数で得たりclGetDeviceInfo関数で見つけたりした使いたいOpenCL デバイスIDを正確に知っていたら、CL_DEVICE_TYPE_USE_IDとデバイスのIDをcl_create_dispatch_queue関数に渡して呼ぶ。
注意：gcl_create_dispatch_queue(...)で作ったディスパッチキューを使い終わったらdispatch_release(...)関数を呼べ。すべてのサンプルコードはこの呼び出しを含んでいる。
いったんキューを作ったら、必要な数のカーネルを投入してよい。あるいは違う特性を持った追加のキューを作ってもよい。
5.3 デバイス上のカーネルの特性を決定する
あるデバイス上でカーネルがローカルとプライベートなメモリをどれだけ消費するかや最適なワークグループのサイズなど、カーネルとデバイスのペアに固有の情報を得たいならgcl_get_kernel_block_workgroup_info関数を使え。この情報は特定のデバイス上で走るカーネルのパフォーマンスのチューニングや、パフォーマンス上の問題をデバグするときに便利である。gcl_get_kernel_block_workgroup_info関数が返した推奨ワークグループサイズは、特定のデバイス上の特定のカーネルのcl_ndrange.local_work_sizeに使うことができる。
注意：gcl_get_kernel_block_workgroup_info関数はgcl_create_dispatch_queue関数で作ったGCDディスパッチキューで実行されるブロックの中で呼ばなければならない。
5.4 サンプルコード：ディスパッチキューの生成
つぎのリスト6-1はカーネルブロックからワークグループ情報を得る方法を示している。ピーク性能を得るために利用できる。
#include <stdio.h>
// OS X v10.7かそれ以降対応
#include <OpenCL/opencl.h>
// この例ではmykernel.cl.hはカーネルブロックの宣言を含んだヘッダである。
// カーネルのソースが違えばファイルの名前は違ってくる。
// このヘッダはXcodeが生成する
#include "mykernel.cl.h"
static void print_device_info(cl_device_id device) {
char name[128];
char vendor[128];
clGetDeviceInfo(device, CL_DEVICE_NAME, 128, name, NULL);
clGetDeviceInfo(device, CL_DEVICE_VENDOR, 128, vendor, NULL);
fprintf(stdout, "%s : %s\n", vendor, name);
}
// グローバルOpenCLコンテクストを得て、
// コンテクストが保持しているデバイスの情報をとる方法のデモ。
// また、デバイスタイプ（CPUかGPU）やOpenCLデバイスディレクトリを指定して
// ディスパッチキューを作る方法のデモ。
static void hello_world_sample1 ()
{
int i;
// グローバルOpenCLコンテクストを調べる
// 注意：特定のデバイスのキューを使うのでなければコンテクストを取ってくる必要はない
cl_context context = gcl_get_context();
// そのコンテクストにどんなデバイスが利用可能か問い合わせる
size_t length;
cl_device_id devices[8];
clGetContextInfo(context,
CL_CONTEXT_DEVICES,
sizeof(devices),
devices,
&length);
// これらのデバイスを渡って基本的な情報をプリントアウトする。
// デバイスに関するどんな情報が利用可能か問い合わせる
fprintf(stdout, "The following devices are available for use:\n");
int num_devices = (int)(length / sizeof(cl_device_id));
for (i = 0; i < num_devices; i++) {
print_device_info(devices[i]);
}
// この作業のために、OpenCLデバイスに関連付けられたディスパッチキューを作る必要がある
// システムにGPU（たぶん唯一のGPU）かあるいはCPUデバイスを返させる
// あるいは、cl_device_idで指定したディスパッチキューを作る。
// このデバイスIDは上のOpenCLコンテクストから来ている。みっつの例が下にある
// 1. GPUベースのキューを尋ねる。ここではデバイスIDを供給しない
// そのかわり、システムは一番使えるGPUを教える
dispatch_queue_t gpu_queue =
gcl_create_dispatch_queue(CL_DEVICE_TYPE_GPU, NULL);
// キューからデバイスを取る。そうするとそれに関するOpenCLの質問ができるようになる
// 戻り値がNULLかどうかを調べることで
// システムにOpenCLで使えるGPUがあったかということをチェックできる
if (gpu_queue != NULL) {
cl_device_id gpu_device =
gcl_get_device_id_with_dispatch_queue(gpu_queue);
fprintf(stdout, "\nAsking for CL_DEVICE_TYPE_GPU gives us:\n");
print_device_info(gpu_device);
} else {
fprintf(stdout, "\nYour system does not contain "
"an OpenCL-compatible GPU\n.");
}
// 2. CL_DEVICE_TYPE_CPUを試してみる。
// すべてのMacはCPUのOpenCLデバイスを持っているので、
// GPUの場合と違ってNULLチェックしなくてもいい
dispatch_queue_t cpu_queue =
gcl_create_dispatch_queue(CL_DEVICE_TYPE_CPU, NULL);
cl_device_id cpu_device =
gcl_get_device_id_with_dispatch_queue(cpu_queue);
fprintf(stdout, "\nAsking for CL_DEVICE_TYPE_CPU gives us:\n");
print_device_info(cpu_device);
// 3. あるいはたぶんコンテクストに見つかったデバイスのリストから特定のデバイスが欲しい場合だろう
// 次のような違いがある：
// CL_DEVICE_TYPE_USE_IDとdevice_idを渡す。この例では
// コンテクストにある最初のデバイスを使う。それがなんだったとしても。
dispatch_queue_t custom_queue =
gcl_create_dispatch_queue(CL_DEVICE_TYPE_USE_ID, devices[0]);
cl_device_id custom_device =
gcl_get_device_id_with_dispatch_queue(custom_queue);
fprintf(stdout,
"\nAsking for CL_DEVICE_TYPE_USE_ID and our own device gives us:\n");
print_device_info(custom_device);
// ということでカーネルを走らせることができる3つのディスパッチキューが手に入った
… // ここでカーネルを走らせる
// GCDのAPIを使ってキューを解放する
dispatch_release(custom_queue);
dispatch_release(cpu_queue);
if (gpu_queue != NULL)
dispatch_release(gpu_queue);
}
次のリスト6-2ではワークグループ情報を得る。
// このリストはワークグループ情報を得る方法を示している
// カーネルブロックからピークパフォーマンスを得るのに役に立つ
static void hello_world_sample2() {
// 2乗するカーネルを走らせるGPUのキューを得る
dispatch_queue_t queue =
gcl_create_dispatch_queue(CL_DEVICE_TYPE_GPU, NULL);
// それはGPUだった？ ダメだったらCPUを使う
if (queue == NULL) {
gcl_create_dispatch_queue(CL_DEVICE_TYPE_GPU, NULL);
}
// どっちにしても使うデバイスをプリントアウトする
fprintf(stdout, "\nExamining workgroup info for square_kernel on device ");
print_device_info(gcl_get_device_id_with_dispatch_queue(queue));
// その特定のデバイスでそのカーネルを実行するときにOpenCLが一番いいと判断した
// ワークグループサイズを調べる
// このメソッドはOpenCL用に作ったディスパッチキューを使ったブロックの中で
// 実行しないといけない
dispatch_sync(queue,
^{
size_t wgs, preferred_wgs_multiple;
cl_ulong local_memsize, private_memsize;
// 次のふたつの関数呼び出しは、ローカルとプライベートのメモリを
// カーネルがどのくらい使うかを教えてくれる
gcl_get_kernel_block_workgroup_info(square_kernel,
CL_KERNEL_LOCAL_MEM_SIZE,
sizeof(local_memsize),
&local_memsize, NULL);
fprintf(stdout, "Local memory size: %lld\n", local_memsize);
gcl_get_kernel_block_workgroup_info(square_kernel,
CL_KERNEL_PRIVATE_MEM_SIZE,
sizeof(private_memsize),
&private_memsize, NULL);
fprintf(stdout,
"Private memory size: %lld\n", private_memsize);
// このデバイスでカーネルにとっていちばんいいワークグループサイズを
// OepnCLに提案してもらう
gcl_get_kernel_block_workgroup_info(square_kernel,
CL_KERNEL_WORK_GROUP_SIZE,
sizeof(wgs), &wgs, NULL);
fprintf(stdout, "Workgroup size: %ld\n", wgs);
// 最後にワークグループサイズの倍数をOpenCLに尋ねる
// これはパフォーマンスヒントになる
// Finally, you can ask OpenCL for a workgroup size multiple.
// This is a performance hint.
gcl_get_kernel_block_workgroup_info(square_kernel,
CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE,
sizeof(preferred_wgs_multiple),
&preferred_wgs_multiple, NULL);
fprintf(stdout, "Preferred workgroup size multiple: %ld\n",
preferred_wgs_multiple);
// このワークグループサイズを使って
// カーネルを実行するときに望ましいcl_ndrange構造体を作ることができる
});
dispatch_release(queue);
}
int main(int argc, const char* argv[]) {
hello_world_sample1();
hello_world_sample2();
}
5.5 サンプルコード：カーネルのワークグループサイズを得る
リスト6-1にあるように、ホストはOpenCL用に作ったディスパッチキュー上のブロックの中でgcl_get_kernel_block_workgroup_infoを呼ぶことで、ローカルメモリサイズを要求できる。gcl_get_kernel_block_workgroup_info( square_kernel, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(local_memsize), &local_memsize, NULL);リスト6-2にあるように、gcl_get_kernel_block_workgroup_info関数を使ってそのデバイスのそのカーネルでの最適なワークグループサイズを返させる。
gcl_get_kernel_block_workgroup_info( square_kernel, CL_KERNEL_WORK_GROUP_SIZE, sizeof(workgroup_size), &workgroup_size, NULL); fprintf(stdout, "Workgroup size: %ld\n", workgroup_size);最後に、デバイスの性能に従ってワークグループサイズの倍数を gcl_get_kernel_block_workgroup_info関数で調べる。
gcl_get_kernel_block_workgroup_info( square_kernel, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, sizeof(preferred_workgroup_size_multiple), &preferred_workgroup_size_multiple, NULL);このワークグループサイズを使って実行するカーネルにとって最適なcl_ndrangeを決めることができる。
cl_ndrange range = {
1, // 次元の数
{0, 0, 0}, // 各ディメンジョンのオフセット。
// 全部のデータを処理したければすべて0にする。
// ワークグループがみっつの次元よりも小さいとしても
// みっつの次元のオフセットを常に渡すように
{NUM_VALUES, 0, 0}, // グローバルレンジ。
// これはそれぞれの次元でいくつ処理をしたいかを指定する
// これも3次元よりも少なくてもみっつ指定する
{workgroup_size, 0, 0 } // それぞれのワークグループのローカルサイズ
// これはワークグループごとにいくつワークアイテムがあるかを指定する
// ワークグループの数は（グローバルサイズ / ローカルサイズ）になる
// この場合は、（NUM_VALUE / workgroup_size）個のワークグループがあることになる
// これも3次元指定する
};
2015-05-17 22:09
nice!(0)
コメント(0)
トラックバック(0)
コメント 0