忍者ブログ

Memeplexes

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

C#でOpenCL入門 チュートリアルその６　プログラム

 前回はGPU内にバッファと呼ばれるメモリ領域を作り、そこでデータを読み書きしました。
しかしそれだけではいけません。
GPUで計算するというのが本来の目的です。
データを書き込んだり読み込んだりするだけではやりたい事の半分しか終わっていません。
バッファにセットしたデータを使ってなにか計算しなければいけないのです。

さてこの計算は一気にできるものではありません。
２つのオブジェクトが関係しているのです。
プログラムとカーネルがその２つですが、今回はプログラムだけを扱います。
ここでいうプログラムとはC#で書くようなプログラムのことではありません。
OpenCLでプログラムと言ったとき、それはGPUで動くプログラムのことです。
OpenCL Cという、C言語を拡張した言語で書いたプログラムです。

ちなみにカーネルとは、プログラム中に記述された関数のことです。
詳しくは次回以降に。

プログラム

OpenCLでGPUを動かすときには、OpenCL Cという独自の言語を使います。
DirectXで言うとHLSLのような言語です。
この言語によって、GPU特有の命令を実行するのです。
OpenCL Cで書いたプログラムはファイルとして持っておき、それを今回C#側から読み出し、コンパイルします。
（そして実行するのです。今回は実行まではいかず、コンパイルまでですが。）。

OpenCLで書いたプログラムは例えば次のようになります。
拡張子は.clです。

__kernel void myKernelFunction(__global float* numbers)
{
    numbers[0] = 3;
    numbers[1] = 4;
    numbers[2] = 5;
}
このプログラムは実はあまりGPUの並列処理っぽくはない、シングルスレッドなプログラムです。
バッファに対して値をセットしています（シングルスレッドで）。
もちろん、書き方によってはGPUの性能をフルに生かした並列処理なプログラムを書くことは可能です（ちょっと複雑になるのでここには書きませんが）。

さてここで見慣れないキーワードが出てきました。
__kernelと__globalです。
これらはC言語にはない、OpenCL C特有のキーワードです。

__kernelは関数が、GPUで実行されるということを意味します（GPUというか、正確にはCPUでも設定次第で実行できるのでデバイスというのが正しいのでしょうが、めんどくさいので以後GPUとします）。
__globalは変数がGPUのメインメモリに確保されるということを意味します。

では変数がGPUのメインメモリに確保されない場合があるのでしょうか？
そうです。
GPUと一口にいっても、変数がどこに確保されるかは全部で次の４つが考えられます。
以下の４つは、基本的に最初のものほど制限が緩く、最後の方に行くほど制限がきつく（しかし高速に）なっていきます。

__globalは変数がGPUのメインメモリにあり、自由に読み書き可能です。
ランダムアクセス可能なのでGPGPUの計算結果はここに格納することになります。
DirectXでいうとUnorderedAccessViewに相当すると言えるでしょうか。

__constantは変数がGPUのメインメモリにある・・・ところまでは__globalと同じなのですが、読み取り専用です。
読み取り専用にして何が嬉しいのかということですが、NVIDIAのGPUではたしかこちらのほうがパフォーマンスが高かったはずです。
読み取り専用なのでキャッシュを有効活用できるわけですね。

__localは、GPUの共有メモリです。
共有メモリとはなんんでしょう？
これが__globalや__constantと何が違うのかというと、これは複数あるということです。（複数あるという意味で、オブジェクト指向言語で言う非staticフィールドに似ています）
GPUには複数の演算ユニットがあり、__globalや__constantはそれら全てで共有されていました。
__localはそれらのうち一部で共有されるのです。

ですから、__localのついた変数は、実際には演算ユニットのグループの数だけ存在することになります。
グループの中で共有されるのです。
計算結果はここにあってもCPUからは読み込めません。（staticなメソッドから非staticなフィールドを直接読み込めないのと同じようなものです）
計算途中のデータをここに格納するといいでしょう。

__privateは、GPUのレジスタです。
これは一つの演算ユニットに特有の変数です。
これはデフォルトです。
なので、変数の前に何も付けなかった場合、その変数は__privateだとみなされます。
普通に関数内でint a;というような宣言をすると、aは__privateとなります。

プログラムの生成

 OpenCL Cで書いたプログラムは、C#中でプログラムオブジェクトになります。
プログラムオブジェクトを作るには、clCreateProgramWithSource関数を使います。
参考

cl_program clCreateProgramWithSource(
    cl_context context,
    cl_uint count,
    const char **strings,
    const size_t *lengths,
    cl_int *errcode_ret
)

contextは生成したプログラムを使うコンテキストです。
countはプログラムのソースコードの数です。stringsはこの個数の長さの配列です。
stringsはプログラムのソースコード（複数可）です。
lengthsはプログラムのソースコードの文字列の長さです。これにNULLを渡せば、stringsに入っているのはNULLで終わる文字列として解釈されます。
errcode_retはこの関数のエラーコードです。

プログラムの破棄

生成したプログラムオブジェクトは最後に破棄しなくてはいけません。
プログラムオブジェクトの破棄には、clReleaseProgram関数を使います。
参考

cl_int clReleaseProgram(cl_program program)

programは破棄するプログラムです。

プログラムのコンパイル

プログラムは、生成したばかりの状態ではまだ使えません。
一度使うGPU向けにコンパイルする必要があります。
コンパイルにはclBuildProgram関数を使います。
参考

cl_int clBuildProgram(
    cl_program program,
    cl_uint num_devices,
    const cl_device_id *device_list,
    const char *options,
    void (CL_CALLBACK *pfn_notify)(cl_program program, void *user_data),
    void *user_data
)

programはビルドするプログラムです。
num_devicesはdevice_listに入っているデバイスの数です。
device_listはprogramと関連付けるデバイスです。NULLでも構いません。NULLにすると、既に関連付けられている全てのデバイスに対してビルドします。
optionsはビルドオプションです。
pfn_notifyは非同期に呼び出すときのコールバック関数です。NULLだと、ビルドが終了してからこの関数は終了します。

pfn_notifyには次の２つの引数があります：
programはビルドしたプログラムです。
user_dataは、clBuildProgramの引数として渡したuser_dataです。

user_dataは、pfn_notifyの引数として渡される、ユーザーのためのデータです。NULLでも構いません。

プログラムを使うコード

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

        Buffer buffer = Buffer.FromCopiedHostMemory(context, new float[] { 1, 2, 3 });

        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;

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

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


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

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


このプログラムは、OpenCLプログラムのソースコードをロードし、コンパイルし、終了します。
残念ながら出力になにか面白い結果を表示したりはしません。





拍手[0回]

PR