このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。
※なるべく一般的なOpenCLについて記述するよう心がけていますが、試した環境がAMDのGPU/APUのため、一部デバイス・プラットフォーム固有なことが書かれているかもしれません。
OpenCLでは基本的にホストで用意したデータをデバイスで計算してもらい、デバイスで計算した結果をホストが受け取るという形式になっています。
この時、デバイス側で用意する(確保する)メモリ領域のことをバッファーと呼びます。
このような、デバイス・ホスト間の通信ですが、よくあるCPUとGPUの構成だとCPU<->GPU間のPCIeの速度に律速してしまいます。
そのため、速くしたいのであれば、どんなバッファーをどのように使うかが重要になってきます。
現時点でのOpenCL(2.0)のバッファー領域をおおまかに分類すると
に分けることができます。
ホスト・デバイスそれぞれのみアクセスできる領域とは、例えばそれぞれのRAMそのもので、それぞれのシステムが管轄している領域です。
よくある構成だと、
といったところでしょうか。
基本的には、OpenCLのアプリケーションを使う場合は、上記のホストのみorデバイスのみがアクセスできる領域を使うことになります。
そして、デバイスからホスト、あるいはホストからデバイスのメモリを参照したい場合は、clReadBuffer
やclWriteBuffer
を使うことで、ホスト・デバイス間でデータ転送できるようになっています。
しかしここで問題になるのは、このホスト・デバイス間の転送は、素直に転送されるとは限りません。
というのも、DMAエンジンの要件(アライメントや最小転送サイズ)などの関係上、ホスト領域<->(ホスト側)DMA用領域<->デバイス領域となっていることがあるからです(実装依存)。
そうなると、転送時に毎回DMA用領域へのコピーが入ったりして遅くなる原因となってしまいます。
ただ、そのような場合、実はこのDMA用領域は、ホストもデバイスも両方が直接アクセスできる領域となっているはずなので、ホスト側が最初からこの領域を使えば良いということになります。そのように無駄な領域を介さずに直接転送できることを「Zero Copyできる」と呼んだりしていて、Zero Copyにすることこそが速度向上につながることになるというわけです。
しかし一方で、実際に本当にそのような中間領域があるのか、それはどんなものかは先述の通り実装依存なので、細かく正確に知りたい場合は、各プラットフォームやドライバのマニュアルを読むしかありません。
例えばAMDの場合は、『OpenCL Programming Optimization Guide』の1.3あたりに書いてあって、まとめると
が該当します(Host-Visible device memoryもあるのですが、AMDのベンダー拡張機能なので今回は割愛)。
これらを制御するために、OpenCLではバッファーの作成・割り当てに、以下のような方法があります。
clCreateBuffer
して、clEnqueue(Write|Read)Buffer
する方法USE_HOST_PTR
COPY_HOST_PTR
ALLOC_HOST_PTR
+ clEnqueueMapBuffer
今回はこれらをそれぞれ説明していきます。
それぞれどのように使えるのか、サンプルコードを作りました。make && ./svm
で実行できます。
このサンプルは、n個の点が正方形領域に存在している時
というものです。
今回、このサンプルコードを以下の2環境で動作させてみました。
実行した結果から先に言うと、概ね以下の様な結果となりました。
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]
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]
この結果から分かることは
といったところです。
以下、このサンプルを参考にしながら、バッファーの種類をそれぞれを紹介していきます。
サンプルコードでの関数は以下の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で最速だということもあります(あった))。
サンプルコードでは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
で指定したホストのポインタから直接転送できず、間に転送用領域が噛んでいることも多いです。
そのため、このままだと転送速度に律速してしまうことも多いです。
サンプルコードで対応する関数は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
が解放された後だと、アクセスできなくなるかもしれません。
アクセスできなくなるどころか、今回の環境だと、アクセスした後にデバイスから応答がなくなって、再起動しないと復活しないという状況になったりしたので、よく気をつけないと大変なことになります。
サンプルコードで対応する関数は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と組み合わせると、初期値を指定できるようになるので、そのような使い方をするのが良いと思います。
サンプルコードで対応する関数は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の中身をホスト側で明示的にコピーしなくても、メモリプールの中身をそのままデバイス側に渡せるようになりました。
対応するサンプルコードは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をインストールするときなどはきれいさっぱりさせておいたほうが無難そうです。
対応するサンプルコードは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には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なものは存在していませんでした。遠くない未来にはこのような機能が使えるようになるのかもしれません。
double val[2];
を含む構造体は普通に8byteアラインされるのですが、OpenCL Cでdouble2 val;
を含む構造体は16byteアラインされます。なので、今回で言うとListItemのサイズが変わるので、ホスト側でアラインを揃えるorパディングするか、OpenCL C側で__attribute__ ((packed))
をつけて圧縮するか、cl_double2
を使う必要があります。boost::simple_segregated_storage::add_block
を使うと良いです。
keisuke.kimura in Livox Mid-360をROS1/ROS2で動かしてみた
Sorry for the delay in replying. I have done SLAM (FAST_LIO) with Livox MID360, but for various reasons I have not be...
Miya in ウエハースケールエンジン向けSimulated Annealingを複数タイルによる並列化で実装しました
作成されたプロファイラがとても良さそうです :) ぜひ詳細を書いていただきたいです!...
Deivaprakash in Livox Mid-360をROS1/ROS2で動かしてみた
Hey guys myself deiva from India currently i am working in this Livox MID360 and eager to knwo whether you have done the...
岩崎システム設計 岩崎 満 in Alveo U50で10G Ethernetを試してみる
仕事の都合で、検索を行い、御社サイトにたどりつきました。 内容は大変参考になりま...
Prabuddhi Wariyapperuma in Livox Mid-360をROS1/ROS2で動かしてみた
This issue was sorted....