Est-ce que CUDA CURAND est susceptible aux conflits de données?

Question

Je génère 1 bloc de 256 threads à partir de mon noyau Setup() pour configurer un tableau RNGstates avec 256 états CURAND :

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

Maintenant, je génère 1000 blocs de 256 threads à partir de mon noyau Generate() pour remplir le tableau résultat avec 256 000 nombres aléatoires. Cependant, je le fais en n'utilisant que les 256 états de RNGstates, de sorte que chaque état sera accédé par 1000 threads (un de chaque bloc) :

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

Je sais que l'appel à curand_uniform() met à jour les états d'une manière ou d'une autre, donc je suppose qu'une opération d'écriture a lieu.

Devrais-je m'inquiéter des conflits de données qui se produisent lorsque les 1000 threads mappés à chacun des 256 états CURAND tentent de mettre à jour l'état implicitement via curand_uniform()? Cela impactera-t-il la qualité de mes nombres aléatoires (par exemple, obtenir des valeurs dupliquées fréquemment)?

Merci beaucoup.

Answer 1

Je pense que le partage des états aura certainement un impact sur la qualité. Les valeurs dupliquées sont la meilleure situation pour partager les états. Les conflits de données pourraient totalement ruiner les états.

Vous pourriez garder un état pour chacun de vos threads.

Lors de l'utilisation de 1000 blocs, 256 000 états sont nécessaires pour votre cas. Le code devrait ressembler à

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

et

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

Pour réduire les besoins en mémoire pour plusieurs blocs, vous pourriez limiter votre #bloc à un petit nombre, et générer plusieurs nombres aléatoires par thread, au lieu de 1 nombre aléatoire par thread.

__global__ void generate_uniform_kernel(curandState *state, 
                                unsigned int *result)
{
    int id = threadIdx.x + blockIdx.x * 64;
    unsigned int count = 0;
    float x;
    /* Copier l'état dans la mémoire locale pour plus d'efficacité */
    curandState localState = state[id];
    /* Générer des nombres pseudo-aléatoires uniformes */
    for(int n = 0; n < 10000; n++) {
        x = curand_uniform(&localState);
        /* Vérifier si > .5 */
        if(x > .5) {
            count++;
        }
    }
    /* Copier l'état dans la mémoire globale */
    state[id] = localState;
    /* Stocker les résultats */
    result[id] += count;
}

Consultez la section Exemples d'API pour les périphériques dans le manuel de référence de cuRAND pour des exemples complets sur la façon de gérer plusieurs blocs.

Answer 2

Vous pouvez également utiliser curandStateMtgp32_t là où vous n'avez besoin que d'un par bloc (si les blocs ont au maximum 256 threads chacun) http://docs.nvidia.com/cuda/curand/device-api-overview.html#bit-generation-1