Comment obtenir un numéro" aléatoire " dans OpenCL

j'ai résolu ce problème "non aléatoire" ces derniers jours et j'ai trouvé trois approches différentes:

1. Xorshift - j'ai créé generator à partir de celui-ci. Tout ce que vous avez à faire est d'en fournir un uint2 nombre (seed) pour le grain entier et chaque élément de travail calculera son propre nombre de rand

   // 'randoms' is uint2 passed to kernel
   uint seed = randoms.x + globalID;
   uint t = seed ^ (seed << 11);  
   uint result = randoms.y ^ (randoms.y >> 19) ^ (t ^ (t >> 8));
   

2. Java aléatoire - j'ai utilisé le code de .next(int bits) méthode pour générer un nombre aléatoire. Cette fois, vous avez pour fournir un ulong nombre de graines.

   // 'randoms' is ulong passed to kernel
   ulong seed = randoms + globalID;
   seed = (seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1);
   uint result = seed >> 16;
   

3. il suffit de générer tout sur CPU et de le passer au noyau dans un gros tampon.

j'ai testé les trois approches (générateurs) dans mon algorithme d'évolution calculant Minimum dominant ensemble dans les graphiques.

j'aime les nombres générés par le premier, mais il semble que mon algorithme d'évolution ne le fasse pas.

le deuxième générateur génère des nombres qui ont un motif visible mais mon algorithme d'évolution l'aime comme ça de toute façon et tout fonctionne un peu plus vite qu'avec le premier générateur.

mais la troisième approche montre qu'il est absolument correct de simplement fournir tous les numéros de l'hôte (cpu). D'abord j'ai pensé que générer (dans mon cas) 1536 numéros int32 et les passer à GPU dans chaque appel du noyau serait trop cher (pour calculer et transférer à GPU). Mais il s'avère, il est aussi rapide que mes précédentes tentatives. Et la charge CPU reste sous 5%.

BTW, j'ai aussi essayé MWC64X Random mais après avoir installé le nouveau pilote GPU la fonction mul_hi commence à causer l'échec de la compilation (même L'Analyseur de noyau AMD entier s'est écrasé).

voici l'algorithme utilisé par java.util.Random classe selon le doc:

(seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1)

voir la documentation pour ses différentes implémentations. Passer l'identification de l'ouvrier pour la graine et la boucle un peu de temps devrait produire le hasard décent

ou un autre metod serait d'avoir certains des opérations aléatoires se produisent qui sont assez ceratain à débordement:

 long rand= yid*xid*as_float(xid-yid*xid);
 rand*=rand<<32^rand<<16|rand;
 rand*=rand+as_double(rand);

xid=get_global_id(0); et yid= get_global_id(1);

j'implémente actuellement un traceur de chemin en temps réel. Vous savez peut-être déjà que le traçage d'un chemin nécessite de nombreux nombres aléatoires.

Avant de générer des nombres aléatoires sur le GPU, je les ai simplement générés sur le CPU (en utilisant rand (), ce qui craint) et les ai passés au GPU.

Qui est rapidement devenu un goulot d'étranglement.

Maintenant, je générais les nombres aléatoires sur le GPU avec le Générateur de nombres Pseudorandom Park-Miller (PRNG).

Il est extrêmement simple à mettre en œuvre et réalise de très bons résultats.

J'ai pris des milliers d'échantillons (dans la plage de 0,0 à 1,0) et en moyenne ensemble.

La valeur résultante était très proche de 0.5 (ce qui est normal). Entre les différents passages, la divergence par rapport à 0.5 était d'environ 0.002. Il a donc une distribution très uniforme.



Voici un papier décrivant le algorithm:

http://www.cems.uwe.ac.uk/~irjohnso/coursenotes/ufeen8-15-m/p1192-parkmiller.pdf

Et voici un article sur l'algorithme ci-dessus optimisé pour CUDA (qui peut facilement être porté sur OpenCL): http://www0.cs.ucl.ac.uk/staff/ucacbbl/ftp/papers/langdon_2009_CIGPU.pdf



Voici un exemple de comment je l'utilise:

int rand(int* seed) // 1 <= *seed < m
{
    int const a = 16807; //ie 7**5
    int const m = 2147483647; //ie 2**31-1

    *seed = (long(*seed * a))%m;
    return(*seed);
}

kernel random_number_kernel(global int* seed_memory)
{
    int global_id = get_global_id(1) * get_global_size(0) + get_global_id(0); // Get the global id in 1D.

    // Since the Park-Miller PRNG generates a SEQUENCE of random numbers
    // we have to keep track of the previous random number, because the next
    // random number will be generated using the previous one.
    int seed = seed_memory[global_id];

    int random_number = rand(&seed); // Generate the next random number in the sequence.

    seed_memory[global_id] = *seed; // Save the seed for the next time this kernel gets enqueued.
}

le code sert d'exemple. Je n'ai pas tester.

Le tableau "seed_memory" est rempli avec rand() qu'une seule fois avant la première exécution du noyau. Après cela, toute la génération de nombres aléatoires se produit sur le GPU. Je pense qu'il est également possible d'utiliser simplement l'id du noyau au lieu d'initialiser le tableau avec rand().

la fermeture de la boucle sur un AMD devgurus de la discussion [lien mort] et annonce:

https://github.com/clMathLibraries/clRNG

il semble Qu'OpenCL n'offre pas une telle fonctionnalité. Cependant, certaines personnes ont fait des recherches sur cela et fournir le code sous licence BSD pour produire de bons nombres aléatoires sur GPU.

les GPU n'ont pas de bonnes sources d'aléatoire, mais cela peut être facilement surmonté en semant un grain avec une graine aléatoire de l'hôte. Après cela, vous avez juste besoin d'un algorithme qui peut fonctionner avec un grand nombre de threads simultanés.

ce lien décrit une implémentation Mersenne Twister utilisant OpenCL: Parallèle Mersenne Twister. Vous pouvez également trouver une implémentation dans le SDK NVIDIA.

j'ai eu le même problème. www.thesalmons.org/john/random123/papers/random123sc11.pdf

vous pouvez trouver la documentation ici. http://www.thesalmons.org/john/random123/releases/latest/docs/index.html

Vous pouvez télécharger la bibliothèque ici: http://www.deshawresearch.com/resources_random123.html

Ceci est ma version de bruit pseudo-aléatoire de flotteur OpenCL, en utilisant la fonction trigonométrique

//noise values in range if 0.0 to 1.0
static float noise3D(float x, float y, float z) {
    float ptr = 0.0f;
    return fract(sin(x*112.9898f + y*179.233f + z*237.212f) * 43758.5453f, &ptr);
}

__kernel void fillRandom(float seed, __global float* buffer, int length) {
    int gi = get_global_id(0);
    float fgi = float(gi)/length;
    buffer[gi] = noise3D(fgi, 0.0f, seed);
}

Vous pouvez générer 1D ou 2D noize en passant à noise3D normalisé index les coordonnées comme premier paramètre, et la graine aléatoire (générée sur CPU par exemple) comme dernier paramètre.

voici quelques images de bruit générées avec ce noyau et différentes graines:

pourquoi pas? vous pouvez simplement écrire un noyau qui génère des nombres aléatoires, difficile qui aurait besoin de plus d'appels du noyau et éventuellement passer les nombres aléatoires comme argument à votre autre noyau qui en a besoin

vous ne pouvez pas générer de nombres aléatoires dans le noyau, la meilleure option est de générer le nombre aléatoire dans l'hôte (CPU) et de le transférer dans le GPU à travers des tampons et de l'utiliser dans le noyau.