CUDAでプログラミングする前にやっておきたいこと

こちらの記事もご確認下さい。

CUDA6.0用にCPUとGPUの速度比較コードを修正 - 株式会社CFlatの明後日スタイルのブログ

GPUとCPUの速度比較をしたい

以前CUDAでのプログラミングが完成した際に、CPUとの速度比較を行いたいという当たり前の要望が上がりました。 そこでGPUとCPUを切り替えるコードを作成したのですが、どうせなら最初から作っておくべきだったと反省しました。 今回はその際に作成したGPUとCPUの速度比較コードを公開します。

OpenMPの設定

GPUとCPUの速度比較をする際にCPUの方も並列化しないと不公平です。CPUの並列化には設定が簡単なOpenMPを使います。 CUDAプロジェクトでOpenMPを使用できるようにするには、プロジェクトのプロパティを開き、構成プロパティ/CUDA C/C++/Command Lineの追加オプションに-Xcompiler "/openmp"と記載します。 あとは並列化したいfor文の前に#pragma omp parallel forを付け加えるだけです。 CUDAでのプログラミングが初めての方はこの記事も参考にして下さい。

GPUとCPUの切り替えを行うラッピング

下記にあるSwitchableCPUGPU.cuhとTimer.cuh内で宣言されたマクロを呼ぶと、USE_GPUの切り替えでGPUとCPUを切り替えることができます。

SwitchableCPUGPU.cuh

#pragma once
#include "Timer.cuh"

#define USE_GPU

#ifdef USE_GPU

#define SWITCHABLE_DEVICE __device__
#define SWITCHABLE_GLOBAL __global__
#define SWITCHABLE_TIMER CudaEventTimer

#else

#define SWITCHABLE_DEVICE  
#define SWITCHABLE_GLOBAL  
#define SWITCHABLE_TIMER Timer

#endif

template <typename T>
void SwitchableCudaMalloc( T& val, int size )
{
#ifdef USE_GPU
  cudaMalloc( (void**)&val, sizeof(T)*size );
#endif
}

template <typename T>
void SwitchableCudaFree( T* val )
{
#ifdef USE_GPU
  cudaFree( val );
#endif
}

template <typename T>
void SwitchableCudaMemcpyHostToDevice( const T* const host, T* const device, int size )
{
#ifdef USE_GPU
  cudaMemcpy( device, host, sizeof(T)*size, cudaMemcpyHostToDevice );
#endif
}

template <typename T>
void SwitchableCudaMemcpyDeviceToHost( const T* const device, T* const host, int size )
{
#ifdef USE_GPU
  cudaMemcpy( host, device, sizeof(T)*size, cudaMemcpyDeviceToHost);
#endif
}

Timer.cuh

#pragma once
#include <string>
#include <time.h>

class CudaEventTimer
{
public :
  CudaEventTimer( const std::string& message ) : m_message( message )
  {
    cudaEventCreate(&m_start);
    cudaEventCreate(&m_end);
    cudaEventRecord( m_start, 0 );
  }
  ~CudaEventTimer()
  {
    cudaEventRecord( m_end, 0 );
    cudaEventSynchronize( m_end );

    float time;
    cudaEventElapsedTime( &time, m_start, m_end );
    printf("%s = %f sec.\n",m_message.c_str(), time*0.001);

    cudaEventDestroy( m_start );
    cudaEventDestroy( m_end );
  }

private:
  cudaEvent_t m_start;
  cudaEvent_t m_end;
  std::string m_message;
};

class Timer
{
public :
  Timer( const std::string& message ) : m_message( message )
  {
    m_start = clock();
  }
  ~Timer()
  {
    m_end = clock();
    printf("%s = %f sec.\n",m_message.c_str(), (double)(m_end - m_start)/CLOCKS_PER_SEC);
  }

private:
  clock_t m_start;
  clock_t m_end;
  std::string m_message;
};

行列計算でGPUとCPUを比較

下記が行列の掛け算を行うテストコードです。CUDA関連の処理をラッピングされた関数から呼ぶことによって、GPUとCPUの切り替えを容易にしています。

#include <stdio.h>
#include "SwitchableCPUGPU.cuh"

SWITCHABLE_GLOBAL void Calculate( float* matrixA, float* matrixB, float* matrixC, int iLength, int col = 0, int row = 0 )
{
#ifdef USE_GPU
  row = blockIdx.x * blockDim.x + threadIdx.x;
  col = blockIdx.y * blockDim.y + threadIdx.y;

  if ( row > iLength || col > iLength ) return;
#endif

  float target = 0.0f;

  for ( int i = 0 ; i < iLength ; ++i ) {
    target += matrixA[row*iLength + i] * matrixB[i*iLength + col];
  }
  matrixC[row*iLength + col] = target;
}

int main()
{
  // 行列のサイズ決定
  const int iLength = 1024;
  const int iSize = iLength * iLength;

  // CPU側の変数初期化
  float* matrixA = (float*)malloc(sizeof(float)*iSize);
  float* matrixB = (float*)malloc(sizeof(float)*iSize);
  float* matrixC = (float*)malloc(sizeof(float)*iSize);

  for ( int col = 0; col < iLength ; ++col ){
    for ( int row = 0; row < iLength ; ++row ){
      matrixA[col*iLength + row] = rand() % (1000);
      matrixB[col*iLength + row] = rand() % (1000);
      matrixC[col*iLength + row] = 0.0f;
    }
  }

  // ここから時間計測
  SWITCHABLE_TIMER t("time");

  // GPU側の変数初期化
  float* d_matrixA;
  float* d_matrixB;
  float* d_matrixC;
  SwitchableCudaMalloc( d_matrixA, iSize );
  SwitchableCudaMalloc( d_matrixB, iSize );
  SwitchableCudaMalloc( d_matrixC, iSize );

  SwitchableCudaMemcpyHostToDevice( matrixA, d_matrixA, iSize );
  SwitchableCudaMemcpyHostToDevice( matrixB, d_matrixB, iSize );

  // 行列計算
#ifdef USE_GPU
  const int iThread = 16;
  dim3 thread( iThread, iThread );
  const int iBlock = ( iLength + iThread - 1 )/iThread;
  dim3 block( iBlock, iBlock );

  Calculate<<<block, thread>>>( d_matrixA, d_matrixB, d_matrixC, iLength );
  cudaThreadSynchronize();

#else
#pragma omp parallel for
  for ( int i = 0 ; i < iLength ; ++i ) {
    for ( int j = 0 ; j < iLength ; ++j ) {
      Calculate( matrixA, matrixB, matrixC, iLength, i, j );
    }
  }
#endif

  // 後処理
  SwitchableCudaMemcpyDeviceToHost( d_matrixC, matrixC, iSize );

  free( matrixA );
  free( matrixB );
  free( matrixC );
  SwitchableCudaFree( d_matrixA );
  SwitchableCudaFree( d_matrixB );
  SwitchableCudaFree( d_matrixC );

  return 0;
}

実行結果

実行結果は次のようになりました。参考までに並列化しなかったCPUの結果も載せておきます。 GPUには苦手な計算もあるので、OpenMPと比べて明らかに遅い場合はGPUの使用を検討した方が良いかもしれません。

GPU:0.152510s
CPU OpenMP:1.220000s
CPU 並列化無し:7.254000s

GPU計算の高速化

今回のプログラムでは全く行っていませんが、GPU計算には色々な高速化手法があります。 興味のある方は下記の書籍等を参考にして下さい。

はじめてのCUDAプログラミング―驚異の開発環境[GPU+CUDA]を使いこなす! (I・O BOOKS)

CUDA by Example 汎用GPUプログラミング入門