[PR]
×
[PR]上記の広告は3ヶ月以上新規記事投稿のないブログに表示されています。新しい記事を書く事で広告が消えます。
プログラミング、3DCGとその他いろいろについて
[PR]上記の広告は3ヶ月以上新規記事投稿のないブログに表示されています。新しい記事を書く事で広告が消えます。
こちらも合わせてお読みください。
今回のテーマはプログラムです。
「プログラム?
前回もプログラム書いたじゃないか」という話ですが、そういう意味のプログラムではあありません。
OpenCLでプログラムといった時、それはGPU(正確に言うとデバイス)で動くプログラムのことです。
ここまではCPU側で動くプログラムしか書いて来ませんでしたからね。
OpenCLでGPUで動くプログラムを動かすときには、OpenCL Cという言語を使います。
OpenCLでGPUを動かすときには、OpenCL Cという独自の言語を使います。
DirectXで言うとHLSLのような言語です。
この言語によって、GPU特有の命令を実行するのです。
OpenCL Cで書いたプログラムはファイルとして持っておき、それを今回C#側から読み出し、コンパイルします。
(そして実行するのです。今回は実行まではいかず、コンパイルまでですが。)。
OpenCLで書いたプログラムは例えば次のようになります。
拡張子は.clです。
__kernel void myKernelFunction(__global float* numbers){numbers[0] = 3;numbers[1] = 4;numbers[2] = 5;}このプログラムは実はあまりGPUの並列処理っぽくはない、シングルスレッドなプログラムです。
バッファに対して値をセットしています(シングルスレッドで)。
もちろん、書き方によってはGPUの性能をフルに生かした並列処理なプログラムを書くことは可能です(ちょっと複雑になるのでここには書きませんが)。
さてここで見慣れないキーワードが出てきました。
__kernelと__globalです。
これらはC言語にはない、OpenCL C特有のキーワードです。
__kernelは関数が、GPUで実行されるということを意味します(GPUというか、正確にはCPUでも設定次第で実行できるのでデバイスというのが正しいのでしょうが、めんどくさいので以後GPUとします)。
__globalは変数がGPUのメインメモリに確保されるということを意味します。
では変数がGPUのメインメモリに確保されない場合があるのでしょうか?
そうです。
GPUと一口にいっても、変数がどこに確保されるかは全部で次の4つが考えられます。
以下の4つは、基本的に最初のものほど制限が緩く、最後の方に行くほど制限がきつく(しかし高速に)なっていきます。
__globalは変数がGPUのメインメモリにあり、自由に読み書き可能です。
ランダムアクセス可能なのでGPGPUの計算結果はここに格納することになります。
DirectXでいうとUnorderedAccessViewに相当すると言えるでしょうか。
__constantは変数がGPUのメインメモリにある・・・ところまでは__globalと同じなのですが、読み取り専用です。
読み取り専用にして何が嬉しいのかということですが、NVIDIAのGPUではたしかこちらのほうがパフォーマンスが高かったはずです。
読み取り専用なのでキャッシュを有効活用できるわけですね。
__localは、GPUの共有メモリです。
共有メモリとはなんんでしょう?
これが__globalや__constantと何が違うのかというと、これは複数あるということです。(複数あるという意味で、オブジェクト指向言語で言う非staticフィールドに似ています)
GPUには複数の演算ユニットがあり、__globalや__constantはそれら全てで共有されていました。
__localはそれらのうち一部で共有されるのです。
ですから、__localのついた変数は、実際には演算ユニットのグループの数だけ存在することになります。
グループの中で共有されるのです。
計算結果はここにあってもCPUからは読み込めません。(staticなメソッドから非staticなフィールドを直接読み込めないのと同じようなものです)
計算途中のデータをここに格納するといいでしょう。
__privateは、GPUのレジスタです。
これは一つの演算ユニットに特有の変数です。
これはデフォルトです。
なので、変数の前に何も付けなかった場合、その変数は__privateだとみなされます。
普通に関数内でint a;というような宣言をすると、aは__privateとなります。
まずプログラムを生成する必要があります。
OpenCL Cで書いたプログラムを、C#中でプログラムオブジェクトにします。
プログラムオブジェクトを作るには、Context.CreateProgramWithSourceメソッドを使います。
public Program CreateProgramWithSource(string source);
sourceはOpenCL Cで書いたプログラムのソースです。(ファイル名ではありません。ソースの文字列そのものです)
こうしてC#で扱えるようになったプログラムオブジェクトは、一度ビルドする必要があります。
Program.Buildメソッドを使うのです。
public void Build();
ではコードを見ていきましょう。
今回はプログラムを生成し、ビルドするだけで意味のある計算などはしません。
using OpenCLNet; class Program { static void Main() { var platform = OpenCL.GetPlatform(0); var context = platform.CreateDefaultContext(); var program = context.CreateProgramWithSource( System.IO.File.ReadAllText("myKernelProgram.cl") ); program.Build(); program.Dispose(); context.Dispose(); } }
このC#のプログラムは"myKernelProgram.cl"というファイルを読み込んで、プログラムオブジェクトを生成し、ビルドします。
残念ながら何か面白い結果を表示したりはしません。
__kernel void myKernelFunction(__global float* numbers) { numbers[0] = 3; numbers[1] = 4; numbers[2] = 5; }
このOpenCL Cのプログラムは数字をシングルスレッドで書き込みます。