このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。
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 の使いかたでした。試した印象だと、かなり使いづらいのではないかという手ごたえがあります。
どうなんでしょうね…
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....