Compute Unified Device Architecture

Cet article est une ébauche concernant l’informatique. Vous pouvez partager vos connaissances en l’améliorant (comment ?) selon les recommandations des projets correspondants.

Schéma de principe de l'architecture CUDA

CUDA (Compute Unified Device Architecture) est une technologie de GPGPU (General-Purpose Computing on Graphics Processing Units), c'est-à-dire qu'on utilise un processeur graphique (GPU) pour exécuter des calculs généraux habituellement exécutés par le processeur central (CPU). CUDA permet de programmer des GPU en C. Cette technologie a été développée par NVIDIA pour leurs cartes graphiques GeForce 8 Series, et utilise un pilote unifié utilisant une technique de streaming (flux continu). NVIDIA s'engage à ce que ses futures cartes graphiques restent compatibles avec CUDA.

Le premier kit de développement pour CUDA a été publié le 15 février 2007.

Utilisations de CUDA

Architecture Tesla

L'architecture Tesla, qui offre selon NVidia le pouvoir de calcul d'un superordinateur (4 téraflops en simple précision, 80 gigaflops en double précision) pour une somme de 10 000 dollars, est construite sur CUDA.

Boîte à outils Pyrit

Pyrit est une boîte à outils Open Source utilisant le parallélisme massif de calcul de CUDA pour craquer des clés Wi-Fi WPA-PSK et WPA2-PSK en gagnant un facteur 20 sur le même calcul réalisé par le processeur seul. Cela ne signifie pas nécessairement « le crépuscule du Wi-Fi » comme cela a été parfois annoncé, mais diminue tout de même de trois ordres de grandeur le degré de protection escompté pour les communications.

Avantages

CUDA présente plusieurs avantage par rapport aux autres API proposant d'effectuer des calculs génériques sur GPU:

• Le code peut lire à n'importe quelle adresse mémoire.
• Mémoire locale. Un pool de threads (un bloc) peut partager une zone mémoire (jusqu'à 48k par multiprocesseur selon les modèles de carte) qui présente des temps d'accès très court (Idem OpenCL)
• Plus de souplesse de codage qu'OpenCL, et garantit que les binaires s'executent sur les plateformes compatibles.

Limitations

• Les transfert entre le host et le device peuvent être un goulot d'étranglement (peut être évité par des copies asynchrones).
• Regroupement des threads en groupe de 32 pour des questions de performances (les warps). Au sein d'un warp, les divergence (du aux executions conditionnelles) peuvent nuire grandement aux performances. Il s'agit d'une limitation dûe au modèle SIMD. De ce fait, tous les algorithmes ne gagnent pas à être portés sur CUDA, et plus généralement sur un modèle SIMD.
• Un code C valide peut être rejeté à cause des contraintes matérielle (manque de mémoire, nombre de threads).
• Les première versions (1.x) de CUDA ne supportent pas la récursion, les pointeurs de fonctions, et d'autres limitations qui tendent à disparaitre.
• La précision des calculs. La double-précision n'est disponible que depuis la version 1.3; et la norme IEEE 754 n'est que partiellement atteinte: en double-précision, le seul mode d'arrondi est round-to-nearest-even. En simple précision, les nombres ne sont pas dénormalisés, le signal NaN est absenti, et les 2 seuls modes d'arrondi sont chop and round-to-nearest even, etc.

Exemples

Exemple avec émulation de carte

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
 
 
__global__ void mykernel(float * A1, float * A2, float * R)
{
        int p = threadIdx.x;
        R[p] = A1[p] + A2[p];
}
 
int main()
{
        float A1[]={1,2,3,4,5,6,7,8,9};
        float A2[]={10,20,30,40,50,60,70,80,90};
        float R[9];
 
        //9 additions, aucune boucle !
        mykernel<<<1,9>>>(A1,A2,R);
 
        //sortie à l'ecran
        for(int i=0; i<9 ; i++)
                printf("%f\n",R[i]);
 
}

Cet exemple fonctionne seulement si on émule la carte graphique car on ne recopie pas les données sur la carte.

Compilation par :

nvcc -deviceemu -o run prog.cu

Exemple avec une carte graphique NVidia

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <cuda_runtime.h>
 
 
__global__ void mykernel(float * A1, float * A2, float * R)
{
        int p = threadIdx.x;
        R[p] = A1[p] + A2[p];
}
 
int main()
{
        float A1[]={1,2,3,4,5,6,7,8,9};
        float A2[]={10,20,30,40,50,60,70,80,90};
        float R[9];
        int taille_mem=sizeof(float) * 9;
        // on alloue de la memoire sur la carte graphique
        float * a1_device;
        float * a2_device;
        float * r_device;
        cudaMalloc ( (void**) &a1_device, taille_mem);
        cudaMalloc ( (void**) &a2_device, taille_mem);
        cudaMalloc ( (void**) &r_device, taille_mem);
        // on copie les donnees sur la carte
        cudaMemcpy( a1_device,A1,taille_mem,cudaMemcpyHostToDevice);
        cudaMemcpy( a2_device,A2,taille_mem,cudaMemcpyHostToDevice);
 
        //9 additions, aucune boucle !
        mykernel<<<1,9>>>(a1_device,a2_device,r_device);
 
        // on recupere le resultat
        cudaMemcpy(R,r_device,taille_mem,cudaMemcpyDeviceToHost);
        //sortie à l'ecran
        for(int i=0; i<9 ; i++)
                printf("%f\n",R[i]);
 
}

Compilation par :

nvcc -o add_cuda add_cuda.cu

Voir aussi

installation de CUDA en fonction des systemes d'exploitation. points spécifiques

• sous Mandriva

Produits concurrents

• OpenCL (API/langage créé à l'origine par Apple, sous licence BSD)
• Larrabee de Intel (encore en projet)
• ATI Stream.
• Compute Shaders de Direct3D 11

Articles connexes

• GPU
• Tesla
• HMPP

Liens externes

• (en) Page sur CUDA du site de NVIDIA
• (fr) Article sur CUDA Tom's Hardware
• (fr) Explications sur CUDA du site Hardware.fr
• (fr) Introduction à CUDA sur le site Developpez.com