C#でOpenCL入門 (Cloo版) グループID、ローカルID

グループIDとローカルID

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

今回はOpenCL Cに用意された関数、get_group_id()とget_local_id()について解説します。

get_group_id()とget_local_id()

以前私の書いた文章から引用します：

今回は、OpenCL C言語に用意された関数、get_group_id()とget_local_id()についてメモします。
get_group_id()
get_local_id()

size_t get_group_id(uint dimIndex)

size_t get_local_id(uint dimIndex)

dimIndexは、次元の数です。0はx、1はy、2はzですね。
戻り値はそれぞれのIDを意味します。

どういう事かというと、こういうことです：

GPGPUには、スレッドがたくさんあるのですが、ただだらーっと並べられるのではなく、スレッドは３次元の構造を取ります。




そしてそれがいくつかのグループにまとめられているのです。



get_group_idは、それを実行するスレッドの所属するグループが、何グループ目なのかを表します。
get_local_idは、それを実行するスレッドが、グループの中で何番目のスレッドなのかをあらわします。
上の２つの関数のdimIndexは、このうちXかYかZかを表すのです。

…というわけで、Clooを使ってこの２つの関数の働きを見てみましょう。

サンプルコード

Program.cs

using Cloo;
using System.Linq;

struct MyItem
{
    public float GlobalID;
    public float GroupID;
    public float LocalID;

    public override string ToString()
    {
        return string.Format("{0}, {1}, {2}", GlobalID, GroupID, LocalID);
    }
}

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);
        const int elementCount = 6;
        ComputeBuffer<MyItem> buffer = new ComputeBuffer<MyItem>(
            context, 
            ComputeMemoryFlags.ReadWrite,
            elementCount
            );

        ComputeKernel kernel = program.CreateKernel("myKernelFunction");
        kernel.SetMemoryArgument(0, buffer);
        ComputeCommandQueue commandQueue = new ComputeCommandQueue(
            context,
            devices[0],
            ComputeCommandQueueFlags.None
            );

        commandQueue.Execute(
            kernel,
            null, 
            new long[] { elementCount },
            new long[] { 2 }, 
            null
            );

        var dataFromGpu = new MyItem[elementCount];
        commandQueue.ReadFromBuffer(
            buffer,
            ref dataFromGpu,
            true,
            null
            );

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

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

myKernelProgram.cl

typedef struct MyItem_s
{
	float GlobalID;
	float GroupID;
	float LocalID;
} MyItem;

__kernel void myKernelFunction(__global MyItem* items)
{
	int globalThreadID = get_global_id(0);

	items[globalThreadID].GlobalID = globalThreadID;
	items[globalThreadID].GroupID = get_group_id(0);
	items[globalThreadID].LocalID = get_local_id(0);
}

このプログラムは、６スレッド作り、それを１グループ２スレッドでグループ分けしています。

（スレッド０、スレッド１）、（スレッド２、スレッド３）、（スレッド４、スレッド５）

そしてそれぞれのスレッドの、各IDを出力しています。 グローバルID、グループID、ローカルIDを、です。

その結果はこうなります。

0, 0, 0
1, 0, 1
2, 1, 0
3, 1, 1
4, 2, 0
5, 2, 1

つまり、

	スレッド０	スレッド１	スレッド２	スレッド３	スレッド４	スレッド５
グローバルID	0	1	2	3	4	5
グループID	0	0	1	1	2	2
ローカルID	0	1	0	1	0	1

というわけですね。