OpenCL2.0 の機能を Broadwell で試す (SVM編)

Broadwell、いわゆる 5th Generation Intel(R) Core(TM) Processor がリリースされてから大分経ちましたが、皆さん、Broadwellは使っていますか？

Broadwell の GPU(Gen8) は、OpenCL 2.0をサポートしており、GPUでもOpenCL 2.0の機能を試すことができます。今回は、SVM(“Shared Virtual Machine” の略で Support Vector Machine とは一切関係無いです)を試してみます。

SVM は、GPU プログラミングを簡単にする(可能性のある)物体です。
GPUプログラミングが面倒である原因はいくらかありますが、その原因の一つに、GPUとホストでメモリが分かれているというのがあります。
SVM を使うと、この問題がいくらか緩和されます。

OpenCL では、ハードウェア、ドライバの制限によって、SVM の機能は4段階にかわります。SVMの機能の分類を以下に示します。

下に行くほど使いやすいSVM になります。今のGen8のドライバでは、3番目 の Fine Grain Buffer with atomicがサポートされています。(これがハードウェアからくる制限なのかドライバからくる制限なのかはよくわからないですが)

https://bitbucket.org/fixstars/blog/src/master/opencl-sharedvm/svm.cpp?at=master が、Fine Grain Buffer SVMを使うサンプルです。ホストで作ったリンクリストに含まれる値の合計をGPUで求めます。

(cl.hpp は、まだ OpenCL 2.0 に対応していないので、C++とCが混ざっています)

まず、デバイスのSVMの機能を調べます。CL_DEVICE_SVM_CAPABILITIES で取得できます。

gpus[di].getInfo(CL_DEVICE_SVM_CAPABILITIES, &svm_capability);
printf("svm capability(%s) : %s%s%s%s\n",
       dev_name.c_str(),
       (svm_capability&CL_DEVICE_SVM_COARSE_GRAIN_BUFFER?"coarse grain, ":""),
       (svm_capability&CL_DEVICE_SVM_FINE_GRAIN_BUFFER?"fine grain buffer, ":""),
       (svm_capability&CL_DEVICE_SVM_FINE_GRAIN_SYSTEM?"fine grain system, ":""),
       (svm_capability&CL_DEVICE_SVM_ATOMICS?"atomic, ":"")
    );

対応は、

となります。

clSVMAlloc でSVMを割り当てます

struct List2
    :public List
{
    void *operator new (size_t count) {
        cl_svm_mem_flags alloc_flags = CL_MEM_READ_WRITE|CL_MEM_SVM_ATOMICS|CL_MEM_SVM_FINE_GRAIN_BUFFER|CL_MEM_SVM_ATOMICS;
        void *ptr = clSVMAlloc(gpu_context(),
                               alloc_flags,
                               count * sizeof(List2),
                               0);

        svm_pointers.push_back(ptr);

        return ptr;
    }

    void operator delete (void *ptr) {
        clSVMFree(gpu_context(), ptr);
    }
};

Fine Grain Buffer SVM では、カーネル内で使うメモリを事前に決めておく必要があります。ここではサンプルなので、とりあえず全部svm_buffersに入れています。プログラムの規模が大きくなってくると、この制限は辛い場合があるかもしれません。

一旦メモリを割り当てたあとは、ホストでの読み書きは通常のポインタと同じように扱えます

for (int i=1; i<=10; i++) {
    List2 *l2 = new List2();
    l2->value = i;
    l2->chain = l;

    l = l2;
}

SVM は、cl_mem ではないので、 clSetKernelArgSVMPointer を使ってカーネルの引数に設定します。
カーネル内で使えるSVMは、clSetKernelArgSVMPointer で指定したポインタだけになります。引数に設定していないSVMを使う場合は、別途 clSetKernelExecInfo で設定しておく必要があります(このためにsvm_pointersにSVMを入れました)。

clSetKernelExecInfo(kernel(),
                    CL_KERNEL_EXEC_INFO_SVM_PTRS,
                    svm_pointers.size() * sizeof(void*),
                    svm_pointers.data());

clSetKernelArgSVMPointer(kernel(), 0, ret_val);
clSetKernelArgSVMPointer(kernel(), 1, l);

カーネル側では、通常の __global ポインタと同じように使えます。

__kernel void f(__global int *sum_ret, __global struct List *p) {
    int sum = 0;
    while (p) {
        sum += p->value;
        p = p->chain;
    }
    *sum_ret = sum;
}

以上が Fine Grain Buffer SVM の使いかたでした。試した印象だと、かなり使いづらいのではないかという手ごたえがあります。

• カーネル実行前に使うポインタを列挙しておく必要がある
• ホスト-デバイス間はAPI使わないとメモリ見えない (詳細は3.3.6.4 Memory Ordering Rules: Host-side and Device-side Commandsにあります。要約できる文章ではないので各自で確認してください)

どうなんでしょうね…

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください。