OpenCLメモ (その3) - work sizeの調整

今回は、この前の簡単な配列加算のプログラムの速度を確認し、work sizeの調整で実行速度が変わることを確認してみる。


OpenCLのkernelがどのように実行されているのか確認したい。

まあ、vtuneがあれば悩まなくていいのだが、さすがにあれはお値段がお高すぎる。とりあえずいろいろ調べると、簡単にはIntel Graphics Performance Analyzersを使えばよいようだ。

(なんか昔も調べた気がするので、たぶん調べるの2回目…)

まず、Intel Graphics Monitorを起動する。「Analyze Application」から、起動する実行ファイルのパスを指定する。



正しくプロファイルを取得するには、トレースする時間を適切に指定する必要がある。[Managae Profiles...]から、[Tracing]タブのTracing Durationを調整する。



以下のような順序になるよう調整しないと、うまくトレースが取得できなかった。

1. 起動
2. OpenCLが実行され、完了
3. 起動してから、Tracing Durationの時間が経過し、トレースが終了
4. プログラムが終了

特にはまるのは、3と4が逆転してしまうと、トレースが取得できなかったことで、プログラムのほうにOpenCL実行後、


Sleep(5000);


を入れて調整した。なんかもっといいやり方がありそうだけど…。

さて、取得したトレースはIntel Platform Analyzerで見ることができて、こんな感じ。



キューにタスクが積まれ、転送→転送→Add→転送の順序で実行している様子とその時間が確認できる。



さて、この前の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];
}


このkernelを以下のように実行していた。


//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指定では、1つのwork itemを持つwork groupがwidth * heightだけ作られることになる。しかし、前回見たGPUの構造から、この指定は以下の理由で効率が悪いことがわかる。

・複数のwork itemをまとめて実行するSIMD-32等が効果を発揮できない
同じwork groupのwork itemしかまとめることはできない

・異なるwork groupだとどのsubsliceで実行されるわからない
メモリアクセスがばらばらで非効率

そこで、通常第0次元のwork itemを複数まとめて実行するよう指定する。このとき、一般的に32, 64など2の累乗がよいとされる。また、OpenCL 1.xでは、globalWorkSizeはlocalWorkSizeで割り切れる必要があることに注意して、localWorkSizeの倍数に切り上げる。

ここでは、work groupあたり32のwork itemを実行するよう指定する。


int ceil_int_div(int i, int div) {
return (i + div - 1) / div;
}

//iをdivで切り上げ
int ceil_int(int i, int div) {
return ceil_int_div(i, div) * div;
}

cl_uint ExecuteAddKernel(ocl_args_d_t *ocl, cl_uint width, cl_uint height) {
cl_int err = CL_SUCCESS;

size_t localWorkSize[2] = { 32, 1 };
size_t globalWorkSize[2] = {
ceil_int(width, localWorkSize[0]),
ceil_int(height, localWorkSize[1])
};

(...略...)


切り上げを行うことで、work item数が実際の数を上回ることがあり得る。確保していないメモリ領域にアクセスしないよう、kernelコードを以下のように変更して、widthとheightは引数で与えるようにする。


__kernel void Add(__global int* pA, __global int* pB, __global int* pC, int width, int height) {
const int x = get_global_id(0);
const int y = get_global_id(1);

if (x < width && y < height) {
const int id = y * width + x;
pC[id] = pA[id] + pB[id];
}
}




kenrnel実行の所要時間を確認する。

work groupあたり1つのwork item = 33.11ms




work groupあたり32のwork item = 1.64 ms



というわけで劇的に高速化していることが確認できる。このように、work sizeの調整は重要であるが、これをOpenCL runtimeに自動的に調整させることも可能である。その際は、clEnqueueNDRangeKernel()のlocal_work_sizeにNULLを指定すればよい。



今回の変更はこちら



続き>>



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)の使用


スポンサーサイト

コメントの投稿

非公開コメント

No title

AMDのAPUで実験は行わないんですか?

あちらのほうが(iGPU)性能は高いみたいですしdGPU版RADEONにも流用出来てお得な気がするんですが・・

Re: No title

AMD APUはいま持ってないので…。

ZEN + iGPUが出たら面白そうなんですが。

No title

あらら、あのAPU機は解体されちゃってたんですか・・

AMD APUはコンパクトに作って緩く使うぶんには良いんですが、ゴリゴリ使うにはもう少しCPU性能が欲しいところですね。
プロフィール

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