CUDAで画像合成

CUDA Mac

今回はCUDAを使って画像合成（アルファブレンディング）をやってみた。元のソースコードは下記を参照。

テスト環境は下記の通り。画像ファイルの読み込みには、MacPorts経由でインストールしたOpenCVを利用した。

まずは(1)CPUでの合成を行い、(2)次にGPUでの合成処理を行った。また「CUDAメモリ転送のパフォーマンス」でも記載したように、通常のメモリを転送元に指定するより、'pinned memory'を使う方が速いので、(3)一旦pinned memoryにデータをコピーした後、次にdevice上のメモリへ転送する処理も行ってみた。この時のCPU側のコードはこんな感じ。

  uchar *h_data1 = NULL;
  uchar *h_data2 = NULL;
  uchar *h_data3 = NULL;
  cutilSafeCall( cudaHostAlloc( (void**)&h_data1, len3, 0 ) );
  cutilSafeCall( cudaHostAlloc( (void**)&h_data2, len3, 0 ) );
  cutilSafeCall( cudaHostAlloc( (void**)&h_data3, len3, 0 ) );
	
  double t1 = getMsecTime();

  // main memory -> pinned memory
  memcpy(h_data1, img1->imageData, len3 );
  memcpy(h_data2, img2->imageData, len3 );
	
  // pinned memory -> device memory
  cutilSafeCall(cudaMemcpy( m_srcBuffer1, h_data1, len3, cudaMemcpyHostToDevice));
  cutilSafeCall(cudaMemcpy( m_srcBuffer2, h_data2, len3, cudaMemcpyHostToDevice));

  // compose images
  composeImages1(m_gridSize, m_blockSize, m_dstBuffer22, m_srcBuffer1, m_srcBuffer2, imageWidth, imageHeight, alpha);
  cutilSafeCall(cudaThreadSynchronize());
  cutilCheckMsg("kernel failed.");
	
  // device memory -> main memory
  cutilSafeCall(cudaMemcpy( image22->imageData, m_dstBuffer22, len3, cudaMemcpyDeviceToHost));
	
  double t2 = getMsecTime();
  std::cout << "Time: " << t2 - t1 << " msec" << std::endl;
	
  cutilSafeCall( cudaFreeHost(h_data1) );
  cutilSafeCall( cudaFreeHost(h_data2) );
  cutilSafeCall( cudaFreeHost(h_data3) );

GPU側で処理するコードは下記。

 __global__
void render1(uchar3 *dstPtr, uchar3 *srcPtr1, uchar3 *srcPtr2, uint width, uint height, int w)
{
  const uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
  const uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
  const int n = y * width + x;

  int a = 255 - w;
  int b = (srcPtr1[ n ].x * w + srcPtr2[ n ].x * a)>>8;
  int g = (srcPtr1[ n ].y * w + srcPtr2[ n ].y * a)>>8;
  int r = (srcPtr1[ n ].z * w + srcPtr2[ n ].z * a)>>8;

  dstPtr[ n ] = make_uchar3((uchar)b, (uchar)g, (uchar)r);
}

extern "C"
void composeImages1(dim3 gridSize, dim3 blockSize, uchar3 *dstPtr, uchar3 *srcPtr1, uchar3 *srcPtr2, uint width, uint height, int w)
{
 render1<<<gridSize, blockSize>>>(dstPtr, srcPtr1, srcPtr2, width, height, w);
}

実行結果は下記の通り。

# 処理方法 時間(msec) 速度比
(1) CPU 162.34 1
(2) GPU 63.0 2.57
(3) GPU(pinned) 41.2 3.94

この結果から、次の事が分かる。

  1. アルファブレンディングの処理をCPUからGPUに持って行くと、約2.6倍の性能が出た。
  2. 画像転送元のメモリとして'pinned memory'を使うと、約3.9倍の性能が出た。

GPUでの処理はもう少し速くなるかと思っていたが、SSEよりは少し速い程度という結果になった。やはり計算そのものよりも、デバイス側へのメモリ転送がネックになっているのだろう。実際、(3)のように直接デバイス側にデータ転送するのではなく、一旦pinned memoryを経由させてデータを渡す方が速くなると言うのも面白い。但し、(2)に比べて(3)の処理は安定していないようで、動作速度のばらつきが大きかった。また、処理結果をメインメモリへ戻す際には、pinned memoryを使っても処理は向上しなかった。