はじめに

GPGPUを上手く使うと処理を大幅に高速化できるのはご存知かと思います。MacやRadeonを積んでいるマシンの場合、OpenCLを使うのがメジャーだと思います。使い方はC++でOpenCL(使ってみよう編)などを参考にすれば良いのですが、デバイスの初期化など定型的な長い処理が多いのに気づくと思います。いろいろと面倒ですので、シンプルなライブラリを書いてみました。

少し前に書いたのですが、しばらくコーディングから離れていたため、作った自分でも使い方を忘れてしまいました。。。そこで、備忘録をかねて使い方を書きます。MacBook Pro(13inch)とAMD A10-7870K(Windows10)で動作確認しています。

当初、公開することを考えていなかったので、コードの質等についてはご了承ください。

使い方

ライブラリのソースはGitHubに置いてあります。

サンプルコードと処理の流れは下の通りです。

sample.cpp

#include <iostream>
#include <string>
#include <vector>
#include "opencl_helper.hpp"

int main(int argc, char *argv[])
{
    const int bufSize = 1024;
    try
    {
        OpenCLController controller(CL_DEVICE_TYPE_GPU); //OpenCLで使用するデバイス確保などの準備

        //ゼロコピーバッファの準備(対応していないハードではOpenCLDeviceBufferを使用してください)
        OpenCLUnifiedBuffer<float> buf1(controller, bufSize, CL_MEM_READ_WRITE);
        OpenCLUnifiedBuffer<float> buf2(controller, bufSize, CL_MEM_READ_WRITE);
        OpenCLUnifiedBuffer<float> buf3(controller, bufSize, CL_MEM_READ_WRITE);

        OpenCLProgram program(controller, std::string("kernel.cl")); //カーネルファイルの指定, コンパイル
        OpenCLKernel product(program, std::string("product")); //指定したカーネルの関数をロード

        //ホスト側でバッファを書き換える準備(OpenCLDeviceBufferではデバイス->ホストのデータ転送)
        buf1.MapToHost();
        buf2.MapToHost();

        for (int i = 0; i < bufSize; i++)
        {
            buf1[i] = static_cast<float>(i);
            buf2[i] = static_cast<float>(bufSize - i - 1);
        }

        //バッファの書き換えが終了したことをデバイスに通知(OpenCLDeviceBufferではホスト->デバイスのデータ転送)
        buf1.Unmap();
        buf2.Unmap();


        //カーネルの引数をセット
        product.SetArgument(0, buf1.GetMemObject()); //GetMemObjcet()でバッファのハンドルであるcl_memを渡す
        product.SetArgument(1, buf2.GetMemObject());
        product.SetArgument(2, buf3.GetMemObject());
        product.SetArgument(3, bufSize); //普通の引数ならそのまま渡せばよい

        product.Execute(1, std::vector<size_t>{static_cast<size_t>(bufSize)}); //カーネル実行

        //計算結果の回収
        buf1.MapToHost();
        buf2.MapToHost();
        buf3.MapToHost();

        for (int i = 0; i < bufSize; i++)
        {
            std::cout << buf1[i] << ", " << buf2[i] << ", " << buf3[i] << std::endl;
        }


    }
    catch (OpenCLException oe)
    {
        std::cerr << oe << std::endl;
    }
    return 0;
}

OpenCLのカーネルファイルは以下の通りです、サンプルプログラムのバイナリと同じところにソースを置いてください。OpenCLProgramのインスタンスを作るときにコンパイルされます。

kernel.cl


__kernel
void product(global float *a, global float *b, global float *c, int n)
{
    int i = get_global_id(0);
    if(i < n)
    {
        c[i] = a[i] * b[i];
    }
}

デバイスやカーネルの解放はデストラクタで行います。エラー時には例外を投げるので、OpenCLを使う部分をtry-catchで囲んでください。

一次元データの処理であれば、このサンプルを少しいじるだけでいろいろなことに使えると思います。二次元データは取り扱ったことがないので、よくわかりません。

メモリの取り扱い方

GPGPUでボトルネックになるのがホスト-デバイス間のデータ転送です。iGPUを使っている場合、メインメモリをGPUと共有しているため、データ転送する必要がありません。OpenCLUnifiedBufferでは、CL_MEM_ALLOC_HOST_PTRによってデバイスメモリをホストと同じポインタ空間にアロケートします。おそらくdGPUでは明示的にデータ転送しなければならないと思いますので(未確認)、OpenCLDeviceBufferを使ってください。OpenCLUnifiedBufferと同じ使い方ですが、MapToHost()でデバイスからホストに、Unmap()でホストからデバイスにデータを転送します。

OpenCLControllerのコンストラクタでユニファイドメモリをサポートしているか情報収集しているので、OpenCLUnifiedBufferを使えない環境では例外を投げるようにしています(サポートしていないハードがないので動作未確認)。

やり残していること

namespaceをわける。
OpenCLControllerなどコピーするとNGなものをコピー禁止にする。
エラーコードを網羅していない。
一番目のGPUしか使えないのをどうにかする。
コメントを含めてドキュメント不足。

C++からOpenCLをラクに使うためにライブラリを書いてみた

はじめに

使い方

メモリの取り扱い方

やり残していること