OpenCLメモ (その7) [終] - reductionとshared local memory(SLM)の使用

最後に、配列の総和計算でshared local memory(SLM)の使用方法を確認する。


配列の総和の計算は、CPUだと非常に単純で、

float sum = 0.0f;
for (int i = 0; i < nSize; i++) {
sum += ptrArray[i];
}

とまあ、これだけなのだが、このままのループの形だと並列性がないので、GPUでやろうとするとやけに面倒だったりする。

こういったいわゆるreductionの計算は、なんとか並列性をひねりだし、work group内で共有できるメモリ(shared local memory, SLM)を使うと効率的に計算できることが知られている。

計算の仕方としては、もう図にしたほうが早いので、下のような感じ。(図はwork groupあたり16 work itemの例)
CUDAでもお馴染みの方法だったりする(最近はshuffleとかも使うらしいが)

まあ、ひとつづつ足すのではなく、なるべく並列に計算できるところを見つけて、GPUを活かそう、というもの。無限に演算器があるなら、計算のオーダーをO(n)からO(log(n))に落とすイメージ。

通常、GPUでは異なるスレッドの計算結果を得ることはできないが、__localメモリを使えば、work group内の計算結果を取得できることを利用している。



このようにkernelを1回起動することで、work group内のwork item分の和を取ることができる。つまり、work groupあたり16work itemの設定なら、1回で要素数を1/16にすることができ、これを繰り返すことで総和を得ることができる。もちろん、work groupあたりのwork item数が大きいほど一気に要素数を減らせるので、実際にはwork groupあたり16work itemなどにする。

実際のkernelコードはこんな感じ。


__kernel void reduce_add(__global float* pIn, __global float* pOut, int nSize) {
const int lid = get_local_id(0);
const int gid = get_global_id(0);
__local float shared[GROUP_SIZE];

shared[lid] = (gid < nSize) ? pIn[gid] : 0;
barrier(CLK_LOCAL_MEM_FENCE);
for (int offset = get_local_size(0) >> 1;
offset > 0; offset >>= 1) {
if (lid < offset) {
shared[lid] += shared[lid + offset];
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (lid == 0) {
pOut[get_group_id(0)] = shared[0];
}
}


shared local memoryは、kernelコードでは__localで取ることができる。

__localに書き込んだ値を他のスレッドで使う場合、かならず

barrier(CLK_LOCAL_MEM_FENCE);

によりwork group内のスレッドの同期をとる必要がある。work group内のスレッドはすべてが同時に実行されるわけではないので、きちんとすべてのスレッドが書き込みを終えてから次に進むようにしないと、おかしな値を使用してしまう可能性があるからである。

実際に実行している様子はこちら。



よく見えないけど、reduction_addが2回実行されていて4M要素の総和を取っている。(4M → 16K → 64)



この方法、GPUでもある程度高速に総和を取れてよいのだけど、加算の計算順序が変わるので、計算結果がCPUと多少変わるのが面倒なところ。浮動小数点演算は数学の計算と違って可換ではないのだ…。

コードはこちら



OpenCLは勉強自体はこれまで少しづつやっていたことだが、時間がなかったのでこれまでまとめたりできなかった。この3連休でやっとまとめることができたので、これできっともう忘れないだろう(棒)

だいたい通常の計算と、image、__localの使い方がわかったし、これでやっとOpenCLの基本の基本ぐらいはわかったように思う。おおまかにCUDAとの対応関係もつかめたし、多少はCUDAの知識も使えそうな感じ。

あとは、実際にやりたいこと次第、になるのだと思う。



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