C#でOpenCL入門 (OpenCLNet版) データ並列

データ並列

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

前回はついにGPUを使ってGPUのメモリに書き込みをしました。
しかしGPUを使っているというのにひとつのスレッドしか使っていませんでした。
GPUは並列計算をしてこそ意味があります。

という訳で今回は、OpenCL Cで書いたプログラムを並列に実行してみましょう。

データ並列で実行

カーネルを並列で実行するのに、前回のEnqueueTaskメソッドは使えません。
代わりにCommandQueue.EnqueueNDRangeKernelメソッドを使います。

public void EnqueueNDRangeKernel(Kernel kernel, int workDim, int[] globalWorkOffset, int[] globalWorkSize, int[] localWorkSize);

kernelは並列で実行するカーネル。
workDimは実行するスレッド群の次元です。1からDevice.MaxWorkItemDimensionsまでの値を取ることができます。
globalWorkOffsetはグローバルIDのオフセットです。nullでも構いません。
globalWorkSizeは全スレッドの量を表します。
localWorkSizeは1グループのスレッドの量を表します。

さて次元とかグループとか言った単語の解説をしたいのですが、ここは以前私が書いた解説をそのまま引用させて頂きます：

さて次元だとかグループだとかよくわからない話が出てきました。
これはどういう事でしょうか？

DirectX11の時にも話したのですが、GPGPUをするときにはスレッドがこのようにたくさん用意されます。





実行されるスレッドは３次元の箱のように配置されます。
そしてこれが次のように分割されます：



GPUの演算ユニットは、いくつかのグループをなしています。
そのためスレッドもグループをなすことになります。
ちなみにOpenCL Cの__local変数は、このグループの中で共有される変数だというわけです。

この関数の次元とは、スレッドが何次元に配置されるかを表します。
DirectXでは3次元でしたが、OpenCLでは１次元もありです（つまり横に並べただけ）。
その次元数をwork_dim変数にセットするのです。

お分かりいただけたでしょうか。
ちなみにwork_dimはここで言うworkDimのことです。

サンプルコード

Program.cs

using OpenCLNet;

class Program
{
    static void Main()
    {
        var platform = OpenCL.GetPlatform(0);
        var context = platform.CreateDefaultContext();
        var commandQueue = context.CreateCommandQueue(platform.QueryDevices(DeviceType.DEFAULT)[0]);

        var program = context.CreateProgramWithSource(
            System.IO.File.ReadAllText("myKernelProgram.cl")
            );
        program.Build();

        var kernel = program.CreateKernel("myKernelFunction");
        const int elementCount = 3;
        var buffer = context.CreateBuffer(MemFlags.READ_WRITE, elementCount * sizeof(float));
        kernel.SetArg(0, buffer);
        commandQueue.EnqueueNDRangeKernel(
            kernel,
            1, 
            null,
            new[] { elementCount },
            new[] { 1 }
            );

        float[] readBack = new float[3];
        buffer.Read(commandQueue, 0, readBack, 0, readBack.Length);

        foreach (var number in readBack)
        {
            System.Console.WriteLine(number);
        }

        buffer.Dispose();
        kernel.Dispose();
        program.Dispose();
        commandQueue.Dispose();
        context.Dispose();
    }
}

myKernelProgram.cl

__kernel void myKernelFunction(__global float* numbers)
{
	int globalThreadID = get_global_id(0);
    numbers[globalThreadID] = globalThreadID;
}

このプログラムを実行すると、次のように表示します：

0
1
2

このプログラムは、まずGPUのメモリにfloatが３つあるバッファを作ります。
（この段階ではバッファの中身は適当で、特に決まっていないようです。普通のでたらめな数値が入っていることもあれば、NaNとか0とかが入っている可能性もあります）
そしてGPUが3スレッド使って、それぞれの要素にグローバルID（スレッドのIDのこと）を代入します。
スレッドのIDはシンプルに0, 1, 2です。
それらが代入されて、バッファは{0, 1, 2}となります。
最後にそれがCPU側に読み戻されて、0 1 2という出力になるのです。