トップ > CUDA

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

関連エントリ

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

CUDA

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

関連エントリ

GF100のキャッシュの性能

CUDA

Motivation

GTX480に代表されるGF100およびGTX460などのGF104にはL1, L2キャッシュが搭載され、メモリアクセスに関する制限が緩和されました。何も考えずに書くとGT200系では性能が出なかった場合でも、十分な性能が発揮されることがあります。

そこで、定数がテーブルに入っていて、それを参照して計算するというコードを例に、キャッシュの性能を測定します。

続きを読む

CUDAでFunctorを使う 捕捉

CUDA

CUDAでFunctorを使う - fjnlの生存記録のような何かに対する捕捉です。

当初、CUDAはデバイスコードへのパラメータをshared memory経由で渡していました。しかし、Compute Capability 2.0*1からはconstant memory経由で渡されるように変更されています。shared memoryに比べればconstant memoryは容量に余裕がある事が多いため、これは嬉しい変更であると思います。

さて、メンバのないfunctorは、以下のようにPTX上では1バイトの変数として扱われています。Compute Capability 2.0以前向けのコードでshared memoryをカツカツに使っている場合には、functorを使いたくても使えないという場合があるかもしれません。(レアケースすぎるかなぁ…)

.param .align 1 .b8 __cudaparm__Z22element_wise_operationI3mulEvPdPKdS3_T__f[1]

なお、あるデバイス関数がどれだけのメモリを使っているかは、nvccのコンパイルオプションに-Xptxas -vを付けるとコンパイル時に表示されます。

サンプルコードを-arch=sm_13と-arch=sm_20でコンパイルした時の出力は以下です。
sm_20からは、パラメータを渡す時にshared memoryではなくconstant memoryが使われていることがわかります。

$ nvcc func.cu -arch=sm_20 -Xptxas -v
ptxas info    : Compiling entry function '_Z22element_wise_operationI3mulEvPdPKdS3_T_' for 'sm_20'
ptxas info    : Used 9 registers, 57 bytes cmem[0]
ptxas info    : Compiling entry function '_Z22element_wise_operationI3addEvPdPKdS3_T_' for 'sm_20'
ptxas info    : Used 9 registers, 57 bytes cmem[0]

$ nvcc func.cu -arch=sm_13 -Xptxas -v
ptxas info    : Compiling entry function '_Z22element_wise_operationI3mulEvPdPKdS3_T_' for 'sm_13'
ptxas info    : Used 5 registers, 25+16 bytes smem
ptxas info    : Compiling entry function '_Z22element_wise_operationI3addEvPdPKdS3_T_' for 'sm_13'
ptxas info    : Used 5 registers, 25+16 bytes smem

*1:GTX470やGTX480など

CUDAでFunctorを使う

CUDA

CUDAの.cuファイルはC++として解釈されるので、templateといった記法やfunctorといったデザインをデバイスコード上でも使うことができます。nvcc 3.0で動作を確認しました。もう2.xの環境が手元になかったので、2.xで使えるかはわかりません。

struct add {
    template <class T>
    __device__
    T operator ()(T const a, T const b) const {
        return a + b;
    }
};

struct mul {
    template <class T>
    __device__
    T operator ()(T const a, T const b) const { 
        return a * b;
    }
};

template <class F>
__global__
void element_wise_operation(double* out, double const* a, double const* b, F f) {
    size_t const i = blockIdx.x * blockDim.x + threadIdx.x;
    out[i] = f(a[i], b[i]);
}

サンプルとして、このようなfunctor classと、それを利用するコードを書きました。を定義しました。operator ()を__device__修飾することにより、__device__や__global__のコードから呼ぶことができるようになります。

呼び出す時は、このようにホスト側でfunctorを生成して渡すことができます。

element_wise_operation<<<1, 10>>>(out, a, b, add());
element_wise_operation<<<1, 10>>>(out, a, b, mul());

また、このように書くことで、パフォーマンスが心配になりますが、functorの呼び出しはインライン展開されるため、普通に書き下した場合と同等のPTXコードが生成されるようです。