(cache) OpenCL 2.0勉強会#1：Shared Virtual MemoryなどのOpenCLのバッファー関連まとめ ~ /proc/cpuinfo

Fixstarsでは半年ごとに2つの社内勉強会を開催しているのですが、今は、もうすぐ確定するはずのOpenCL 2.1に向けてOpenCL 2.0の勉強会を開催しています。先週の発表者は私で、タイトルに書いたようなことを話したのですが、せっかくなのでここでまとめたいと思います。

※なるべく一般的なOpenCLについて記述するよう心がけていますが、試した環境がAMDのGPU/APUのため、一部デバイス・プラットフォーム固有なことが書かれているかもしれません。

OpenCLのバッファーとは

OpenCLでは基本的にホストで用意したデータをデバイスで計算してもらい、デバイスで計算した結果をホストが受け取るという形式になっています。この時、デバイス側で用意する（確保する）メモリ領域のことをバッファーと呼びます。

このような、デバイス・ホスト間の通信ですが、よくあるCPUとGPUの構成だとCPU<->GPU間のPCIeの速度に律速してしまいます。そのため、速くしたいのであれば、どんなバッファーをどのように使うかが重要になってきます。

概要

現時点でのOpenCL(2.0)のバッファー領域をおおまかに分類すると

ホストのみアクセスできる領域
ホストとデバイスの両方が直接アクセスできる領域
デバイスのみアクセスできる領域

に分けることができます。

ホスト・デバイスそれぞれのみアクセスできる領域とは、例えばそれぞれのRAMそのもので、それぞれのシステムが管轄している領域です。よくある構成だと、

ホスト(CPU)：DDR上にあり、OSの仮想メモリの管轄
デバイス(GPU)：GPUボードの上にあるGDDR上にあり、ドライバやFWの管轄

といったところでしょうか。

基本的には、OpenCLのアプリケーションを使う場合は、上記のホストのみorデバイスのみがアクセスできる領域を使うことになります。そして、デバイスからホスト、あるいはホストからデバイスのメモリを参照したい場合は、clReadBufferやclWriteBufferを使うことで、ホスト・デバイス間でデータ転送できるようになっています。

しかしここで問題になるのは、このホスト・デバイス間の転送は、素直に転送されるとは限りません。というのも、DMAエンジンの要件（アライメントや最小転送サイズ）などの関係上、ホスト領域<->（ホスト側）DMA用領域<->デバイス領域となっていることがあるからです（実装依存）。そうなると、転送時に毎回DMA用領域へのコピーが入ったりして遅くなる原因となってしまいます。

ただ、そのような場合、実はこのDMA用領域は、ホストもデバイスも両方が直接アクセスできる領域となっているはずなので、ホスト側が最初からこの領域を使えば良いということになります。そのように無駄な領域を介さずに直接転送できることを「Zero Copyできる」と呼んだりしていて、Zero Copyにすることこそが速度向上につながることになるというわけです。

しかし一方で、実際に本当にそのような中間領域があるのか、それはどんなものかは先述の通り実装依存なので、細かく正確に知りたい場合は、各プラットフォームやドライバのマニュアルを読むしかありません。例えばAMDの場合は、『OpenCL Programming Optimization Guide』の1.3あたりに書いてあって、まとめると

ホストのみアクセス可能：Unpinned Host Memory
Pinned Host Memory / Device-Visible Host Memory
デバイスのみアクセス可能：Device Memory

が該当します（Host-Visible device memoryもあるのですが、AMDのベンダー拡張機能なので今回は割愛）。

これらを制御するために、OpenCLではバッファーの作成・割り当てに、以下のような方法があります。

なにもしない（フラグなしでclCreateBufferして、clEnqueue(Write|Read)Bufferする方法
USE_HOST_PTR
COPY_HOST_PTR
ALLOC_HOST_PTR + clEnqueueMapBuffer
Coarse Grain SVM
Fine Grain SVM

今回はこれらをそれぞれ説明していきます。

使ってみた結果

それぞれどのように使えるのか、サンプルコードを作りました。make && ./svmで実行できます。このサンプルは、n個の点が正方形領域に存在している時

正方領域を、m x m個のブロックを持つ格子分割して
各ブロック内の中心座標から、そのブロック内に属している点の距離の相乗平均（幾何平均）を計算して
最後にその平均値の最大と最大となるブロックの番号を取得する

というものです。

今回、このサンプルコードを以下の2環境で動作させてみました。

GPU: Hawaii

CPU: Core i7-2600K
OS: Ubuntu 14.04 LTS
AMD-APP SDK: 3.0β (1642.5)

APU: Kaveri

OS: Ubuntu 15.10
AMD-APP SDK: 3.0β (1642.5)

実行した結果から先に言うと、概ね以下の様な結果となりました。

GPU: Hawaii (Radeon R9 290X)

CPU: Core i7-2600K
Host               : 3536[ms]
Host Pool          : 4746[ms]
Host OpenMP        : 1555[ms]
OpenCL Normal      : 1479[ms]
OpenCL UseHostPtr  : 1496[ms]
OpenCL CopyHostPtr : 1544[ms]
OpenCL MapBuffer   : 1547[ms]
OpenCL MapPool     : 1157[ms]
OpenCL SVMCoarse   : 658[ms]
OpenCL SVMFine     : 759[ms]

APU: Kaveri (A10-7850K)

Host               : 10755[ms]
Host Pool          : 8857[ms]
Host OpenMP        : 3082[ms]
OpenCL Normal      : 4773[ms]
OpenCL UseHostPtr  : 4385[ms]
OpenCL CopyHostPtr : 5107[ms]
OpenCL MapBuffer   : 4556[ms]
OpenCL MapPool     : 2651[ms]
OpenCL SVMCoarse   : 1679[ms]
OpenCL SVMFine     : 1702[ms]

この結果から分かることは

まず前提として、元のCPU（リファレンス）がリストを使う構造となっており、素直に移植するとGPU（特にOpenCL 1.2以前）に不向きな方法をとっています。
普通に何も考えずにやると、OpenCLデバイスに投げやすいようにホスト側でデータ変換＆コピーが入ってしまい、だいぶ遅くなります。
OpenCL 1.2までの範囲だと、使用するメモリ領域全体をALLOC_HOST_PTRしたMapPoolで、なんとかCPUより速いぐらいでした。
OpenCL 2.0でSVMが使えると、ポインタをそのまま使うことができるため、余計なコピーなどが入らず、OpenMPで全コア使った時より2倍ぐらいは速くできました。
しかし、そもそもGPUではポインタを使わない実装にしたほうが良く、今回のサンプルもいくらでももっと速くできる方法はあるはずです（もっと最適化してみるとか含めて）。ただしメモリの消費量の関係などで実装はなるべく元のCPUと合わせたままにしたい、という場合もあるかもしれないので、今回はそういう場合だということにしました。

といったところです。

解説

以下、このサンプルを参考にしながら、バッファーの種類をそれぞれを紹介していきます。

ホスト側

サンプルコードでの関数は以下の3つです。

MaxGeoMeanHost
MaxGeoMeanHostPool
MaxGeoMeanHostOmp

1つ目は素直に実装したものです。

しかし、リストのアイテムを追加するのに毎回newしていると遅いため、2つ目のようにメモリプールを使うと速くすることができます。

static boost::object_pool<ListItem<Point>> pool(n);

struct Allocator
{
    static ListItem<Point>* Allocate(const Point& val)
    {
        // return new ListItem<Point>(val);
        return pool.construct(val);
    }
    static void Deallocate(ListItem<Point>* const)
    {
        // delete ptr;
        // 最終的にpoolが消える時に全て消えるので不要
    }
};

void Add(const T val)
{
    ListItem<T>* const t = ALLOCATOR::Allocate(val);
    if (first == nullptr)
    {
        first = t;
        last = first;
    }
    else
    {
        last->Append(*t);
        last = t;
    }
    count++;
}

また、メモリプールを使って割り当てる領域を決めておくと、実は後でデバイスに投げる時も割り当てる領域を固定できるので有利です。本来は、1つ目と同様にstd::listを使って、std::listのallocatorを指定するのが一番正解に近いかもしれないのですが、後で OpenCLに投げる時のことを考えて自前でリストを用意しました（Hawaiiの方でプールしたほうが遅いのは、実は自作したのが遅いせいかもしれません）。

3つ目は、メモリプールを使った版をOpenMPで並列化したものです。

#pragma omp parallel for
for (int i = 0; i < gridCount; i++)
{

グリッドへの分割は簡単には並列化できないので、各グリッド内で幾何平均を求める部分だけ並列化しています。以下のOpenCLデバイスで実行する部分も、同じ場所だけ並列化することにします。

ここでは3つ目がCPUで最速なので、OpenCLにした場合にここからどれだけ速くできるかが指標になります。なお、余談ですが、計算内容などによっては、実はOpenMPよりもOpenCLで書いてデバイスにCPUを指定する方が、CPUで最速だということもあります（あった））。

clEnqueue(Write|Read)Buffer

サンプルコードではMaxGeoMeanOclNormalに対応します。

この方法は、太古の昔から使われている方法です。特に何も指定せずcl::Bufferを作成してデバイス側に領域を確保して、必要そうな箇所でアプリケーション側からclEnqueue(Write|Read)Bufferを叩いてデータ転送します。

// バッファー作成
cl::Buffer bufferDst(context, CL_MEM_WRITE_ONLY, sizeof(cl_double)*gridTotalCount);
cl::Buffer bufferGrid(context, CL_MEM_READ_ONLY, sizeof(Point)*gridTotalCount*gridStride);
cl::Buffer bufferGridCount(context, CL_MEM_READ_ONLY, sizeof(cl_ulong)*gridTotalCount);

// ホスト->デバイス転送
queue.enqueueWriteBuffer(bufferGrid, CL_FALSE, 0, sizeof(Point)*gridTotalCount*gridStride, gridBuffer.get());
queue.enqueueWriteBuffer(bufferGridCount, CL_FALSE, 0, sizeof(cl_ulong)*gridTotalCount, count.get());

// 略

// デバイス->ホスト転送
queue.enqueueReadBuffer(bufferDst, CL_TRUE, 0, sizeof(cl_double)*gridTotalCount, dst.get());

転送が発生する箇所が明確になっているのが利点で、この利点のおかげか、みんなこれを使っていることが多いです。

しかし、先述の通り、実際にはclEnqueue(Write|Read)Bufferで指定したホストのポインタから直接転送できず、間に転送用領域が噛んでいることも多いです。そのため、このままだと転送速度に律速してしまうことも多いです。

USE_HOST_PTR

サンプルコードで対応する関数はMaxGeoMeanOclUseHostPtrです。

CL_MEM_USE_HOST_PTRを指定してcl::Buffer作ると、アプリケーション側でいちいちどこでWrite/Readするか指定しなくても、OpenCLランタイム側でよろしくやってくれる方法です。

cl::Buffer bufferDst(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, sizeof(cl_double)*gridTotalCount, dst.get());

OpenCL 1.2以前では、おそらく使い勝手は一番良いと思います。

ただし、どこでどう転送されるのか分からない（実装依存な）ので、速度を気にしたくなってきた時にはだんだん使いづらくなります。また、間で転送用領域があってコピーされることがあるかもしれない、というのは変わりません。

注意点として、いつどこでどう転送されるか分からないということは、ホスト側のポインタが生きていてもcl::Bufferが解放された後だと、アクセスできなくなるかもしれません。アクセスできなくなるどころか、今回の環境だと、アクセスした後にデバイスから応答がなくなって、再起動しないと復活しないという状況になったりしたので、よく気をつけないと大変なことになります。

COPY_HOST_PTR

サンプルコードで対応する関数はMaxGeoMeanOclCopyHostPtrです。

CL_MEM_COPY_HOST_PTRを指定してcl::Bufferを作ると、作った時に、指定したポインタからデバイス側バッファーにコピーが作成されます。

cl::Buffer bufferGridCount(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_ulong)*gridTotalCount, count.get());

ただ、今回のようにCOPY_HOST_PTRだけ単独では、あまり使っても意味がないし、むしろ遅くなることも多いです。後述のALLOC_HOST_PTRと組み合わせると、初期値を指定できるようになるので、そのような使い方をするのが良いと思います。

ALLOC_HOST_PTR

サンプルコードで対応する関数は2つあります。

MaxGeoMeanOclMapBuffer
MaxGeoMeanOclMapPool

この方法は、CL_MEM_ALLOC_HOST_PTRを指定してcl::Bufferを作ったあと、、そのバッファーに対応するホスト側領域のポインタがclEnqueueMapBufferで取得できるというものです。

cl::Buffer bufferGridCount(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_ulong)*gridTotalCount, count.get());

つまり、確実に転送用領域がない（デバイスから直接転送されてくるZero Copyな）メモリを確保できるものです。取得した後は、そこへ普通のポインタと同じようにRead/Writeすることが可能です。

ただ、ホスト側のキャッシュと同期したりするために、触り終わったらclEnqueueUnmapBufferが必要なので少し面倒ではあります。

queue.enqueueUnmapMemObject(bufferGrid, gridBuffer);

また、ほとんどの環境で、clEnqueueMapBufferで確保できるサイズは限られているので（サイズは実装依存）、なんでもかんでも全部のホスト側バッファーを置き換えればいいというわけでもないです。

サンプルコードの前者と後者の違いは、後者は先に書いたメモリプール（ListItemの確保先）もMapBufferして、デバイス側からZeroCopyアクセスできるようにしている点です。そのおかげで、後者では、ListItemの中身をホスト側で明示的にコピーしなくても、メモリプールの中身をそのままデバイス側に渡せるようになりました。

Coarse Grain SVM

対応するサンプルコードはMaxGeoMeanOclSVMCoarseです。

SVM(Shared Virtual Memory≠Supoort Vector Machine)は、OpenCL 2.0の新機能のひとつです。名前の通り、ホストとデバイスで仮想メモリ空間を共有できるようになります。メモリ空間を共有するので、ポインタをそのまま使うことができるようになります。

SVMの一番基本的な機能はCoarse Grain SVMと呼ばれるもので、サポートされているかどうかは以下のようにデバイス情報から取得できます

cl_device_svm_capabilities svmCap;
device.getInfo(CL_DEVICE_SVM_CAPABILITIES, &svmCap);
const bool coarseSupported = ((svmCap&CL_DEVICE_SVM_COARSE_GRAIN_BUFFER) != 0);

SVMを使う場合、領域の確保はclSVMAlloc、解放はclSVMFreeを使います。

// 確保
ListItem<Point>* const poolSource = static_cast<ListItem<Point>*>(clSVMAlloc(context(), CL_MEM_READ_WRITE, sizeof(ListItem<Point>)*n, 0));

// 解放
clSVMFree(context(), poolSource);

Coarse Grain SVMでは、ホストが触る前にclEnqueueSVMMap、使い終わったらclEnqueueSVMUnmapが必要です。

// ホストからアクセスする前に実行
clEnqueueSVMMap(queue(), CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, poolSource, sizeof(ListItem<Point>)*n, 0, nullptr, nullptr);

// ここではホストからpoolSourceに好きにアクセスできる

// ホストが終わってデバイスからアクセスする前に実行
clEnqueueSVMUnmap(queue(), poolSource, 0, nullptr, nullptr);

// ここからデバイスがアクセスできる

なお、Unmapが完了する前にEnqueueNDRangeKernelなどでカーネルの実行をキューに入れてはいけないので、必ずキューに入れる前に完了を待つ必要があります。

そして、カーネルにSVM領域を引数に渡す場合はclSetKernelArgではなくclSetKernelArgSVMPointerになります。

// kernel.setArg(2, bufferPool);
clSetKernelArgSVMPointer(kernel(), 2, poolSource);

ただし今回のように、引数にはポインタだけ渡して、実体（メモリプール）は引数に渡さない（渡す必要がない）場合は、clSetKernelExecInfoを使うことでカーネル側からそのSVM領域にアクセスすることができるようになります。

clSetKernelExecInfo(kernel(), CL_KERNEL_EXEC_INFO_SVM_PTRS, sizeof(poolSource), &poolSource);

このclSetKernelExecInfoの引数は「clSVMAllocで確保されたメモリ領域へのポインタの配列」で、つまり「clSVMAllocで確保されたメモリ領域の個数」なので注意してください。

void* arrayDouble = clSVMAlloc(context(), CL_MEM_READ_WRITE, sizeof(double)*10, 0);
void* arrayFloat = clSVMAlloc(context(), CL_MEM_READ_WRITE, sizeof(float)*10, 0);
void* arrayInt = clSVMAlloc(context(), CL_MEM_READ_WRITE, sizeof(int)*10, 0);
void* svmPtrs[] = {arrayDouble, arrayFloat, arrayInt};
clSetKernelExecInfo(kernel(), CL_KERNEL_EXEC_INFO_SVM_PTRS, sizeof(void*)*3, svmPtrs);
// clSetKernelExecInfo(kernel(), CL_KERNEL_EXEC_INFO_SVM_PTRS, sizeof(double*)*10, arrayDouble); とかではない

注意点としては、SVMを使う場合、clBuildProgramの時に"-cl-std=CL2.0"オプションを付ける必要があります。今回の環境(AMD-APP SDK 3.0β 1642.5)だと、カーネル実行後にSVM領域を読みに行くと「領域外アクセス」もしくは「nullptr dereference」の例外が起きたのですが、なんでホスト側が領域外アクセスなのか全く意味不明なので、気づくのにかなり時間をつかいました・・・。

また、AMD-APP SDK 2.9やyum install clinfoなどでlibOpenCL.soなどが2.9と3.0混在している環境を作ってしまうと、clBuildProgramが「internal link error」と言われたりするので、最初にSDKをインストールするときなどはきれいさっぱりさせておいたほうが無難そうです。

Fine Grain SVM

対応するサンプルコードはMaxGeoMeanOclSVMFineです。

Fine Grain SVMがサポートされているかどうかは以下のようにデバイス情報から取得できます

cl_device_svm_capabilities svmCap;
device.getInfo(CL_DEVICE_SVM_CAPABILITIES, &svmCap);
const bool fineSupported = ((svmCap&CL_DEVICE_SVM_FINE_GRAIN_BUFFER) != 0);

Fine Grainになると、clEnqueueSVMMap/clEnqueueSVMUnmapすら不要になって、確保したら後は好きにホストからでもデバイスからでもアクセスできるようになります。

// 確保
ListItem<Point>* const poolSource = static_cast<ListItem<Point>*>(clSVMAlloc(context(), CL_MEM_READ_WRITE, sizeof(ListItem<Point>)*n, 0));

// 不要：clEnqueueSVMMap(queue(), CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, poolSource, sizeof(ListItem<Point>)*n, 0, nullptr, nullptr);
// 不要　clEnqueueSVMUnmap(queue(), poolSource, 0, nullptr, nullptr);

// いつでもホストからでもデバイスからでもアクセスできる

// 解放
clSVMFree(context(), poolSource);

// ここではもちろんアクセスできない

ただし、ホスト側の(CPU)キャッシュと整合性を取らなければならなくなるからか、少し遅くなります。

その他のSVM

今回紹介しませんでしたが、SVMにはCoarseとFine以外もあります。

例えば、先ほどFineの場合はホストとデバイスが好きにアクセスできると述べましたが、そうなると「ホストとデバイス同時にアクセスした場合どうなるの」という疑問が湧きます。そのあたりについてはCL_DEVICE_SVM_ATOMICSやCL_MEM_SVM_ATOMICSという機能が該当します。が、同時にアクセスしたい場合が思いつかなかったので今回は割愛しました。

また、clSVMAlloc/clSVMFreeではなくて。普通のnew/deleteやmalloc/freeしたものに（ホストはもちろん）デバイスからもアクセスできるようになるFine Grain System（もはやSVMではない）という最上位の機能も存在します。ただし本稿執筆時点で主要なGPU/APUでFine Grain Systemなものは存在していませんでした。遠くない未来にはこのような機能が使えるようになるのかもしれません。

まとめ

一般論としては、基本的にALLOC_HOST_PTRを使うのが最速なはずです。
ただし今回のようにポインタを使いたい（使ったほうが速い）場合は、SVMを使うと更に速くかつ簡潔なコードが書けます。
SVMの使い分けは、ホスト側からアクセスする領域が明確なのであればCoarse Grain SVM（のほうが速そう）が良いです。その辺りの制御が難しいならFine Grain SVMを使いましょう。

小ネタ

gccのdouble val[2];を含む構造体は普通に8byteアラインされるのですが、OpenCL Cでdouble2 val;を含む構造体は16byteアラインされます。なので、今回で言うとListItemのサイズが変わるので、ホスト側でアラインを揃えるorパディングするか、OpenCL C側で__attribute__ ((packed))をつけて圧縮するか、cl_double2を使う必要があります。
OpenCLとはちょっと違う話ですが、今回のようにあらかじめ確保しておいた領域をメモリプールとして使いたい場合は、boost::simple_segregated_storage::add_blockを使うと良いです。

/proc/cpuinfo

2015年6月18日木曜日

OpenCL 2.0勉強会#1：Shared Virtual MemoryなどのOpenCLのバッファー関連まとめ

OpenCLのバッファーとは

概要

使ってみた結果

解説

ホスト側

clEnqueue(Write|Read)Buffer

USE_HOST_PTR

COPY_HOST_PTR

ALLOC_HOST_PTR

Coarse Grain SVM

Fine Grain SVM

その他のSVM

まとめ

小ネタ

0 コメント:

コメントを投稿

About us

Popular Posts

Contributors

Blog Archive

Subscribe Via Email