0

curand などを使用して、各スレッドで同じになる 0 から 49k までの疑似乱数を 1 つ生成する最良の方法は何かと考えていました。

一度に 1 つ生成する必要がありますが、約 10,000 回生成する必要があるため、カーネル内で乱数を生成することを好みます。

また、0.0 から 1.0 の間の浮動小数点数を使用することもできますが、ほとんどの投稿と例ではスレッドごとに異なる PRN を使用する方法が示されているため、PRN をすべてのスレッドで使用できるようにする方法がわかりません。

ありがとう

4

1 に答える 1

8

おそらく、 curand のドキュメント、特にデバイス APIを勉強する必要があるだけでしょう。各スレッドで同じシーケンスを取得するための鍵は、各スレッドの状態を作成し (ほとんどの例ではこれを行います)、同じシーケンス番号を各スレッドの init 関数に渡すことです。curand_initでは、パラメーターのシーケンスは次のとおりです。

curand_init(seed, subsequence number, offset, state)

各 init 呼び出しのシードを同じに設定することで、各スレッドに同じシーケンスを生成します。サブシーケンス番号とオフセット番号を同じに設定することで、スレッドごとに、そのシーケンス内で同じ開始値を選択します。

デモ用のコードは次のとおりです。

// compile with: nvcc -arch=sm_20 -lcurand -o t89 t89.cu
#include <stdio.h>
#include <curand.h>
#include <curand_kernel.h>

#define SCALE 49000
#define DSIZE 5000
#define nTPB 256

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__device__ float getnextrand(curandState *state){

  return (float)(curand_uniform(state));
}

__device__ int getnextrandscaled(curandState *state, int scale){

  return (int) scale * getnextrand(state);
}


__global__ void initCurand(curandState *state, unsigned long seed){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    curand_init(seed, 0, 0, &state[idx]);
}

__global__ void testrand(curandState *state, int *a1, int *a2){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    a1[idx] = getnextrandscaled(&state[idx], SCALE);
    a2[idx] = getnextrandscaled(&state[idx], SCALE);
}

int main() {

    int *h_a1, *h_a2, *d_a1, *d_a2;
    curandState *devState;

    h_a1 = (int *)malloc(DSIZE*sizeof(int));
    if (h_a1 == 0) {printf("malloc fail\n"); return 1;}
    h_a2 = (int *)malloc(DSIZE*sizeof(int));
    if (h_a2 == 0) {printf("malloc fail\n"); return 1;}
    cudaMalloc((void**)&d_a1, DSIZE * sizeof(int));
    cudaMalloc((void**)&d_a2, DSIZE * sizeof(int));
    cudaMalloc((void**)&devState, DSIZE * sizeof(curandState));
    cudaCheckErrors("cudamalloc");



     initCurand<<<(DSIZE+nTPB-1)/nTPB,nTPB>>>(devState, 1);
     cudaDeviceSynchronize();
     cudaCheckErrors("kernels1");
     testrand<<<(DSIZE+nTPB-1)/nTPB,nTPB>>>(devState, d_a1, d_a2);
     cudaDeviceSynchronize();
     cudaCheckErrors("kernels2");
     cudaMemcpy(h_a1, d_a1, DSIZE*sizeof(int), cudaMemcpyDeviceToHost);
     cudaMemcpy(h_a2, d_a2, DSIZE*sizeof(int), cudaMemcpyDeviceToHost);
     cudaCheckErrors("cudamemcpy");
     printf("1st returned random value is %d\n", h_a1[0]);
     printf("2nd returned random value is %d\n", h_a2[0]);

     for (int i=1; i< DSIZE; i++){
       if (h_a1[i] != h_a1[0]) {
         printf("mismatch on 1st value at %d, val = %d\n", i, h_a1[i]);
         return 1;
         }
       if (h_a2[i] != h_a2[0]) {
         printf("mismatch on 2nd value at %d, val = %d\n", i, h_a2[i]);
         return 1;
         }
       }
     printf("thread values match!\n");

}
于 2013-03-06T16:03:24.107 に答える