Introduction à CUDA: Une Démarche Simplifiée Ce billet deblog s'adresse aux programmeurs C++ souhaitant se familiariser avec CUDA, la plateforme et le modèle de programmation parallèle de NVIDIA. Cette mise à jour résume les avancées et simplifications apportées par CUDA, notamment dans le contexte de l'intelligence artificielle et de l'apprentissage profond (Deep Learning). Pour bien suivre ce tutoriel, vous aurez besoin d'un ordinateur équipé d'une GPU compatible CUDA (sous Windows, WSL ou Linux en 64 bits), ou d'une instance cloud disposant de GPU (AWS, Azure, Google Colab, etc.). Le téléchargement gratuit de la CUDA Toolkit est également requis. Démarrer Simplement L'objectif initial est d'ajouter les éléments de deux tableaux, chacun contenant un million d'éléments. Code Initial Voici un exemple de code C++ basique qui réalise cette opération : ```cpp include include void add(int a, int b, int* c, int n) { for (int i = 0; i < n; i++) { c[i] = a[i] + b[i]; } } int main() { const int N = 1000000; int a = new int[N]; int b = new int[N]; int* c = new int[N]; // Initialisation des tableaux for (int i = 0; i < N; i++) { a[i] = -i; b[i] = i * i; } add(a, b, c, N); // Vérification de l'exactitude for (int i = 0; i < N; i++) { assert(c[i] == i * i - i); } std::cout << "Pas d'erreur!" << std::endl; delete[] a; delete[] b; delete[] c; return 0; } ``` Compilez et exécutez ce code : bash g++ -o add add.cpp ./add Passage à CUDA Pour exécuter cette opération en parallèle sur une GPU, commencez par transformer la fonction add en un noyau CUDA : cpp __global__ void add(int* a, int* b, int* c, int n) { for (int i = 0; i < n; i++) { c[i] = a[i] + b[i]; } } La spécification __global__ indique que la fonction est exécutée sur la GPU et peut être appelée depuis la CPU. Allocation de Mémoire Utilisez la mémoire unifiée (Unified Memory) pour simplifier l'allocation de mémoire accessible par la GPU et la CPU : cpp int* a; int* b; int* c; cudaMallocManaged(&a, N * sizeof(int)); cudaMallocManaged(&b, N * sizeof(int)); cudaMallocManaged(&c, N * sizeof(int)); Libération de Mémoire cpp cudaFree(a); cudaFree(b); cudaFree(c); Lancement du Noyau Utilisez la syntaxe <<< >>> pour lancer le noyau : cpp add<<<1, 1>>>(a, b, c, N); Cette ligne lance une thread unique pour exécuter add(). Pour synchroniser le CPU avec le GPU, utilisez cudaDeviceSynchronize() avant la vérification finale : cpp cudaDeviceSynchronize(); Code Complet de Base ```cpp include include global void add(int a, int b, int* c, int n) { for (int i = 0; i < n; i++) { c[i] = a[i] + b[i]; } } int main() { const int N = 1000000; int a; int b; int* c; cudaMallocManaged(&a, N * sizeof(int)); cudaMallocManaged(&b, N * sizeof(int)); cudaMallocManaged(&c, N * sizeof(int)); // Initialisation des tableaux for (int i = 0; i < N; i++) { a[i] = -i; b[i] = i * i; } add<<<1, 1>>>(a, b, c, N); cudaDeviceSynchronize(); // Vérification de l'exactitude for (int i = 0; i < N; i++) { assert(c[i] == i * i - i); } std::cout << "Pas d'erreur!" << std::endl; cudaFree(a); cudaFree(b); cudaFree(c); return 0; } ``` Optimisation avec le Parallélisme Parallélisme au Niveau des Thread Modifiez la boucle pour utiliser plusieurs threads : ```cpp global void add(int a, int b, int* c, int n) { int index = threadIdx.x; int stride = blockDim.x; for (int i = index; i < n; i += stride) { c[i] = a[i] + b[i]; } } ``` Lancez le noyau avec 256 threads par bloc : cpp add<<<1, 256>>>(a, b, c, N); Parallélisme au Niveau des Blocs Augmentez le nombre de blocs pour tirer parti de tous les processeurs parallèles (SMs) de la GPU : cpp int blockSize = 256; int numBlocks = (N + blockSize - 1) / blockSize; add<<<numBlocks, blockSize>>>(a, b, c, N); Préfetching de la Mémoire Unifiée La migration de la mémoire peut entraîner des latences significatives. Pour l'optimiser, préchargez la mémoire : cpp cudaMemPrefetchAsync(a, N * sizeof(int), 0); cudaMemPrefetchAsync(b, N * sizeof(int), 0); cudaMemPrefetchAsync(c, N * sizeof(int), 0); Résultats et Comparaison Voici une comparaison des temps d'exécution et des vitesses obtenues avec les différentes versions du programme : | Version | Temps (ns) | Accélération vs. Thread Unique | Bande Passante (GB/s) | |---------|------------|-----------------------------|---------------------| | Thread Unique | 91,811,206 | 1x | 137 | | Bloc Unique (256 threads) | 2,049,034 | 45x | 6 | | Multiple Blocs | 47,520 | 1932x | 265 | Avec la précharge de la mémoire, le temps de noyau est réduit à moins de 50 microsecondes, atteignant une bande passante de 265 GB/s, soit plus de 80% du pic de bande passante de l'NVIDIA T4. Exercices Pratiques Documentation : Explorez la documentation de la CUDA Toolkit, y compris les guides de démarrage rapide, de programmation et de bonnes pratiques. Expérimentation avec printf() : Essayez d'afficher les valeurs de threadIdx.x et blockIdx.x pour certains ou tous les threads. Sont-elles affichées dans l'ordre prévu ? Utilisation de threadIdx et blockIdx en 2D et 3D : Affichez les valeurs de threadIdx.y, threadIdx.z, blockIdx.y et blockIdx.z pour comprendre pourquoi elles existent et comment leur assigner des valeurs autres que 0 ou 1. Perspective et Ressources Ce tutoriel offre une introduction simplifiée à CUDA, mais il reste d'autres aspects à découvrir. Pour approfondir vos connaissances, consultez les cours en ligne offerts par NVIDIA DLI, notamment : Getting Started with Accelerated Computing in Modern CUDA C++ : Cours complet avec des ressources dédiées, un environnement de programmation sophistiqué, et des présentations détaillées. Fundamentals of Accelerated Computing with CUDA Python : Pour les programmeurs Python. Accelerated Computing Section of the NVIDIA DLI Self-Paced Catalog : Matériaux intermédiaires et avancés pour divers niveaux de compétence. En bref, CUDA permet une accélération massive des applications computationnelles et à haute consommation de bande passante, ouvrant la voie à des performances exceptionnelles dans des domaines tels que l'IA, l'apprentissage profond, le traitement d'images et de signaux, et les simulations physiques. Commencer par cette introduction simplifiée est un excellent point de départ pour explorer ses possibilités.