忍者ブログ

Memeplexes

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

C#でOpenCL入門 (OpenCLNet版) スレッドとグループの個数

スレッドとグループの個数

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


GPGPUでは普通複数のスレッドを使います。
複数のスレッドで仕事をバーっとやってもらうのです。

そのスレッドの個数はCPU側で設定します。
ときに、GPU側でもスレッドの個数がいくつかを知りたいことがあります。

それには、get_global_size()関数を使います。
この関数は、今のコマンドで、スレッドが全部でいくつ実行されているかを取得します。
参考

size_t get_global_size(uint dimIndex)

dimIndexは、次元を表すインデックスです。X軸方向のスレッド個数を取得したいなら0,Y軸方向なら1,Zなら2です。

サンプルコード

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 = 6;
        var buffer = context.CreateBuffer(
            MemFlags.READ_WRITE, 
            elementCount * sizeof(float)
            );
        kernel.SetArg(0, buffer);

        commandQueue.EnqueueNDRangeKernel(
            kernel,
            2,
            null,
            new[] { 2, 3 },
            new[] { 1, 1 });

        var readBack = new float[elementCount];
        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* items)
{
	items[get_global_id(0)] = 
		get_global_size(0);
		//get_global_size(1);
}

このプログラムは、スレッドを(2×3)個作ります。 そして、バッファの先頭に、0番目の次元のスレッド個数を書き込みます。 ここでは0番目の次元を書き込んでいるので、2という数字(2×3の2)を書き込みます。

実行結果はこうなります。

2
2
0
0
0
0

myKernelProgram.clのコメントアウト位置を変えると、今度は1番目の次元の個数を書き込みます。 そうすると、バッファの先頭には3が書きこまれます(2×3の3です)。

3
3
0
0
0
0

グループのサイズ

以上では、スレッドの総数をGPU側で取得しました。
しかし、1グループ内のスレッド数を取得したい場合もあります。
それにはget_local_size関数を使います。
参考

size_t get_local_size(uint dimIndex)

dimIndexは取得するサイズの次元です。

たとえばスレッドを{{スレッド1,スレッド2、スレッド3}、{スレッド4、スレッド5,スレッド6}}と実行したとします。
全部で6つ。
1グループ3スレッドです。
そいう言う場合は、get_local_size(0)が3になります。
1グループ内のスレッド数を取得するのです。
サンプルを次のように改変します:

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 = 24;
        var buffer = context.CreateBuffer(
            MemFlags.READ_WRITE, 
            elementCount * sizeof(float)
            );
        kernel.SetArg(0, buffer);

        commandQueue.EnqueueNDRangeKernel(
            kernel,
            2,
            null,
            new[] { 4, 6 },
            new[] { 2, 3 });

        var readBack = new float[elementCount];
        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* items)
{
	items[get_global_id(0)] = 
		get_local_size(0);
		//get_local_size(1);
}

これを実行すると次のようになります:

2
2
2
2
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0

このプログラムはまず4×6このスレッドを作り、2×3のスレッドを1グループにして、グループ分けしています。 そしてバッファの先頭に、ローカルサイズ(1グループのスレッド数)を書き込みます。 この場合は1グループにX方向には2ある(2×3の2)ので、2が書きこまれています。

グループの個数

スレッドが次のように実行されているとしましょう:
{{スレッド0、スレッド1}、{スレッド2、スレッド3}、{スレッド4、スレッド5}}
このとき、グループの個数は3個です。

グループの個数を得るには、get_num_groups()関数を使います。 参考

size_t get_num_groups(uint dimIndex)

dimIndexは次元のインデックスです。X軸方向のグループの数を得たい時には0。Y軸方向なら1。Zなら2です。

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 = 24;
        var buffer = context.CreateBuffer(
            MemFlags.READ_WRITE, 
            elementCount * sizeof(float)
            );
        kernel.SetArg(0, buffer);

        commandQueue.EnqueueNDRangeKernel(
            kernel,
            2,
            null,
            new[] { 2, 12 },
            new[] { 2, 3 });

        var readBack = new float[elementCount];
        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* items)
{
	items[get_global_id(0)] = 
		get_num_groups(0);
		//get_num_groups(1);
}

このプログラムは、2×12個のスレッドを作り、それを2×3のスレッドを持つグループでグループ分けしています。 そうすると、グループの数は1×4となります。 結果、このプログラムは次のような文字列を出力します:

1
1
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0
0

グループの数は1×4なので、その最初の1を出力したのです。 もしmyKernelProgram.clのコメントアウト位置を変えれば、4を出力するでしょう。

拍手[0回]

PR