C#でOpenCL入門 (OpenCLNet版) グループ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かを表すのです。

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

サンプルコード

Program.cs

using OpenCLNet;
using System.Runtime.InteropServices;

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()
    {
        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 = 6;
        Mem buffer = context.CreateBuffer(
            MemFlags.READ_WRITE, 
            elementCount * System.Runtime.InteropServices.Marshal.SizeOf(typeof(MyItem))
            );
        kernel.SetArg(0, buffer);
        commandQueue.EnqueueNDRangeKernel(
            kernel,
            1,
            null,
            new[] { elementCount },
            new[] { 2 });

        var readBack = new MyItem[elementCount];
        ReadBuffer(commandQueue, buffer, readBack);

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

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

    static void ReadBuffer<T>(CommandQueue queue, Mem buffer, T[] systemBuffer) where T :struct
    {
        GCHandle handle = GCHandle.Alloc(systemBuffer, GCHandleType.Pinned);

        queue.EnqueueReadBuffer(
            buffer,
            true,
            0,
            Marshal.SizeOf(typeof(T)) * systemBuffer.Length,
            handle.AddrOfPinnedObject()
            );

        handle.Free();
    }
}

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

というわけですね。