Xperia acroの予約が始まり、予約しようかしまいか迷っている人です。

それはどうでもいいとして。

今回はCUDAでちょっと遊んでみる・・・というより実験してみた記録です。

なぜ今CUDAなのか？というと、「日経ソフトウェアでGPUプログラミングの特集をやっていたのでついでにこっちでもやってみよう！」という感じです。

そのままプログラムを組んで実行結果を見てもしょうがないのでもうちょっと(私的に)実用的なルーチンでいろいろと試してみることにしました。

あらかじめ書いておきますと、グラフィックはCUDAのバージョンとしてはかなり古い(CUDA1.1)ものを使っているのでちょっといろいろと微妙ですが、実験レベルなら問題ないだろう、との考えです。

ちなみに、CUDAをコンパイルできる環境を整えるのに一苦労しました。32bit版と64bit版で微妙に付属してるライブラリが異なっているので妙なマージを行ってみたり、

予約語のハイライトがどこにあるのかわからなくて結局サンプルSDKに入っていることをインストール後に知ったりとしっちゃかめっちゃかでした。

VisualStudioとは2005,2008,2010のどれかで統合できるのでそれらを持っているなら環境を整えてみるのも一興かと思います。

試しで書いてみたのはCUDAを使ったRGB<=>YUV変換ルーチンです。(A)RGB32と(A)YUVを相互変換するルーチンです。

これについては動画変換プログラムを作ったときにCPU側でできる限り最適になるようにSSE2を使ってルーチンを組みましたが、それのCUDAバージョンと言ったところです。

CUDAの処理速度をみる分にはある意味最適なルーチン(CPU側で高速な処理ができるので)ですね。

CUDAを使ったRGB<=>YUV変換

CUDAのルーチン自体は普通に変換部分を書けばいいだけなのでものすごい単純です。

__global__ void RGBToYUVCuda(const uchar4 *input,uchar4 *output,const int width,const int pitch)
{
	float y,u,v,r,g,b; int idx,i; uchar4 src,dst;
	idx = blockDim.x * blockIdx.x + threadIdx.x;
	input = (uchar4 *)((unsigned char *)input + idx * pitch);
	output = (uchar4 *)((unsigned char *)output + idx * pitch);
	for(i = 0;i < width;i++){
		src = input&#91;i&#93;; b = src.x; g = src.y; r = src.z;
		y = r *  0.299f + g *  0.587f + b *  0.114f;
		u = r * -0.168f + g * -0.331f + b *  0.500f;
		v = r *  0.500f + g * -0.419f + b * -0.081f;
		/* u = 0.564 * (b - y); */
		/* v = 0.713 * (r - y); */
		dst = make_uchar4((unsigned char)(v + 128.0f),(unsigned char)(u + 128.0f),(unsigned char)y,src.w);
		output&#91;i&#93; = dst;
	}
}
__global__ void YUVToRGBCuda(const uchar4 *input,uchar4 *output,const int width,const int pitch)
{
	float y,u,v,r,g,b; int idx,i; uchar4 src,dst;
	idx = blockDim.x * blockIdx.x + threadIdx.x;
	input = (uchar4 *)((unsigned char *)input + idx * pitch);
	output = (uchar4 *)((unsigned char *)output + idx * pitch);
	for(i = 0;i < width;i++){
		src = input&#91;i&#93;; v = (float)src.x - 128.0f; u = (float)src.y - 128.0f; y = src.z;
		r = y				+ v *  1.4026f;
		g = y + u * -0.3444f + v * -0.7114f;
		b = y + u *  1.7330f			   ;
		
		dst = make_uchar4((unsigned char)b,(unsigned char)g,(unsigned char)r,src.w);
		output&#91;i&#93; = dst;
	}
}
&#91;/cpp&#93;
<p>全部浮動小数でやっていますが、整数演算にしても大して速度は変わりませんでした。</p>
<br>
<p>で、制御部は、</p>

void ConvRGBAndYUVCuda(void *lpDest,const void *lpSrc,uint32_t width,uint32_t height)
{
	uchar4 *src,*dst,*tmp; size_t len,pitch,lines; cudaError_t err; cudaDeviceProp cdp; int threads,blocks;
	cudaGetDeviceProperties(&cdp,0);
	pitch = width * 4; len = pitch * height;
	err = cudaMalloc((void **)&src,len);
	err = cudaMalloc((void **)&dst,len);
	err = cudaMalloc((void **)&tmp,len);
	while(!(width & 0x01)){ width /= 2; pitch /= 2; height *= 2; } /* pitch分割処理 */
	
	err = cudaMemcpy(src,lpSrc,len,cudaMemcpyHostToDevice);
	for(blocks = height,threads = 1;blocks > cdp.maxThreadsPerBlock;blocks /= 2,threads *= 2);
	lines = blocks * threads;
	RGBToYUVCuda<<<blocks,threads>>>(src,tmp,width,pitch);
	if(lines != height){ RGBToYUVCuda<<<height - lines,1>>>((uchar4 *)((unsigned char *)src + lines * pitch),(uchar4 *)((unsigned char *)tmp + lines * pitch),width,pitch); }
	YUVToRGBCuda<<<blocks,threads>>>(tmp,dst,width,pitch);
	if(lines != height){ YUVToRGBCuda<<<height - lines,1>>>((uchar4 *)((unsigned char *)tmp + lines * pitch),(uchar4 *)((unsigned char *)dst + lines * pitch),width,pitch); }
	err = cudaMemcpy(lpDest,dst,len,cudaMemcpyDeviceToHost);
	cudaFree(src); cudaFree(dst); cudaFree(tmp);
}

相変わらずエラーを受け取ってもその処理は書いていません。

ちょっと変わったことをしているのは、処理の分割数をできるだけ多くするために擬似的に縦方向に長いような絵に見せかける演算を行っています。

この処理がないと分割数が多くならず処理が早くなりません。

GPUプログラミングで肝となる部分だと思うのですが、

「できるだけ処理を分割して単純な処理を大量に発行する」

と処理がものすごく早くなります。分割数が少なすぎて早くならない場合はどうにかして分割数を多くなるように変形しましょう。

もう一つの注意点は、「メモリの確保と転送が(対象がデバイスとなるため)メインメモリより処理が遅くなることがあるので発行回数をできるだけ減らす」ですか。

正しいテストを行ったわけではないのでこれが完全に正しいかどうかは不明ですが…。

CPUvsGPUその結果は？

CPUとの純粋な演算速度勝負をやってみた例です。メモリ転送については繰り返しの最初のみ元の絵をコピーして繰り返しの最後のみ結果の絵をコピーするようにしています。

CPUが1コアのみ使用なので、スレッド化するとどうなるかが微妙に見物なのですが・・・。時間はGetTickCountで計測したものを表記しています。

対象の絵は800×600 32bit PNGを使っています。

処理繰り返し回数	1	512	1024	2048	4096	8192
CPU処理時間(Core2Quad Q8400)[ms]	0	1841	3650	7301	14586	29188
GPU処理時間(GeForce 9600GT)[ms]	156	920	1670	3182	6209	12215
GPU処理時間(pitch分割停止)[ms]	156	2996	5850	11513	22870	45568

1回のみの場合はメモリのセットアップにかなりの時間がとられるのでどうしてもCPUのほうが早いです。

が、繰り返しの回数が多くなるとさすがにGPUの面目躍如といった感じの結果が得られました。これについては日経ソフトウェアの記事のほうが詳しいです。

また、ちょっとした注意点なのが、CPUで浮動小数演算を行ったときとGPUで浮動小数演算を行ったときの結果が多少違うことがあります。

この辺はある意味誤差レベルなのでそれほど気にする必要はないはずですが、気にしたい人は気に留めておくといいかもです。

また、分割数が少なくなる(pitch分割処理をコメントアウトした)場合にはCPUにすら速度的に負けてしまうことになります。

これについては作っていく段階でテストを行ってどれくらいがいいのか見極める必要があるという面ではなかなか難しいものがあるのではないかと思います。

次は動画変換時の動き検索とかに使えないかな～と思いながら

ちょっとプログラムを組んでいく予定です。簡単に早くなるとは思いませんが、これはこれでおもしろいということも収穫ですし、がんばればいろんな処理が高速化できるんだろうな～というのもよくわかりました。

ちなみに、GPUのプログラムだと内部的にはスレッドやら何やらで処理が細切れになっているのですが、それが完璧に隠されているのがこちらとしては考えやすいです。

同期が不必要な処理かつ大量の同じ処理の繰り返しなら考えてみるのもいいと思います。

ADVシステムの方にもいろいろと修正が必要な箇所が累積しているのが自分の頭で考えてわかっているのでその辺も組みながら、というところでしょうか。

どうでもよくないですが、システム管理の仕事をしているのですが、管理の手間賃を結局もらっていないのですが、これってどうなのでしょうか？