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

データ並列

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

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

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

データ並列で実行

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

public void Execute(ComputeKernel kernel, long[] globalWorkOffset, long[] globalWorkSize, long[] localWorkSize, ICollection<ComputeEventBase> events);

kernelは並列で実行するカーネル。
globalWorkOffsetはグローバルIDのオフセットです。nullでも構いません。
globalWorkSizeは全スレッドの量を表します。
localWorkSizeは1グループのスレッドの量を表します。
eventsはとりあえずnullにしときましょう。

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

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

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





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



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

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

お分かりいただけたでしょうか。
ちなみにwork_dimはここで言うglobalWorkSize.Length(あるいはlocalWorkSize.Length)のことです。

サンプルコード

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);
        ComputeKernel kernel = program.CreateKernel("myKernelFunction");
        const int elementCount = 3;
        ComputeBuffer<float> buffer = new ComputeBuffer<float>(
            context,
            ComputeMemoryFlags.ReadWrite,
            elementCount);
        kernel.SetMemoryArgument(0, buffer);

        ComputeCommandQueue commandQueue = new ComputeCommandQueue(
            context, 
            devices[0],
            ComputeCommandQueueFlags.None
            );
        commandQueue.Execute(
            kernel, 
            null, 
            new long[] { elementCount },
            new long[] { 1 },
            null
            );
        float[] dataFromGpu = new float[elementCount];
        commandQueue.ReadFromBuffer(buffer, ref dataFromGpu, true, null);

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

        commandQueue.Dispose();
        buffer.Dispose();
        kernel.Dispose();
        program.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という出力になるのです。