2012年12月31日

OpenCLの使い方について その3

いつの間にか大みそかですが前回の続きです。
[1]の内容を適当にメモを入れながらまとめていきます。
この資料の後半で、前回見ていたOpenCL Postprocessing Exampleが例題として解説されています。
ただ、NVIDIA GPU Computing SDK 4.2ではPBOを使用していますが、[1]ではRenderBufferを使用していますので、内容が少し古いのかもしれません。

OpenCLカーネルとGLSLはどちらがいいのか?
・OpenCLは、GLSLに準備されていないいくつかの機能を持っています。
 ・バラバラな書込み
 ・ローカルメモリ
 ・スレッドの同期
 ・アトミックなメモリ操作(あるスレッドの読書き中はロックできるなど)

OpenCL/GLの相互関係
・OpenGLは、OpenCLとデータを共有することができます。
 ・バッファ(Vertex/Pixelbuffer)←VBOやPBO
 ・Texture
 ・Renderbuffer
・マッピング
 ・OpenCL image→OpenGL texture,renderbuffer
 ・OpenCL buffer→OpenGL buffer

共有データ
・OpenCLのメモリオブジェクトは、OpenGLオブジェクトから作成されます。
 ・OpenGLのオブジェクトが変更されると無効になります。
 ・OpenGLのオブジェクトが削除されても有効なままです。
・使用の前/後で確保/解放する必要があります。
 ・同期するためのAPIsが必要です。
・最良の方法
 ・GLリソースを解放するに、CLリソースを解放することです。

共有するためのAPI
・GLオブジェクトからCLオブジェクトを作成するAPI
 ・clCreateFromGLBuffer
 ・clCreateFromGLTexture2D
 ・clCreateFromGLTexture3D
 ・clCreateFromGLRenderbuffer


・GL vertex bufferの共有
GLuint vbo;
cl_mem vbo_cl;
// create buffer object
glGenBuffers(1, &vbo);
glBindBuffer(GL_ARRAY_BUFFER, vbo);

// initialize buffer obejct
unsigned int size = mesh_width*mesh_height*4*sizeof(float);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);

// create OpenCL buffer from GL VBO
vbo_cl = clCreateFromGLBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, vbo, NULL);

共有するためのAPI
・OpenCLで使用するためのオブジェクトのロック
 ・clEnqueueAcquireGLObjects
 ・clEnqueueReleaseGLObjects
・追加で、同期させるために必要となるAPI
 ・clFinish, clWaitForEvents
 ・glFinish, glFlush


・確保と解放
glFinish();
// すべてのGLコールが終わるので、安全にCLのバッファを確保できます
clEnqueueAcquireGLObjects(cqCommandQueue, 1, vbo_cl, 0,0,0);

// OpenCL manipulates the buffer ...

clEnqueueReleaseGLObjects(cqCommandQueue, 1, vbo_cl, 0,0,0);
clFinish(cqCommandQueue);
// すべてのCLコールが終わるので、安全にGLはバッファを使用できます


OpenCL images
・オプションです:すべてのOpenCLデバイスでサポートされているわけではありません
 ・CL_DEVICE_IMAGE_SUPPORTをチェックしてください。
・OpenGLのテクスチャに似ています。
・読込み可能OR書込み可能です。
・サンプラを介して読込めます。
 ・補間(最近傍、バイリニア)
 ・正規化/非正規化座標系
 ・境界処理(Clamp, Repeat)

OpenCL sampler
・サンプラは、ホスト上で作成でき、カーネル引数として渡すことができます。
 ・clCreateSampler
 ・clGetSamplerInfo
・サンプラは、カーネルコード内でconstとして定義することもできます。
const sampler_t constSampler =
    CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

OpenCL C画像関数
・画像はカーネル引数として渡します。
 ・__read_onlyまたは__write_onlyのどちらかで
__kernel void someKernel(__read_only image2d_t inputImage)
・画像にアクセスするための関数
 ・read_imagei/ui/f(image, sampler, coord);
 ・write_image(image, coord, value);

OpenCL images
・例:
const sampler_t sampler = CLK_NORMALIZED_COORDS_TRUE |
    CLK_ADDRESS_CLAMP | CLK_FILTER_BILINEAR;

__kernel void imageKernel(__read_only image2d_t inputImage,
    __write_only image2d_t outputImage, int width, int height)
{
    int x = get_global_id(0); int y = get_global_id(1);

    float2 normalizedCoord = (float2)((x+0.5f)/width, (y+0.5f)/height);
    uint4 value = read_imageui(inputImage, sampler, normalzedCoord);

    int2 unnormalizedCoord = (int2)(x,y);
    write_imageui(outputImage, unnormalizedCoord, value);
}

OpenCL Postprocessing Example
・OpenGLでレンダリングしたシーンに後から処理を行います。
 ・2Dボックスフィルタは、2パスに分けたフィルタとして実装されています。
 ・最後のパスのハイライトを促進します。
・FrameBufferObjectにシーンをレンダリングします。
 ・RenderBufferはColorとDepth用の2種類です。
・OpenCLカーネルは、OpenGLのテクスチャに書き込みます。
 ・OpenGLはスクリーンぴったりの四角形にテクスチャを描きます。
1.OpenGLで描画
2.OpenCLで縦方向のフィルタを掛ける
3.OpenCLで横方向のフィルタを掛ける
・FBOのRendertaget
// create and bind the FBO
glGenFramebuffersEXT(1, &fbo);
glBindFramebufferEXT(GL_FRAMEBUFFER_EXT, fbo);

// create a RGBA8 render buffer
glGenRenderbuffersEXT(1, &rb_color);
glBindRenderbufferEXT(GL_RENDERBUFFER_EXT, rb_color);
glRenderbufferStorageEXT(GL_RENDERBUFFER_EXT, GL_RGBA8, width, height);

// attach it as color attachment to the FBO
glFramebufferRenderbufferEXT(GL_FRAMEBUFFER_EXT,
    GL_COLOR_ATTACHMENT0_EXT, GL_RENDERBUFFER_EXT, rb_color);

// Do the same for the depth attachment
// ...
・FBOカラーアタッチメントからCL Imageを準備
// Create the CL image form the color renderbuffer - will read from this in the kernel
cl_mem cl_scene;
cl_scene = clCreateFromGLRenderbuffer(cxGPUContext, CL_MEM_READ_ONLY, rb_color, 0);

// CL can query properties on this image as with normal CL images
cl_image_format image_format;
clGetImageInfo(cl_texture, CL_IMAGE_FORMAT, sizeof(cl_image_format),
    &image_format, NULL);

// image_format will be GL_UNSIGNED_INT8, CL_BGRA
・最後のレンダリングパスのためのGLテクスチャを準備
// create GL texture
glGenTextures(1, &tex_screen);
glBindTexture(GL_TEXTURE_2D, tex_screen);

// Set Texture Parameters
// ...

// Setup data storage
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA8, size_x, size_y, 0, GL_RGBA,
    GL_UNSIGNED_BYTE, NULL);

// create CL image from Screen Texture - the CL kernel will write to this
cl_mem cl_screen;
cl_screen = clCreateFromGLTexture2D(cxGPUContext, CL_MEM_WRITE_ONLY,
    GL_TEXTURE_2D, 0, tex_screen, 0);

1. OpenGLでシーンのレンダリング。出力:GL Renderbuffer→CL Image
2.OpenCLで縦方向のフィルタを掛ける。入力:CL Image、出力:CL Buffer
3.OpenCLで横方向のフィルタを掛ける。入力:CL Buffer、出力:GL Texture

以下略。

FBOの場合も大体同じ流れのようですが、速度的にはどちらがいいんでしょうね。
ついでにFBOも色々忘れていたので復習です[2]〜[4]。
一応FBOやTextureでも共有できましたが、clSetKernelArgでエラーがでて少し詰まりました。
Renderbufferを渡すときにバッファで受けるのかimage2d_tで受け取るのかに注意が必要だったようです[5]。

[1] http://sa09.idav.ucdavis.edu/docs/SA09_GL_interop.pdf
[2] http://wiki.livedoor.jp/mikk_ni3_92/d/FrameBufferObject
[3] http://oshiro.bpe.es.osaka-u.ac.jp/people/staff/imura/computer/OpenGL/framebuffer_object/disp_content
[4] http://www.songho.ca/opengl/gl_fbo.html
[5] OpenCL/OpenGL Interop on OSX

web拍手 by FC2
posted by シンドラー at 21:32 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

2012年12月25日

OpenCLの使い方について その2

今回はそのうち必要になりそうなOpenGLとの連携方法についてのメモです。

NVIDIA GPU Computing SDKを使用する場合には、SDKのOpenCLフォルダにサンプルプログラムがあります。
Visual StudioのバージョンによってOpenCLのユーティリティも再コンパイルする必要がありますので、一度はここのoclRelease_vs○○.slnを開いてコンパイルすることになるのではないかと思います。

上記ソリューションの中身は、サンプルプログラムとユーティリティになっており、今回の内容に関連がありそうなものは以下の通りです。

NVIDIA Computing SDK OpenCL Samples
-2_Graphics_Rendering
--OpenCLを使ったマーチンキューブス法のサンプル: oclMarchingCubes
--OpenCL+OpenGLでVBOを使用したサンプル: oclSimpleGL
--OpenCLで3Dテクスチャを使用したサンプル: oclSimpleTexture3D
--OpenCLでボリュームレンダリングのサンプル: oclVolumeRender
-3_Image_Video
--OpenCLで再帰的にガウスフィルタを掛けるサンプル: oclRecursiveGaussian
--OpenCLでOpenGLの描画結果(PBO)に画像処理を施すサンプル: oclPostprocessGL

おそらくVBOかPBOを使ってOpenCLとバッファの共有を行うと思いますので、oclSimpleGLとoclPostprocessGLが参考になると思います。
今回は後者のPBOを使ったOpenCLとの連携方法についてです。

以下Windowsで64ビットでglew,glutを使う場合です。oclPostprocessGLの一部を抜粋改変しています。

1) glutで作成するウィンドウのハンドルを保持する変数を準備して、ウィンドウ作成時に取得します。
int iGLUTWindowHandle;
iGLUTWindowHandle = glutCreateWindow("OpenCL/OpenGL post-processing");

2) レンダリング結果を転送するPBOと、画像処理結果を保存するPBOを準備します。
GLuint pbo[2];
// float型でRGBAの場合
GLuint size_tex_data = image_width*image_height*4*sizeof(float);
glGenBuffers(2, pbo);
for (int i=0; i<2; i++)
{
    glBindBuffer(GL_ARRAY_BUFFER, pbo[i]);
    glBufferData(GL_ARRAY_BUFFER, size_tex_data, NULL, GL_DYNAMIC_DRAW);
    glBindBuffer(GL_ARRAY_BUFFER, 0);
}

3) 最終的なレンダリングのためのテクスチャを準備します。
GLuint texId;
glGenTextures(1, &texId);
glBindTexture(GL_TEXTURE_2D, texId);
//set basic parameters
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST);
// buffer data
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, image_width, image_height, 0, GL_RGBA, GL_FLOAT, NULL);

4) OpenCLの準備をします。(以下通常の方法との変更点のみ記述)
通常のclCreateContextの第一引数はNULLでいいみたいですが、GLUTと連携する場合には下記のように設定します。
cl_context_properties props[] =
{
    CL_GL_CONTEXT_KHR, (cl_context_properties)wglGetCurrentContext(),
    CL_WGL_HDC_KHR, (cl_context_properties)wglGetCurrentDC(),
    CL_CONTEXT_PLATFORM, (cl_context_properties)m_platformID,
    0
};
//コンテキストの作成 (第一引数にpropsを指定するのが違い)
cxGPUContext = clCreateContext(props, m_deviceCount, m_devices, NULL, NULL, &ciErrNum);

//PBO用のメモリ領域を確保
cl_mem cl_pbos[2] = {0, 0};
cl_pbos[0] = clCreateFromGLBuffer(cxGPUContext, CL_MEM_READ_ONLY, pbo[0], &ciErrNum);
cl_pbos[1] = clCreateFromGLBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, pbo[1], &ciErrNum);

5) display関数で描画します。
5-1) Frame bufferへのレンダリング
普通のレンダリングです。
renderScene();

5-2) レンダリング結果をPBOに転送
// PBOをPACKモードでバインド
glBindBuffer(GL_PIXEL_PACK_BUFFER_ARB, pbo[0]);
// PBOにフレームバッファの内容を転送
glReadPixels(0, 0, image_width, image_height, GL_BGRA, GL_FLOAT, NULL);

5-3) OpenCLで画像処理
// OpenGLでの処理を終了させる
glFlush();
// PBOをOpenCL側でバインド
clEnqueueAcquireGLObjects(cqCommandQueue, 2, cl_pbos, 0, NULL, NULL);

// カーネルの実行(省略。pbo[1]の方に処理結果が保存される)
executeKernel();

// OpenCL側からPBOのバインドを解除
clEnqueueReleaseGLObjects(cqCommandQueue, 2, cl_pbos, 0, NULL, NULL);
// OpenCLでの処理を終了させる
clFinish(cqCommandQueue);

5-4) PBOの内容をテクスチャに転送
// PBOをUNPACKモードでバインド
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo[1]);
glBindTexture(GL_TEXTURE_2D, texId);
// テクスチャに転送
glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, image_width, image_height, GL_BGRA, GL_FLOAT, NULL);

5-5) 画面いっぱいにテクスチャを表示
普通のレンダリングです。
displayImage();
glutSwapBuffers();
glutPostRedisplay();

6) 各種終了処理(省略)

PBOの復習は参考サイト[1]で行いました。PACK、UNPACKの違いとかテクスチャへの転送とか色々忘れてて困ったものです。
他にも[2]などにテクスチャ(FBO?)を使用する方法やVBOを使用する方法もありますので、気が向いたらまとめたいと思います。

[1] OpenGL de プログラミング
[2] http://sa09.idav.ucdavis.edu/docs/SA09_GL_interop.pdf

web拍手 by FC2
posted by シンドラー at 20:27 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

2012年12月08日

OpenCLの使い方について その1

何となく試してみたくなったのでOpenCLの使い方についてまとめていきます。

まず最初の環境設定ですが、参考サイト[1]を参考に行いました。
OpenCLのライブラリ等はCUDA等をインストールするとNVIDIA GPU Computing SDKフォルダに一緒に入ってきます。
ここは既にインストール済みというとこで省略します。

以下Visual Studio 2010 Express Editionでx64のお話です。
1.新しいプロジェクトを作成します。
2.[1]を参考にヘッダファイルとLibファイルの設定をします。
 追加のインクルードディレクトリ
  $(NVSDKCOMPUTE_ROOT)\OpenCL\common\inc;$(NVSDKCOMPUTE_ROOT)\shared\inc;
 コード生成→ランタイムライブラリ
  マルチスレッド(/MT)
 追加のライブラリディレクトリ
  $(NVSDKCOMPUTE_ROOT)\OpenCL\common\lib\x64;$(NVSDKCOMPUTE_ROOT)\shared\lib\x64;
 追加の依存ファイル
  oclUtils64.lib OpenCL.lib shrUtils64.lib
3.カーネルのプログラム(.cl)を作成
4.前準備をして実行

今回は画像の勾配強度を計算するプログラムを作成してみました。
手順としては、画像読込→イメージオブジェクトの作成→カーネルプログラムで計算→結果の保存、です。

サンプルカーネルプログラム(勾配強度の計算)
const sampler_t s_linear = CLK_FILTER_LINEAR | CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE;

__kernel void Test2d(__read_only image2d_t im, __global float4 *Ary)
{
    const unsigned int idx = get_global_id(0);
    const unsigned int idy = get_global_id(1);
    const unsigned int index = idy*get_image_width(im)+idx;
    float4 dx, dy;
    dx = read_imagef(im, s_linear, (float2)(idx+1.0, idy))-read_imagef(im, s_linear, (float2)(idx-1.0, idy));
    dy = read_imagef(im, s_linear, (float2)(idx, idy+1.0))-read_imagef(im, s_linear, (float2)(idx, idy-1.0));
    Ary[index] = hypot(dx, dy);
}
hypot関数は、sqrt(dx*dx+dy*dy)を計算してくれる関数です。この辺の情報は[4]が見やすかったです。

実行結果
1024x1024のLenna画像で約13ms(Core i7-2600K, GTX580, データ転送時間含む)でした。
lenna_1024_result.jpg

以下メモ
・イメージオブジェクトを使用すればGPUの補間の恩恵が得られる
・__localはグループ内で共有するCUDAでいうシェアードメモリみたいな感じなので注意
・get_global_idでグローバルなIDの取得ができるのでCUDAよりわかりやすいかも
・CUDAとの関係などのPPT資料[5]
・clEnqueueNDRangeKernelの引数に注意
 (グローバルワークサイズとローカルワークサイズの関係がまだよくわからない)
・OpenGLとの連携は、CUDAと同様にPBOやVBOでバッファの共有ができる[6][7]

[1] OpenGL de プログラミング
[2] OpenCL - 画像処理(αブレンディング) チューニング
[3] 株式会社フィックスターズ, "OpenCL入門 マルチコアCPU・GPUのための並列プログラミング", 株式会社インプレスジャパン(2010)
[4] OpenCL API 1.1 Quick Reference Card
[5] http://www.seas.upenn.edu/~cis565/Lectures2011/FLecture15_OpenCL.ppt
[6] http://kmotk.jugem.jp/?eid=89
[7] http://www.steradian.co.jp/2009/12/siggraph-asia-2009---1-h.html

web拍手 by FC2
posted by シンドラー at 19:47 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

広告


この広告は60日以上更新がないブログに表示がされております。

以下のいずれかの方法で非表示にすることが可能です。

・記事の投稿、編集をおこなう
・マイブログの【設定】 > 【広告設定】 より、「60日間更新が無い場合」 の 「広告を表示しない」にチェックを入れて保存する。