Différences entre versions de « Cuda hello world »
(13 versions intermédiaires par le même utilisateur non affichées) | |||
Ligne 16 : | Ligne 16 : | ||
[[Fichier:Cuda program workflow.png|centré]] | [[Fichier:Cuda program workflow.png|centré]] | ||
− | + | = Les bases = | |
− | |||
L'API CUDA permet d'exécuter une fonction sur le GPU grâce à la notation chevrons ''<<<...>>>'' : | L'API CUDA permet d'exécuter une fonction sur le GPU grâce à la notation chevrons ''<<<...>>>'' : | ||
<source lang="c"> | <source lang="c"> | ||
− | fonction<<< | + | fonction<<<NB_BLOCK, NB_THREAD>>>(); |
</source> | </source> | ||
* fonction : correspond au nom de la fonction | * fonction : correspond au nom de la fonction | ||
− | * | + | * NB_BLOCK : correspond au nombre de blocs utilisés |
− | * | + | * NB_THREAD : correspond au nombre de threads par bloc |
La fonction appelée doit être précédée du spécificateur ''__global__'' pour préciser au compilateur que sont exécution se fait sur le GPU: | La fonction appelée doit être précédée du spécificateur ''__global__'' pour préciser au compilateur que sont exécution se fait sur le GPU: | ||
<source lang="c"> | <source lang="c"> | ||
Ligne 35 : | Ligne 34 : | ||
# nvcc hello.cu -o hello | # nvcc hello.cu -o hello | ||
</pre> | </pre> | ||
− | == | + | == Programmation CPU vs. GPU == |
Ci-dessous l'exemple classique du ''Hello World'' utilisant le CPU: | Ci-dessous l'exemple classique du ''Hello World'' utilisant le CPU: | ||
<source lang="c"> | <source lang="c"> | ||
Ligne 63 : | Ligne 62 : | ||
} | } | ||
</source> | </source> | ||
− | + | ||
− | + | = Synchronisation CPU / GPU = | |
+ | Le programme précédent affichera le résultat ci-dessous : | ||
<pre> | <pre> | ||
Hello World du CPU | Hello World du CPU | ||
</pre> | </pre> | ||
− | Le GPU n'a | + | Le GPU n'a simplement rien fait. Il n'était pas informé de quoique ce soit car les instructions n'ont pas été copiées sur la mémoire du GPU. Cela se fait grâce à la fonction ''cudaDeviceSynchronize'' qu'il faut ajouter juste après l'instruction chevron: |
<source lang="c"> | <source lang="c"> | ||
... | ... | ||
Ligne 75 : | Ligne 75 : | ||
... | ... | ||
</source> | </source> | ||
− | + | Et nous avons bien le résultat attendu: | |
<pre> | <pre> | ||
Hello World du CPU | Hello World du CPU | ||
Hello World du GPU | Hello World du GPU | ||
</pre> | </pre> | ||
− | == Parallélisation == | + | |
+ | = Manipulation mémoire = | ||
+ | == Gestion manuellement == | ||
+ | Le programme ci-dessous permet d'afficher le nom passé en paramètre : | ||
+ | <source lang="c"> | ||
+ | #include <stdio.h> | ||
+ | #include <stdlib.h> | ||
+ | #include <cuda_runtime.h> | ||
+ | #include <cuda.h> | ||
+ | |||
+ | __global__ void say_my_name(char * name){ | ||
+ | printf("GPU says : %s\n", name); | ||
+ | } | ||
+ | |||
+ | int main(int argc, char * argv[]) { | ||
+ | if(argc < 2){ | ||
+ | printf("I don't know your name !"); | ||
+ | return EXIT_FAILURE; | ||
+ | } | ||
+ | printf("CPU says : %s\n", argv[1]); | ||
+ | say_my_name<<<1,1>>>(argv[1]); | ||
+ | cudaDeviceSynchronize(); | ||
+ | return EXIT_SUCCESS; | ||
+ | } | ||
+ | </source> | ||
+ | Le programme précédent affichera : | ||
+ | <pre> | ||
+ | # name.bin tala | ||
+ | CPU says : tala | ||
+ | </pre> | ||
+ | Il manque, encore une fois, l’exécution du code côté GPU... Ce qui est normal car la variable passée en paramètre ''argv[1]'' est dans la mémoire centrale donc inaccessible côté GPU ! | ||
+ | |||
+ | Pour pouvoir l'utiliser, il faut copier le contenu de cette variable dans la mémoire du GPU. | ||
+ | |||
+ | Cela se fait en trois étapes: | ||
+ | # on fait l'allocation mémoire avec ''cudaMalloc'' | ||
+ | # on procède à la copie avec ''cudaMemCpy'' | ||
+ | # on procède à la libération mémoire avec ''cudaFree'' | ||
+ | ===allocation=== | ||
+ | ''cudaMalloc'' prend en paramètre un pointeur côté CPU et la taille de l'allocation : | ||
+ | <source lang="c"> | ||
+ | cudaMalloc(void ** ptr, size_t size); | ||
+ | </source> | ||
+ | Le pointeur ''ptr'' permet de faire le lien entre les variables dans la mémoire centrale CPU et celle dans la mémoire graphique. | ||
+ | ===copie=== | ||
+ | ''cudaMemCpy'' prend en paramètre : | ||
+ | * un pointeur source | ||
+ | * un pointeur destination | ||
+ | * la taille de la copie | ||
+ | * le sens de la copie | ||
+ | <source lang="c"> | ||
+ | cudaMemCpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind direction); | ||
+ | </source> | ||
+ | Notez que le paramètre ''direction'' de type ''cudaMemcpyKind'' peut prendre les valeurs suivantes: | ||
+ | * cudaMemcpyHostToDevice (CPU → GPU) | ||
+ | * cudaMemcpyDeviceToHost (GPU → CPU) | ||
+ | ===libération=== | ||
+ | ''cudaFree'' prend en paramètre le pointeur côté GPU à libérer | ||
+ | <source lang="c"> | ||
+ | cudaFree(void *ptr); | ||
+ | </source> | ||
+ | |||
+ | Ce qui donne les modifications suivantes : | ||
+ | <source lang="c"> | ||
+ | ... | ||
+ | printf("CPU says : %s\n", argv[1]); | ||
+ | char * name; | ||
+ | int size = sizeof(char) * strlen(argv[1]) + 1; | ||
+ | cudaMalloc(&name, size); | ||
+ | cudaMemcpy(name, argv[1], size, cudaMemcpyHostToDevice); | ||
+ | say_my_name<<<1,1>>>(name); | ||
+ | cudaDeviceSynchronize(); | ||
+ | cudaFree(name); | ||
+ | ... | ||
+ | </source> | ||
+ | Pour avoir le résultat attendu: | ||
+ | <pre> | ||
+ | # ./name.bin tala | ||
+ | CPU says : tala | ||
+ | GPU says : tala | ||
+ | </pre> | ||
+ | == Gestion unifiée == | ||
+ | [https://fr.wikipedia.org/wiki/Maxwell_(architecture_de_carte_graphique) L'architecture Maxwell] a introduit le support de la mémoire unifiée. Les variables allouées grâce à la fonction ''cudaMallocManaged'' voit leurs espaces mémoire transiter entre CPU et GPU au besoin. Cella permet d'écrire les programmes plus simplement, sans avoir à appeler explicitement la fonction cudaMemcpy. | ||
+ | |||
+ | Dans l'exemple suivant on créer une chaîne de caractère que l'on modifie dans le kernel GPU avant de l'afficher avec le CPU: | ||
+ | <source lang="c"> | ||
+ | #include <stdio.h> | ||
+ | #include <stdlib.h> | ||
+ | #include <cuda_runtime.h> | ||
+ | #include <cuda.h> | ||
+ | |||
+ | __global__ void say_hello(char *str){ | ||
+ | printf("CPU said : %s\n", str); | ||
+ | memcpy(str, "world", 5); | ||
+ | } | ||
+ | |||
+ | int main(int argc, char * argv[]) { | ||
+ | char *str; | ||
+ | cudaMallocManaged(&str, 5 * sizeof(char)); | ||
+ | say_hello<<<1,1>>>(str); | ||
+ | cudaDeviceSynchronize(); | ||
+ | printf("GPU said : %s\n", str); | ||
+ | return EXIT_SUCCESS; | ||
+ | } | ||
+ | </source> | ||
+ | La sortie permet de vérifier que l'espace mémoire est géré correctement: | ||
+ | <pre> | ||
+ | CPU said : hello | ||
+ | GPU said : world | ||
+ | </pre> | ||
+ | |||
+ | = Affichage des erreurs = | ||
+ | L'exemple précédent met en évidence le manque de visibilité sur ce qui se passe côté GPU. Pour cela on peut utiliser la fonction ''gpuAssert'' pour afficher les erreurs lors de l’exécution du code sur le GPU: | ||
+ | <source lang="c"> | ||
+ | #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } | ||
+ | |||
+ | inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) { | ||
+ | if (code != cudaSuccess) { | ||
+ | fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); | ||
+ | if (abort) exit(code); | ||
+ | } | ||
+ | } | ||
+ | </source> | ||
+ | Le code précédent est à mettre au dessus de la fonction ''main'' et avant les fonctions chevrons. Si on reprend l'exemple précédent cela donnerait : | ||
+ | <source lang="c"> | ||
+ | ... | ||
+ | __global__ void say_my_name(char * name){ | ||
+ | printf("GPU says : %s\n", name); | ||
+ | } | ||
+ | |||
+ | int main(int argc, char * argv[]) { | ||
+ | if(argc < 2){ | ||
+ | printf("I don't know your name !"); | ||
+ | return EXIT_FAILURE; | ||
+ | } | ||
+ | printf("CPU says : %s\n", argv[1]); | ||
+ | say_my_name<<<1,1>>>(argv[1]); | ||
+ | gpuErrchk(cudaDeviceSynchronize()); | ||
+ | return EXIT_SUCCESS; | ||
+ | } | ||
+ | </source> | ||
+ | Comme l'allocation n'est pas faite avec ''cudaMalloc'', cela affiche l'erreur suivante: | ||
+ | <pre> | ||
+ | # ./name.bin tala | ||
+ | CPU says : tala | ||
+ | GPUassert: an illegal memory access was encountered name.cu 30 | ||
+ | </pre> | ||
+ | |||
+ | Il est aussi possible d'utiliser la fonction ''cudaPeekAtLastError'' pour afficher la dernière erreur : | ||
+ | <source lang="c"> | ||
+ | gpuErrchk(cudaPeekAtLastError()); | ||
+ | </source> | ||
+ | |||
+ | = Parallélisation = | ||
+ | Le gros avantage du GPU c'est sa capacité à paralléliser l'exécution du code et pour cela on doit s'attarder sur la syntaxe chevron. On peut mixer le nombre de blocs ainsi que le nombre de thread par blocs jusqu'à un total de 1024, comme expliqué [https://forums.developer.nvidia.com/t/maximum-number-of-threads-on-thread-block/46392 ici] : | ||
+ | * on peut faire <<<1024, 1>>> ou <<<128,8>>> (128 * 8 = 1024) | ||
+ | * on ne peut pas faire <<<256, 10>>> car 256 * 10 = 2560 > 1024 | ||
+ | |||
+ | Il faut comprendre que le code exécuté par les différentes instances est le même, il faudra donc modifier son comportement grâce au numéro de groupe ou de thread. Ces numéros sont accessible grâce aux variables : | ||
+ | * threadIdx : id du thread | ||
+ | * blockIdx : id du bloc | ||
+ | * blockDim : dimension de chaque bloc | ||
+ | * gridDim : dimension de la grille (nombre de blocs) | ||
+ | |||
+ | On peut schématiser la parallélisation effectuée par le GPU comme la matrice suivante: | ||
+ | [[Fichier:Cuda parallel matrix.png|centré]] | ||
+ | |||
+ | Dans l'exemple suivant on affiche les différents indexes : | ||
+ | |||
+ | <source lang="c"> | ||
+ | #include <stdio.h> | ||
+ | #include <stdlib.h> | ||
+ | #include <cuda_runtime.h> | ||
+ | #include <cuda.h> | ||
+ | |||
+ | __global__ void whoami() { | ||
+ | printf("gridDim[%d], blockDim [%d], blockIdx [%d], threadIdx[%d]\n", gridDim.x, blockDim.x, blockIdx.x, threadIdx.x); | ||
+ | } | ||
+ | |||
+ | int main(int argc, char *argv[]) { | ||
+ | whoami<<<2,3>>>(); | ||
+ | cudaDeviceSynchronize(); | ||
+ | return EXIT_SUCCESS; | ||
+ | } | ||
+ | </source> | ||
+ | Cette exemple donne la sortie suivante: | ||
+ | <pre> | ||
+ | gridDim[2], blockDim [3], blockIdx [0], threadIdx[0] | ||
+ | gridDim[2], blockDim [3], blockIdx [0], threadIdx[1] | ||
+ | gridDim[2], blockDim [3], blockIdx [0], threadIdx[2] | ||
+ | gridDim[2], blockDim [3], blockIdx [1], threadIdx[0] | ||
+ | gridDim[2], blockDim [3], blockIdx [1], threadIdx[1] | ||
+ | gridDim[2], blockDim [3], blockIdx [1], threadIdx[2] | ||
+ | </pre> | ||
+ | On peut remarquer que: | ||
+ | * gridDim.x correspond au nombre de blocs (premier paramètre de la fonction chevron) | ||
+ | * blockDim.x correspond au nombre de threads par blocs (deuxième paramètre de la fonction chevron) | ||
+ | * blockIdx.x correspond à l'index du bloc (premier paramètre de la fonction chevron) | ||
+ | * threadIdx.x correspond à l'index du thread (compris entre 0 et blockDim.x - 1) | ||
+ | |||
+ | On peut utiliser l'astuce suivante pour calculer un index unique par thread : | ||
+ | <source lang="c"> | ||
+ | __global__ void whoami() { | ||
+ | int index = blockDim.x * blockIdx.x + threadIdx.x; | ||
+ | int total = gridDim.x * blockDim.x; | ||
+ | printf("I am %d over %d\n", index, total); | ||
+ | } | ||
+ | </source> | ||
+ | Ce qui donne le résultat suivant: | ||
+ | <pre> | ||
+ | I am 0 over 6 | ||
+ | I am 1 over 6 | ||
+ | I am 2 over 6 | ||
+ | I am 3 over 6 | ||
+ | I am 4 over 6 | ||
+ | I am 5 over 6 | ||
+ | </pre> |
Version actuelle datée du 23 juin 2021 à 08:09
Concepts
Le framework CUDA permet d'utiliser le GPU pour faire des calculs généraux, habituellement réalisés par le CPU. CUDA est une extension en C/C++ qui propose une API pour gérer le GPU et elle distingue deux entités:
- l'entité host ou hôte pour le CPU
- l'entité device ou périphérique pour le GPU
De manière générale on essayera de faire le traitement sérialisé sur le CPU et de décharger les calculs parallélisés sur le GPU.
La difficulté réside dans le faite que CPU et GPU utilisent deux espaces mémoire séparés... qu'il va falloir gérer par le biais de pointeurs et fonctions spécifiques !
Les programmes utilisant le GPU se déroule en 5 étapes:
- instanciation des instructions et variables dans la mémoire centrale (RAM)
- copie des instruction et variables dans la mémoire du GPU
- exécution des instruction et lecture des variables par le GPU
- modification des variables dans la mémoire du GPU
- rapatriement des variables modifiées par le GPU dans la mémoire centrale
Les bases
L'API CUDA permet d'exécuter une fonction sur le GPU grâce à la notation chevrons <<<...>>> :
fonction<<<NB_BLOCK, NB_THREAD>>>();
- fonction : correspond au nom de la fonction
- NB_BLOCK : correspond au nombre de blocs utilisés
- NB_THREAD : correspond au nombre de threads par bloc
La fonction appelée doit être précédée du spécificateur __global__ pour préciser au compilateur que sont exécution se fait sur le GPU:
__global__ void fonction(){
// choses à faire
}
Les fichiers contenant du code CUDA doivent impérativement porter l'extension .cu et être compilés avec le compilateur CUDA : nvcc
# nvcc hello.cu -o hello
Programmation CPU vs. GPU
Ci-dessous l'exemple classique du Hello World utilisant le CPU:
#include <stdio.h>
int main(void) {
printf("Hello World du CPU\n");
return EXIT_SUCCESS;
}
Si nous voulons utiliser le GPU il faut sortir l'instruction à exécuter (ici le printf) dans une fonction __global__ puis lancer son exécution grâce à la notation chevrons:
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>
__global__ void cuda_hello(){
printf("Hello World du GPU\n");
}
int main(void) {
printf("Hello World du CPU\n");
cuda_hello<<<1,1>>>();
return EXIT_SUCCESS;
}
Synchronisation CPU / GPU
Le programme précédent affichera le résultat ci-dessous :
Hello World du CPU
Le GPU n'a simplement rien fait. Il n'était pas informé de quoique ce soit car les instructions n'ont pas été copiées sur la mémoire du GPU. Cela se fait grâce à la fonction cudaDeviceSynchronize qu'il faut ajouter juste après l'instruction chevron:
...
cuda_hello<<<1,1>>>();
cudaDeviceSynchronize();
...
Et nous avons bien le résultat attendu:
Hello World du CPU Hello World du GPU
Manipulation mémoire
Gestion manuellement
Le programme ci-dessous permet d'afficher le nom passé en paramètre :
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>
__global__ void say_my_name(char * name){
printf("GPU says : %s\n", name);
}
int main(int argc, char * argv[]) {
if(argc < 2){
printf("I don't know your name !");
return EXIT_FAILURE;
}
printf("CPU says : %s\n", argv[1]);
say_my_name<<<1,1>>>(argv[1]);
cudaDeviceSynchronize();
return EXIT_SUCCESS;
}
Le programme précédent affichera :
# name.bin tala CPU says : tala
Il manque, encore une fois, l’exécution du code côté GPU... Ce qui est normal car la variable passée en paramètre argv[1] est dans la mémoire centrale donc inaccessible côté GPU !
Pour pouvoir l'utiliser, il faut copier le contenu de cette variable dans la mémoire du GPU.
Cela se fait en trois étapes:
- on fait l'allocation mémoire avec cudaMalloc
- on procède à la copie avec cudaMemCpy
- on procède à la libération mémoire avec cudaFree
allocation
cudaMalloc prend en paramètre un pointeur côté CPU et la taille de l'allocation :
cudaMalloc(void ** ptr, size_t size);
Le pointeur ptr permet de faire le lien entre les variables dans la mémoire centrale CPU et celle dans la mémoire graphique.
copie
cudaMemCpy prend en paramètre :
- un pointeur source
- un pointeur destination
- la taille de la copie
- le sens de la copie
cudaMemCpy(void *dst, const void *src, size_t count, enum cudaMemcpyKind direction);
Notez que le paramètre direction de type cudaMemcpyKind peut prendre les valeurs suivantes:
- cudaMemcpyHostToDevice (CPU → GPU)
- cudaMemcpyDeviceToHost (GPU → CPU)
libération
cudaFree prend en paramètre le pointeur côté GPU à libérer
cudaFree(void *ptr);
Ce qui donne les modifications suivantes :
...
printf("CPU says : %s\n", argv[1]);
char * name;
int size = sizeof(char) * strlen(argv[1]) + 1;
cudaMalloc(&name, size);
cudaMemcpy(name, argv[1], size, cudaMemcpyHostToDevice);
say_my_name<<<1,1>>>(name);
cudaDeviceSynchronize();
cudaFree(name);
...
Pour avoir le résultat attendu:
# ./name.bin tala CPU says : tala GPU says : tala
Gestion unifiée
L'architecture Maxwell a introduit le support de la mémoire unifiée. Les variables allouées grâce à la fonction cudaMallocManaged voit leurs espaces mémoire transiter entre CPU et GPU au besoin. Cella permet d'écrire les programmes plus simplement, sans avoir à appeler explicitement la fonction cudaMemcpy.
Dans l'exemple suivant on créer une chaîne de caractère que l'on modifie dans le kernel GPU avant de l'afficher avec le CPU:
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>
__global__ void say_hello(char *str){
printf("CPU said : %s\n", str);
memcpy(str, "world", 5);
}
int main(int argc, char * argv[]) {
char *str;
cudaMallocManaged(&str, 5 * sizeof(char));
say_hello<<<1,1>>>(str);
cudaDeviceSynchronize();
printf("GPU said : %s\n", str);
return EXIT_SUCCESS;
}
La sortie permet de vérifier que l'espace mémoire est géré correctement:
CPU said : hello GPU said : world
Affichage des erreurs
L'exemple précédent met en évidence le manque de visibilité sur ce qui se passe côté GPU. Pour cela on peut utiliser la fonction gpuAssert pour afficher les erreurs lors de l’exécution du code sur le GPU:
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort = true) {
if (code != cudaSuccess) {
fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
Le code précédent est à mettre au dessus de la fonction main et avant les fonctions chevrons. Si on reprend l'exemple précédent cela donnerait :
...
__global__ void say_my_name(char * name){
printf("GPU says : %s\n", name);
}
int main(int argc, char * argv[]) {
if(argc < 2){
printf("I don't know your name !");
return EXIT_FAILURE;
}
printf("CPU says : %s\n", argv[1]);
say_my_name<<<1,1>>>(argv[1]);
gpuErrchk(cudaDeviceSynchronize());
return EXIT_SUCCESS;
}
Comme l'allocation n'est pas faite avec cudaMalloc, cela affiche l'erreur suivante:
# ./name.bin tala CPU says : tala GPUassert: an illegal memory access was encountered name.cu 30
Il est aussi possible d'utiliser la fonction cudaPeekAtLastError pour afficher la dernière erreur :
gpuErrchk(cudaPeekAtLastError());
Parallélisation
Le gros avantage du GPU c'est sa capacité à paralléliser l'exécution du code et pour cela on doit s'attarder sur la syntaxe chevron. On peut mixer le nombre de blocs ainsi que le nombre de thread par blocs jusqu'à un total de 1024, comme expliqué ici :
- on peut faire <<<1024, 1>>> ou <<<128,8>>> (128 * 8 = 1024)
- on ne peut pas faire <<<256, 10>>> car 256 * 10 = 2560 > 1024
Il faut comprendre que le code exécuté par les différentes instances est le même, il faudra donc modifier son comportement grâce au numéro de groupe ou de thread. Ces numéros sont accessible grâce aux variables :
- threadIdx : id du thread
- blockIdx : id du bloc
- blockDim : dimension de chaque bloc
- gridDim : dimension de la grille (nombre de blocs)
On peut schématiser la parallélisation effectuée par le GPU comme la matrice suivante:
Dans l'exemple suivant on affiche les différents indexes :
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <cuda.h>
__global__ void whoami() {
printf("gridDim[%d], blockDim [%d], blockIdx [%d], threadIdx[%d]\n", gridDim.x, blockDim.x, blockIdx.x, threadIdx.x);
}
int main(int argc, char *argv[]) {
whoami<<<2,3>>>();
cudaDeviceSynchronize();
return EXIT_SUCCESS;
}
Cette exemple donne la sortie suivante:
gridDim[2], blockDim [3], blockIdx [0], threadIdx[0] gridDim[2], blockDim [3], blockIdx [0], threadIdx[1] gridDim[2], blockDim [3], blockIdx [0], threadIdx[2] gridDim[2], blockDim [3], blockIdx [1], threadIdx[0] gridDim[2], blockDim [3], blockIdx [1], threadIdx[1] gridDim[2], blockDim [3], blockIdx [1], threadIdx[2]
On peut remarquer que:
- gridDim.x correspond au nombre de blocs (premier paramètre de la fonction chevron)
- blockDim.x correspond au nombre de threads par blocs (deuxième paramètre de la fonction chevron)
- blockIdx.x correspond à l'index du bloc (premier paramètre de la fonction chevron)
- threadIdx.x correspond à l'index du thread (compris entre 0 et blockDim.x - 1)
On peut utiliser l'astuce suivante pour calculer un index unique par thread :
__global__ void whoami() {
int index = blockDim.x * blockIdx.x + threadIdx.x;
int total = gridDim.x * blockDim.x;
printf("I am %d over %d\n", index, total);
}
Ce qui donne le résultat suivant:
I am 0 over 6 I am 1 over 6 I am 2 over 6 I am 3 over 6 I am 4 over 6 I am 5 over 6