忍者ブログ

Memeplexes

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

C#でOpenCL入門 チュートリアルその７　カーネル

 今回のメモは、カーネルについてです。
カーネルとは、GPUで動かす関数のことです。

前回、「プログラム」という似た概念が出てきました。
プログラムは複数（１つ以上）のカーネルをまとめたものです。
プログラムはGPU版dllといったところでしょうか。

前回はプログラムのロード、つまりカーネルの塊をロードしたのです。
今回は、その中からカーネルを一つ取り出し、それをGPUで実行します。

今回することを３つにまとめると次のようになります：
１．プログラムからカーネル生成
２．カーネルに引数のセット
３．カーネルの実行

カーネルの生成

カーネルの生成にはclCreateKernel関数を使います。
参考

cl_kernel clCreateKernel(
    cl_program program,
    const char *kernel_name,
    cl_int *errcode_ret
)
programはカーネルを持つプログラムです。
kernel_nameはカーネル関数の名前です。関数には__kernelが付いていなければいけませんん。
errcode_retはこの関数のエラーコードです。NULLでもかまいません。

戻り値は生成されたカーネルです。
 カーネルの破棄

生成したカーネルは使い終わったら破棄しなければいけません。
破棄にはclReleaseKernel関数を使います。
参考

cl_int clReleaseKernel(cl_kernel kernel)
kernelは破棄するカーネルです。

 カーネルにパラメーターをセット

生成したカーネルは実行する前にパラメーターというか、引数をセットする必要があります。
ここで言う引数とは、計算に必要なデータや、計算結果を格納するバッファのことです。

引数のセットにはclSetKernelArg関数を使います。
参考

cl_int clSetKernelArg(
    cl_kernel kernel,
    cl_uint arg_index,
    size_t arg_size,
    const void *arg_value
)

kernelは引数をセットするカーネルです。
arg_indexは引数の位置。例えば最初の引数なら0です。
arg_sizeは、arg_valueに渡したメモリのサイズです。ただし、引数が__localな場合には、単にサイズだけをセットし、arg_valueにはNULLを入れてやります。
arg_valueはセットする引数です。ここにはバッファやイメージをセットすることが出来ますし、単なるintをセットすることも出来ます。これはNULLでも構いません。どういうときにNULLになるかというと、たとえば引数が__localな場合です。__localが付いている場合には必ずNULLでなくてはいけません。

 カーネルの実行

下準備が全て終わったカーネルは、clEnqueueTask関数で実行できます。
参考

cl_int clEnqueueTask(
    cl_command_queue command_queue,
    cl_kernel kernel,
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event
)

command_queueはカーネルの実行を受け持つコマンドキューです。（なので、すぐには実行されません。キューに入れられるだけです）
kernelは実行するカーネルです。
num_events_in_wait_listはevent_wait_listに入っているイベントの数です。
event_wait_listはこの実行コマンドを実行する前に完結しなければならないイベント群です。
eventはこの実行が終わったことを教えるイベントオブジェクトを返します。

実はこの関数、カーネルをシングルスレッドで実行します。
あまりGPUっぽくありませんね。
なぜこんな関数を使うかというと、そのほうが単純だからです。
よりGPUらしい並列計算を行うには、もう少し複雑な別の関数、clEnqueueNDRangeKernelを使います（後述）。
そちらは、何十スレッドでも何百スレッドでも一気に実行できます。

サンプルコード

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, 2, 3 });
        kernel.SetArgument(0, buffer);

        commandQueue.EnqueueTask(kernel);

        float[] 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 EnqueueTask(Kernel kernel)
    {
        OpenCLFunctions.clEnqueueTask(
            InternalPointer,
            kernel.InternalPointer,
            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 clEnqueueTask(
        IntPtr commandQueue,
        IntPtr kernel,
        int countOfEventsInWaitList,
        IntPtr[] eventWaitList,
        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)
}

myKernelProgram.cl
__kernel void myKernelFunction(__global float* numbers)
{
    numbers[0] = 3;
    numbers[1] = 4;
    numbers[2] = 5;
}


このプログラムの実行結果はこうなります：

3
4
5

このプログラムが何をしているのかというと、まず{1, 2, 3}というバッファをGPUのメモリに作ります。
そしてそれをGPUが、{3, 4, 5}というデータに書き換えているのです。
バッファはCPU側に読み戻し、Console.WriteLineで出力しています。











拍手[0回]

PR