CPU使用率を抑えたCUDA実行

出典: トータル・ディスクロージャ・サイト(事実をありのままに)

目次

Kernel終了実行待ちのCPU使用率削減

CUDAの仕様では、実行中のCUDA Contextの数がホストの論理CPU数以下の場合、CUDA Kernelの実行終了を待つ際にCPU使用率100 %でポーリングする動作がデフォルトとなっている。忙しくポーリングするおかげで、実行終了がホストに伝わるレイテンシは短くなるが、GPU動作中にホストで行いたい計算が在ったり、CPUをスリープさせて消費電力を抑える必要が在ったり、この挙動を変えたい場合も多い。

この挙動については、cudaSetDeviceFlagsというAPIでGPUごとに制御できる。このAPIは引数を1つだけとり、cudaDeviceScheduleSpinでContext数や論理CPU数に関係無くポーリングを強制、逆にcudaDeviceScheduleYieldでスリープを強制、cudaDeviceScheduleAutoでContext数と論理CPU数を比べるデフォルトの挙動となる。なお、このAPIは、実際にGPUが何かしらの処理を始める前という、かなり限られた時期でしか呼び出すことができないので、できるだけプログラムの処理の早い段階で呼んでおきたい。

メモリ転送のCPU使用率削減

上記のAPIを使っても、メモリ転送ではやはりCPUがフル回転してしまう。このメモリ転送も実はCPUを使わずに行う方法が在る。何らかの形でPinnedホストメモリを確保し、これを非同期転送すれば、CPUを介さない転送となる。ただし、非同期APIの同期待ちは、cudaDeviceScheduleYieldを指定してもCPU使用率100 %でポーリングする。この同期待ちでCPUの使用率を下げるには、cudaDeviceScheduleBlockingSyncというフラグ(前節のフラグと同時指定可)を指定する必要が在る。

一方、Streamが処理中かどうかの状態をクエリするAPI、cudaStreamQueryはCPUをほとんど使わずに実行できる。このAPIがcudaSuccessを返せば処理終了であるので転送終了とみなして処理を続行でき、cudaErrorNotReadyを返せばまだ処理中であることが分かる。適切なタイミングで定期的にこのAPIを呼んで終了を確認し、その合間にホスト側の処理を挟むやり方も在るので、覚えておくと良いだろう。

CUBLASを使った行列積計算を、できるだけCPUを使わずに行うサンプルである。実行環境は以下のようなものである。

CPU: Core i7 2600K 3.4 GHz
GPU: GeForce GTX 580
OS: Ubuntu 12.04
CUDA: CUDA 5.0
GPUドライバ: バージョン304.64

cudaDeviceScheduleBlockingSyncを使った例

#include <unistd.h>
#include <cublas_v2.h>

static const size_t n = 8000;

int
main()
{
  cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync);
  cudaStream_t stream;
  cudaStreamCreate(&stream);
  cublasHandle_t handle;
  cublasCreate(&handle);
  cublasSetStream(handle, stream);
  float *d_a, *d_b, *d_c, *h_a, *h_b, *h_c;
  size_t d_a_pitch, d_b_pitch, d_c_pitch;
  cudaHostAlloc((void **)&h_a, n * n * sizeof(float), cudaHostAllocDefault);
  cudaHostAlloc((void **)&h_b, n * n * sizeof(float), cudaHostAllocDefault);
  cudaHostAlloc((void **)&h_c, n * n * sizeof(float), cudaHostAllocDefault);
  cudaMallocPitch((void **)&d_a, &d_a_pitch, n * sizeof(float), n);
  cudaMallocPitch((void **)&d_b, &d_b_pitch, n * sizeof(float), n);
  cudaMallocPitch((void **)&d_c, &d_c_pitch, n * sizeof(float), n);
  cudaMemsetAsync(&h_a, 0, n, stream);
  cudaMemsetAsync(&h_b, 0, n, stream);
  cudaMemcpy2DAsync(d_a, d_a_pitch, h_a, n * sizeof(float),
                    n * sizeof(float), n,
                    cudaMemcpyHostToDevice, stream);
  cudaMemcpy2DAsync(d_b, d_b_pitch, h_b, n * sizeof(float),
                    n * sizeof(float), n,
                    cudaMemcpyHostToDevice, stream);
  const float alpha = 1.0f;
  const float beta = 0.0f;
  cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n,
              &alpha,
              d_a, d_a_pitch / sizeof(float),
              d_b, d_b_pitch / sizeof(float),
              &beta,
              d_c, d_c_pitch / sizeof(float));
  cudaMemcpy2DAsync(h_c, n * sizeof(float), d_c, d_c_pitch,
                    n * sizeof(float), n,
                    cudaMemcpyDeviceToHost, stream);
  cudaStreamSynchronize(stream);
  cudaFree(d_c);
  cudaFree(d_b);
  cudaFree(d_a);
  cudaFreeHost(h_c);
  cudaFreeHost(h_b);
  cudaFreeHost(h_a);
  cublasDestroy(handle);
  cudaStreamDestroy(stream);
  return 0;
}

cudaSetDeviceFlags(cudaDeviceScheduleBlockingSync)でCPU使用率を減らすフラグを最初に立て、その後Pinnedメモリ上の入力を非同期転送APIでGPU上に転送し、Streamと紐付けたCUBLASのcublasSgemmで計算、計算結果も非同期転送APIでPinnedメモリ上に受け取る、という流れである。最後に、cudaStreamSynchronizeで同期を行うが、cudaDeviceScheduleBlockingSyncフラグの効果で同期処理もCPU使用率が少なくなる。

このプログラムをtimeコマンド付きで実行する。

$ time ./a.out
real    0m2.427s
user    0m0.064s
sys     0m0.568s

全実行時間2.427秒のうち、CPUが動いていたのは0.632秒である。この0.632秒の大部分はCUBLASの初期化とPinnedホストメモリの確保であるので、これらの処理を一度で済ませて再利用すれば、よりCPU使用率を下げることができるだろう。

cudaStreamQueryを使った例

#include <unistd.h>
#include <cublas_v2.h>

static const size_t n = 8000;

int
main()
{
  cudaStream_t stream;
  cudaStreamCreate(&stream);
  cublasHandle_t handle;
  cublasCreate(&handle);
  cublasSetStream(handle, stream);
  float *d_a, *d_b, *d_c, *h_a, *h_b, *h_c;
  size_t d_a_pitch, d_b_pitch, d_c_pitch;
  cudaHostAlloc((void **)&h_a, n * n * sizeof(float), cudaHostAllocDefault);
  cudaHostAlloc((void **)&h_b, n * n * sizeof(float), cudaHostAllocDefault);
  cudaHostAlloc((void **)&h_c, n * n * sizeof(float), cudaHostAllocDefault);
  cudaMallocPitch((void **)&d_a, &d_a_pitch, n * sizeof(float), n);
  cudaMallocPitch((void **)&d_b, &d_b_pitch, n * sizeof(float), n);
  cudaMallocPitch((void **)&d_c, &d_c_pitch, n * sizeof(float), n);
  cudaMemsetAsync(&h_a, 0, n, stream);
  cudaMemsetAsync(&h_b, 0, n, stream);
  cudaMemcpy2DAsync(d_a, d_a_pitch, h_a, n * sizeof(float),
                    n * sizeof(float), n,
                    cudaMemcpyHostToDevice, stream);
  cudaMemcpy2DAsync(d_b, d_b_pitch, h_b, n * sizeof(float),
                    n * sizeof(float), n,
                    cudaMemcpyHostToDevice, stream);
  const float alpha = 1.0f;
  const float beta = 0.0f;
  cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, n, n,
              &alpha,
              d_a, d_a_pitch / sizeof(float),
              d_b, d_b_pitch / sizeof(float),
              &beta,
              d_c, d_c_pitch / sizeof(float));
  cudaMemcpy2DAsync(h_c, n * sizeof(float), d_c, d_c_pitch,
                    n * sizeof(float), n,
                    cudaMemcpyDeviceToHost, stream);
  while (cudaStreamQuery(stream) == cudaErrorNotReady) {
    sleep(1);
  }
  cudaFree(d_c);
  cudaFree(d_b);
  cudaFree(d_a);
  cudaFreeHost(h_c);
  cudaFreeHost(h_b);
  cudaFreeHost(h_a);
  cublasDestroy(handle);
  cudaStreamDestroy(stream);
  return 0;
}

同様にtimeコマンドを使って実行する。

$ time ./a.out

real    0m3.329s
user    0m0.056s
sys     0m0.584s

CPUが動いていたのは0.640秒と、上の例とほとんど変わらない。全実行時間が伸びているが、これはおそらくsleep関数が1秒単位であるための切り上げであろう。


この記事へのコメントをお寄せください

  • サイトへの書き込みに差し支えございましたら トータルディスクロージャーサイトサポート係へメールをお送りください
  • トータル・ディスクロージャ・サイトに投稿された文章と画像は、すべてその著作権がHPCシステムズ株式会社に帰属し、HPCシステムズ株式会社が著作権を所有することに同意してください。
  • あなたの文章が他人によって自由に編集、配布されることを望まない場合は、投稿を控えてください。
  • コメントを書き込む場合は名前にひらがなを織り交ぜてください。
  • あなたの投稿する文章と画像はあなた自身によって書かれたものであるか、パブリック・ドメインかそれに類する自由なリソースからの複製であることを約束してください。あなたが著作権を保持していない作品を許諾なしに投稿してはいけません!

<comments hideform="false" />


Comments

ノート:CPU使用率を抑えたCUDA実行

個人用ツール