SHOEISHA iD

※旧SEメンバーシップ会員の方は、同じ登録情報(メールアドレス&パスワード)でログインいただけます

CodeZine編集部では、現場で活躍するデベロッパーをスターにするためのカンファレンス「Developers Summit」や、エンジニアの生きざまをブーストするためのイベント「Developers Boost」など、さまざまなカンファレンスを企画・運営しています。

特集記事

CUDA:「超並列コンピューティング」はじめました(その2)
~高速化TipsとOpenGLでのグラフィクスへの応用


  • X ポスト
  • このエントリーをはてなブックマークに追加

メモリ・アクセスを速くする

 お次はGPU内で動くデバイス側コード(kernel)の高速化。cudaMemcpyでコピー元/先に指定されるのはデバイス内メモリ、すなわちビデオ・カード上のビデオ・メモリです。kernel呼び出し時に引数に渡しているのはビデオ・メモリ上の位置を指すポインタで、kernelはGPUの"外部"にあるビデオ・メモリをアクセスしています。

 GPUはチップ"内部"にもメモリを持っていて、GPUの外にあるビデオ・メモリよりずっと高速。kernelがこのGPUチップ内メモリを使うよう仕向けてやればメモリ・アクセスの頻度が多いほど高速になるはずです。

 簡単な例としてSMA(Simple Moving Average:単純移動平均)を考えます。SMAは株価や気温といった均等間隔に並ぶデータの平滑化、デコボコを埋めて滑らかにする計算です。

 例えば毎日正午の気温を測定した値がsize日分、float input[size]に納められているとしましょう。input[0]が本日、input[i]がi日前の正午の気温データです。そうするとi日前における過去N日間分の平均気温 result[i] は、

result[i] = (input[i] + input[i+1] + …… + input[i+N-1]) / N

ですね。これを i = 0..size-N について計算しようと。

 まず、素直に書いたkernelはこんなところでしょうか。

list02
__global__ void kernel_SMA(float* result, const float* input, size_t size, size_t N) {
  size_t i = blockDim.x * blockIdx.x + threadIdx.x;
  if ( i <= size - N ) {
    float sum = 0.0f;
    for ( size_t t = 0; t < N; ++t ) {
      sum += input[i+t];
    }
    result[i] = sum / N;
  }
}

 入力データinputの各要素はN回読み込まれます。いったん読みだしてGPU内部のメモリに置いとけば、それ以降より高速な読み出しができるハズ。

 ・高速化の方法1つめ、read-onlyキャッシュを紹介しましょう、こんなコードです。

list03
__global__ void kernel_SMA_ROCache(float* result, const float* __restrict__ input, size_t size, size_t N) {
  size_t i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i <= size - N) {
    float sum = 0.0f;
      for (size_t t = 0; t < N; ++t) {
        sum += input[i + t];
    }
    result[i] = sum / N;
  }
}

 素直に書いたコードとの違いは引数:const float* inputが__restrict__で修飾されているところ、それだけ。__restrict__は"このポインタが指す先の値は変化しない"ことをコンパイラに教えてあげるキーワードです。これによってinput[~]は変化しない、すなわちread-only であるとみなされ、GPU内のread-onlyキャッシュに置いて最適化するヒントとなります。

 ただし、__restrict__が効果を現すのはGPUのCompute Capabilityが3.5以降のもの。有効にするにはコンパイル時のオプションを調整しておかなくてはなりません。

 ・高速化の方法2つめはshared-memory(共有メモリ)の利用です。

 kernelコードでshared-memoryを利用するには、変数を__shared__修飾します。例えば、

__shared__ int data;
__shared__ float buffer[100];

のように。__shared__修飾された変数はkernel呼び出し時に指定したブロックに属する全スレッドで共有されます(ブロックをまたいだ共有はできません)。1つのスレッドで書いたbuffer[]を同一ブロック内の他のスレッドが読めるわけ。

 shared-memoryを利用したSMAはこうなりました。

list04
__global__ void kernel_cached_SMA(float* result, const float* input, size_t size, size_t N) {
  size_t gid = blockDim.x * blockIdx.x + threadIdx.x;

  extern __shared__ float cache[]; // block内の全threadで共有される

  if (gid < size) // 値が存在すれば書いてよい
    cache[threadIdx.x] = input[gid];

  // blockDim.xで足りない分、オフセットを与えて読み込む。
  if ((threadIdx.x  < N - 1) && (gid + blockDim.x < size))
    cache[threadIdx.x + blockDim.x] = input[gid + blockDim.x];

  syncthreads(); // block内の全threadがここで待ち合わせ

  if (gid < size - N + 1) {
    // cacheを入力データとしてSMAを求める
    float sum = 0.0f;
    for ( size_t t = 0; t < N; ++t )
      sum += cache[threadIdx.x + t];
    result[gid] = sum / N;
  }
}

 extern __shared__ float cache[];に要素数が書かれていません。extern __shared__としておくと、その大きさ(バイト数)を動的に外から与えることができます。呼び出し時、kernel_cached_SMA<<<>>>の3つめのパラメータがextern __shared__したブロック内共有領域のバイト数となります。

 extern __shared__ float cache[]の中身を各スレッドで手分けして埋めています。syncthreads()は、ブロック内の全スレッドがこの地点で待ち合わせるため。キャッシュが全部埋まるのを待たないと、その後に続く計算結果がおかしくなっちゃいますからね。

 500回ほど繰り返し、キャッシュの効果を調べてみたところ……。

list05
enum KernelType {
  SMA = 0,
  SMA_ROCache,
  cached_SMA,
};

chrono::milliseconds device_SMA(float* result, const float* input, size_t size, size_t N, KernelType kernelType, int iteration) {
  chrono::high_resolution_clock::time_point start, stop;

  float* dinput  = nullptr;
  float* dresult = nullptr;

  CUERRCHK(cudaMalloc((void**)&dinput,  size*sizeof(float)));
  CUERRCHK(cudaMalloc((void**)&dresult, size*sizeof(float)));

  CUERRCHK(cudaMemcpy(dinput, input, size*sizeof(float), cudaMemcpyHostToDevice));
  CUERRCHK(cudaMemset(dresult, 0, size*sizeof(float))); // テストのため。0に初期化しておく。

  start = chrono::high_resolution_clock::now();
  const unsigned int block = 256; // 1-blockあたりのthread数
  assert( block > N );
  const unsigned int grid = (size + block -1) / block; // 1-grid あたりの block数
  for ( int i = 0; i < iteration; ++i ) {
    switch (kernelType) {
    case SMA:
      kernel_SMA <<<grid, block >>>(dresult, dinput, size, N);
      break;
    case SMA_ROCache:
      kernel_SMA_ROCache<<<grid, block >>>(dresult, dinput, size, N);
      break;
    case cached_SMA:
      kernel_cached_SMA<<<grid, block, (block + N - 1) * sizeof(float) >>>(dresult, dinput, size, N);
      break;
    }
  }
  CUERRCHK(cudaGetLastError());
  cudaError_t err = cudaDeviceSynchronize();
  assert(err == cudaSuccess);

  stop = chrono::high_resolution_clock::now();

  CUERRCHK(cudaMemcpy(result, dresult, size*sizeof(float), cudaMemcpyDeviceToHost));

  CUERRCHK(cudaFree(dinput));
  CUERRCHK(cudaFree(dresult));
  return chrono::duration_cast<chrono::milliseconds>(stop - start);
}

 なるほど、確かにGPU内メモリが効いてますね。処理の過程で特定の領域を何度もアクセスしなきゃいけないときは、かなりの効果が期待できそうです。

次のページ
5年ぶりのマンデルブロ集合

この記事は参考になりましたか?

  • X ポスト
  • このエントリーをはてなブックマークに追加
特集記事連載記事一覧

もっと読む

この記事の著者

επιστημη(エピステーメー)

C++に首まで浸かったプログラマ。Microsoft MVP, Visual C++ (2004.01~2018.06) "だった"りわんくま同盟でたまにセッションスピーカやったり中国茶淹れてにわか茶...

※プロフィールは、執筆時点、または直近の記事の寄稿時点での内容です

この記事は参考になりましたか?

この記事をシェア

  • X ポスト
  • このエントリーをはてなブックマークに追加
CodeZine(コードジン)
https://codezine.jp/article/detail/8480 2023/09/25 15:04

おすすめ

アクセスランキング

アクセスランキング

イベント

CodeZine編集部では、現場で活躍するデベロッパーをスターにするためのカンファレンス「Developers Summit」や、エンジニアの生きざまをブーストするためのイベント「Developers Boost」など、さまざまなカンファレンスを企画・運営しています。

新規会員登録無料のご案内

  • ・全ての過去記事が閲覧できます
  • ・会員限定メルマガを受信できます

メールバックナンバー

アクセスランキング

アクセスランキング