CURANDライブラリを使う [Device API]

CUDA

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

関連エントリ