/* * Ce fichier CUDA doit être compilé comme suit: * * nvcc -O3 (-arch sm_13) -lrt -o AddVec AddVec.cu * * Le paramètre entre parenthèses, -arch sm_13, n'est nécessaire que si vous employez des "doubles" dans votre code CUDA. De * plus, seules les cartes avec la capacité de calcul dite "1.3" (et plus) permettent le calcul des "doubles". L'option "-lrt", * nécessaire ici, permet d'utiliser le chronomètre en nanosecondes. * * Le fichier généré, "AddVec", est le programme à exécuter. * * * Le programme additionne deux vecteurs aléatoires de 10 000 000 de composantes avec le CPU et avec le GPU. Le temps que met * chacune des deux additions est comptabilisé et affiché à l'écran. */ #include /* Contient les fonctions permettant d'afficher à l'écran, entre autres. */ #include /* Contient le générateur de nombres aléatoires. */ #include /* Pour chronométrer les fonctions. */ /* * La fonction qui suit additionne les vecteurs "v1" et "v2" et place le résultat dans "res". Cette fonction, de type * "__global__" est démarrée par le CPU dans le code principal mais s'exécute sur le GPU. Ce dernier démarrera autant de copies * de cette fonction qu'il y aura de "threads" à exécuter. Les éléments à additionner sont indicés par "i" et celui-ci est * obtenu en connaissant la position du "thread" dans le bloc (threadIdx), la position du bloc dans la grille (blockIdx) ainsi * qu'en connaissant le format de la grille (gridDim) et des blocs (blockDim). * La condition "if" doit être ajoutée pour être certain de ne pas lire et écrire en dehors de la mémoire prévue pour les * vecteurs "v1", "v2" et "res". En effet, le dernier bloc de la grille pourrait contenir moins de "threads" que les autres * blocs. */ __global__ void AddVecGPU(unsigned int TailleVec, float* v1, float* v2, float* res) { unsigned int i = blockDim.x * blockIdx.x + threadIdx.x; if (i < TailleVec) res[i] = v1[i] + v2[i]; } /* * La fonction suivante additionne deux vecteurs avec le CPU. */ void AddVecCPU(unsigned int TailleVec, float* v1, float* v2, float* res) { unsigned int i = 0; while(i < TailleVec) { res[i] = v1[i] + v2[i]; i++; } } /* * Code principal. */ int main() { unsigned int N = 10000000; /* Le nombre de composantes des vecteurs à additionner. */ unsigned int i; /* * Les deux structures qui suivent vont contenir, respectivement, les temps de départ et de fin; leur différence sera le temps, * en nanosecondes, pris par le processus. */ struct timespec leTemps1; struct timespec leTemps2; /* * Allocation des vecteurs d'entrée "h_v1" et "h_v2" et du vecteur de résultats "h_res" définis dans la mémoire du CPU (host). */ float *h_v1 = (float*) malloc( N*sizeof(float) ); float *h_v2 = (float*) malloc( N*sizeof(float) ); float *h_res = (float*) malloc( N*sizeof(float) ); /* * Initialisation des vecteurs aléatoires dont les composantes sont comprises entre 0 et 1. */ i = 0; while(i < N) { h_v1[i] = (float) rand()/RAND_MAX; h_v2[i] = (float) rand()/RAND_MAX; i++; } /* * Définition et allocation des vecteurs "d_v1", "d_v2" et "d_res" sur la mémoire du GPU (device); l'initialisation sera faite * plus loin à partir des vecteurs "h_v1" et "h_v2". */ float *d_v1; float *d_v2; float *d_res; cudaMalloc( (void**) &d_v1, N*sizeof(float) ); /* "cudaMalloc" est conçue pour allouer des espaces mémoire 2d, d'où */ cudaMalloc( (void**) &d_v2, N*sizeof(float) ); /* le besoin du "cast" (void**) pour un vecteur.*/ cudaMalloc( (void**) &d_res, N*sizeof(float) ); /* * On copie maintenant les vecteurs "h_v1" et "h_v2" de la mémoire du CPU (host) dans l'emplacement alloué par "cudaMalloc" sur * la mémoire du GPU (device) pour "d_v1" et "d_v2", respectivement. De cette façon, les deux ensembles de vecteurs seront * identiques. */ cudaMemcpy(d_v1, h_v1, N*sizeof(float), cudaMemcpyHostToDevice); cudaMemcpy(d_v2, h_v2, N*sizeof(float), cudaMemcpyHostToDevice); /* * Avant d'appeler la fonction "__global__", il faut choisir le format de chaque bloc de "threads" ainsi que le nombre de * "threads" les composant. Ensuite, il faut déterminer le format de la grille ainsi que le nombre de blocs qui la * constitueront. Ici, les blocs ont été choisi en format unidimensionnel puisque les données d'entrée sont des vecteurs. Le * format de la grille est déterminé comme suit: si le nombre de blocs nécessaires est supérieur à 65535, alors la grille est * choisie bidimensionnelle. Sinon, le programme sélectionne une grille unidimensionnelle. Ceci est nécessaire puisque CUDA ne * permet pas qu'une des dimensions de la grille (gridDim.x, gridDim.y et gridDim.z) dépasse 65535. * Par exemple, ici, 10 000 000 / 256 = 39062.5, donc 39063 blocs sont nécessaires et une grille unidimensionnelle suffit. Mais * si vous souhaitez choisir N > 65535*256 = 16 776 960, alors une grille bidimensionnelle sera nécessaire. */ dim3 DimenBlocs(256, 1, 1); /* Blocs unidimensionnels de 256 threads. */ dim3 DimenGrille; unsigned int NbBlocsNec = (N + DimenBlocs.x - 1)/DimenBlocs.x; if(NbBlocsNec > 65535) /* Grille bidimensionnelle nécessaire. */ { DimenGrille.x = 65535; DimenGrille.y = (NbBlocsNec + 65535 - 1)/ 65535; DimenGrille.z = 1; } else /* Grille unidimensionnelle suffisante. */ { DimenGrille.x = NbBlocsNec; DimenGrille.y = 1; DimenGrille.z = 1; } /* * Tout juste avant de lancer la fonction "__global__", on démarre le chronomètre (en nanosecondes). */ clock_gettime(CLOCK_REALTIME, &leTemps1); AddVecGPU<<>>(N, d_v1, d_v2, d_res); cudaThreadSynchronize(); /* Cette fonction attend que le GPU ait terminé. */ clock_gettime(CLOCK_REALTIME, &leTemps2); /* Le chronomètre est stoppé. */ printf("Temps GPU: %f s\n", (float) (leTemps2.tv_nsec-leTemps1.tv_nsec) / 1000000000); clock_gettime(CLOCK_REALTIME, &leTemps1); AddVecCPU(N, h_v1, h_v2, h_res); clock_gettime(CLOCK_REALTIME, &leTemps2); printf("Temps CPU: %f s\n", (float) (leTemps2.tv_nsec-leTemps1.tv_nsec) / 1000000000); /* * Libération de la mémoire CPU (host) et GPU (device). */ free(h_v1), free(h_v2), free(h_res); cudaFree(d_v1), cudaFree(d_v2), cudaFree(d_res); }