2014年05月05日

レイキャスティングと重心空間を用いたデプスバッファの計算 その2

良く考えたらポリゴンの内部かどうかは別として、レイキャスティングの時点で深度値は計算できてますよね…。それにスクリーン座標系の深度値って非線形だった気がするのですが、深度値をZ座標として平面計算してしまって問題ないのでしょうか…。

とりあえずOpenCLで実装をしてみました。GlobalWorkSize(size_t gws[2])の設定を久々だったので間違えていて無駄に時間がかかりましたのでメモしておきます。

LocalWorkSizeはsize_t lws[2] = {16,16};にしておきます。GlobalWorkSizeは、LocalWorkSizeの定数倍になっていないと駄目なようです。2次元の画像の場合、widthとheightを16の倍数にする必要があります。そのために、16で割ってceil関数で整数にしてもう一度16を掛けます。
void setGlobalWorkSize2D(const int width, const int height, 
const size_t lws[2], size_t gws[2])
{
gws[0] = lws[0]*ceil((double)width/(double)lws[0]);
gws[1] = lws[1]*ceil((double)height/(double)lws[1]);
}
各頂点、各面のデータに対して処理する場合、基本的には1次元の頂点数、面数を使えばいいのだと思いますが、面倒なので?2次元に分割して使用することにします。
void setGlobalWorkSize1D(const int dataNum, 
const size_t lws[2], size_t gws[2])
{
// データ数の平方根を横幅にする
gws[0] = lws[0]*ceil(sqrt((double)dataNum)/(double)lws[0]);
// 縦幅
gws[1] = lws[1]*ceil(ceil(((double)dataNum/(double)gws[0]))/(double)lws[1]);
}
普通に1次元でやった方が楽ですかね。

cl側での参照の仕方を書いていなかったので追記します。(2014.5.6)
上の方法で1次元データを2次元で処理する場合、横幅はGlobalWorkSizeと一致しますので、それを使って判定できます。領域外かどうかは、データ数を渡してあげる必要はあります。(必要なのかどうかはわかりませんが)
__kernel void Test1d(__global float *vertexPos, __global float *screenPos, int vertexNum)
{
// x座標
const unsigned int idx = get_global_id(0);
// y座標
const unsigned int idy = get_global_id(1);
// get_global_size(0)で横幅を取得して一次元に直す
unsigned int index = get_global_size(0)*idy + idx;
// データ数を渡しておいて範囲外の時は処理しない
if (index >= vertexNum) return;
...
}

2次元画像の場合、GlobalWorkSizeと画像のサイズは異なっていますので、領域外かどうか判定するためにget_image_dim関数を使用して画像の大きさを取得します。
__kernel void Test1d(__read_only image2d_t im_in, __write_only image2d_t im_out)
{
// x座標
const unsigned int idx = get_global_id(0);
// y座標
const unsigned int idy = get_global_id(1);
// 画像の大きさを取得して範囲外の時は処理しない
const int2 imageSize = get_image_dim(im_out);
if (idx >= imageSize.x || idy >= imageSize.y) return;
...
}

実行結果

640x480の画像で2秒ぐらいかかるのでやはり使えなさそうですね。メモリ的にもギリギリみたいですし。

center_gravity_003.png

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

2013年11月20日

OpenCLでSSAO+シャドウマッピング

以前もSSAOをやっていたのですが、CPUで遅かったのでOpenCLで試してみることにしました。
SSAOについては、[1]がとても参考になりました。
以前やっていたのはHorizon-based Ambient Occlusionもどきで、今回はStarCraft II Ambient Occlusionもどきになります。
サンプリング点の生成方法は、[2]を参考にさせていただいて、中心の方にサンプリング点が増える方式を使わせていただき、サンプリング点100個として計算しました。

シャドウマッピングについては、ガウスフィルタでぼかすことを前提にしてますので、一番基本的なものにしています。時間があればバイラテラルフィルタとかもOpenCLで実装しておきたいところです。

手順としては以下のようになります。

(1) 視点及び光源からレンダリングして深度値や法線ベクトルなどをFBOに出力
(2) FBOをOpenCLから参照してSSAO+シャドウマッピング
(3) ↑の結果にガウスフィルタ×5回
(4) 通常のレンダリング結果×↑の画像
(5) FBOをOpenGLに戻す

実行結果

通常のレンダリング結果というかテクスチャの色を出力しただけのもの
opencl_ssao_001.png

(2)〜(3)を計算した結果
opencl_ssao_002.png

↑の結果をレンダリング結果に掛け合わせたもの+ガンマ補正
opencl_ssao_003.png

(2)〜(3)で色を(SSAO, シャドウマッピング, 1.0, 1.0)にして出力した結果
opencl_ssao_004.png

データ転送を含まない(2)〜(4)の処理が約30msぐらいで、リアルタイムにはちょっとぎりぎり間に合わないぐらいですがやっぱり並列処理は早いですね。

一応OpenCLのソースコードだけ置いておきます。このソースコードを参考にして何か問題が起こったとしても責任はとりません。
gpuMain.cl

[1] http://ambientocclusion.hatenablog.com/entry/2013/11/07/152755
[2] http://marina.sys.wakayama-u.ac.jp/~tokoi/?date=20101122
web拍手 by FC2
posted by シンドラー at 23:08 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

2013年11月07日

LU分解で行列式と逆行列の計算 その3

clBLASを使えるようになったので、LU分解で逆行列の計算を試してみることにしました。
といっても参考サイト[1][2]の遂次用っぽいアルゴリズムをそのままOpenCL用に書いただけなので、全然並列化効率とか上がっていないと思います。
(計算結果に依存する部分があるため、ループを回して何度もカーネル関数を実行しています。)

LU分解のやり方は大体以前の記事「LU分解で行列式と逆行列の計算 その2」の方法で、AをMAにしてMAをLUにしてA-1=U-1L-1Mで逆行列を計算しています。本来はLUx=bとした場合、Ly=b、Ux=yとしてxを求めるべきなのですが、上記のように三角行列の逆行列U-1とL-1を計算して逆行列を求めています。[1]によると連立1次方程式を解くのに逆行列は無用の長物らしいので、あまり良くないやり方なんだとは思います。

掛け算については前回のclBLASを使用させていただいております。

実行結果

適当に作った1024x1024の行列の逆行列を計算したところ、最初から最後まで合わせても約1秒で計算できました。ちゃんと並列化された逆行列計算のライブラリなら一瞬で終わると思いますので自分で作る意味は無いと思いますが。

せっかく作ったので一応clのコードだけ公開しておきます。万が一参考にして何か問題が起こっても責任はとりません。
gpuMatrixTest.cl

[1] http://www.math.meiji.ac.jp/~mk/labo/text/linear-eq-1.pdf
[2] http://www.kata-lab.itc.u-tokyo.ac.jp/OpenLecture/SP20120612.pdf
web拍手 by FC2
posted by シンドラー at 01:35 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

2013年10月31日

clMath(clBLAS)の使い方について

行列計算の並列処理でCUDAを使ったものはありますが、OpenCLを使ったものはどうかなと探してみると、clMathというものがありました。
clMathは、現在clBLASとclFFTがあり、行列計算や高速フーリエ変換をOpenCLを使ってできるようです。

今回は、clBLASの使い方について簡単にメモっておきます。

OpenCLはインストール済みと仮定します。
githubにclMathのデータがあります。

[1] https://github.com/clMathLibraries

ユーザ登録が必要かどうかはわかりませんが、自分の場合はgithubにユーザ登録して、clBLASのDownload ZIPをダウンロードしました。

clblas-develop.zipを適当なフォルダに解凍して、CMake(cmake-gui)を立ち上げます。

clBLAS-develop/srcフォルダをソースフォルダに設定し、適当にclBLAS-develop/buildフォルダを作成してそこをビルドフォルダに設定します。
Configureを押すと、ACMLが見つからないというエラーと、Boostが見つからないというエラーがでていました。

ACMLとBoostは、サンプルプログラム?等に必要な雰囲気でなくてもかまわない気がしましたが、一応入れることにしました。

ACMLは、下記のページからacml5.3.1-win64.exe PGI fortran用をダウンロードしてインストールしました。

[2] http://developer.amd.com/tools-and-sdks/cpu-development/amd-core-math-library-acml/acml-downloads-resources/

ACML_ROOTにF:/AMD/acml5.3.1/win64といったインストールフォルダを指定すればACMLのエラーは消えます。
Boostは、省略します。

Generateを押せば、buildフォルダにclBLAS.slnができていますので、ビルドすればエラーなく成功しました。

ライブラリを使用するときは、clBLAS-develop/src及びclBLAS-develop/src/includeにあるヘッダファイルと、clBLAS-develop/build/library/ReleaseフォルダにあるclBLAS.libとclBLAS.dllがあれば良さそうです。

一番基本となる行列同士の掛け算は、[3]の下のほうにサンプルソースコードがあります。

[3] https://github.com/clMathLibraries/clBLAS

APIの説明は、[4]にあります。

[4] http://clmathlibraries.github.io/clBLAS/

[3]では、clblasSgemmという関数を使って行列の掛け算を計算しています。

Sはsingle precision、gemmはGeneral matrix-matrix multiplicationの略で、double型で計算をしたい場合はclblasDgemmという関数になります。

で、引数がかなりわかりにくかったです。APIドキュメントには書かれていますが、ひとつの関数で下記のような計算を行うことができるようになっています。

gemm.png
[4]より引用

α、βは複素数のバージョンもあるのですが、上記の2関数は実数用のものになっています。
C=αAB+βCで、AやBは引数によって転置させて計算することも可能となっています。
つまり、普通掛け算というとC=ABだと思いますが、これを行うためにはα=1.0、β=0.0に設定しないとおかしな結果になってしまいます。

それ以外は、引数の説明に気をつけながら値を渡せば特に問題はないかと思います。

実行結果

自作行列掛け算と、clBLASを使用した場合(clblasDgemm)の計算時間を比較してみました。
1024x1024の行列A,B(値は-1.0〜1.0の一様乱数)を掛けて、行列Cに代入します。

自作行列掛け算の時間:9629[ms]
clBLAS(転送時間含む):55[ms]

自作行列クラスでデータ生成して、それを一時配列にコピーして、GPU上にメモリを作成して、CPU->GPUにデータ転送して、掛け算計算して、GPU->CPUにデータ転送して、自作クラスにデータコピーして結果を返すまでの時間が上記の結果でした。
一応計算結果をファイルに出力して比較してみましたが、どちらも同じ結果になっていたので合っていると思います。
無駄なデータ転送時間を入れてもこの差はさすが並列処理ですね。まぁ比較対象がひどすぎるというのもありますが。

とりあえず後はGEMV - General matrix-Vector multiplicationとかを試してみようとは思います。
○○solveとかあるので逆行列とかも計算できそうな気がしますが、まだやり方はわかっていません。
web拍手 by FC2
posted by シンドラー at 19:18 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

2013年01月14日

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

OpenCLとOpenGLの連携ができるようになりましたので、少し色々なレンダリングを試してみたいと思います。
今回は以前も少しやっていましたが、異方性反射です。

異方性反射を実現するためには、法線ベクトルだけではなく、接線、従接線ベクトルが必要になります。
この接線、従接線ベクトルの計算は、テクスチャのUV座標を使用する方法や、適当に指定した軸と法線ベクトルの外積を利用して計算する方法などがあるようです[1][2]。
今回は、上方向ベクトルと外積を使って適当に計算してみました。

後は異方性を考慮したレンダリング方法ということで、[3]を参考にAshikhminモデルを試してみました。
OpenCL側では、FBOでレンダリングした結果にガウスフィルタを掛けたりして最終的な画像を生成しているのですが、引数で渡したFBOの画像は必ず参照しないとエラーがでてしまうようです。
(文法的に間違っていないのにコンパイルエラーで-30が返ってきて無駄に詰まりました。)

実行結果

前回と同様にあぴミクさんを使用させていただいております。
[4]でいうところのnuとnvという2つのパラメータを変化させることで結果が変わります。
(見えにくいですがタイトルバーに数値を表示しています)
最初は10-10で広く浅く、100-100ぐらいでフォンシェーディングのような効果、10-2000や2000-10で縦方向もしくは横方向に伸びた光沢が得られています。



本当は部位毎に接線ベクトル方向やこのパラメータを変更した方がいいのだとは思います。

[1] http://marupeke296.com/DXPS_No12_CalcTangentVectorSpace.html
[2] http://lucille.sourceforge.net/blog/archives/000013.html
[3] http://asura.iaigiri.com/XNA_GS/xna19.html
[4] http://www.cs.utah.edu/~shirley/papers/jgtbrdf.pdf
web拍手 by FC2
posted by シンドラー at 03:00 | Comment(0) | TrackBack(0) | OpenCL | このブログの読者になる | 更新情報をチェックする

2013年01月07日

PMXファイルの表示

MMD用に素晴らしいミクさんがたくさん公開されていますね。
最近のミクさんはPMX形式になっているものが殆どですので、PMXファイルを読み込んで表示するようにしました。
PMXファイルについては、VPVP wikiに解説[1]がありますし、読込み可能なライブラリの紹介[2]もありますし、PMXエディタに仕様書も入っています。

自分で作る必要はないのですが、自分で作っていて困った点のメモです。
1.UTF16やUTF8がよくわからない
2.pngの透過色はOpenCVで昔は読み込んでもらえなかった

1.については、とりあえずUTF16だけに対応しました。
Windowsでは、charとかTCHARとかWCHAR色々あってよくわからない文字コードですが、TCHARはVisual Studioのプロパティ設定で振舞いが変わる汎用的な型で、WCHARの方がワイドバイト文字だそうです[3]。
PMXファイルのエンコード方式がUTF16の場合、マルチバイト文字のようですので、WCHARとして読み込んでwcstombs_s関数などでchar型に変換すればいいようです。

2.については、いつの間にか(2.3.1ぐらいから)pngのα値を読み込めるようになっていたようです。
注意点としては、画像読込関数の第2引数に-1を指定する必要があるみたいです[4]。

[1] http://www6.atwiki.jp/vpvpwiki/pages/284.html
[2] http://www6.atwiki.jp/vpvpwiki/pages/288.html
[3] http://note.phyllo.net/?eid=1106043
[4] http://opencv.jp/cookbook/opencv_img.html#png

実行結果

FBOに描画→OpenCLで3x3のガウシアンフィルタをかけた結果です。
ローカルメモリなどを使用していないので遅いですが、FPSは60弱ぐらいなのでまぁいいですかね。


Appearance Miku あぴミクさんを使用させていただいております。
以下ガウシアンフィルタのカーネル関数です。フィルタは普通はCPUから渡すと思いますが、面倒だったので直書きです。
const sampler_t s_linear = CLK_FILTER_LINEAR | CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE;
__kernel void postprocess(__read_only image2d_t im_in, __write_only image2d_t im_out, int imgw, int imgh)
{
    const unsigned int x = get_global_id(0);
    const unsigned int y = get_global_id(1);
    const int index = y * imgw + x;

    const float k[9] = {1.0/16.0, 2.0/16.0, 1.0/16.0,
        2.0/16.0, 4.0/16.0, 2.0/16.0,
        1.0/16.0, 2.0/16.0, 1.0/16.0};

    if( x >= imgw || y >= imgh ) return;

    float4 output;
    output = k[0]*read_imagef(im_in, s_linear, (float2)(x-1, y-1));
    output += k[1]*read_imagef(im_in, s_linear, (float2)(x-1, y ));
    output += k[2]*read_imagef(im_in, s_linear, (float2)(x-1, y+1));
    output += k[3]*read_imagef(im_in, s_linear, (float2)(x , y-1));
    output += k[4]*read_imagef(im_in, s_linear, (float2)(x , y ));
    output += k[5]*read_imagef(im_in, s_linear, (float2)(x , y+1));
    output += k[6]*read_imagef(im_in, s_linear, (float2)(x+1, y-1));
    output += k[7]*read_imagef(im_in, s_linear, (float2)(x+1, y ));
    output += k[8]*read_imagef(im_in, s_linear, (float2)(x+1, y+1));

    write_imagef(im_out, (int2)(x, y), output);
}

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

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 | このブログの読者になる | 更新情報をチェックする