CUDAによるバンディング低減フィルタの高速化（2）

前回の続きです。メモリアクセス周りをもう少し見てみましょう。

カーネル内でメモリアクセスしてるデータは、乱数表と入出力画像です。このうち、読み書きデータ量の多い入出力画像を見てみます。画像のデータ型はPIXEL_YCです。これはAviUtlのフィルタ処理におけるデータ型で、以下のように定義されています。

typedef struct {
    short y;
    short cb;
    short cr;
} PIXEL_YC;

画像データはこれの配列なので、構造体の配列（Array of Structures, AoS）です。AoSはCUDAが苦手なデータ構造です。なぜならコアレスアクセスができないからです。本来、教科書通りのやり方なら、ここでAoSをSoA（Structure of Arrays）に変更することになりますが、それだとコード変更が多くなるので、ここではもっと簡単なやり方で最適化してみます。

まず、AoSだとなぜダメなのか、見てみましょう。コード中でのアクセスは以下のようになっています。（デフォルト値sample_mode==1の場合）

            PIXEL_YC ref_p = src[offset + ref];
            PIXEL_YC ref_m = src[offset - ref];

構造体1個分をまとめて1行で読んでいます。実は、CUDAがこの構造体1個分をまとめて読み込んでくれれば、パフォーマンスロスは発生しません。しかし、実際にはまとめて読み込んでくれません。これはNVCCコンパイラから出力されるPTXコードを見れば分かります。

        ...
	ld.global.nc.u16 	%rs3, [%r65];
	ld.global.nc.u16 	%rs4, [%r65+2];
	ld.global.nc.u16 	%rs5, [%r65+4];
	cvt.s32.s16	%r69, %rs3;
	ld.global.nc.u16 	%rs6, [%r68];
	ld.global.nc.u16 	%rs7, [%r68+2];
	ld.global.nc.u16 	%rs8, [%r68+4];
        ...

命令レベルで見たら、まぁそうなるよね、という感じでしょうか。

しかし、実はCUDAには、2要素、4要素をまとめて読み込む命令があります。CUDAの組み込み型short2やshort4を使うと、まとめて読み込む命令を吐いてくれます。ここではshort4を使って、まとめて読み込めるようにコードを変更してみます。short4は、名前から想像できると思いますが、4要素のshort型を持った型です。PIXEL_YCは3要素なので、ちょっとデータ構造の変更が必要です。short4で読み込めるように、以下のようなPIXEL_YCAというのを導入します。

struct PIXEL_YCA {
    short    y;
    short    cb;
    short    cr;
    short    a; // 使用しない
};

y,cb,crの値はPIXEL_YCからコピーして最後の要素aは使用しません。データ量が4/3に増えるのがデメリットとしてあります。これを処理するカーネルは、PIXEL_YCをshort4に置換すれば完成です。

template <int sample_mode, bool blur_first>
__global__ void kl_reduce_banding(BandingParam prm,
    short4* __restrict__ dst,
    const short4* __restrict__ src,
    const uint8_t* __restrict__ rand)
{
    …
}

これで、このカーネルの実行時間を測ると、以下のようになります。

劇的という訳ではありませんが、20%程高速化しました。読み書きデータ量が4/3に増えているにもかかわらず、こちらの方が速いようです。4要素をまとめて読み込んでくれていることは、PTXで確認できます。

	ld.global.nc.v4.u16 	{%rs3, %rs4, %rs5, %rs6}, [%r65];
	cvt.s32.s16	%r69, %rs3;
	ld.global.nc.v4.u16 	{%rs10, %rs11, %rs12, %rs13}, [%r68];

実は、short3という組み込み型もあり、今回のPIXEL_YC構造体は3要素なので、short3だとピッタリフィットするのですが、short3だと速くなりません。

template <int sample_mode, bool blur_first>
__global__ void kl_reduce_banding_short3(BandingParam prm,
    short3* __restrict__ dst,
    const short3* __restrict__ src,
    const uint8_t* __restrict__ rand)
{
    …
}

また、short4を使わないでPIXEL_YCAのままアクセスしても、今回の場合は速くなりませんでした。

template <int sample_mode, bool blur_first>
__global__ void kl_reduce_banding_YCA(BandingParam prm,
    PIXEL_YCA* __restrict__ dst,
    const PIXEL_YCA* __restrict__ src,
    const uint8_t* __restrict__ rand)
{
    …
}

どちらも、PTXを見ると、各short要素を１つずつ読んでいました。short3は3要素を読み込む命令がないためと考えられます。PIXEL_YCA経由アクセスの場合は、アラインメントが保証できないからでしょう。short4で読み込む場合、データが2×4=8バイトにアラインしていないといけませんが、PIXEL_YCAが型としては8バイトアラインになっていないからです。

というわけで、今回、2回に分けて、バンディング低減フィルタのCUDA化と高速化を行いました。CUDAは、元がGPUなだけあって画像処理との相性は良く、今回のように割りと簡単にCUDA化できるものは多いです。

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください。