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_INTENSITYCL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, CL_SNORM_INT16, CL_HALF_FLOAT, CL_FLOAT
CL_LUMINANCECL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, CL_SNORM_INT16, CL_HALF_FLOAT, CL_FLOAT
CL_DEPTHCL_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_STENCILCL_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 RGBCL_RGB, CL_RGBx
CL_UNORM_SHORT_555規格化された x-5-5-5 4-channel xRGBCL_RGB, CL_RGBx
CL_UNORM_INT_101010規格化された x-10-10-10 4-channel xRGBCL_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_imagefCL_UNORM_INT8
CL_UNORM_INT16
規格化された値 [-1.0 … 1.0]
read_imagefCL_SNORM_INT8
CL_SNORM_INT16
規格化された値 [0.0 … 1.0]
read_imagefCL_HALF_FLOAT
CL_FLOAT
入力と同じ値
read_imageiCL_SIGNED_INT8
CL_SIGNED_INT16
CL_SIGNED_INT32
入力と同じ値
read_imageuiCL_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)の使用


スポンサーサイト

コメントの投稿

非公開コメント

プロフィール

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