0

正規分布に従って生成される確率変数を生成する必要がある cuda プログラムを作成しています。確率変数の値を 0 から 8 の間に制限したいので、確率変数をカーネル関数内で生成し、確率変数の結果をさらに使用するために使用します。その目的で cuRAND ライブラリを使用する予定です。curand_normal デバイス API を使用して値を生成しようとしましたが、成功しませんでした。誰かがカーネル関数のコードを提供してくれると非常に助かります。ご協力ありがとうございました。

以下に提供するコードは、 gpu で探しているものの cpu 実装です。

  #include "stdafx.h"
    #include <iostream>
    #include <random>

    using namespace std;
    int _tmain(int argc, _TCHAR* argv[])
    {
        const int nrolls=10000;  // number of experiments
        const int nstars=100;    // maximum number of stars to distribute
        int i;
        default_random_engine generator;
        normal_distribution<double> distribution(0.0,3);


       for (i=0;i<=nstars;i++)
       {   int number = distribution(generator);
           printf("%d\n\n",number);
        }


        return 0;
    }

私が C++ を知らず、他のサイトで見た他のコードに従ってこのプログラムを作成したことを追加したいと思います。ありがとう。

4

1 に答える 1

4

これは、約 0 から 8 の間の離散値を取ることができる、ほぼ「正規」分布の乱数セットを生成するこのコードの適応です。0 から 8 の範囲を持つというコメントの要求がわかりません0の平均で。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <curand_kernel.h>
#include <math.h>
#define SCALE 2.0
#define SHIFT 4.5
#define DISCRETE
#define BLOCKS 1024
#define THREADS 512

#define CUDA_CALL(x) do { if((x) != cudaSuccess) { \
    printf("Error at %s:%d\n",__FILE__,__LINE__); \
    return EXIT_FAILURE;}} while(0)

__global__ void setup_kernel(curandState *state)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    /* Each thread gets different seed, a different sequence
       number, no offset */
    curand_init(7+id, id, 0, &state[id]);
}



__global__ void generate_normal_kernel(curandState *state,
                                int *result)
{
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    float x;
    /* Copy state to local memory for efficiency */
    curandState localState = state[id];
    /* Generate pseudo-random uniforms */
    for(int n = 0; n < 10; n++) {
        x = (curand_normal(&localState) * SCALE)+SHIFT;
        /* Discretize */
#if defined DISCRETE
        x = truncf(x);
#endif
    }
    /* Copy state back to global memory */
    state[id] = localState;
    /* Store last generated result per thread */
    result[id] = (int) x;
}


int main(int argc, char *argv[])
{
    int i;
    unsigned int total;
    curandState *devStates;
    int *devResults, *hostResults;
    int device;
    struct cudaDeviceProp properties;

    CUDA_CALL(cudaGetDevice(&device));
    CUDA_CALL(cudaGetDeviceProperties(&properties,device));


    /* Allocate space for results on host */
    hostResults = (int *)calloc(THREADS * BLOCKS, sizeof(int));

    /* Allocate space for results on device */
    CUDA_CALL(cudaMalloc((void **)&devResults, BLOCKS * THREADS *
              sizeof(int)));
    /* Set results to 0 */
    CUDA_CALL(cudaMemset(devResults, 0, THREADS * BLOCKS *
              sizeof(int)));

    /* Allocate space for prng states on device */
    CUDA_CALL(cudaMalloc((void **)&devStates, THREADS * BLOCKS *
                  sizeof(curandState)));

    /* Setup prng states */
    setup_kernel<<<BLOCKS, THREADS>>>(devStates);


    /* Generate and use uniform pseudo-random  */
    generate_normal_kernel<<<BLOCKS, THREADS>>>(devStates, devResults);

    /* Copy device memory to host */
    CUDA_CALL(cudaMemcpy(hostResults, devResults, BLOCKS * THREADS *
        sizeof(int), cudaMemcpyDeviceToHost));

    /* Show result */
    if (THREADS*BLOCKS > 20){
      printf("First 20 stored results:\n");
      for (i=0; i<20; i++)
        printf("%d\n", hostResults[i]);
      }

    total = 0;
    for(i = 0; i < BLOCKS * THREADS; i++) {
        total += hostResults[i];
    }
    printf("Results mean = %f\n", (total/(1.0*BLOCKS*THREADS)));



    /* Cleanup */
    CUDA_CALL(cudaFree(devStates));
    CUDA_CALL(cudaFree(devResults));
    free(hostResults);
    return EXIT_SUCCESS;
}

このコードを簡単に変更して、(float の) 連続値の正規分布も生成できます。

正規分布の 2 つのパラメーターは、平均と標準偏差です。これらは、SHIFT および SCALE パラメータを使用して表されます。SHIFT は平均をゼロから移動します。SCALE は標準偏差を変更します (1.0 から SCALE が示す値まで)。したがって、SHIFT パラメーターと SCALE パラメーターをいじって、必要な分布を得ることができます。乱数発生器の実数値出力の切り捨ては、統計に影響することに注意してください。これは、SCALE または SHIFT を調整することで調整できます。またはtruncf()、 から丸めのフレーバーに切り替えることもできます。

これを次のようにコンパイルできます。

nvcc -arch=sm_20 -o uniform uniform.cu

cc2.0 以上の GPU を使用していると仮定します。

そうでない場合は、次のようにコンパイルしても問題ありません。

nvcc -o uniform uniform.cu

この場合、double が float に降格されているというコンパイラの警告は無視しても問題ありません。

THREADSBLOCKSマシンの制限内で任意の選択です。これらは、独自のコードの特定の起動構成に合わせて変更できます。

于 2013-01-12T07:40:26.227 に答える