1

カーネルから256 個のスレッドの 1 ブロックを生成して、256 個の CURAND 状態を持つ配列をセットアップします。Setup()RNGstates

__global__ void Setup(curandState *RNGstates, long seed) {
    int tid = threadIdx.x;
    curand_init(seed, tid, 0, &RNGstates[tid]);
}

ここで、カーネルから256 スレッドの 1000 ブロックを生成して、配列を 256,000 の乱数で埋めます。ただし、各状態が 1000 スレッド (各ブロックから 1 つ) によってアクセスされるように、 の 256 状態のみを使用してそうします。Generate()resultRNGstates

__global__ void Generate(curandState *RNGstates, float *result) {
    int tid = blockIdx.x*blockDim.x + threadIdx.x;
    float rnd = curand_uniform(&RNGstates[threadIdx.x]);
    result[tid] = rnd;
}

呼び出すcurand_uniform()と何らかの形で状態が更新されることはわかっているので、何らかの書き込み操作が行われていると思います。

では、256 の CURAND 状態のそれぞれにマップされた 1000 のスレッドが を介して暗黙的に状態を更新しようとするときに発生するデータ競合について心配する必要がありますcurand_uniform()か? これは乱数の品質に影響を与えますか (例: 値が頻繁に重複する)?

どうもありがとう。

4

2 に答える 2

3

共有状態は間違いなく品質に影響すると思います。重複する値は、状態を共有するための最良の状況です。データ競争は州を完全に滅ぼす可能性があります。

スレッドごとに 1 つの状態を維持できます。

1000 ブロックを使用する場合、ケースには 256,000 ステートが必要です。コードは次のようになります

__global__ void Setup(curandState *RNGstates, long seed) {
  int tid = blockIdx.x*blockDim.x + threadIdx.x;
  curand_init(seed, tid, 0, &RNGstates[tid]);
}

と

__global__ void Generate(curandState *RNGstates, float *result) {
  int tid = blockIdx.x*blockDim.x + threadIdx.x;
  float rnd = curand_uniform(&RNGstates[tid]);
  result[tid] = rnd;
}

複数のブロックのメモリ要件を減らすには、#block を少数に制限し、スレッドごとに 1 つの乱数ではなく、スレッドごとに複数の乱数を生成することができます。

__global__ void generate_uniform_kernel(curandState *state, 
                                unsigned int *result)
{
    int id = threadIdx.x + blockIdx.x * 64;
    unsigned int count = 0;
    float x;
    /* Copy state to local memory for efficiency */
    curandState localState = state[id];
    /* Generate pseudo-random uniforms */
    for(int n = 0; n < 10000; n++) {
        x = curand_uniform(&localState);
        /* Check if > .5 */
        if(x > .5) {
            count++;
        }
    }
    /* Copy state back to global memory */
    state[id] = localState;
    /* Store results */
    result[id] += count;
}

複数のブロックを処理する方法の完全な例については、cuRAND ref マニュアルのデバイス API の例のセクションを参照してください。

于 2013-01-16T07:03:18.073 に答える
2

また、ブロックごとに 1 つだけ必要な curandStateMtgp32_t を使用することもできます (ブロックがそれぞれ最大 256 スレッドの場合) http://docs.nvidia.com/cuda/curand/device-api-overview.html#bit-generation-1

于 2013-08-22T11:28:10.633 に答える