CURANDライブラリを使う [Host API]
CUDA Toolkit 3.2から,CURANDという疑似乱数ジェネレータライブラリが導入されました.GPUによる乱数の生成を行えるライブラリです.
CURANDは,2つのAPIカテゴリから構成されています.
- Host API
- ホストからの使用を目的としたAPI群.Device APIの上に構築されている.
- Device API
- デバイスからの使用を目的としたAPI群.要するに__device__修飾された関数群.
今回のエントリではHost API+Pseudo Random Generatorの組み合わせのみを扱います.Quasirandom Generatorも作れるそうですが,乱数について詳しいわけではないので割愛します.
Host APIの流れは,
- 初期化(curandCreateGenerator)
- Seed設定(curandSetPseudoRandomGeneratorSeed)
- 乱数列生成(curandGenerate*)
- 解放(curandDestroyGenerator)
となっています.CPU上でのライブラリでも似たような手順を踏む作りが多いと思います.
生成関数(curandGenerate*)は,いくつかの種類があるので,詳しくは公式ドキュメントを参照してください.今回の例で使用しているcurandGenerateUniformは0から1の間の乱数列を生成します.
なお,疑似乱数列は,生成を分割しても同じ乱数列が生成されます.1万個の乱数を得るのに1000個の生成を10回しても,10個の生成を1000回しても,結果として得られる乱数列は同じものとなります.ただ,手元の環境では生成関数の呼び出し関数を少なくした方が高速でした.
以下,サンプルコード.エラー処理の類は省略しています.
#include <iostream> #include <iterator> #include <curand.h> int main() { curandGenerator_t g; cudaSetDeviceFlags(cudaDeviceMapHost); curandCreateGenerator(&g, CURAND_RNG_PSEUDO_DEFAULT); curandSetPseudoRandomGeneratorSeed(g, 0); int const n = 100; float* p; float* dp; cudaHostAlloc(&p, n * sizeof(float), cudaHostAllocMapped); cudaHostGetDevicePointer(&dp, p, 0); curandGenerateUniform(g, dp, n); cudaThreadSynchronize(); std::copy(p, p + n, std::ostream_iterator<float>(std::cout, "\n")); curandDestroyGenerator(g); cudaFreeHost(p); return 0; }
CURANDライブラリを使う [Device API]
CUDA Toolkit 3.2から,CURANDという疑似乱数ジェネレータライブラリが導入されました.GPUによる乱数の生成を行えるライブラリです.
CURANDは,2つのAPIカテゴリから構成されています.
- Host API
- ホストからの使用を目的としたAPI群.Device APIの上に構築されている.
- Device API
- デバイスからの使用を目的としたAPI群.要するに__device__修飾された関数群.
Host APIについてはCURANDライブラリを使う [Host API] - fjnlの生存記録のような何かを参照してください.
以下,サンプルコード.モンテカルロ法による円周率を求めるプログラムです.エラー処理の類は省略しています.
// -*- c++ -*- #include <numeric> #include <iostream> #include <curand_kernel.h> __global__ void monte_pi(unsigned long seed, int n_try, float* out) { int id = blockIdx.x * blockDim.x + threadIdx.x; curandState s; curand_init(seed, id, 0, &s); int inside = 0; for (int i = 0; i < n_try; ++i) { float x = curand_uniform(&s); float y = curand_uniform(&s); if (x * x + y * y < 1) { inside++; } } out[id] = inside / static_cast<float>(n_try); } int main() { cudaSetDeviceFlags(cudaDeviceMapHost); int const n_thread = 256; int const n_block = 256; int const n = n_thread * n_block; int const n_try = 1000; float* out; float* d_out; cudaHostAlloc(&out, n * sizeof(float), cudaHostAllocMapped); cudaHostGetDevicePointer(&d_out, out, 0); monte_pi<<<n_block, n_thread>>>(0, n_try, d_out); cudaThreadSynchronize(); std::cout << 4.0 * std::accumulate(out, out + n, 0.0) / n << std::endl; cudaFreeHost(out); return 0; }