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

2015年3月25日

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の機能の分類を以下に示します。

名前 説明
Coarse Grain ホストから読み書きする場合、clEnqueueSVMMap, clEnqueueSVMUnmap が必要なSVM。機能的にはOpenCL 1.2のCL_MEM_ALLOC_HOST_PTRと大体同じ
Fine Grain Buffer clSVMAlloc で割り当てたポインタをホスト、デバイス両方から読み書きし、API経由でバッファ単位で順序付けられる。
Fine Grain Buffer with atomic clSVMAlloc で割り当てたポインタをホスト、デバイス両方から読み書きし、カーネル内でバイト単位で順序付けられる。
Fine Grain System malloc したポインタをGPUでも読み書きできる。

下に行くほど使いやすい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, ":"")
    );

対応は、

名前 capability
Coarse Grain CL_DEVICE_SVM_COARSE_GRAIN_BUFFER
Fine Grain Buffer CL_DEVICE_SVM_FINE_GRAIN_BUFFER
Fine Grain Buffer with atomic support CL_DEVICE_SVM_FINE_GRAIN_BUFFER|CL_DEVICE_SVM_ATOMICS
Fine Grain System CL_DEVICE_SVM_FINE_GRAIN_SYSTEM

となります。

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にあります。要約できる文章ではないので各自で確認してください)

どうなんでしょうね…

Tags

About Author

nakamura

Leave a Comment

メールアドレスが公開されることはありません。 が付いている欄は必須項目です

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

Recent Comments

Social Media