rigayaの日記兼メモ帳 2016年12月

C91 3日目

C91参加された方、お疲れ様でした。

冬コミ3日目も終わると1年が終わった感じがしますね(病気)

朝はこんな感じで寒かったけど、昼は暖かい感じでよかった。

このページのトップへ

2016年を振り返る (エンコ編)

今年のエンコード周りとか、自分の作ったものとかを振り返る。

このページのトップへ

C91 2日目

今日も参加された方はお疲れ様でした。

今日は日が出て暖かい…と思いきや、ちょっと風が強めで寒かったかも。

今日はあまり初手でならばなきゃ…みたいなとこがなかったので、のんびり行ったら東駐車場に9:45だった。入れたのは10:50ぐらいかなあ。

このページのトップへ

この前は、自動フィールドシフトではいろいろやってだいぶ速くなったのだけど、これ以上は速くするのが難しくなってきているという話だった。

ただ、実はまだ方法が残っているといえば残っていて、実際にコメントを頂いたこともあった。そこで、もう少しやってみたというのが今回。やったのは2つで、

・YUY2フィルタモードへの対応して高速化。
・通常(YC48)モードで高速な簡易解析モードを追加。

どういうことなのかをちょっと書いていく。

2016-12-29(21:14) :
afs :
コメント 8 :
トラックバック 0

このページのトップへ

C91 1日目

C91 1日目参加された方、お疲れ様でした。

朝はかなり寒かったけど、昼間は多少日もさしてよかったかも。

このページのトップへ

自動フィールドシフト高速化 (振り返り編)

自動フィールドシフトはaji様によるとても素晴らしいAviutlプラグインである。なんならrigayaがAviutlを使い始めたのは自動フィールドシフトがあったからだったりする。

ただ残念ながら唯一の欠点として、処理速度が遅いという問題があったので、これをいろんな方法で高速化してきた。まあ試行錯誤とかでだいぶ時間も突っ込んできたし、自動フィールドシフト高速化は(おそらくあまり期待してる人はいないと思うけど)結構いろいろやってみているもののひとつ。

今回は、これを振り返ってみる。

続きを閉じる

自動フィールドシフトの処理の概要

前も書いた気がするけど、画像処理の高速化では、演算をSIMD化してしまうと、演算は十分速くなってしまい、演算そのものよりもその処理でどのくらいのメモリアクセスがあって、どういったアクセスパターンになっているかを把握することが重要になってくる。

自動フィールドシフトも重い処理はほぼSIMD化されているので、ここでは、主にどういう重い処理があって、その処理が読むデータと書き出すデータはどうなっているかを確認すると、まあおおまかにはこんな感じ。

get_ycp_cahce
入力プラグインがデコードしたフレームを受け取り、フレームバッファにコピーする。基本的には、フレームのデコードとメモリコピーにかかる時間。

scan_frame
最も重い処理。スレッド並列化されていて、設定の「スレッド数」はここのスレッド数。フレームバッファの2枚のフレームを参照して縞や動きがあるかを解析し、解析データとして書き出す。

count_motion
解析データから、動きありとして判定された画素の数をフィールドごとに数える。解除パターン決定に使われる。

analyze_frame
2枚の解析データから、解析マップを作成する。

analyze_map_filter
解析マップにぼかしフィルタをかけ、細かなノイズによると思われる判定ノイズを取り除く。

count_stripe
解析マップから、縞ありとして判定された画素の数を数える。解除パターン決定に使われる。

blend
フレームバッファと解析マップから、インタレ解除されたフレームを再構成する。

というわけで、まあいろいろな処理をやるので、当然遅いのだ…。

計算時間の確認

計測環境

OS	Win10 x64	Win10 x64
CPU	i7 4770K (4C/8T)	i7 6700K (4C/8T)
CPU世代	Haswell	Skylake
Core	4.0GHz	4.3GHz
UnCore	4.0GHz	4.1GHz
キャッシュ	L3=8MB	L3=8MB
メモリ	DDR3-1600, 2ch	DDR4-2933, 2ch
CL	8-8-8-24-2	16-18-18-34-2

その他計測環境・計測条件
Aviutl 1.00
入力プラグイン: L-SMASH Works r917 (POP氏ビルド)
入力: MPEG2 1920x1080 29.97fps 10240フレーム
一発勝負

自動フィールドシフトのオリジナル( = r0 )と、これまでのいくつかの高速化版で、どのくらい速くなったのか測定した。

i7 4770K 測定結果

i7 6700K 測定結果

まあ、最初は劇的に、その後は少しずつ速くなっていっているのが分かる。オリジナル(=r0)からくらべると、まあ、だいたい3倍ぐらい速くなっているのが分かる。

それぞれどういう高速化をしたか振り返ると、

r0 (オリジナル) → r5
SSE化とSIMD化されていない箇所のSIMD化
・もともとのMMX命令(64bit幅)からSSE命令(128bit幅)に更新。
もともとはasmだったのだけど、ちょっとアセンブラを書く知識はなかったので、intrinsicで書いてみた。それなりに性能が出せた。

・count_motionなどSIMD化されていなかった処理をSIMD化。
count_motionのSIMD化は劇的な効果があった(15倍以上)。

まあ、これを


  for(pos_y = top; pos_y < scan_h - bottom - ((scan_h - top - bottom) & 1); pos_y++){
    sip = sp->map + pos_y * si_w + left;
    if(is_latter_field(pos_y, sp->tb_order)){
      for(pos_x = left; pos_x < scan_w - right; pos_x++){
        lf_motion += ~*sip & 0x40;
        sip++;
      }
    }else{
      for(pos_x = left; pos_x < scan_w - right; pos_x++){
        ff_motion += ~*sip & 0x40;
        sip++;
      }
    }
  }

こうする古典的なSIMD化である。


#if USE_POPCNT
#define popcnt32(x) _mm_popcnt_u32(x)
#else
#define popcnt32(x) popcnt32_c(x)
#endif

const __m128i xMotion = _mm_set1_epi8(0x40);
for (int pos_y = sp->clip.top; pos_y < y_fin; pos_y++) {
    BYTE *sip = sp->map + pos_y * si_w + sp->clip.left;
    const int is_latter_feild = is_latter_field(pos_y, sp->tb_order);
    const int x_count = scan_w - sp->clip.right - sp->clip.left;
    BYTE *sip_fin = sip + (x_count & ~31);
    for ( ; sip < sip_fin; sip += 32) {
        x0 = _mm_loadu_si128((__m128i*)(sip +  0));
        x1 = _mm_loadu_si128((__m128i*)(sip + 16));
        x0 = _mm_andnot_si128(x0, xMotion);
        x1 = _mm_andnot_si128(x1, xMotion);
        x0 = _mm_cmpeq_epi8(x0, xMotion);
        x1 = _mm_cmpeq_epi8(x1, xMotion);
        DWORD count0 = _mm_movemask_epi8(x0);
        DWORD count1 = _mm_movemask_epi8(x1);
        motion_count[is_latter_feild] += popcnt32(((count1 << 16) | count0));
    }
    if (x_count & 16) {
        x0 = _mm_loadu_si128((__m128i*)sip);
        x0 = _mm_andnot_si128(x0, xMotion);
        x0 = _mm_cmpeq_epi8(x0, xMotion);
        DWORD count0 = _mm_movemask_epi8(x0);
        motion_count[is_latter_feild] += popcnt32(count0);
        sip += 16;
    }
    sip_fin = sip + (x_count & 15);
    for ( ; sip < sip_fin; sip++) {
        motion_count[is_latter_feild] += ((~*sip & 0x40) >> 6);
    }
}

いまとなっては、別にわざわざmovemask + popcntをしなくても途中まではSIMDレジスタで和を取っていってもいい気がするが…実はr16あたりでscan_frameに吸収されたので使われていないコードだったりする。

r5 → r8
メモリアクセス最適化
やはりメモリアクセスの最適化は重要で、かなり高速化した。
・scan_frameで分かれていた関数を統合してメモリアクセスを低減。
・scan_frameなどのフレームを縦にスキャンする部分で、メモリに対してとびとびにアクセスしたのを、なるべく連続アクセスができるように変更。

scan_frameでフレームを縦にスキャンするというのは、まあ図にするとこんな感じで

上下のラインについて計算を行い、加えて一時変数がさらに下のラインの計算結果に影響するので、縦方向に依存関係が発生する。つまり最内ループが縦方向になっている。縦方向にデータを読むと、メモリ上では飛び飛びの位置を読むことになり、メモリアクセスがかなり遅くなる。オリジナル版では上の図のようにMMXレジスタ(64bit幅)を使って、なるべく一度に読み込もうとしていた。

r5では、SSEレジスタ(128bit幅)を使用したことで、一度に処理する量が広がった。

とはいえ、まだわずか128bit( = 16byte)なので、まだまだ縦に飛び飛びのアクセスという感じ。

そこで、r8では、縦方向のループのさらに内側に、横方向にループを追加し、横方向に一度に処理する量を飛躍的に増やした(最大1536byte x4)。

一時変数はレジスタには収まらないので、メモリに読み書きすることになるけど、1.5KB x4 = 6KBなので、L1-Dキャッシュ(32KB)に余裕を持って収まる量、になっている。

こうした最適化はループ構造が無駄に複雑になって、ちゃんと動くコードにするのが大変なのだけど、それなりに成果は出る。

r8 → r10
AVX2対応
・HaswellのAVX2に対応した。256bit整数演算がついに可能になり、さぞかし速くなるだろう…と思いきや、shuffle命令のスループットが半減したり、そもそも一部のshuffle命令がなんじゃそりゃという仕様だったり(vpalignrとか、vpshufbとか、unpack系とか)、期待したよりは効果がなかった。まあメモリ帯域で律速してしまっていると演算速度はあまり関係ないというのもある。

スレッド分割法の改善によるロードアンバランス(負荷不均衡)の抑制
・この時点では、scan_frameだけがスレッド並列化されて高速化されていた。ただ、基本的に並列化では、仕事を各スレッドに分配して計算させて並列に計算させて速くするのだけど、すべてのスレッドが計算を終了しないと、次の計算に進むことができない。

r10まではスレッド間の負荷のバランスがとれていなくて、仕事の割り当ての多い1つのスレッドが終わるまで次に進めず、全体の速度が下がってしまっていたので、分割方法を見直して負荷の不均衡を改善したことで、それなりに高速化に繋がった。r8 → r10での速度向上は、AVX2の効果よりも、こちらの効果のほうが大きい。

r10 → r12
・ソフトウェアプリフェッチ再投入により高速化。
r8で連続アクセスが多くなったのでいらないかなと思ってやめていたprefetch命令を再投入した。Intelのマニュアルとかにはハードウェアプリフェッチがよしなにやってくれるからソフトウェアプリフェッチはいらない的なことを書いてるけど、afsの場合にはそんなことはないみたい。まあ、ハードウェアプリフェッチは4Kページ境界を越えられないので、次の行のアクセスなどは明示的にプリフェッチしておくと速くなるということかも。あるいは、D-TLBキャッシュミスも隠ぺいできるのかもしれない。

・フレームを最後に再構成する際にメモリアライメントを考慮して書き出すことで高速化。
アライメントが取れていないメモリアクセスは遅いということなので、とにかくアライメントを調整すれば速くなる。アライメント大事。

r12 → r16
・scan_frameにcount_motionを統合。
メモリアクセスをさらに減らして、微々たるものだけど高速化。速くするためなら何でもするという方針のもと、0.1ms短縮するために、えらい労力をかけた例。おかげで、count_motion分の0.1msを消すことができた。

・scan_frame以外もスレッド並列を導入(サブスレッド)して高速化。
ただし、対象の処理がほとんどメモリ律速なためか、増やすとすぐ遅くなってしまうので、2スレッドまで。まあ5960Xとかの4chだと、もう少しスレッドを増やしたほうが速くなったりする。

r16 → r19
一部の計算にMMX命令を使ったり(MMXレジスタが使えるので…)、こまごまとした最適化をした。まあ、こんなことではあまり速くならない…。

と、執念深くいろいろやってきた。まあ、忘れたのも多いけどこれ以外に試したことはたくさんあって、いろんなネタが没になった…。コードを書いては捨て、書いては捨て…。

もっと確実に速くなるネタをぱっと思いつけるようになりたいけど、まあ、わたしの知識不足というのがあるのだと思う。

ともかく、ここまでいろいろやってしまうと、もう限界かな、と思うようになってきた。とにかく自動フィールドシフトは参照するデータ量が多く、メモリ帯域との戦いをやっている感じなので、辛い…。

で、もうこうなったら「同じ計算をより高速に」は一度おいておいて、「処理をケチって高速に」という方向性を考え始めた。

(続く>>)

2016-12-28(22:36) :
afs :
コメント 0 :
トラックバック 0

このページのトップへ

2016年を振り返る (PC編)

そろそろ2016年も終わりそうなので、今年を振り返ってみた。 (PC編)

2016-12-27(23:31) :
PC :
コメント 0 :
トラックバック 0

このページのトップへ

FFXV クリア感想

3連休でついにFFXVをクリアしたのでその感想とか。

ネタバレは極力避けて書いたつもり。

このページのトップへ

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)の使用 (いまここ)

2016-12-25(21:01) :
OpenCL :
コメント 0 :
トラックバック 0

このページのトップへ

OpenCLメモ (その6) - imageの使用

今回はconvolution計算を例に、imageNd(image1d/image2d/image3d)の使い方を確認した。

続きを閉じる

簡単な例として、2次元畳みこみ演算を考える。


float *input;  //入力配列
float *output; //出力配列
int width, pitch, height; //2Dサイズ
//重みは適当
const float weight[3][3] = {
    { 1.0f, 2.0f, 1.0f },
    { 2.0f, 4.0f, 2.0f },
    { 1.0f, 2.0f, 1.0f },
};
for (int y = 0; y < height; y++) {
    for (int x = 0; x < width; x++) {
        float sump = 0.0f;
        float sumw = 0.0f;
        for (int j = -1; j <= 1; j++) {
            int yj = CLAMP(y+j, 0, height-1);
            for (int i = -1; i <= 1; i++) {
                int xi = CLAMP(x+i, 0, width-1);
                float w = weight[j+1][i+1];
                sump += w * input[yj * pitch + xi];
                sumw += w;
            }
        }
        output[y * pitch + x] = sump / sumw;
    }
}

これをGPUのkernelに置き換えると、


__kernel void convolve(
    __global float* pIn,
    __global float* pOut,
    int width, int pitch, int height) {
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    if (x < width && y < height) {
        const float weight[] = {
            1.0f, 2.0f, 1.0f,
            2.0f, 4.0f, 2.0f,
            1.0f, 2.0f, 1.0f,
        };
        float sump = 0.0f;
        float sumw = 0.0f;
        int iw = 0;
        for (int j = -1; j <= 1; j++) {
            int yj = clamp(y+j, 0, height-1);
            for (int i = -1; i <= 1; i++) {
                int xi = clamp(x+i, 0, width-1);
                float w = weight[iw];
                sumw += w;
                sump += w * pIn[yj * pitch + xi];
                iw++;
            }
        }
        pOut[y * pitch + x] = sump / sumw;
    }
}

本来は重みも引数で渡すべきだけど、今回はそこは本題ではないので、気にしないことにする。

まあ、この畳みこみは非常に簡単な計算で、自分とその周辺について重みをかけて足しこむというものである。

このとき、入力フレームの画素は、重複して何度もロードされることになる。こういうアクセスパターンの場合、OpenCLのimageを使うとテクスチャキャッシュが効いて速い…と聞いたことがあるので、とりあえずやってみた。

image2dは、cl_image_format構造体とcl_image_desc構造体に適切にパラメータを設定して、clCreateImage()によって作ることができる。

ホスト(CPU)のポインタからimageNd(image1d/image2d/image3d)を作る場合には、clCreateImageによって直接imageNdを作る方法と、一度OpenCLバッファを経由してからclCreateImageで作る方法がある。ただ、後者の方法で作成しないとZeroCopyにならず、転送が無駄に発生してしまう点に注意する必要がある。(わりと罠)

直接image2dを作る
ホスト側のポインタをclCreateImageのhost_ptr引数に指定している。ZeroCopyにならない。


//パラメータを適切に設定する
cl_image_format format;
format.image_channel_order = CL_R;         //チャンネル数
format.image_channel_data_type = CL_FLOAT; //データ型

cl_image_desc img_desc;
img_desc.image_type = CL_MEM_OBJECT_IMAGE2D; //2D
img_desc.image_width = arrayWidth;   //サイズ
img_desc.image_height = arrayHeight; //サイズ
img_desc.image_depth = 0;            //サイズ
img_desc.image_array_size = 0;
img_desc.image_row_pitch = arrayPitch * sizeof(cl_float);
img_desc.image_slice_pitch = 0;
img_desc.num_mip_levels = 0;
img_desc.num_samples = 0;
img_desc.buffer = 0;
img_desc.mem_object = 0;

ocl->srcImg = clCreateImage(ocl->context,
    CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
    &format, &img_desc,
    input, //ホスト側のポインタ
    &err);

一度OpenCLバッファを経由する
この方法はこちらで触れられている。

一度OpenCLバッファを作成し、そのバッファをcl_image_desc構造体のmem_objectに指定してclCreateImage()を呼ぶ。この場合は、もとのOpenCLバッファがZeroCopyなら、作成したimage2dもZeroCopyとなる。


//これまでと同じようにOpenCLバッファを作成する
cl_uint nSize = sizeof(cl_float) * arrayPitch * arrayHeight;
ocl->srcMem = clCreateBuffer(ocl->context,
    CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
    nSize,
    input, //ホスト側のポインタ
    &err);

//パラメータを適切に設定する
cl_image_format format;
format.image_channel_order = CL_R;         //チャンネル数
format.image_channel_data_type = CL_FLOAT; //データ型

cl_image_desc img_desc;
img_desc.image_type = CL_MEM_OBJECT_IMAGE2D; //2D
img_desc.image_width = arrayWidth;   //サイズ
img_desc.image_height = arrayHeight; //サイズ
img_desc.image_depth = 0;
img_desc.image_array_size = 0;
img_desc.image_row_pitch = arrayPitch * sizeof(cl_float);
img_desc.image_slice_pitch = 0;
img_desc.num_mip_levels = 0;
img_desc.num_samples = 0;
img_desc.buffer = 0;
img_desc.mem_object = ocl->srcMem; //作成したOpenCLバッファをここに指定する

ocl->srcImg = clCreateImage(ocl->context,
    CL_MEM_READ_ONLY,
    &format, &img_desc,
    nullptr,
    &err);

作成したimageの解放は、OpenCLバッファと同じように行えばよい。


err = clReleaseMemObject(srcImg);

cl_image_format構造体で、使用するimageのformat(チャンネルのフォーマットとデータ型)を指定したが、imageで使用可能なformatは様々なものが用意されている。

まず、チャンネル数については、下記から選択できるが、フォーマットによっては特定のデータ型との組み合わせのみ可能なものがある。

フォーマット	データ型の制約
CL_R CL_Rx CL_A
CL_INTENSITY	CL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, CL_SNORM_INT16, CL_HALF_FLOAT, CL_FLOAT
CL_LUMINANCE	CL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, CL_SNORM_INT16, CL_HALF_FLOAT, CL_FLOAT
CL_DEPTH	CL_UNORM_INT16, CL_FLOAT
CL_RG CL_RGx CL_RA
CL_RGB CL_RGBx	CL_UNORM_SHORT_565, CL_UNORM_SHORT_555, CL_UNORM_INT101010
CL_RGBA
CL_sRGB CL_sRGBx CL_sRGBA CL_sBGRA	CL_UNORM_INT8
CL_ARGB CL_BGRA CL_ABGR	CL_UNORM_INT8, CL_SNORM_INT8, CL_SIGNED_INT8, CL_UNSIGNED_INT8
CL_DEPTH_STENCIL	CL_UNORM_INT24, CL_FLOAT

また、データ型は以下から選ぶことになる。選んだデータ型によってkernelコードで使用すべき関数と、kernelで取り出せる値のデータ型が変わる。また、データ型によっては、特定のフォーマットとの組み合わせのみ可能なものがある。

データ型	kernelで取り出せるデータ型	formatの制約
CL_SNORM_INT8	規格化された符号付き8bit整数
CL_SNORM_INT16	規格化された符号付き16bit整数
CL_UNORM_INT8	規格化された符号なし8bit整数
CL_UNORM_INT16	規格化された符号なし16bit整数
CL_UNORM_SHORT_565	規格化された 5-6-5 3-channel RGB	CL_RGB, CL_RGBx
CL_UNORM_SHORT_555	規格化された x-5-5-5 4-channel xRGB	CL_RGB, CL_RGBx
CL_UNORM_INT_101010	規格化された x-10-10-10 4-channel xRGB	CL_RGB, CL_RGBx
CL_SIGNED_INT8	符号付き8bit整数
CL_SIGNED_INT16	符号付き16bit整数
CL_SIGNED_INT32	符号付き32bit整数
CL_UNSIGNED_INT8	符号なし8bit整数
CL_UNSIGNED_INT16	符号なし16bit整数
CL_UNSIGNED_INT32	符号なし32bit整数
CL_HALF_FLOAT	半精度浮動小数点
CL_FLOAT	単精度浮動小数点
CL_UNORM_INT24	規格化された符号なし24bit整数

imageNdをカーネルで使う際には、すこしコードを変える必要があり、
・引数としては、imageNd_t型
・取り出す際にread_imagef(floatで取り出す場合)、read_imagei/read_imageui(整数型で取り出す場合)を使用
の2点変更する必要がある。

read_imageの取り出し方にはさまざまな方法があるが、まずデータ型によって使用すべき関数が決まっている。

使用すべき関数	データ型	値域
read_imagef	CL_UNORM_INT8 CL_UNORM_INT16	規格化された値 [-1.0 … 1.0]
read_imagef	CL_SNORM_INT8 CL_SNORM_INT16	規格化された値 [0.0 … 1.0]
read_imagef	CL_HALF_FLOAT CL_FLOAT	入力と同じ値
read_imagei	CL_SIGNED_INT8 CL_SIGNED_INT16 CL_SIGNED_INT32	入力と同じ値
read_imageui	CL_UNSIGNED_INT8 CL_UNSIGNED_INT16 CL_UNSIGNED_INT32	入力と同じ値

座標値の指定は、image1d/image2dであればint2型/float2型で、image3dであればint4型/float4型で指定する。

座標の指定方法と値の取り出し方について、3つのオプションがあり、これを指定することができる。

・取り出す座標を指定する際に、0.0～1.0に規格化された座標を用いるかどうか
　- CLK_NORMALIZED_COORDS_TRUE (規格化された浮動小数点の座標で指定)
　- CLK_NORMALIZED_COORDS_FALSE (通常のインデックスで指定)

・領域外の座標を指定した場合にどのように処理するか
　- CLK_ADDRESS_NONE (なにもしない > どんな値が取得されるかは未定義)
　- CLK_ADDRESS_CLAMP (imageの境界の色が採用される)
　- CLK_ADDRESS_CLAMP_TO_EDGE (imageの端の値が採用される)
　- CLK_ADDRESS_REPEAT (imageの反対側の端から繰り返す)
　- CLK_ADDRESS_MIRRORED_REPEAT (imageを折り返して繰り返す)

・座標値に合わせて値の補間を行うか
　- CLK_FILTER_NEAREST (補間を行わず、最も近い座標の値をそのまま使用する)
　- CLK_FILTER_LINEAR (線形補間)

また、関数によりimageNdから取得された値は、float4, half4, int4, short4, char4などの4要素のベクトル型であるが、imageのフォーマット(チャンネル)によって、何番目の要素にどのデータが入っているかが異なる。

チャンネル	data4の中身
CL_R	(r, 0, 0, 1)
CL_A	(0, 0, 0, a)
CL_RG	(r, g, 0, 1)
CL_RA	(r, 0, 0, a)
CL_RGB	(r, g, b, 1)
CL_RGBA CL_BGRA CL_ARGB	(r, g, b, a)
CL_INTENSITY	(I, I, I, I)
CL_LUMINANCE	(L, L, L, 1)

これらを踏まえて、kernelを書き直すと、以下のようになる。


__kernel void convolve(
    __read_only image2d_t imgIn,
    __global float* pOut,
    int width, int pitch, int height) {
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    if (x < width && y < height) {
        const float weight[] = {
            1.0f, 2.0f, 1.0f,
            2.0f, 4.0f, 2.0f,
            1.0f, 2.0f, 1.0f,
        };
        float sump = 0.0f;
        float sumw = 0.0f;
        int iw = 0;
        for (int j = -1; j <= 1; j++) {
            int yj = clamp(y+j, 0, height-1);
            for (int i = -1; i <= 1; i++) {
                int xi = clamp(x+i, 0, width-1);
                float w = weight[iw];
                sumw += w;
                sump += w * read_imagef(imgIn, (int2)(xi,yj)).x;
                iw++;
            }
        }
        pOut[y * pitch + x] = sump / sumw;
    }
}

kernelのほうは引数を変更したが、ホスト(CPU)からkernel引数を渡す場合は特に変更は必要なく、普通にclSetKernelArgで渡せばよい。


clSetKernelArg(ocl->kernel, //対象のkernel
    0,                     //第0引数
    sizeof(cl_mem),        //引数のサイズ
    (void *)&ocl->srcImg); //引数へのポインタ

とすればよい。

計算速度の確認

image2dを使用しない場合

image2dを使用した場合

というわけでkernelの実行時間をみると1.91ms → 2.94msと、image2dを使うことで逆に遅くなってしまった。原因は残念ながらよくわからない…。コードの書き方の問題だろうか…?

まあ、とりあえずOpenCLのimageの使い方の確認ができたのと、(コードの書き方が悪いのかもしれないが)遅くなることもある、ということがわかった。

コードはこちら。

続き>>

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

2016-12-25(16:27) :
OpenCL :
コメント 0 :
トラックバック 0

このページのトップへ

	i7 5960X	GTX1060	M/B交換
	i7 6700K		M/B交換
	i7 4770K
新	Celeron J3710		録画用
	Celeron N3150
	VAIO Tap 11
放出		GTX960
放出		GTX970
放出		R7 360