CURAND et les kernelx, où générer?

Ma motivation: J’utilise un algorithme pour modéliser la dynamic des populations et je souhaite utiliser CUDA afin de pouvoir prendre en compte un grand nombre de nœuds dans les simulations numériques. Bien que ce soit la première fois que j’utilise du code sur un processeur graphique, les résultats semblent prometteurs jusqu’à présent.

Contexte: Je dois prendre en compte le bruit stochastique, qui joue un rôle crucial dans l’évolution des systèmes complexes que je souhaite étudier. Autant que je sache, la génération de nombres aléatoires dans CUDA peut être assez gênante comparée à une opération similaire sur un processeur. Dans la documentation, je vois qu’il faut stocker les états du RNG et continuer à en informer le kernel (fonction globale) qui doit (générer et) utiliser les nombres aléatoires. J’ai trouvé ces exemples assez instructifs. Peut-être y a-t-il autre chose que vous me recommandiez de lire à ce sujet?

Question: Quel est l’avantage de générer n valeurs de départ, de les stocker dans un tableau de la mémoire globale du périphérique, puis de les alimenter au kernel, ce qui génère deux nombres aléatoires à utiliser, par opposition à la génération de 2n nombres aléatoires, les dans la mémoire globale du périphérique, puis les alimentant directement au kernel qui doit les utiliser? Il me manque quelque chose de vraiment important ici, car il me semble que l’on économiserait des ressources dans le deuxième cas (qui n’est jamais utilisé dans des exemples). Il semble également que l’on serait plus sûr en ce qui concerne la dissortingbution des nombres générés.

Mon code est plutôt long mais j’ai essayé de faire un petit exemple de ce dont j’ai besoin. C’est ici:

Mon code:

#include  #include  #include  #include  #include  __global__ void update (int n, float *A, float *B, float p, float q, float *rand){ int idx = blockIdx.x*blockDim.x + threadIdx.x; int n_max=n*n; int i, j; i=idx/n; //col j=idx-i*n; //row float status; //A, B symmesortingc //diagonal untouched, only need 2 random numbers per thread //ie n*(n-1) random numbers in total int idx_rand = (2*n-1-i)*i/2+j-1-i; if(idxi){ if(rand[idx_rand]<p){ status=A[idx]; if(status==1){ if(rand[idx_rand+n*(n-1)/2] < q){ B[idx]=-1.0f; B[i+n*j]=-1.0f; } } else if(status==0){ if(rand[idx_rand+n*(n-1)/2] < q){ B[idx]=1.0f; B[i+n*j]=1.0f; } } } } } __global__ void fill(float *A, int n, float num){ int idx = blockIdx.x*blockDim.x + threadIdx.x; if(idx<n){ A[idx]=num; } } void swap(float** a, float** b) { float* temp = *a; *a = *b; *b = temp; } int main(int argc, char* argv[]){ int t, n, t_max, seed; seed = atoi(argv[1]); n = atoi(argv[2]); t_max = atoi(argv[3]); int blockSize = 256; int nBlocks = n*n/blockSize + ((n*n)%blockSize == 0?0:1); curandGenerator_t prng; curandCreateGenerator(&prng, CURAND_RNG_PSEUDO_DEFAULT); curandSetPseudoRandomGeneratorSeed(prng, (unsigned long long) seed); float *h_A = (float *)malloc(n * n * sizeof(float)); float *h_B = (float *)malloc(n * n * sizeof(float)); float *d_A, *d_B, *d_rand; cudaMalloc(&d_A, n * n * sizeof(float)); cudaMalloc(&d_B, n * n * sizeof(float)); cudaMalloc(&d_rand, n * (n-1) * sizeof(float)); fill <<>> (d_A, n*n, 0.0f); fill <<>> (d_B, n*n, 0.0f); for(t=1; t<t_max+1; t++){ //generate random numbers curandGenerateUniform(prng, d_rand, n*(n-1)); //update B update <<>> (n, d_A, d_B, 0.5f, 0.5f, d_rand); //do more stuff swap(&d_A, &d_B); } cudaMemcpy(h_A, d_A, n*n*sizeof(float),cudaMemcpyDeviceToHost); //print stuff curandDestroyGenerator(prng); cudaFree(d_A); cudaFree(d_B); cudaFree(d_rand); free(h_A); free(h_B); return 0; } 

J’aimerais que vous me disiez ce qui ne va pas (et quelques astuces pour le réparer). Si les experts pouvaient me dire ce que je pouvais espérer économiser (en temps d’exécution) dans le meilleur des cas, après tous les ajustements de performance auxquels ils pourraient penser, ce serait formidable, car plusieurs tâches sont actuellement à ma charge et mon coût est élevé. Le bénéfice en termes de “temps d’étude” est donc très important.

Et ça y est, merci pour la lecture!

Pour mémoire, mes spécifications matérielles sont ci-dessous. J’ai toutefois l’intention d’utiliser Amazon EC2 pour cela à un moment donné.

Mon matériel (actuel):

  Device 0: "GeForce 8800 GTX" CUDA Driver Version / Runtime Version 5.5 / 5.5 CUDA Capability Major/Minor version number: 1.0 Total amount of global memory: 768 MBytes (804978688 bytes) (16) Multiprocessors, ( 8) CUDA Cores/MP: 128 CUDA Cores GPU Clock rate: 1350 MHz (1.35 GHz) Memory Clock rate: 900 Mhz Memory Bus Width: 384-bit Maximum Texture Dimension Size (x,y,z) 1D=(8192), 2D=(65536, 32768), 3D=(2048, 2048, 2048) Maximum Layered 1D Texture Size, (num) layers 1D=(8192), 512 layers Maximum Layered 2D Texture Size, (num) layers 2D=(8192, 8192), 512 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 16384 bytes Total number of registers available per block: 8192 Warp size: 32 Maximum number of threads per multiprocessor: 768 Maximum number of threads per block: 512 Max dimension size of a thread block (x,y,z): (512, 512, 64) Max dimension size of a grid size (x,y,z): (65535, 65535, 1) Maximum memory pitch: 2147483647 bytes Texture alignment: 256 bytes Concurrent copy and kernel execution: No with 0 copy engine(s) Run time limit on kernels: Yes Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: No Alignment requirement for Surfaces: Yes Device has ECC support: Disabled Device supports Unified Addressing (UVA): No Device PCI Bus ID / PCI location ID: 7 / 0 

En général, la génération de nombres aléatoires est un processus qui peut être mis en parallèle sur le GPU et peut donc tirer parti du GPU pour générer des nombres plus rapidement, dans de nombreux cas, qu’ils ne peuvent être générés sur le CPU. C’est la principale motivation pour utiliser une API / bibliothèque comme CURAND.

Quel est l’avantage de générer n valeurs de départ, de les stocker dans un tableau de la mémoire globale du périphérique, puis de les alimenter vers le kernel, ce qui génère deux nombres aléatoires à utiliser, par opposition à la génération de 2n nombres aléatoires, stockés dans la mémoire globale du périphérique, puis les alimenter directement au kernel qui doit les utiliser?

Les deux sont des approches valables et peuvent tirer parti de l’accélération du GPU: soit générer les numéros à l’avance et les stocker, soit les générer à la volée.

Parmi les raisons pour lesquelles vous voudrez peut-être envisager une approche plutôt qu’une autre, citons:

  1. Générer les nombres à l’avance n’est utile que si vous savez combien (ou une limite supérieure sur le nombre) dont vous aurez besoin. Si votre algorithme est assez variable (peut-être en présence de différents ensembles de données), cela peut être difficile à déterminer.
  2. Le stockage des numéros générés peut poser problème. Pour certains types d’algorithmes (par exemple, la simulation de Monte Carlo), il peut être nécessaire de générer autant de nombres aléatoires qu’il peut être prohibitif de le faire à l’avance et de tous les stocker. Dans ce cas, leur génération à la volée peut vous permettre d’éviter le recours à de grandes quantités de stockage de nombres aléatoires.
  3. Il serait peut-être possible d’obtenir une utilisation légèrement meilleure de la machine en générant les numéros à la volée, afin d’éviter la surcharge d’un appel du kernel supplémentaire pour générer les numéros avant leur utilisation.

Encore une fois, le principal avantage de CURAND est la performance. Si la génération de nombres aléatoires ne représente qu’une petite partie du coût de calcul global de votre application, l’approche utilisée ou peu importe l’utilisation de CURAND (par exemple à la place d’une méthode de RNG basée sur CPU ordinaire) peut être indifférente.

Lorsque vous utilisez des API hôtes cuRand comme indiqué dans votre code, vous devez d’abord générer des nombres aléatoires, les stocker dans un fichier global, puis les charger dans le kernel de travail. Le stockage et le chargement vont coûter un peu, ce qui peut nuire à la performance.

Toutefois, si vous utilisez les API de périphérique comme indiqué dans votre lien, vous pouvez économiser la bande passante pour le stockage et le chargement de ces nombres aléatoires.

Dans les deux cas, vous devrez charger et stocker la même quantité de données d’état RNG.