La génération de nombre aléatoire à l'intérieur de Cuda noyau dans une plage variant

Je suis en train de générer de nombres aléatoires de nombres aléatoires à l'intérieur de la cuda noyau. Je souhaite générer des nombres aléatoires à partir d'une distribution uniforme et dans la forme entière, à partir de 1 jusqu'à 8. Les nombres aléatoires serait différent pour chacun des threads. La gamme jusqu'à laquelle nombre aléatoire peut être généré serait également varier d'un thread à l'autre. Le maximum de l'échelle dans un thread peut être aussi bas que 2 ou dans l'autre thread, il peut être élevé que 8, mais pas plus que cela. Donc, je suis en fournissant un exemple ci-dessous de la façon dont je veux que les chiffres générée : 

In thread#1 --> maximum of the range is 2 and so the random number should be between 1 and 2
In thread#2 --> maximum of the range is 6  and so the random number should be between 1 and 6
In thread#3 --> maximum of the range is 5 and so the random number should be between 1 and 5

et ainsi de suite...

InformationsquelleAutor duttasankha | 2013-08-29
  1. 16

    EDIT: j'ai édité ma réponse à corriger certaines des lacunes souligné dans les autres réponses (@tudorturcu) et les commentaires.

    1. Utilisation CURAND pour générer un uniforme
      distribution
      entre 0.0 et 1.0
    2. Puis multipliez ce chiffre par la gamme désirée (la plus grande valeur plus petite
      valeur + 0.999999).
    3. Puis ajouter le décalage (+ la plus petite valeur).
    4. Puis tronquer un nombre entier.

    Quelque chose comme ceci dans votre code de périphérique:

    int idx = threadIdx.x+blockDim.x*blockIdx.x;
//assume have already set up curand and generated state for each thread...
//assume ranges vary by thread index
float myrandf = curand_uniform(&(my_curandstate[idx]));
myrandf *= (max_rand_int[idx] - min_rand_int[idx] + 0.999999);
myrandf += min_rand_int[idx];
int myrand = (int)truncf(myrandf);

    Vous devez:

    #include <math.h>

    pour truncf

    Voici un exemple développé:

    $ cat t527.cu
#include <stdio.h>
#include <curand.h>
#include <curand_kernel.h>
#include <math.h>
#include <assert.h>
#define MIN 2
#define MAX 7
#define ITER 10000000
__global__ void setup_kernel(curandState *state){
int idx = threadIdx.x+blockDim.x*blockIdx.x;
curand_init(1234, idx, 0, &state[idx]);
}
__global__ void generate_kernel(curandState *my_curandstate, const unsigned int n, const unsigned *max_rand_int, const unsigned *min_rand_int,  unsigned int *result){
int idx = threadIdx.x + blockDim.x*blockIdx.x;
int count = 0;
while (count < n){
float myrandf = curand_uniform(my_curandstate+idx);
myrandf *= (max_rand_int[idx] - min_rand_int[idx]+0.999999);
myrandf += min_rand_int[idx];
int myrand = (int)truncf(myrandf);
assert(myrand <= max_rand_int[idx]);
assert(myrand >= min_rand_int[idx]);
result[myrand-min_rand_int[idx]]++;
count++;}
}
int main(){
curandState *d_state;
cudaMalloc(&d_state, sizeof(curandState));
unsigned *d_result, *h_result;
unsigned *d_max_rand_int, *h_max_rand_int, *d_min_rand_int, *h_min_rand_int;
cudaMalloc(&d_result, (MAX-MIN+1) * sizeof(unsigned));
h_result = (unsigned *)malloc((MAX-MIN+1)*sizeof(unsigned));
cudaMalloc(&d_max_rand_int, sizeof(unsigned));
h_max_rand_int = (unsigned *)malloc(sizeof(unsigned));
cudaMalloc(&d_min_rand_int, sizeof(unsigned));
h_min_rand_int = (unsigned *)malloc(sizeof(unsigned));
cudaMemset(d_result, 0, (MAX-MIN+1)*sizeof(unsigned));
setup_kernel<<<1,1>>>(d_state);
*h_max_rand_int = MAX;
*h_min_rand_int = MIN;
cudaMemcpy(d_max_rand_int, h_max_rand_int, sizeof(unsigned), cudaMemcpyHostToDevice);
cudaMemcpy(d_min_rand_int, h_min_rand_int, sizeof(unsigned), cudaMemcpyHostToDevice);
generate_kernel<<<1,1>>>(d_state, ITER, d_max_rand_int, d_min_rand_int, d_result);
cudaMemcpy(h_result, d_result, (MAX-MIN+1) * sizeof(unsigned), cudaMemcpyDeviceToHost);
printf("Bin:    Count: \n");
for (int i = MIN; i <= MAX; i++)
printf("%d    %d\n", i, h_result[i-MIN]);
return 0;
}
$ nvcc -arch=sm_20 -o t527 t527.cu -lcurand
$ cuda-memcheck ./t527
========= CUDA-MEMCHECK
Bin:    Count:
2    1665496
3    1668130
4    1667644
5    1667435
6    1665026
7    1666269
========= ERROR SUMMARY: 0 errors
$
    • Je l'ai probablement fait quelque chose comme cela. Pouvez-vous il suffit de mettre ceci dans le code afin que je puisse comparer les deux. Merci encore.
    InformationsquelleAutor Robert Crovella
  2. 4

    @Robert n'est pas de générer une parfaitement distribution uniforme (bien que tous les numéros de la plage sont générés et tous les nombres générés sont dans la gamme). À la fois la plus petite et la plus grande valeur de 0,5, la probabilité d'être choisi pour le reste de la les numéros de la plage.

    À l'étape 2, vous devez multiplier par le nombre de valeurs dans la gamme: (la plus grande valeur plus petite valeur + 0.999999). *

    À l'étape 3, le décalage doit être (+ la plus petite valeur) au lieu de (+ plus petite valeur de + 0,5).

    Les étapes 1 et 4 restent les mêmes.

    *Comme @Kamil Czerski noté, 1.0 est inclus dans la distribution. L'ajout de 1.0 au lieu de 0.99999 parfois le résultat dans un certain nombre dehors de la plage souhaitée.

    • Notez que 1.0 est inclus dans curand_uniform. Il ya une petite chance que Vous dessinez exactement 1.0 et après multipliant par (largest_value - smallest_value + 1) ajout (smallest_value) et l'arrondi vers zéro, Vous sortir des limites du terrain.Here est ma version de la génération uniforme des entiers, mais l'idée de base est la même que Robert Crovella. J'utilise 0.999999 au lieu de 1 et de faire exactement comme Vous les proposer.
    • Merci d'avoir remarqué cette erreur. Je trouve la décision d'inclure 1.0 dans la distribution et exclure 0.0 assez étrange. Je vais modifier ma réponse à inclure vos modifications. Vous pouvez la modifier pour y inclure votre code d'échantillon.
    InformationsquelleAutor tudorturcu