忍者ブログ

Memeplexes

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

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

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

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


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

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

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

size_t get_global_size(uint dimIndex)

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

サンプルコード

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
2
0
0
0
0

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

3
3
0
0
0
0

グループのサイズ

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

size_t get_local_size(uint dimIndex)

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

たとえばスレッドを｛｛スレッド１，スレッド２、スレッド３｝、｛スレッド４、スレッド５,スレッド６｝｝と実行したとします。
全部で６つ。
１グループ３スレッドです。
そいう言う場合は、get_local_size(0)が３になります。
１グループ内のスレッド数を取得するのです。
サンプルを次のように改変します：

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

このプログラムはまず４×６このスレッドを作り、２×３のスレッドを１グループにして、グループ分けしています。 そしてバッファの先頭に、ローカルサイズ（１グループのスレッド数）を書き込みます。 この場合は１グループにX方向には２ある（２×３の２）ので、２が書きこまれています。

グループの個数

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

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

size_t get_num_groups(uint dimIndex)

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

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);
}

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

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

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

拍手[1回]

PR