OS XのOpenCL - その13 [OS XのOpenCL]
OS X用のOpenCLのガイド訳も前回で峠は越した。あともうちょっと。今日からはサンプルコードが続く。今日のサンプルコードが理解できればOS XのOpenCLは、おおまかにはわかったと言っていいんではないかな。
8.5 例：バッファオブジェクトを確保して利用して解放する
つぎのリスト9-1は入力と出力に一つずつバッファを作って、入力用を初期化して、各値を2乗するカーネル（リスト9-2）を呼んで、最後に結果をチェックしている。
#include <stdio.h>
#include <stdlib.h>
#include <OpenCL/opencl.h>
#include "kernels.cl.h"
#define COUNT 2048
static void display_device(cl_device_id device)
{
char name_buf[128];
char vendor_buf[128];
clGetDeviceInfo(device, CL_DEVICE_NAME,
sizeof(char)*128, name_buf, NULL);
clGetDeviceInfo(device, CL_DEVICE_VENDOR,
sizeof(char)*128, vendor_buf, NULL);
fprintf(stdout, "Using OpenCL device: %s %s\n",
vendor_buf, name_buf);
}
static void buffer_test(const dispatch_queue_t dq)
{
unsigned int i;
// ホストとOpenCLデバイスの同期を取るためにセマフォを使う
dispatch_semaphore_t dsema = dispatch_semaphore_create(0);
// 入力データをホスト側で作る
cl_float* host_input = (float*)malloc(sizeof(cl_float) * COUNT);
// 初期データで埋める
for (i=0; i<COUNT; i++)
host_input[i] = (cl_float)i;
// この単精度浮動小数点数の2乗をOpenCLを使ってやる
// まず、OpenCLデバイスに入力を保持させる
// この場合出力に同じバッファを使いまわしてもいいが
// 別のものを使うことにする
// メモリ確保 1: 入力を保持できる大きさのバッファを作る
// CL_MEM_COPY_HOST_PTRフラグを使ってホスト側の入力データを渡していることに注意
// これは与えたホストデータでデバイス側のメモリ領域を初期化するように指示することになる
// device-side memory region with the supplied host data.
void* device_input =
gcl_malloc(sizeof(cl_float)*COUNT, host_input,
CL_MEM_COPY_HOST_PTR);
// メモリ確保 2: カーネルが計算した結果を保持するバッファを作る
void* device_results = gcl_malloc(sizeof(cl_float)*COUNT, NULL, 0);
// で、OpenCLに作業を送る準備が揃った
// これはホストアプリとは非同期に実行される
dispatch_async(dq, ^{
cl_ndrange range = {
1, // 1次元である
{0}, // 先頭から始める
{COUNT}, // COUNT個のワークアイテムを実行する
{0} // ワークグループへの分割はOpenCLに任せる
};
square_kernel(
&range, (cl_float*) device_input,
(cl_float*) device_results );
// この時点で計算は終わっているが、
// 結果はまだデバイスの上にある
// ホストで結果をチェックするには
// ホストメモリ空間に結果を書き戻さなくてはならない
// ホスト側の入力バッファを再利用する
gcl_memcpy(host_input, device_results, COUNT * sizeof(cl_float));
// ということで、ディスパッチセマフォを立てて
// ホストに継続できることを知らせる
dispatch_semaphore_signal(dsema);
});
// ここでホストはOpenCLデバイスがカーネルの計算をしているあいだ
// それに無関係な他のことができる
// しかしここではOpenCLが終わるのを待つ
dispatch_semaphore_wait(dsema, DISPATCH_TIME_FOREVER);
// 結果のテスト
int results_ok = 1;
for (i=0; i<COUNT; i++)
{
cl_float truth = (cl_float)i * (cl_float)i;
if (host_input[i] != truth) {
fprintf(stdout,
"Incorrect result @ index %d: Saw %1.4f, expected %1.4f\n\n",
i, host_input[i], truth);
results_ok = 0;
break;
}
}
if (results_ok)
fprintf(stdout, "Buffer results OK!\n");
// デバイス側のメモリを解放
gcl_free(device_input);
// ホスト側の解放
free(host_input);
}
int main (int argc, const char * argv[])
{
// CPUベースのディスパッチキューを取得
dispatch_queue_t dq = gcl_create_dispatch_queue(CL_DEVICE_TYPE_CPU, NULL);
if (!dq)
{
fprintf(stdout, "Unable to create a CPU-based dispatch queue.\n");
exit(1);
}
// このディスパッチキューのOpenCLデバイスを表示
display_device(gcl_get_device_id_with_dispatch_queue(dq));
buffer_test(dq);
fprintf(stdout, "\nDone.\n\n");
dispatch_release(dq);
}
つぎのリスト9-2は入力配列を2乗するカーネルである
// 入力配列を2乗する簡単なカーネル
// 結果は別のバッファに保持するが、
// デベロッパがそうしたければ入力配列に保存しても良い
// 入力と結果は'global'に宣言されていることに注意
// これはデバイスのグローバルメモリに確保されていることを示している
kernel void square( global float* input, global float* results )
{
// 入力された単精度浮動小数点数をそれぞれのワークアイテムで2乗するようなカーネルを
// ホスト側のコードで起動している
// それぞれのワークアイテムが処理すべきアイテムはグローバルワークアイテムIDに対応している
size_t index = get_global_id(0);
float val = input[index];
results[index] = val * val;
}
9 OpenCLとOpenGLでデータを共有する
OpenGL（Open Graphics Library）は2次元あるいは3次元のコンピュータグラフィクスを生成するアプリを書くためのAPIである。OpenCLとOpenGLは同時に利用できるように設計されている。とくにOpenCLとOpenGLはデータを共有することができ、そうすることでオーバーヘッドを減らすことができる。例えば、OpenGLオブジェクトから作られたOpenGLオブジェクトとOpenCLメモリオブジェクトは同じメモリをアクセスできる。さらに、GLSL（OpenGL Shading Language）シェーダとOpenCLカーネルは共有したデータにアクセスできる。 OpenCLとOpenGLが滑らかに一緒に動作させるためには：- すべての計算とレンダリングをGPUで行うようにプログラムする。これはパフォーマンスを改善できる。なぜならホストとGPUとの間でデータの転送が起こらないからである。
- 効率的にデータが共有されるようにメモリの確保に気をつける
10 GCDを使ってOpenCL/OpenGLの共同作業を制御する
{{decafish : これもさっきの続きみたいなので、このチャプタもスキップする}}2015-07-20 18:16
nice!(0)
コメント(0)
トラックバック(0)
コメント 0