CUDA 5.5の新機能(3): Philox法CURAND

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

1年半前のSC11でBest Paper Awardを受賞した、GPGPU上で特に高速に動く乱数生成アルゴリズム、Philox。CUDA 5.5で、このアルゴリズムがCURANDについに実装された。

目次

テストプログラム

#include <cstdio>

#include <sys/time.h>

#include <curand.h>
#include <thrust/device_ptr.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/reduce.h>

struct func : thrust::unary_function<int, int> {
  double *d_a;
  func(double *d_a_) : d_a(d_a_)
  {}
  __device__ int
  operator()(int i) const
  {
    double x = d_a[i * 2 + 0];
    double y = d_a[i * 2 + 1];
    double r2 = x * x + y * y;
    return (r2 <= 1.0) ? 1 : 0;
  }
};

void
benchmark(const size_t &n, double *d_a, const curandRngType_t &rngtype,
          const char *name)
{
  struct timeval tv0, tv1;
  gettimeofday(&tv0, NULL);
  curandGenerator_t generator;
  curandCreateGenerator(&generator, rngtype);
  curandGenerateUniformDouble(generator, d_a, n * 2);
  thrust::counting_iterator<int> count(0);
  thrust::transform_iterator<func, thrust::counting_iterator<int> >
    trans(count, func(d_a));
  int sum = thrust::reduce(trans, trans + n);
  double pi = 4.0 * double(sum) / double(n);
  curandDestroyGenerator(generator);
  gettimeofday(&tv1, NULL);
  double error = std::abs(pi - M_PI) / M_PI;
  int sec = tv1.tv_sec - tv0.tv_sec;
  int usec = tv1.tv_usec - tv0.tv_usec;
  if (usec < 0) {
    --sec;
    usec += 1000000;
  }
  printf("%s (%d Points)\nPI:\t%.70f\nError:\t%.70f\nTime:\t%d.%06d\n",
         name, n, pi, error, sec, usec);
}

int
main()
{
  const size_t n = 150000000;

  double *d_a;
  cudaMalloc((void **)&d_a, n * 2 * sizeof(double));
  benchmark(n, d_a, CURAND_RNG_PSEUDO_PHILOX4_32_10, "Philox4x32-10");
  benchmark(n, d_a, CURAND_RNG_PSEUDO_PHILOX4_32_10, "Philox4x32-10");
  benchmark(n, d_a, CURAND_RNG_PSEUDO_XORWOW, "XORWOW");
  benchmark(n, d_a, CURAND_RNG_PSEUDO_MRG32K3A, "MRG32k3a");
  benchmark(n, d_a, CURAND_RNG_PSEUDO_MTGP32, "Mersenne Twister");
  benchmark(n, d_a, CURAND_RNG_QUASI_SOBOL64, "Sobol64");
  benchmark(n, d_a, CURAND_RNG_QUASI_SCRAMBLED_SOBOL64, "Scrambled Sobol64");
  cudaFree(d_a);
  return 0;
}

CUDA 4.1のころの検証記事のプログラムにCURAND_RNG_PSEUDO_PHILOX4_32_10を使った行を追加(初期化オーバーヘッドを考慮して、同じ行を2行追加している)。なお、前回の実験環境がTesla M2090だったのに対して、今回はメモリの少ないTesla C2050を使っているので、それに合わせて問題サイズnも小さくなっている。

実験環境

CPU: Quad-Core Xeon W3565 3.20 GHz
GPU: Tesla C2050 × 2
OS: CentOS 6.4
CUDA: CUDA 5.5 RC

実行結果

Philox4x32-10 (150000000 Points)
PI:	3.1415069600000000704653757566120475530624389648437500000000000000000000
Error:	0.0000272771168137047695400644597985007067109108902513980865478515625000
Time:	0.097795
Philox4x32-10 (150000000 Points)
PI:	3.1415069600000000704653757566120475530624389648437500000000000000000000
Error:	0.0000272771168137047695400644597985007067109108902513980865478515625000
Time:	0.090189
XORWOW (150000000 Points)
PI:	3.1416632799999999470230704901041463017463684082031250000000000000000000
Error:	0.0000224810845945060966963920806582777345283830072730779647827148437500
Time:	0.069279
MRG32k3a (150000000 Points)
PI:	3.1415478399999998693203906441340222954750061035156250000000000000000000
Error:	0.0000142646086665754331744950314520004042151413159444928169250488281250
Time:	0.083944
Mersenne Twister (150000000 Points)
PI:	3.1415007466666664548426979308715090155601501464843750000000000000000000
Error:	0.0000292548822399499132483439989282913984425249509513378143310546875000
Time:	0.123347
Sobol64 (150000000 Points)
PI:	3.2915026400000000350587470165919512510299682617187500000000000000000000
Error:	0.0477178307120465702850786726685328176245093345642089843750000000000000
Time:	0.084097
Scrambled Sobol64 (150000000 Points)
PI:	3.9815186666666666503999749693321064114570617675781250000000000000000000
Error:	0.2673567536253046994509929845662554726004600524902343750000000000000000
Time:	0.084249

精度に関しては問題になるような点は見当たらない。性能面ではさすがにMRG32k3aには劣るが、あちらは最初から浮動小数点数を生成するアルゴリズム。Philoxのように生成の対象が整数のアルゴリズムは、浮動小数点数への変換処理も加わって不利になっている可能性は有る。それを差し引いても、元論文では僅差で勝っているはずのXORWOWに対して大敗しているのは気になる点である。

ちなみに、このプログラムを元論文と同じGeForce GTX 580で走らせるとこうなる。

Philox4x32-10 (75000000 Points)
PI:     3.1413323200000000667841959511861205101013183593750000000000000000000000
Error:  0.0000828667553368431507613764419595270283025456592440605163574218750000
Time:   0.020651
XORWOW (75000000 Points)
PI:     3.1416579733333334090161770291160792112350463867187500000000000000000000
Error:  0.0000207919201318650673443802490147902517492184415459632873535156250000
Time:   0.025295

こちらはきちんとPhiloxのほうが速いのであるが、GF100→GF110の間のアーキテクチャのマイナーチェンジがPhiloxに有利に働いているのか、それとも別の要因が在るのか、今後の課題となるであろう。

まとめ

新アルゴリズムPhiloxは、環境によっては従来手法に比べて高い性能を発揮する。事前にきちんとベンチマークを採り、性能の優位を確認した上で使いこなせば、かなり強力な機能となるだろう。


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

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

<comments hideform="false" />


Comments

ノート:CUDA 5.5の新機能(3): Philox法CURAND

個人用ツール