CUDA6.0用にCPUとGPUの速度比較コードを修正

以前の記事でCUDAプログラミングする際に、CPUとGPUを切り替える方法を書いたのですが、 その2日後にCUDA6.0が正式リリースされ、Unified Memoryという 新しいメモリ確保の仕組みが提供されました。

メモリ確保が簡単に

CUDA5.5まではCPUとGPUの両方にメモリ領域を確保して、それぞれのメモリからメモリへデータを転送するという作業をユーザーがする必要がありました。 CUDA6.0からはCPUとGPUで使用するメモリを1変数で管理できるようになりました。当然データを転送する関数をユーザーが呼ぶ必要もありません。

やるべきことはGPUで使用する変数を用意してcudaMallocManaged()で領域を確保し、使い終わったらcudaFree()で解放するだけです。 メモリ管理が簡単になったので、今回は下記のようなコンストラクタとデストラクタでメモリの確保と解放を行うCpuGpuDataクラスを作ってみました。

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 <class T>
class CpuGpuData {
public:
  CpuGpuData( const int iSize )
  {
    cudaMallocManaged( &m_data, sizeof(T)*iSize );
  }
  ~CpuGpuData()
  {
    cudaFree( m_data );
  }

  T* m_data;
};

Timerクラスは前回と同じです。

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;
};

コードもすっきり

下記が前回と同じ処理を行うテストコードです。メモリの確保と解放をクラスが行っているのですっきりしています。

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

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;

  CpuGpuData<float> matrixA( iSize );
  CpuGpuData<float> matrixB( iSize );
  CpuGpuData<float> matrixC( iSize );

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

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

  // 行列計算
#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>>>( matrixA.m_data, matrixB.m_data, matrixC.m_data, iLength );
  cudaDeviceSynchronize();

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

  return 0;
}

実行結果

実行結果も載せておきます。このプログラムでは前回の結果とほぼ同じです。 Unified Memoryの実行速度についてはこちら で詳しく調査されています。メモリの転送が自動で行われてしまうので、注意しないと遅くなるようです。

GPU:0.149291s
CPU OpenMP:1.231000s
CPU 並列化無し:7.213000s

GPU計算の高速化

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

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

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