OpenCLメモ (その1) - かんたんな計算

最近のCPUにはたいていGPUがついてきていて、

skylake_die_gpus.png
The Compute Architecture of Intel Processor Graphics Gen9より

まあ、大雑把にだいたい30~40%ぐらいGPUだったりする。このGPUを有効利用するには、OpenCLを使って、GPUで計算するのが良いとされるので、年末も近づき、暇になってきたので少し試してみることにした。

OpenCLは以前にもすこしだけやったことがあるが、すぐに忘れてしまうのでメモ書きを残すことにする。

まずはVisual Studio 2015 Communityをインストールし、さらにこのへんからダウンロードしたIntel OpenCL SDKをインストールする。

インストールすると、VC++のプロジェクトテンプレートの中にOpenCLのテンプレートが登場する。

opencl_project_template_01.png

Code Builder Project... を選ぶと、設定が出る。

ここでは、"Use explicit Local Work Group Size"にチェックを入れた。

opencl_project_template_02.png

テンプレートは配列加算のサンプルだが、その入出力の配列のメモリに関する設定。デフォルトは"Use Host Memory"になっているが、これはZero Copyを使って、CPU-GPU間転送を省略するものであり、OpenCLの基本とはいえない。そこで、"Allocate & Copy Memory"に変更した。

opencl_project_template_03.png

Finishとすると、OpenCLのサンプルコードがもう動かす事ができる形で自動生成される。

opencl_project_template_04.png

OpenCLといえば、まず動かすまでに長すぎる呪文詠唱で嫌になるという印象が強かったが、最近では便利になったものだ。



とはいえ、どういう呪文を詠唱すればよいのか確認していく。見通しをよくするため、自動生成されたサンプルコードからエラー処理の一切を省き、ある程度簡略化してある。

まず"platform"の取得と選択 FindOpenCLPlatform()

//今回はIntelのGPUを探したいので、
//deviceType = CL_DEVICE_TYPE_GPU
//preferredPlatform = "Intel"
cl_platform_id FindOpenCLPlatform(
const char* preferredPlatform, cl_device_type deviceType) {
//platformの数を取得
cl_uint numPlatforms = 0;
clGetPlatformIDs(0, NULL, &numPlatforms);

//実際にplatformのリストを取得
std::vector platforms(numPlatforms);
clGetPlatformIDs(numPlatforms, &platforms[0], NULL);

for (cl_uint i = 0; i < numPlatforms; i++) {
//platformの名前を取得する
char buf[4096];
clGetPlatformInfo(platform, CL_PLATFORM_NAME,
sizeof(buf), buf, NULL);

//ここではIntelのplatformが見つかったら
if (strstr(buf, preferredPlatform) != 0) {
//そのplatformがdeviceType(ここではGPU)を持つか、確認する
cl_int err = clGetDeviceIDs(
platforms[i], deviceType, 0, NULL, &numDevices);
if (err == CL_SUCCESS && numDevices > 0) {
return platforms[i];
}
}
}
return NULL;
}


つぎに取得したplatformを使って、opencl "context"を作成し、その"device"を取得する。さらに、そのdeviceに命令を投入するためのキューを作成する。


//先程の関数でplatformを作成
cl_platform_id platformId
= FindOpenCLPlatform("Intel", deviceType);

//コンテキスト情報
cl_context_properties contextProperties[] =
{ CL_CONTEXT_PLATFORM, (cl_context_properties)platformId, 0 };

//コンテキストを作成
ocl->context = clCreateContextFromType(
contextProperties, deviceType, NULL, NULL, &err);

//コンテキストのdeviceを取得
err = clGetContextInfo(ocl->context, CL_CONTEXT_DEVICES,
sizeof(cl_device_id), &ocl->device, NULL);

//プロファイリングを有効にしてキューを作成
cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
ocl->commandQueue = clCreateCommandQueue(
ocl->context, ocl->device, properties, &err);


次に、kernel(GPUで実行する関数)をビルドして実行可能にする必要がある。通常、kernelはソースコードを文字列として、clCreateProgramWithSource()関数に渡してビルドする。

Intelのサンプルコードのように、ソースファイルからそのまま文字列を読み込んでもいいが、面倒なので、リソースとして実行ファイルに埋め込むことを考える。

OpenCLTest.rcを新規作成し、ここに


CLDATA KERNEL_DATA DISCARDABLE "Template.cl"


と書いてプロジェクトに登録し、このリソースを対象としてビルドを行い、関数名を指定してkernelを取得する。リソースを対象とする場合、#includeなどは使えないという制約が出るが、そこは…気にしないことにする。


//リソースからOpenCLのカーネルファイルを取得
HMODULE hmdl = GetModuleHandle(NULL);
HRSRC hrsrc = nullptr;
HGLOBAL hdata = nullptr;
const char *source = nullptr;
size_t src_size = 0;
if ( nullptr == (hrsrc = FindResource(hmdl, "CLDATA", "KERNEL_DATA"))
|| nullptr == (hdata = LoadResource(hmdl, hrsrc))
|| nullptr == (source = (const char *)LockResource(hdata))
|| 0 == (src_size = SizeofResource(hmdl, hrsrc))) {
return CL_INVALID_VALUE;
}

//ソースから"program"を作成
ocl->program = clCreateProgramWithSource(ocl->context,
1, (const char**)&source, &src_size, &err);

//ビルドを実行
clBuildProgram(ocl->program, 1, &ocl->device, "", NULL, NULL);

//ビルドした関数のうち、"Add"関数をkernelとして取得
ocl->kernel = clCreateKernel(ocl->program, "Add", &err);


さて、配列加算 C = A + B を行うが、そのサイズを width * height とする。この時、それぞれの配列について、CPU側の配列とは別に、GPU側にバッファを作成してやる必要がある。

ここで、自動生成されたsampleプログラムは、clEnqueueMapBufferを使った効率的な方法を用いていたが、ここではより基本的な、明示的にバッファへの転送を行うclEnqueueWriteBuffer / clEnqueueReadBuffer を用いた方法を見ていく。


int CreateBufferArguments(
ocl_args_d_t *ocl,
cl_int* inputA, //CPU側の入力配列へのポインタ
cl_int* inputB, //CPU側の入力配列へのポインタ
cl_int* outputC, //CPU側の出力配列へのポインタ
cl_uint width, cl_uint height //配列のサイズ
) {
cl_int err = CL_SUCCESS;
ocl->srcA = clCreateBuffer(ocl->context,
CL_MEM_READ_ONLY, //GPUからはこのメモリに対し、読み込みのみ行う
sizeof(cl_uint) * width * height, //配列のサイズ
nullptr,
&err); //エラー情報を受け取る

ocl->srcB = clCreateBuffer(ocl->context,
CL_MEM_READ_ONLY, //GPUからはこのメモリに対し、読み込みのみ行う
sizeof(cl_uint) * width * height, //配列のサイズ
nullptr,
&err); //エラー情報を受け取る

ocl->dstMem = clCreateBuffer(ocl->context,
CL_MEM_ALLOC_HOST_PTR, ////GPUからはこのメモリに対し、書き込みのみ行う
sizeof(cl_uint) * width * height, //配列のサイズ
nullptr, //CPUから転送するものはない
&err);

//inputAから作成したバッファに転送
err = clEnqueueWriteBuffer(ocl->commandQueue, //
ocl->srcA, //転送先
CL_FALSE, //転送が終了するまで待機するか -> しない
0, //オフセット
sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
inputA, //転送元
0, //この関数が待機すべきeventの数
nullptr, //この関数が待機すべき関数のリストへのポインタ
nullptr); //この関数の返すevent

//inputBから作成したバッファに転送
err = clEnqueueWriteBuffer(ocl->commandQueue, //
ocl->srcB, //転送先
CL_FALSE, //転送が終了するまで待機するか -> しない
0, //オフセット
sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
inputB, //転送元
0, //この関数が待機すべきeventの数
nullptr, //この関数が待機すべき関数のリストへのポインタ
nullptr); //この関数の返すevent

return err;
}


ここからいよいよ関数を実行…の前に、関数の引数を一つずつ設定する必要がある。引数がたくさんあると、まあ、簡単に悲惨なコードが出来上がるわけで、正直、このあたりがOpenCLの、CUDAと比べてなかなかに面倒なところだ。

引数を積み終わった後、clEnqueueNDRangeKernelにより、指定のカーネルをGPUで実行するよう指示する。


//関数引数を設定
clSetKernelArg(ocl->kernel, 0, sizeof(cl_mem), (void *)&ocl->srcA);
clSetKernelArg(ocl->kernel, 1, sizeof(cl_mem), (void *)&ocl->srcB);
clSetKernelArg(ocl->kernel, 2, sizeof(cl_mem), (void *)&ocl->dstMem);

//work sizeの指定
//ここでは1要素に対して1 work item
//またグループあたり1 work item (実は効率的でない)
//width * heightの2次元でwork itemを作成
size_t globalWorkSize[2] = { width, height };
size_t localWorkSize[2] = { 1, 1 };

//タスクをキューに積む
clEnqueueNDRangeKernel(
ocl->commandQueue, //タスクを投入するキュー
ocl->kernel, //実行するカーネル
2, //work sizeの次元
NULL, //NULLを指定すること
globalWorkSize, //全スレッド数
localWorkSize, //1グループのスレッド数
0, //この関数が待機すべきeventの数
NULL, //この関数が待機すべき関数のリストへのポインタ
NULL); //この関数の返すevent


work size(work itemの数)で、並列処理する数を指定し、3次元まで設定できる。

work itemは階層構造になっており、通常は複数のwork itemがwork groupを形成し、work groupが集まってwork item全体を構成する。OpenCL 1.2までは各次元のglobalWorkSizeは各次元のlocalWorkSizeで割りきれる必要があることに注意する。

この時、work group内の個々のwork itemのIDがlocal idで、groupによらない一意に決まるIDをglobal id という。下記の図はwork itemが1次元の場合である。

opencl_work_group.png

実際には、work itemは3次元まで設定可能になっていて、サンプルコードは2次元の例になっている。

同じwork groupのwork itemは、GPUで実行される際に必ず同じ演算ブロック(Intelでいうsubslice、NVIDIAでいうSM)で実行され、__localメモリ(IntelでいうShared Local Memory、NVIDIAでいうsharedメモリ)を共有する。また、連続する8, 16, または32のwork itemは同時に実行されるという特徴があり、こうした特徴はメモリアクセス効率を考える際に重要となる。

GPU上で実行するkernelコードはこのようになっていて、単なる配列の加算である。


__kernel void Add(__global int* pA, __global int* pB, __global int* pC) {
const int x = get_global_id(0); //0次元目についてglobal idを取得
const int y = get_global_id(1); //1次元目についてglobal idを取得
const int width = get_global_size(0);

const int id = y * width + x;

pC[id] = pA[id] + pB[id];
}


width * height 分の各work itemがkernelを実行することになるから、2次元のwork itemの指定により、CPUでいうと以下のような2重ループを並列化したことに等しい。


for (int y = 0; y < height; y++) { //<- 1次元目の並列化
for (int x = 0; x < width; x++) { //<- 0次元目の並列化
const int id = y * width + x;
pC[id] = pA[id] + pB[id];
}
}


さて、clEnqueueNDRangeKernelにより、指定のカーネルをGPUで実行するよう指示したが、これはタスクをキューに積んだだけなので、GPUが実際にこのkernelを実行し、完了するには時間がかかる。キューのすべてのタスクが終了するまで待機するには、キューを指定してclFinish()を呼ぶ。


clFinish(ocl->commandQueue);


その後、GPUの計算結果を回収する。出力情報のあるGPUバッファをclEnqueueMapBuffer()によりmappingし、mappingされたCPU側の配列へのポインタが得られる。内部でGPUメモリ→CPUメモリへの転送が行われる。


//inputBから作成したバッフからoutputCに転送
err = clEnqueueReadBuffer(ocl->commandQueue, //
ocl->dstMem, //転送元
CL_FALSE, //転送が終了するまで待機するか -> しない
0, //オフセット
sizeof(cl_uint) * arrayWidth * arrayHeight, //転送サイズ
outputC, //転送先
0, //この関数が待機すべきeventの数
nullptr, //この関数が待機すべき関数のリストへのポインタ
nullptr); //この関数の返すevent

//clEnqueueMapBufferの終了を待機
clFinish(ocl->commandQueue);

//outputCに計算結果が入っていて、CPUから操作できる




というわけで、ここまででやっと配列の加算を行うことができた。

まず、deviceを開き、OpenCL環境を初期化するところまでをまとめると、こんな感じ。



そして、OpenCLバッファの確保から、出力の取り出しまでをまとめると以下のようになる。



実際にはエラー処理などもあるので、結構長いコードになってしまう。CUDAと比べると、とにかく呪文(手続き)が長い、というのが汎用性・標準化を重視したOpenCLのつらいところだと思う。



今回のコードはこのへん



続き>>



OpenCLメモリスト
OpenCLメモ (その1) - かんたんな計算
OpenCLメモ (その2) - Intel GPUの構造
OpenCLメモ (その3) - work sizeの調整
OpenCLメモ (その4) - 転送(コピー)の排除 (USE_HOST_PTR)
OpenCLメモ (その5) - 転送(コピー)の排除 (SVM : OpenCL 2.0)
OpenCLメモ (その6) - imageの使用
OpenCLメモ (その7) [終] - reductionとshared local memory(SLM)の使用


スポンサーサイト

コメントの投稿

非公開コメント

プロフィール

rigaya

Author:rigaya
アニメとか見たり、エンコードしたり。
連絡先(@を半角にしてください!)
rigaya34589@live.jp
github

最新記事
最新コメント
カテゴリ
月別アーカイブ
カウンター
検索フォーム
いろいろ
公開中のAviutlプラグインとかのダウンロード

○Aviutlプラグイン
x264guiEx 2.xx (ミラー)
- x264を使用したH264出力
- x264guiExの導入>
- x264.exeはこちら>

x265guiEx (ミラー)
- x265を使用したH.265/HEVC出力
- x265.exeはこちら>

QSVEnc + QSVEncC (ミラー)
- QuickSyncVideoによるH264出力
- QSVEncCはコマンドライン版
- QSVEncC 導入/使用方法>
- QSVEncCオプション一覧>

NVEnc + NVEncC (ミラー)
- NVIDIAのNVEncによるH264出力
- NVEncCオプション一覧>

VCEEnc + VCEEncC (ミラー)
- AMDのVCEによるH.264出力

ffmpegOut (ミラー)
- ffmpeg/avconvを使用した出力

自動フィールドシフト (ミラー)
- SSE2~AVX2による高速化版
- オリジナル: aji様

エッジレベル調整MT (ミラー)
- エッジレベル調整の並列化/高速化
- SSE2~AVX対応
- オリジナル: まじぽか太郎様

バンディング低減MT (ミラー)
- SSE2~AVX2による高速化版
- オリジナル: まじぽか太郎様

PMD_MT (ミラー)
- SSE2~FMA3による高速化版
- オリジナル: スレ48≫989氏

透過性ロゴ (ミラー)
- SSE2~FMA3によるSIMD版
- オリジナル: MakKi氏

AviutlColor (ミラー)
- BT.2020nc向け色変換プラグイン
- BT.709/BT.601向けも同梱

○その他
x264afs (ミラー)
- x264のafs対応版

aui_indexer (ミラー使い方>)
- lsmashinput.aui/m2v.auiの
 インデックス事前・一括生成

auc_export (ミラー使い方>)
- Aviutl Controlの
 エクスポートプラグイン版
 エクスポートをコマンドから

aup_reseter (ミラー)
- aupプロジェクトファイルの
 終了フラグを一括リセット

CheckBitrate (ミラー, 使い方, ソース)
- ビットレート分布の分析(HEVC対応)

チャプター変換 (ミラー使い方>)
- nero/appleチャプター形式変換

エッジレベル調整 (avisynth)
- Avisynth用エッジレベル調整

メモリ・キャッシュ速度測定
- スレッド数を変えて測定

○ビルドしたものとか
L-SMASH (ミラー)
x264 (ミラー)
x265 (ミラー)

○その他
サンプル動画
その他

○読みもの (ミラー)
Aviutl/x264guiExの色変換
動画関連ダウンロードリンク集
簡易インストーラの概要

○更新停止・公開終了
改造版x264gui
x264guiEx 0.xx
RSSリンクの表示
リンク
QRコード
QR