忍者ブログ

Memeplexes

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

[PR]

×

[PR]上記の広告は3ヶ月以上新規記事投稿のないブログに表示されています。新しい記事を書く事で広告が消えます。

C#でOpenCL入門 チュートリアルその１２　スレッドとグループの個数

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

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

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

size_t get_global_size(uint dimIndex)

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

サンプルコード

MyProgram.cs
using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[6]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(2, 3), new MultiDimension(1, 1));

        var readBack = new float[6];
        commandQueue.ReadBuffer(buffer, readBack);

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 

OpenCLWrappers.cs
using System;
using System.Runtime.InteropServices;
using System.Linq;

class Context
{
    public IntPtr InternalPointer { get; private set; }

    public Context(params IntPtr[] devices)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateContext(
            null,
            devices.Length,
            devices,
            null,
            IntPtr.Zero,
            out error
            );
    }

    ~Context()
    {
        OpenCLFunctions.clReleaseContext(InternalPointer);
    }
}

class CommandQueue
{
    public IntPtr InternalPointer { get; private set; }

    public CommandQueue(Context context, IntPtr device)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateCommandQueue(
            context.InternalPointer,
            device,
            0,
            out error
            );
    }

    ~CommandQueue()
    {
        OpenCLFunctions.clReleaseCommandQueue(InternalPointer);
    }

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

        OpenCLFunctions.clEnqueueReadBuffer(
            InternalPointer,
            buffer.InternalPointer,
            true,
            0,
            Math.Min(buffer.SizeInBytes, Marshal.SizeOf(typeof(T)) * systemBuffer.Length),
            handle.AddrOfPinnedObject(),
            0,
            IntPtr.Zero,
            IntPtr.Zero
            );

        handle.Free();
    }

    public void EnqueueRange(Kernel kernel, MultiDimension globalWorkSize, MultiDimension localWorkSize)
    {
        MultiDimension offset = new MultiDimension();
        OpenCLFunctions.clEnqueueNDRangeKernel(
            InternalPointer,
            kernel.InternalPointer,
            globalWorkSize.Dimension,
            ref offset,
            ref globalWorkSize,
            ref localWorkSize,
            0,
            null,
            IntPtr.Zero
            );
    }
}

class Buffer
{
    public IntPtr InternalPointer { get; private set; }
    public int SizeInBytes { get; private set; }

    private Buffer() { }

    ~Buffer()
    {
        OpenCLFunctions.clReleaseMemObject(InternalPointer);
    }

    public static Buffer FromCopiedHostMemory<T>(Context context, T[] initialData) where T : struct
    {
        Buffer result = new Buffer();
        result.SizeInBytes = Marshal.SizeOf(typeof(T)) * initialData.Length;

        int errorCode;
        GCHandle handle = GCHandle.Alloc(initialData, GCHandleType.Pinned);

        result.InternalPointer = OpenCLFunctions.clCreateBuffer(
            context.InternalPointer,
            MemoryFlags.CopyHostMemory,
            result.SizeInBytes,
            handle.AddrOfPinnedObject(),
            out errorCode
            );

        handle.Free();
        return result;
    }
}

class Program
{
    public IntPtr InternalPointer { get; private set; }

    public Program(Context context, params string[] sources)
    {
        int errorCode;

        InternalPointer = OpenCLFunctions.clCreateProgramWithSource(
            context.InternalPointer,
            sources.Length,
            sources,
            null,
            out errorCode
            );
    }

    ~Program()
    {
        OpenCLFunctions.clReleaseProgram(InternalPointer);
    }

    public void Build(params IntPtr[] devices)
    {
        int error = OpenCLFunctions.clBuildProgram(
            InternalPointer,
            devices.Length,
            devices,
            null,
            null,
            IntPtr.Zero
            );

        if (error != 0)
        {
            int paramValueSize = 0;
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                0, 
                null,
                out paramValueSize
                );
            System.Text.StringBuilder text = new System.Text.StringBuilder(paramValueSize);
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                paramValueSize,
                text,
                out paramValueSize);
            throw new Exception(text.ToString());
        }
    }
}

class Kernel
{
    public IntPtr InternalPointer { get; private set; }

    public Kernel(Program program, string functionName)
    {
        int errorCode;
        InternalPointer = OpenCLFunctions.clCreateKernel(
            program.InternalPointer,
            functionName, 
            out errorCode
            );
    }

    ~Kernel()
    {
        OpenCLFunctions.clReleaseKernel(InternalPointer);
    }

    public void SetArgument(int argumentIndex, Buffer buffer)
    {
        IntPtr bufferPointer = buffer.InternalPointer;
        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(IntPtr)),
            ref bufferPointer
            );
    }
}

OpenCLFunctions.cs
using System;
using System.Runtime.InteropServices;

static class OpenCLFunctions
{
    [DllImport("OpenCL.dll")]
    public static extern int clGetPlatformIDs(int entryCount, IntPtr[] platforms, out int platformCount);

    [DllImport("OpenCL.dll")]
    public static extern int clGetDeviceIDs(
        IntPtr platform,
        DeviceType deviceType,
        int entryCount,
        IntPtr[] devices,
        out int deviceCount
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateContext(
        IntPtr[] properties, 
        int deviceCount,
        IntPtr[] devices,
        NotifyContextCreated pfnNotify,
        IntPtr userData,
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseContext(IntPtr context);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateCommandQueue(
        IntPtr context,
        IntPtr device,
        long properties,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseCommandQueue(IntPtr commandQueue);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateBuffer(
        IntPtr context,
        MemoryFlags allocationAndUsage,
        int sizeInBytes,
        IntPtr hostPtr,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseMemObject(IntPtr memoryObject);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueReadBuffer(
        IntPtr commandQueue,
        IntPtr buffer,
        bool isBlocking,
        int offset,
        int sizeInBytes,
        IntPtr result,
        int numberOfEventsInWaitList,
        IntPtr eventWaitList,
        IntPtr eventObjectOut
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateProgramWithSource(
        IntPtr context,
        int count,
        string[] programSources, 
        int[] sourceLengths, 
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clBuildProgram(
        IntPtr program,
        int deviceCount, 
        IntPtr[] deviceList,
        string buildOptions,
        NotifyProgramBuilt notify,
        IntPtr userData
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseProgram(IntPtr program);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateKernel(IntPtr kernel, string functionName, out int errorCode);

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseKernel(IntPtr kernel);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, ref IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueNDRangeKernel(
        IntPtr commandQueue, 
        IntPtr kernel,
        int workDimension,
        ref MultiDimension globalWorkOffset, 
        ref MultiDimension globalWorkSize,
        ref MultiDimension localWorkSize,
        int countOfEventsInWaitList,
        IntPtr[] eventList,
        IntPtr eventObject
        );

    [DllImport("OpenCL.dll")]
    public static extern int clGetProgramBuildInfo(
        IntPtr program, 
        IntPtr device, 
        ProgramBuildInfoString paramName,
        int paramValueSize,
        System.Text.StringBuilder paramValue,
        out int paramValueSizeReturn
        );
}

delegate void NotifyContextCreated(string errorInfo, IntPtr privateInfoSize, int cb, IntPtr userData);
delegate void NotifyProgramBuilt(IntPtr program, IntPtr userData);

enum DeviceType : long
{
    Default = (1 << 0),
    Cpu = (1 << 1),
    Gpu = (1 << 2),
    Accelerator = (1 << 3),
    All = 0xFFFFFFFF
}

enum MemoryFlags : long
{
    ReadWrite = (1 << 0),
    WriteOnly = (1 << 1),
    ReadOnly = (1 << 2),
    UseHostMemory = (1 << 3),
    HostAccessible = (1 << 4),
    CopyHostMemory = (1 << 5)
}

struct MultiDimension
{
    public int X;
    public int Y;
    public int Z;
    public int Dimension;

    public MultiDimension(int x)
    {
        X = x;
        Y = 0;
        Z = 0;
        Dimension = 1;
    }

    public MultiDimension(int x, int y)
    {
        X = x;
        Y = y;
        Z = 0;
        Dimension = 2;
    }
}

enum ProgramBuildInfoString
{
    Options = 0x1182,
    Log = 0x1183
}

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)が３になります。
１グループ内のスレッド数を取得するのです。
サンプルを次のように改変します：


MyProgram.cs
using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[24]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(4, 6), new MultiDimension(2, 3));

        var readBack = new float[24];
        commandQueue.ReadBuffer(buffer, readBack);

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 

myKernelFunction.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なら２です。

この関数の動作サンプルは次のようになります：

MyProgram.cs
using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[24]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(2, 12), new MultiDimension(2, 3));

        var readBack = new float[24];
        commandQueue.ReadBuffer(buffer, readBack);

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 

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のコメントアウト位置を変えれば、４を出力するでしょう。














拍手[0回]

PR

C#でOpenCL入門 チュートリアルその１１　非バッファ引数

OpenCLでは、GPU側にデータを受け渡すのに、関数の引数という形を取ります。
DirectXではグローバル変数という形をとっていたのと対照的ですね。

前回まではGPU側に渡す引数は全てバッファという形をとっていましたが、今回はもっと簡単な引数を渡してみます。
GPUに一つのintを渡してみましょう。

奇妙なことに、intを渡すにせよバッファを渡すにせよ、使う関数は同じ（clSetKernelArg）です。
これまでバッファのハンドルを渡していた引数に、直接intな値を渡しこむのです。
intはCPU側に、バッファの本体はGPU内にあることを考えれば不思議ですね。
もしかすると関数内部で渡すデータの型を判別し、バッファならバッファを、intやその他普通の構造体ならそのデータそのものをGPUのカーネル引数として渡しているような仕組みになっているのかもしれませんね。
あるいは関数のやっていることは両方とも同じで、単にGPU内での引数の扱い方が違うというだけなのかもしれませんが。

サンプルコード

では、GPUにintを一つ渡すサンプルを書きます。

MyProgram.cs
using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[3]);
        kernel.SetArgument(0, buffer);
        kernel.SetArgument(1, 2);

        commandQueue.EnqueueRange(kernel, new MultiDimension(3), new MultiDimension(1));

        var readBack = new float[3];
        commandQueue.ReadBuffer(buffer, readBack);

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 


OpenCLWrappers.cs
using System;
using System.Runtime.InteropServices;
using System.Linq;

class Context
{
    public IntPtr InternalPointer { get; private set; }

    public Context(params IntPtr[] devices)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateContext(
            null,
            devices.Length,
            devices,
            null,
            IntPtr.Zero,
            out error
            );
    }

    ~Context()
    {
        OpenCLFunctions.clReleaseContext(InternalPointer);
    }
}

class CommandQueue
{
    public IntPtr InternalPointer { get; private set; }

    public CommandQueue(Context context, IntPtr device)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateCommandQueue(
            context.InternalPointer,
            device,
            0,
            out error
            );
    }

    ~CommandQueue()
    {
        OpenCLFunctions.clReleaseCommandQueue(InternalPointer);
    }

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

        OpenCLFunctions.clEnqueueReadBuffer(
            InternalPointer,
            buffer.InternalPointer,
            true,
            0,
            Math.Min(buffer.SizeInBytes, Marshal.SizeOf(typeof(T)) * systemBuffer.Length),
            handle.AddrOfPinnedObject(),
            0,
            IntPtr.Zero,
            IntPtr.Zero
            );

        handle.Free();
    }

    public void EnqueueRange(Kernel kernel, MultiDimension globalWorkSize, MultiDimension localWorkSize)
    {
        MultiDimension offset = new MultiDimension();
        OpenCLFunctions.clEnqueueNDRangeKernel(
            InternalPointer,
            kernel.InternalPointer,
            globalWorkSize.Dimension,
            ref offset,
            ref globalWorkSize,
            ref localWorkSize,
            0,
            null,
            IntPtr.Zero
            );
    }
}

class Buffer
{
    public IntPtr InternalPointer { get; private set; }
    public int SizeInBytes { get; private set; }

    private Buffer() { }

    ~Buffer()
    {
        OpenCLFunctions.clReleaseMemObject(InternalPointer);
    }

    public static Buffer FromCopiedHostMemory<T>(Context context, T[] initialData) where T : struct
    {
        Buffer result = new Buffer();
        result.SizeInBytes = Marshal.SizeOf(typeof(T)) * initialData.Length;

        int errorCode;
        GCHandle handle = GCHandle.Alloc(initialData, GCHandleType.Pinned);

        result.InternalPointer = OpenCLFunctions.clCreateBuffer(
            context.InternalPointer,
            MemoryFlags.CopyHostMemory,
            result.SizeInBytes,
            handle.AddrOfPinnedObject(),
            out errorCode
            );

        handle.Free();
        return result;
    }
}

class Program
{
    public IntPtr InternalPointer { get; private set; }

    public Program(Context context, params string[] sources)
    {
        int errorCode;

        InternalPointer = OpenCLFunctions.clCreateProgramWithSource(
            context.InternalPointer,
            sources.Length,
            sources,
            null,
            out errorCode
            );
    }

    ~Program()
    {
        OpenCLFunctions.clReleaseProgram(InternalPointer);
    }

    public void Build(params IntPtr[] devices)
    {
        int error = OpenCLFunctions.clBuildProgram(
            InternalPointer,
            devices.Length,
            devices,
            null,
            null,
            IntPtr.Zero
            );

        if (error != 0)
        {
            int paramValueSize = 0;
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                0, 
                null,
                out paramValueSize
                );
            System.Text.StringBuilder text = new System.Text.StringBuilder(paramValueSize);
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                paramValueSize,
                text,
                out paramValueSize);
            throw new Exception(text.ToString());
        }
    }
}

class Kernel
{
    public IntPtr InternalPointer { get; private set; }

    public Kernel(Program program, string functionName)
    {
        int errorCode;
        InternalPointer = OpenCLFunctions.clCreateKernel(
            program.InternalPointer,
            functionName, 
            out errorCode
            );
    }

    ~Kernel()
    {
        OpenCLFunctions.clReleaseKernel(InternalPointer);
    }

    public void SetArgument(int argumentIndex, Buffer buffer)
    {
        IntPtr bufferPointer = buffer.InternalPointer;
        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(IntPtr)),
            ref bufferPointer
            );
    }

    public void SetArgument<T>(int argumentIndex, T value)where T : struct
    {
        GCHandle handle = GCHandle.Alloc(value, GCHandleType.Pinned);

        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(T)),
            handle.AddrOfPinnedObject()
            );
        handle.Free();
    }
}

OpenCLFunctions.cs
using System;
using System.Runtime.InteropServices;

static class OpenCLFunctions
{
    [DllImport("OpenCL.dll")]
    public static extern int clGetPlatformIDs(int entryCount, IntPtr[] platforms, out int platformCount);

    [DllImport("OpenCL.dll")]
    public static extern int clGetDeviceIDs(
        IntPtr platform,
        DeviceType deviceType,
        int entryCount,
        IntPtr[] devices,
        out int deviceCount
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateContext(
        IntPtr[] properties, 
        int deviceCount,
        IntPtr[] devices,
        NotifyContextCreated pfnNotify,
        IntPtr userData,
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseContext(IntPtr context);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateCommandQueue(
        IntPtr context,
        IntPtr device,
        long properties,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseCommandQueue(IntPtr commandQueue);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateBuffer(
        IntPtr context,
        MemoryFlags allocationAndUsage,
        int sizeInBytes,
        IntPtr hostPtr,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseMemObject(IntPtr memoryObject);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueReadBuffer(
        IntPtr commandQueue,
        IntPtr buffer,
        bool isBlocking,
        int offset,
        int sizeInBytes,
        IntPtr result,
        int numberOfEventsInWaitList,
        IntPtr eventWaitList,
        IntPtr eventObjectOut
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateProgramWithSource(
        IntPtr context,
        int count,
        string[] programSources, 
        int[] sourceLengths, 
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clBuildProgram(
        IntPtr program,
        int deviceCount, 
        IntPtr[] deviceList,
        string buildOptions,
        NotifyProgramBuilt notify,
        IntPtr userData
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseProgram(IntPtr program);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateKernel(IntPtr kernel, string functionName, out int errorCode);

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseKernel(IntPtr kernel);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, ref IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueNDRangeKernel(
        IntPtr commandQueue, 
        IntPtr kernel,
        int workDimension,
        ref MultiDimension globalWorkOffset, 
        ref MultiDimension globalWorkSize,
        ref MultiDimension localWorkSize,
        int countOfEventsInWaitList,
        IntPtr[] eventList,
        IntPtr eventObject
        );

    [DllImport("OpenCL.dll")]
    public static extern int clGetProgramBuildInfo(
        IntPtr program, 
        IntPtr device, 
        ProgramBuildInfoString paramName,
        int paramValueSize,
        System.Text.StringBuilder paramValue,
        out int paramValueSizeReturn
        );
}

delegate void NotifyContextCreated(string errorInfo, IntPtr privateInfoSize, int cb, IntPtr userData);
delegate void NotifyProgramBuilt(IntPtr program, IntPtr userData);

enum DeviceType : long
{
    Default = (1 << 0),
    Cpu = (1 << 1),
    Gpu = (1 << 2),
    Accelerator = (1 << 3),
    All = 0xFFFFFFFF
}

enum MemoryFlags : long
{
    ReadWrite = (1 << 0),
    WriteOnly = (1 << 1),
    ReadOnly = (1 << 2),
    UseHostMemory = (1 << 3),
    HostAccessible = (1 << 4),
    CopyHostMemory = (1 << 5)
}

struct MultiDimension
{
    public int X;
    public int Y;
    public int Z;
    public int Dimension;

    public MultiDimension(int x)
    {
        X = x;
        Y = 0;
        Z = 0;
        Dimension = 1;
    }
}

enum ProgramBuildInfoString
{
    Options = 0x1182,
    Log = 0x1183
}


myKernelProgram.cl
 
__kernel void myKernelFunction(__global float* items, __const int number)
{
	items[get_global_id(0)] = number;
}


このサンプルプログラムは、まずGPU内に{0, 0, 0}というバッファを作ります。
そしてGPUに2という数字(int)を送ります。
GPUを動かして、2をバッファに代入します。
CPUにバッファの内容を読み戻し、表示します。
結果はこうなります。

2
2
2


拍手[0回]

C#でOpenCL入門 チュートリアルその１０　グループID、ローカル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には、スレッドがたくさんあるのですが、ただだらーっと並べられるのではなく、スレッドは３次元の構造を取ります。

gpgpuThreads3D.png


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

gpgpuThreadsGrouped.png

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

サンプルコード

MyProgram.cs


using System;

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 MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        const int bufferSize = 6;
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new MyItem[bufferSize]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(6), new MultiDimension(2));

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

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 

OpenCLWrappers.cs
using System;
using System.Runtime.InteropServices;
using System.Linq;

class Context
{
    public IntPtr InternalPointer { get; private set; }

    public Context(params IntPtr[] devices)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateContext(
            null,
            devices.Length,
            devices,
            null,
            IntPtr.Zero,
            out error
            );
    }

    ~Context()
    {
        OpenCLFunctions.clReleaseContext(InternalPointer);
    }
}

class CommandQueue
{
    public IntPtr InternalPointer { get; private set; }

    public CommandQueue(Context context, IntPtr device)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateCommandQueue(
            context.InternalPointer,
            device,
            0,
            out error
            );
    }

    ~CommandQueue()
    {
        OpenCLFunctions.clReleaseCommandQueue(InternalPointer);
    }

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

        OpenCLFunctions.clEnqueueReadBuffer(
            InternalPointer,
            buffer.InternalPointer,
            true,
            0,
            Math.Min(buffer.SizeInBytes, Marshal.SizeOf(typeof(T)) * systemBuffer.Length),
            handle.AddrOfPinnedObject(),
            0,
            IntPtr.Zero,
            IntPtr.Zero
            );

        handle.Free();
    }

    public void EnqueueRange(Kernel kernel, MultiDimension globalWorkSize, MultiDimension localWorkSize)
    {
        MultiDimension offset = new MultiDimension();
        OpenCLFunctions.clEnqueueNDRangeKernel(
            InternalPointer,
            kernel.InternalPointer,
            globalWorkSize.Dimension,
            ref offset,
            ref globalWorkSize,
            ref localWorkSize,
            0,
            null,
            IntPtr.Zero
            );
    }
}

class Buffer
{
    public IntPtr InternalPointer { get; private set; }
    public int SizeInBytes { get; private set; }

    private Buffer() { }

    ~Buffer()
    {
        OpenCLFunctions.clReleaseMemObject(InternalPointer);
    }

    public static Buffer FromCopiedHostMemory<T>(Context context, T[] initialData) where T : struct
    {
        Buffer result = new Buffer();
        result.SizeInBytes = Marshal.SizeOf(typeof(T)) * initialData.Length;

        int errorCode;
        GCHandle handle = GCHandle.Alloc(initialData, GCHandleType.Pinned);

        result.InternalPointer = OpenCLFunctions.clCreateBuffer(
            context.InternalPointer,
            MemoryFlags.CopyHostMemory,
            result.SizeInBytes,
            handle.AddrOfPinnedObject(),
            out errorCode
            );

        handle.Free();
        return result;
    }
}

class Program
{
    public IntPtr InternalPointer { get; private set; }

    public Program(Context context, params string[] sources)
    {
        int errorCode;

        InternalPointer = OpenCLFunctions.clCreateProgramWithSource(
            context.InternalPointer,
            sources.Length,
            sources,
            null,
            out errorCode
            );
    }

    ~Program()
    {
        OpenCLFunctions.clReleaseProgram(InternalPointer);
    }

    public void Build(params IntPtr[] devices)
    {
        int error = OpenCLFunctions.clBuildProgram(
            InternalPointer,
            devices.Length,
            devices,
            null,
            null,
            IntPtr.Zero
            );

        if (error != 0)
        {
            int paramValueSize = 0;
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                0, 
                null,
                out paramValueSize
                );
            System.Text.StringBuilder text = new System.Text.StringBuilder(paramValueSize);
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                paramValueSize,
                text,
                out paramValueSize);
            throw new Exception(text.ToString());
        }
    }
}

class Kernel
{
    public IntPtr InternalPointer { get; private set; }

    public Kernel(Program program, string functionName)
    {
        int errorCode;
        InternalPointer = OpenCLFunctions.clCreateKernel(
            program.InternalPointer,
            functionName, 
            out errorCode
            );
    }

    ~Kernel()
    {
        OpenCLFunctions.clReleaseKernel(InternalPointer);
    }

    public void SetArgument(int argumentIndex, Buffer buffer)
    {
        IntPtr bufferPointer = buffer.InternalPointer;
        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(IntPtr)),
            ref bufferPointer
            );
    }
}

OpenCLFunctions.cs
using System;
using System.Runtime.InteropServices;

static class OpenCLFunctions
{
    [DllImport("OpenCL.dll")]
    public static extern int clGetPlatformIDs(int entryCount, IntPtr[] platforms, out int platformCount);

    [DllImport("OpenCL.dll")]
    public static extern int clGetDeviceIDs(
        IntPtr platform,
        DeviceType deviceType,
        int entryCount,
        IntPtr[] devices,
        out int deviceCount
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateContext(
        IntPtr[] properties, 
        int deviceCount,
        IntPtr[] devices,
        NotifyContextCreated pfnNotify,
        IntPtr userData,
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseContext(IntPtr context);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateCommandQueue(
        IntPtr context,
        IntPtr device,
        long properties,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseCommandQueue(IntPtr commandQueue);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateBuffer(
        IntPtr context,
        MemoryFlags allocationAndUsage,
        int sizeInBytes,
        IntPtr hostPtr,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseMemObject(IntPtr memoryObject);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueReadBuffer(
        IntPtr commandQueue,
        IntPtr buffer,
        bool isBlocking,
        int offset,
        int sizeInBytes,
        IntPtr result,
        int numberOfEventsInWaitList,
        IntPtr eventWaitList,
        IntPtr eventObjectOut
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateProgramWithSource(
        IntPtr context,
        int count,
        string[] programSources, 
        int[] sourceLengths, 
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clBuildProgram(
        IntPtr program,
        int deviceCount, 
        IntPtr[] deviceList,
        string buildOptions,
        NotifyProgramBuilt notify,
        IntPtr userData
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseProgram(IntPtr program);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateKernel(IntPtr kernel, string functionName, out int errorCode);

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseKernel(IntPtr kernel);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, ref IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueNDRangeKernel(
        IntPtr commandQueue, 
        IntPtr kernel,
        int workDimension,
        ref MultiDimension globalWorkOffset, 
        ref MultiDimension globalWorkSize,
        ref MultiDimension localWorkSize,
        int countOfEventsInWaitList,
        IntPtr[] eventList,
        IntPtr eventObject
        );

    [DllImport("OpenCL.dll")]
    public static extern int clGetProgramBuildInfo(
        IntPtr program, 
        IntPtr device, 
        ProgramBuildInfoString paramName,
        int paramValueSize,
        System.Text.StringBuilder paramValue,
        out int paramValueSizeReturn
        );
}

delegate void NotifyContextCreated(string errorInfo, IntPtr privateInfoSize, int cb, IntPtr userData);
delegate void NotifyProgramBuilt(IntPtr program, IntPtr userData);

enum DeviceType : long
{
    Default = (1 << 0),
    Cpu = (1 << 1),
    Gpu = (1 << 2),
    Accelerator = (1 << 3),
    All = 0xFFFFFFFF
}

enum MemoryFlags : long
{
    ReadWrite = (1 << 0),
    WriteOnly = (1 << 1),
    ReadOnly = (1 << 2),
    UseHostMemory = (1 << 3),
    HostAccessible = (1 << 4),
    CopyHostMemory = (1 << 5)
}

struct MultiDimension
{
    public int X;
    public int Y;
    public int Z;
    public int Dimension;

    public MultiDimension(int x)
    {
        X = x;
        Y = 0;
        Z = 0;
        Dimension = 1;
    }
}

enum ProgramBuildInfoString
{
    Options = 0x1182,
    Log = 0x1183
}


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
というわけですね。



拍手[1回]

C#でOpenCL入門 チュートリアルその９　OpenCL Cのコンパイルエラー捕捉

GPGPUのチュートリアルなのに足踏みばかりしていたこのシリーズですが、前回、ついにGPGPUらしいことをしました。
バッファの中に、数字を代入したのです（マルチスレッドで）。
バッファの要素の数だけGPUでスレッドを起動して、そのスレッド番号を所定の場所に書き込みました。

今回も、少し足踏みをします。
前回のプログラムは、GPU側で動くプログラムに不備（コンパイルエラー）があった場合、それを知ることは出来ませんでした。
妙な値がConsole.WriteLineで出力されるだけで、デバッグしにくいものでした。
今回は、OpenCL Cで書かれたプログラムをビルドするときの、コンパイルエラーを補足する方法をメモします。

コンパイルエラーを見つける

OpenCL C言語で書かれたプログラムのコンパイルエラーは、clGetProgramBuildInfo関数で取得できます。
参考

cl_int clGetProgramBuildInfo(
    cl_program program,
    cl_device_id device,
    cl_program_build_info param_name,
    size_t param_value_size,
    void *param_value,
    size_t *param_value_size_ret
)

programはビルドしたプログラムです。
deviceはビルド情報を得るデバイスです。
param_nameは取得したい情報の種類です。次の３種類があります。

cl_program_build_info 値 param_valueの型
 解説
CL_PROGRAM_BUILD_STATUS 0x1181 cl_build_status ビルドステータスを返します。次の４種類のいずれかを返します。
CL_BUILD_NONE(-1)はまだビルドされていないことを表します。
CL_BUILD_ERROR(-2)はビルドにエラーが起きたことを表します。
CL_BUILD_SUCCESS(0)はビルドが成功したことを表します。
CL_BUILD_IN_PROGRESS(-3)はまだビルド中であることを表します。
CL_PROGRAM_BUILD_OPTIONS 0x1182 char[] clBuildProgram関数のoptions引数に渡されたオプションを表します。
もしプログラムがCL_BUILD_NONEな状態だった場合には、空の文字列が返されます。
CL_PROGRAM_BUILD_LOG 0x1183 char[] ビルドログを表します。
もしプログラムがCL_BUILD_NONEな状態だった場合には、空の文字列が返されます。

param_valueはこの関数の結果（ビルド情報）を格納するメモリへのポインタです。
param_value_sizeはparam_valueで指定されているメモリのサイズです。
param_value_size_retは、取得したいビルド情報のサイズを取得します。NULLでも構いません。

サンプルコード

では早速、エラーの起きそうなOpenCL Cのコードを書いて、コンパイルエラーを検出してみましょう。
まずこんなOpenCL Cのコードを用意します：

myKernelProgram.cl
__kernel void myKernelFunction(__global float* items)
{
	items[0] = 0
}

これはいけませんね。
items[0] = 0の後にセミコロン（；）が足りません。
これではコンパイルエラーが発生しコンパイル出来ません。

エラーを補足するC#側のコードは次のようになります。

MyProgram.cs
using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[1]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(1), new MultiDimension(1));

        var readBack = new float[1];
        commandQueue.ReadBuffer(buffer, readBack);

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 

OpenCLWrappers.cs
using System;
using System.Runtime.InteropServices;
using System.Linq;

class Context
{
    public IntPtr InternalPointer { get; private set; }

    public Context(params IntPtr[] devices)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateContext(
            null,
            devices.Length,
            devices,
            null,
            IntPtr.Zero,
            out error
            );
    }

    ~Context()
    {
        OpenCLFunctions.clReleaseContext(InternalPointer);
    }
}

class CommandQueue
{
    public IntPtr InternalPointer { get; private set; }

    public CommandQueue(Context context, IntPtr device)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateCommandQueue(
            context.InternalPointer,
            device,
            0,
            out error
            );
    }

    ~CommandQueue()
    {
        OpenCLFunctions.clReleaseCommandQueue(InternalPointer);
    }

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

        OpenCLFunctions.clEnqueueReadBuffer(
            InternalPointer,
            buffer.InternalPointer,
            true,
            0,
            Math.Min(buffer.SizeInBytes, Marshal.SizeOf(typeof(T)) * systemBuffer.Length),
            handle.AddrOfPinnedObject(),
            0,
            IntPtr.Zero,
            IntPtr.Zero
            );

        handle.Free();
    }

    public void EnqueueRange(Kernel kernel, MultiDimension globalWorkSize, MultiDimension localWorkSize)
    {
        MultiDimension offset = new MultiDimension();
        OpenCLFunctions.clEnqueueNDRangeKernel(
            InternalPointer,
            kernel.InternalPointer,
            globalWorkSize.Dimension,
            ref offset,
            ref globalWorkSize,
            ref localWorkSize,
            0,
            null,
            IntPtr.Zero
            );
    }
}

class Buffer
{
    public IntPtr InternalPointer { get; private set; }
    public int SizeInBytes { get; private set; }

    private Buffer() { }

    ~Buffer()
    {
        OpenCLFunctions.clReleaseMemObject(InternalPointer);
    }

    public static Buffer FromCopiedHostMemory<T>(Context context, T[] initialData) where T : struct
    {
        Buffer result = new Buffer();
        result.SizeInBytes = Marshal.SizeOf(typeof(T)) * initialData.Length;

        int errorCode;
        GCHandle handle = GCHandle.Alloc(initialData, GCHandleType.Pinned);

        result.InternalPointer = OpenCLFunctions.clCreateBuffer(
            context.InternalPointer,
            MemoryFlags.CopyHostMemory,
            result.SizeInBytes,
            handle.AddrOfPinnedObject(),
            out errorCode
            );

        handle.Free();
        return result;
    }
}

class Program
{
    public IntPtr InternalPointer { get; private set; }

    public Program(Context context, params string[] sources)
    {
        int errorCode;

        InternalPointer = OpenCLFunctions.clCreateProgramWithSource(
            context.InternalPointer,
            sources.Length,
            sources,
            null,
            out errorCode
            );
    }

    ~Program()
    {
        OpenCLFunctions.clReleaseProgram(InternalPointer);
    }

    public void Build(params IntPtr[] devices)
    {
        int error = OpenCLFunctions.clBuildProgram(
            InternalPointer,
            devices.Length,
            devices,
            null,
            null,
            IntPtr.Zero
            );

        if (error != 0)
        {
            int paramValueSize = 0;
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                0, 
                null,
                out paramValueSize
                );
            System.Text.StringBuilder text = new System.Text.StringBuilder(paramValueSize);
            OpenCLFunctions.clGetProgramBuildInfo(
                InternalPointer,
                devices.First(),
                ProgramBuildInfoString.Log,
                paramValueSize,
                text,
                out paramValueSize);
            throw new Exception(text.ToString());
        }
    }
}

class Kernel
{
    public IntPtr InternalPointer { get; private set; }

    public Kernel(Program program, string functionName)
    {
        int errorCode;
        InternalPointer = OpenCLFunctions.clCreateKernel(
            program.InternalPointer,
            functionName, 
            out errorCode
            );
    }

    ~Kernel()
    {
        OpenCLFunctions.clReleaseKernel(InternalPointer);
    }

    public void SetArgument(int argumentIndex, Buffer buffer)
    {
        IntPtr bufferPointer = buffer.InternalPointer;
        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(IntPtr)),
            ref bufferPointer
            );
    }
}

OpenCLFunctions.cs
using System;
using System.Runtime.InteropServices;

static class OpenCLFunctions
{
    [DllImport("OpenCL.dll")]
    public static extern int clGetPlatformIDs(int entryCount, IntPtr[] platforms, out int platformCount);

    [DllImport("OpenCL.dll")]
    public static extern int clGetDeviceIDs(
        IntPtr platform,
        DeviceType deviceType,
        int entryCount,
        IntPtr[] devices,
        out int deviceCount
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateContext(
        IntPtr[] properties, 
        int deviceCount,
        IntPtr[] devices,
        NotifyContextCreated pfnNotify,
        IntPtr userData,
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseContext(IntPtr context);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateCommandQueue(
        IntPtr context,
        IntPtr device,
        long properties,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseCommandQueue(IntPtr commandQueue);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateBuffer(
        IntPtr context,
        MemoryFlags allocationAndUsage,
        int sizeInBytes,
        IntPtr hostPtr,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseMemObject(IntPtr memoryObject);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueReadBuffer(
        IntPtr commandQueue,
        IntPtr buffer,
        bool isBlocking,
        int offset,
        int sizeInBytes,
        IntPtr result,
        int numberOfEventsInWaitList,
        IntPtr eventWaitList,
        IntPtr eventObjectOut
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateProgramWithSource(
        IntPtr context,
        int count,
        string[] programSources, 
        int[] sourceLengths, 
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clBuildProgram(
        IntPtr program,
        int deviceCount, 
        IntPtr[] deviceList,
        string buildOptions,
        NotifyProgramBuilt notify,
        IntPtr userData
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseProgram(IntPtr program);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateKernel(IntPtr kernel, string functionName, out int errorCode);

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseKernel(IntPtr kernel);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, ref IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueNDRangeKernel(
        IntPtr commandQueue, 
        IntPtr kernel,
        int workDimension,
        ref MultiDimension globalWorkOffset, 
        ref MultiDimension globalWorkSize,
        ref MultiDimension localWorkSize,
        int countOfEventsInWaitList,
        IntPtr[] eventList,
        IntPtr eventObject
        );

    [DllImport("OpenCL.dll")]
    public static extern int clGetProgramBuildInfo(
        IntPtr program, 
        IntPtr device, 
        ProgramBuildInfoString paramName,
        int paramValueSize,
        System.Text.StringBuilder paramValue,
        out int paramValueSizeReturn
        );
}

delegate void NotifyContextCreated(string errorInfo, IntPtr privateInfoSize, int cb, IntPtr userData);
delegate void NotifyProgramBuilt(IntPtr program, IntPtr userData);

enum DeviceType : long
{
    Default = (1 << 0),
    Cpu = (1 << 1),
    Gpu = (1 << 2),
    Accelerator = (1 << 3),
    All = 0xFFFFFFFF
}

enum MemoryFlags : long
{
    ReadWrite = (1 << 0),
    WriteOnly = (1 << 1),
    ReadOnly = (1 << 2),
    UseHostMemory = (1 << 3),
    HostAccessible = (1 << 4),
    CopyHostMemory = (1 << 5)
}

struct MultiDimension
{
    public int X;
    public int Y;
    public int Z;
    public int Dimension;

    public MultiDimension(int x)
    {
        X = x;
        Y = 0;
        Z = 0;
        Dimension = 1;
    }
}

enum ProgramBuildInfoString
{
    Options = 0x1182,
    Log = 0x1183
}

OpenCLWrappers.csに注目してください。
ここではビルドしてエラーが出た場合、例外をスローしています。
（もちろん本当はSystem.Exceptionではなく専用の例外クラスを定義してそれをスローするのが望ましいのですが）
 
このプログラムをVisual StudioでF5キーを押し実行すると、このように例外が出てくれます：
OpenCLBuildExceptionUnhandled.JPG


この例外のメッセージには、

<program source>:3:14: error: expected ';' after expression
    items[0] = 0


とあります。
つまり「items[0] = 0のあとに';'がありませんよ。3行目の14文字目です。」と言っているわけですね。
きちんとコンパイルエラーが警告されています。
良かったですね！


















 

拍手[0回]

C#でOpenCL入門 チュートリアルその８　データ並列

 前回はついにGPUを使って計算（というか代入）しました。
しかし、せっかくGPUを使っているというのにシングルスレッドでした。
これではなんのためにGPUを使っているのかわかりません。

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

カーネル（OpenCL Cで書いたGPU用プログラム）を並列に実行するには、clEnqueueNDRangeKernel関数を使います。
参考

cl_int clEnqueueNDRangeKernel(
    cl_command_queue command_queue,
    cl_kernel kernel,
    cl_uint work_dim,
    const size_t *global_work_offset,
    const size_t *global_work_size,
    const size_t *local_work_size,
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event
)

command_queueは実行を受け持つコマンドキューです。（ですから、直ちには実行されません。）
kernelは実行するカーネル関数です。
work_dimは実行するスレッド群の次元です。1~CL_DEVICE_MAX_WORK_ITEM_DIMENSIONSまでの値を取ることが出来ます。
global_work_offsetはグローバルIDのオフセットです。NULLでも構いません。
global_work_sizeは全スレッドの量を表します。
local_work_sizeは１グループのスレッドの量を表します。
num_events_in_wait_listとevent_wait_listは、この実行コマンドを実行にうつす前にすでに起きていなければいけないイベントのリストを表します。
eventは、このコマンドが実行されたことを知らせるイベントを返します。NULLでも構いません。

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

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

gpgpuThreads3D.png



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

gpgpuThreadsGrouped.png

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

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


MyProgram.cs
using System;

class MyProgram
{
    static void Main()
    {
        IntPtr device = getDevices(getPlatforms()[0], DeviceType.Default)[0];
        Context context = new Context(device);
        CommandQueue commandQueue = new CommandQueue(context, device);

        Program program = new Program(context, System.IO.File.ReadAllText("myKernelProgram.cl"));
        program.Build(device);

        Kernel kernel = new Kernel(program, "myKernelFunction");
        const int bufferSize = 3;
        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[bufferSize]);
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueRange(kernel, new MultiDimension(bufferSize), new MultiDimension(1));

        float[] readBack = new float[bufferSize];
        commandQueue.ReadBuffer(buffer, readBack);

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

    private static IntPtr[] getDevices(IntPtr platform, DeviceType deviceType)
    {
        int deviceCount;
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, 0, null, out deviceCount);
       
        IntPtr[] result = new IntPtr[deviceCount];
        OpenCLFunctions.clGetDeviceIDs(platform, deviceType, deviceCount, result, out deviceCount);
        return result;
    }


    private static IntPtr[] getPlatforms()
    {
        int platformCount;
        OpenCLFunctions.clGetPlatformIDs(0, null, out platformCount);

        IntPtr[] result = new IntPtr[platformCount];
        OpenCLFunctions.clGetPlatformIDs(platformCount, result, out platformCount);
        return result;
    }
}
 

OpenCLWrappers.cs
using System;
using System.Runtime.InteropServices;

class Context
{
    public IntPtr InternalPointer { get; private set; }

    public Context(params IntPtr[] devices)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateContext(
            null,
            devices.Length,
            devices,
            null,
            IntPtr.Zero,
            out error
            );
    }

    ~Context()
    {
        OpenCLFunctions.clReleaseContext(InternalPointer);
    }
}

class CommandQueue
{
    public IntPtr InternalPointer { get; private set; }

    public CommandQueue(Context context, IntPtr device)
    {
        int error;
        InternalPointer = OpenCLFunctions.clCreateCommandQueue(
            context.InternalPointer,
            device,
            0,
            out error
            );
    }

    ~CommandQueue()
    {
        OpenCLFunctions.clReleaseCommandQueue(InternalPointer);
    }

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

        OpenCLFunctions.clEnqueueReadBuffer(
            InternalPointer,
            buffer.InternalPointer,
            true,
            0,
            Math.Min(buffer.SizeInBytes, Marshal.SizeOf(typeof(T)) * systemBuffer.Length),
            handle.AddrOfPinnedObject(),
            0,
            IntPtr.Zero,
            IntPtr.Zero
            );

        handle.Free();
    }

    public void EnqueueRange(Kernel kernel, MultiDimension globalWorkSize, MultiDimension localWorkSize)
    {
        MultiDimension offset = new MultiDimension();
        OpenCLFunctions.clEnqueueNDRangeKernel(
            InternalPointer,
            kernel.InternalPointer,
            globalWorkSize.Dimension,
            ref offset,
            ref globalWorkSize,
            ref localWorkSize,
            0,
            null,
            IntPtr.Zero
            );
    }
}

class Buffer
{
    public IntPtr InternalPointer { get; private set; }
    public int SizeInBytes { get; private set; }

    private Buffer() { }

    ~Buffer()
    {
        OpenCLFunctions.clReleaseMemObject(InternalPointer);
    }

    public static Buffer FromCopiedHostMemory<T>(Context context, T[] initialData) where T : struct
    {
        Buffer result = new Buffer();
        result.SizeInBytes = Marshal.SizeOf(typeof(T)) * initialData.Length;

        int errorCode;
        GCHandle handle = GCHandle.Alloc(initialData, GCHandleType.Pinned);

        result.InternalPointer = OpenCLFunctions.clCreateBuffer(
            context.InternalPointer,
            MemoryFlags.CopyHostMemory,
            result.SizeInBytes,
            handle.AddrOfPinnedObject(),
            out errorCode
            );

        handle.Free();
        return result;
    }
}

class Program
{
    public IntPtr InternalPointer { get; private set; }

    public Program(Context context, params string[] sources)
    {
        int errorCode;

        InternalPointer = OpenCLFunctions.clCreateProgramWithSource(
            context.InternalPointer,
            sources.Length,
            sources,
            null,
            out errorCode
            );
    }

    ~Program()
    {
        OpenCLFunctions.clReleaseProgram(InternalPointer);
    }

    public void Build(params IntPtr[] devices)
    {
        OpenCLFunctions.clBuildProgram(
            InternalPointer,
            devices.Length,
            devices,
            null,
            null,
            IntPtr.Zero
            );
    }
}

class Kernel
{
    public IntPtr InternalPointer { get; private set; }

    public Kernel(Program program, string functionName)
    {
        int errorCode;
        InternalPointer = OpenCLFunctions.clCreateKernel(
            program.InternalPointer,
            functionName, 
            out errorCode
            );
    }

    ~Kernel()
    {
        OpenCLFunctions.clReleaseKernel(InternalPointer);
    }

    public void SetArgument(int argumentIndex, Buffer buffer)
    {
        IntPtr bufferPointer = buffer.InternalPointer;
        OpenCLFunctions.clSetKernelArg(
            InternalPointer,
            argumentIndex,
            Marshal.SizeOf(typeof(IntPtr)),
            ref bufferPointer
            );
    }
}

OpenCLFunctions.cs
using System;
using System.Runtime.InteropServices;

static class OpenCLFunctions
{
    [DllImport("OpenCL.dll")]
    public static extern int clGetPlatformIDs(int entryCount, IntPtr[] platforms, out int platformCount);

    [DllImport("OpenCL.dll")]
    public static extern int clGetDeviceIDs(
        IntPtr platform,
        DeviceType deviceType,
        int entryCount,
        IntPtr[] devices,
        out int deviceCount
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateContext(
        IntPtr[] properties, 
        int deviceCount,
        IntPtr[] devices,
        NotifyContextCreated pfnNotify,
        IntPtr userData,
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseContext(IntPtr context);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateCommandQueue(
        IntPtr context,
        IntPtr device,
        long properties,
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseCommandQueue(IntPtr commandQueue);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateBuffer(
        IntPtr context,
        MemoryFlags allocationAndUsage,
        int sizeInBytes,
        IntPtr hostPtr,//
        out int errorCodeReturn
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseMemObject(IntPtr memoryObject);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueReadBuffer(
        IntPtr commandQueue,
        IntPtr buffer,
        bool isBlocking,
        int offset,
        int sizeInBytes,
        IntPtr result,
        int numberOfEventsInWaitList,
        IntPtr eventWaitList,
        IntPtr eventObjectOut
        );

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateProgramWithSource(
        IntPtr context,
        int count,
        string[] programSources, 
        int[] sourceLengths, 
        out int errorCode
        );

    [DllImport("OpenCL.dll")]
    public static extern int clBuildProgram(
        IntPtr program,
        int deviceCount, 
        IntPtr[] deviceList,
        string buildOptions,
        NotifyProgramBuilt notify,
        IntPtr userData
        );

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseProgram(IntPtr program);

    [DllImport("OpenCL.dll")]
    public static extern IntPtr clCreateKernel(IntPtr kernel, string functionName, out int errorCode);

    [DllImport("OpenCL.dll")]
    public static extern int clReleaseKernel(IntPtr kernel);

    [DllImport("OpenCL.dll")]
    public static extern int clSetKernelArg(IntPtr kernel, int argumentIndex, int size, ref IntPtr value);

    [DllImport("OpenCL.dll")]
    public static extern int clEnqueueNDRangeKernel(
        IntPtr commandQueue, 
        IntPtr kernel,
        int workDimension,
        ref MultiDimension globalWorkOffset, 
        ref MultiDimension globalWorkSize,
        ref MultiDimension localWorkSize,
        int countOfEventsInWaitList,
        IntPtr[] eventList,
        IntPtr eventObject
        );
}

delegate void NotifyContextCreated(string errorInfo, IntPtr privateInfoSize, int cb, IntPtr userData);
delegate void NotifyProgramBuilt(IntPtr program, IntPtr userData);

enum DeviceType : long
{
    Default = (1 << 0),
    Cpu = (1 << 1),
    Gpu = (1 << 2),
    Accelerator = (1 << 3),
    All = 0xFFFFFFFF
}

enum MemoryFlags : long
{
    ReadWrite = (1 << 0),
    WriteOnly = (1 << 1),
    ReadOnly = (1 << 2),
    UseHostMemory = (1 << 3),
    HostAccessible = (1 << 4),
    CopyHostMemory = (1 << 5)
}

struct MultiDimension
{
    public int X;
    public int Y;
    public int Z;
    public int Dimension;

    public MultiDimension(int x)
    {
        X = x;
        Y = 0;
        Z = 0;
        Dimension = 1;
    }
}


myKernelProgram.cl
__kernel void myKernelFunction(__global float* numbers)
{
	int globalThreadID = get_global_id(0);
    numbers[globalThreadID] = globalThreadID;
}


このプログラムは実行すると、次のような出力をします：
 
0
1
2

このプログラムは、まずGPUのメモリに{0, 0, 0}というバッファを作ります。
そしてGPUが３スレッド使って、それぞれの要素にグローバルID（スレッドのIDのこと）を代入します。
スレッドのIDは0, 1, 2です。
それが代入されて、バッファは{0, 1, 2}となります。
最後にそれがCPU側に読み戻されて、012という出力になったわけです。
 






拍手[0回]