忍者ブログ

Memeplexes

プログラミング、3DCGとその他いろいろについて

C#でOpenCL入門 (Cloo版) プログラム

プログラム

こちらも合わせてお読みください。

今回のテーマはプログラムです。
「プログラム?
前回もプログラム書いたじゃないか」という話ですが、そういう意味のプログラムではあありません。
OpenCLでプログラムといった時、それはGPU(正確に言うとデバイス)で動くプログラムのことです。
ここまではCPU側で動くプログラムしか書いて来ませんでしたからね。

OpenCLでGPUで動くプログラムを動かすときには、OpenCL Cという言語を使います。


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#中でプログラムオブジェクトにします。
プログラムオブジェクトを作るには、ComputeProgramのコンストラクタを使います。

public ComputeProgram(ComputeContext context, string source);
public ComputeProgram(ComputeContext context, string[] source);
public ComputeProgram(ComputeContext context, IList<byte[]> binaries, IList<ComputeDevice> devices);

contextはプログラムオブジェクトの所属するコンテキストです。

sourceはOpenCL Cで書いたプログラムのソースです。(ファイル名ではありません。ソースの文字列そのものです)

binariesはプログラムのバイナリでしょうね。

devicesはバイナリの行くデバイスでしょう。

プログラムのビルド

こうしてC#で扱えるようになったプログラムオブジェクトは、一度ビルドする必要があります。
ComputeProgram.Buildメソッドを使うのです。

public void Build(ICollection<ComputeDevice> devices, string options, ComputeProgramBuildNotifier notify, IntPtr notifyDataPtr);

devicesはプログラムオブジェクトの行くデバイスでしょうね。

optionsはOpenCL Cのコンパイラへ渡すオプションです。

notifyはコールバックです。ビルドが成功しても失敗しても呼ばれるそうです。notifyはnullでもかまいませんが、nullでなかったとき、ビルドが完了する前に戻る可能性があるそうです。nullだったときには、ビルドが完了するまで戻りません。このコールバックはOpenCLの実装によっては非同期に呼ばれる可能性があるのだそうです。

サンプルコード

ではコードを見ていきましょう。
今回はプログラムを生成し、ビルドするだけで意味のある計算などはしません。

Program.cs

using Cloo;
using System.Linq;

class Program
{
    static void Main()
    {
        ComputePlatform platform = ComputePlatform.Platforms[0];
        ComputeDevice[] devices = platform
            .Devices
            .Where(d => d.Type == ComputeDeviceTypes.Gpu)
            .ToArray();
        ComputeContext context = new ComputeContext(
            devices,
            new ComputeContextPropertyList(platform),
            null, 
            System.IntPtr.Zero
            );
        ComputeProgram program = new ComputeProgram(
            context,
            System.IO.File.ReadAllText("myKernelProgram.cl")
            );
        program.Build(devices, null, null, System.IntPtr.Zero);
        program.Dispose();
        context.Dispose();
    }
}

このC#のプログラムは"myKernelProgram.cl"というファイルを読み込んで、プログラムオブジェクトを生成し、ビルドします。
残念ながら何か面白い結果を表示したりはしません。

myKernelProgram.cl

__kernel void myKernelFunction(__global float* numbers)
{
    numbers[0] = 3;
    numbers[1] = 4;
    numbers[2] = 5;
}

このOpenCL Cのプログラムは数字をシングルスレッドで書き込みます。

拍手[1回]

PR