cudaCommencer avec cuda


Remarques

CUDA est une technologie de programmation parallèle et un langage de programmation NVIDIA propriétaires pour leurs GPU.

Les GPU sont des machines hautement parallèles capables d'exécuter des milliers de threads légers en parallèle. Chaque thread GPU est généralement plus lent dans l'exécution et son contexte est plus petit. D'autre part, le GPU est capable d'exécuter plusieurs milliers de threads en parallèle et même plus simultanément (les nombres précis dépendent du modèle de GPU réel). CUDA est un dialecte C ++ conçu spécifiquement pour l'architecture GPU NVIDIA. Cependant, en raison des différences d’architecture, la plupart des algorithmes ne peuvent pas être simplement copiés-collés à partir de C ++, ils seraient exécutés mais seraient très lents.

Terminologie

  • host - fait référence au matériel normal basé sur le processeur et aux programmes normaux exécutés dans cet environnement
  • device - fait référence à un GPU spécifique exécuté par les programmes CUDA. Un hôte unique peut prendre en charge plusieurs périphériques.
  • kernel - une fonction qui réside sur le périphérique et qui peut être appelée à partir du code hôte.

Structure du processeur physique

Le processeur GPU compatible CUDA a la structure physique suivante:

  • la puce - le processeur entier du GPU. Certains GPU en ont deux.
  • streaming multiprocessor (SM) - chaque puce contient jusqu'à ~ 100 SM, selon le modèle. Chaque SM fonctionne presque indépendamment l'un de l'autre, en utilisant uniquement la mémoire globale pour communiquer entre eux.
  • CUDA core - une unité de calcul scalaire unique d'un SM. Leur nombre précis dépend de l'architecture. Chaque cœur peut gérer quelques threads exécutés simultanément dans une succession rapide (similaire à l'hyperthreading dans le CPU).

De plus, chaque SM comporte un ou plusieurs ordonnanceurs de chaîne . Chaque planificateur envoie une seule instruction à plusieurs cœurs CUDA. Cela amène le SM à fonctionner en mode SIMD à 32 largeurs .

Modèle d'exécution CUDA

La structure physique du GPU a une influence directe sur la manière dont les noyaux sont exécutés sur le périphérique et sur la façon dont ils sont programmés dans CUDA. Le noyau est appelé avec une configuration d'appel qui spécifie le nombre de threads parallèles générés.

  • the grid - représente tous les threads générés lors de l'appel du noyau. Il est spécifié comme un ensemble de blocs de 1 ou 2 dimensions
  • le bloc - est un ensemble semi-indépendant de threads . Chaque bloc est affecté à un seul SM. En tant que tels, les blocs ne peuvent communiquer que par la mémoire globale. Les blocs ne sont synchronisés en aucune façon. S'il y a trop de blocs, certains peuvent être exécutés séquentiellement après d'autres. D'un autre côté, si les ressources le permettent, plusieurs blocs peuvent s'exécuter sur le même serveur de stockage, mais le programmeur ne peut pas en bénéficier (sauf pour l'amélioration des performances évidentes).
  • le thread - une séquence scalaire d'instructions exécutées par un seul cœur CUDA. Les threads sont «légers» avec un contexte minimal, permettant au matériel de les échanger rapidement. En raison de leur nombre, les threads CUDA fonctionnent avec quelques registres qui leur sont assignés et une pile très courte (de préférence aucune du tout!). Pour cette raison, le compilateur CUDA préfère incorporer tous les appels de fonctions pour aplatir le noyau afin qu'il ne contienne que des sauts et des boucles statiques. Les appels de fonction et les appels de méthode virtuels, bien que pris en charge par la plupart des nouveaux périphériques, entraînent généralement une pénalité majeure en termes de performances.

Chaque thread est identifié par un index de bloc blockIdx et un index de thread dans le threadIdx . Ces nombres peuvent être vérifiés à tout moment par n'importe quel thread en cours d'exécution et constituent le seul moyen de distinguer un thread d'un autre.

De plus, les threads sont organisés en chaînes , chacune contenant exactement 32 threads. Les threads au sein d'une même chaîne s'exécutent dans une synchronisation parfaite, en mode SIMD. Les threads provenant de différentes chaînes, mais dans le même bloc, peuvent s'exécuter dans n'importe quel ordre, mais peuvent être forcés de se synchroniser par le programmeur. Les threads provenant de différents blocs ne peuvent pas être synchronisés ou interagir directement de quelque manière que ce soit.

Organisation de la mémoire

Dans la programmation normale du processeur, l’organisation de la mémoire est généralement masquée par le programmeur. Les programmes typiques agissent comme s'il n'y avait que de la RAM. Toutes les opérations de mémoire, telles que la gestion des registres, l'utilisation de la mise en cache L1-L2-L3, la permutation sur disque, etc. sont gérées par le compilateur, le système d'exploitation ou le matériel lui-même.

Ce n'est pas le cas avec CUDA. Alors que les nouveaux modèles de GPU cachent partiellement le fardeau, par exemple via la mémoire unifiée de CUDA 6, il est toujours utile de comprendre l’organisation pour des raisons de performances. La structure de base de la mémoire CUDA est la suivante:

  • Mémoire hôte - la RAM normale. Principalement utilisé par le code hôte, mais les nouveaux modèles de GPU peuvent également y accéder. Lorsqu'un noyau accède à la mémoire de l'hôte, le processeur graphique doit communiquer avec la carte mère, généralement via le connecteur PCIe, ce qui le rend relativement lent.
  • Mémoire de l'appareil / Mémoire globale - la mémoire hors puce principale du GPU, disponible pour tous les threads.
  • Mémoire partagée - située dans chaque SM permet un accès beaucoup plus rapide que global. La mémoire partagée est privée à chaque bloc. Les threads d'un même bloc peuvent l'utiliser pour la communication.
  • Registers - Mémoire la plus rapide, privée et non adressable de chaque thread. En général, ils ne peuvent pas être utilisés pour la communication, mais quelques fonctions intrinsèques permettent de mélanger leur contenu dans une chaîne.
  • La mémoire locale - mémoire privée de chaque fil qui est adressable. Ceci est utilisé pour les déversements de registres et les tableaux locaux avec indexation variable. Physiquement, ils résident dans la mémoire globale.
  • Mémoire de texture, mémoire constante - une partie de la mémoire globale marquée comme immuable pour le noyau. Cela permet au GPU d'utiliser des caches spéciaux.
  • Cache L2 - sur puce, disponible pour tous les threads. Compte tenu de la quantité de threads, la durée de vie attendue de chaque ligne de cache est nettement inférieure à celle du processeur. Il est principalement utilisé des modèles d'accès à la mémoire mal alignés et partiellement aléatoires.
  • Cache L1 - situé dans le même espace que la mémoire partagée. Encore une fois, la quantité est plutôt petite, étant donné le nombre de threads qui l'utilisent, ne vous attendez donc pas à ce que les données y restent longtemps. La mise en cache L1 peut être désactivée.

Versions

Capacité de calcul Architecture Nom de code GPU Date de sortie
1.0 Tesla G80 2006-11-08
1.1 Tesla G84, G86, G92, G94, G96, G98, 2007-04-17
1.2 Tesla GT218, GT216, GT215 2009-04-01
1.3 Tesla GT200, GT200b 2009-04-09
2.0 Fermi GF100, GF110 2010-03-26
2.1 Fermi GF104, GF106 GF108, GF114, GF116, GF117, GF119 2010-07-12
3.0 Kepler GK104, GK106, GK107 2012-03-22
3.2 Kepler GK20A 2014-04-01
3.5 Kepler GK110, GK208 2013-02-19
3.7 Kepler GK210 2014-11-17
5.0 Maxwell GM107, GM108 2014-02-18
5.2 Maxwell GM200, GM204, GM206 2014-09-18
5.3 Maxwell GM20B 2015-04-02
6,0 Pascal GP100 2016-10-01
6.1 Pascal GP102, GP104, GP106 2016-05-27

La date de publication marque la sortie du premier processeur graphique prenant en charge une capacité de calcul donnée. Certaines dates sont approximatives, par exemple la carte 3.2 a été publiée au deuxième trimestre 2014.

Compiler et exécuter les exemples de programmes

Le guide d'installation NVIDIA se termine par l'exécution des exemples de programmes pour vérifier votre installation de CUDA Toolkit, mais n'indique pas explicitement comment. Commencez par vérifier toutes les conditions préalables. Vérifiez le répertoire CUDA par défaut pour les exemples de programmes. S'il n'est pas présent, il peut être téléchargé à partir du site Web officiel de CUDA. Accédez au répertoire où les exemples sont présents.

$ cd /path/to/samples/
$ ls
 

Vous devriez voir une sortie similaire à:

0_Simple     2_Graphics  4_Finance      6_Advanced       bin     EULA.txt
1_Utilities  3_Imaging   5_Simulations  7_CUDALibraries  common  Makefile
 

Assurez-vous que le Makefile est présent dans ce répertoire. La commande make dans les systèmes UNIX générera tous les exemples de programmes. Vous pouvez également accéder à un sous-répertoire dans lequel un autre Makefile est présent et exécuter la commande make partir de là pour créer uniquement cet échantillon.

Exécutez les deux exemples de programmes deviceQuery - deviceQuery et bandwidthTest :

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 
 

La sortie sera similaire à celle ci-dessous:

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 950M"
  CUDA Driver Version / Runtime Version          7.5 / 7.5
  CUDA Capability Major/Minor version number:    5.0
  Total amount of global memory:                 4096 MBytes (4294836224 bytes)
  ( 5) Multiprocessors, (128) CUDA Cores/MP:     640 CUDA Cores
  GPU Max Clock rate:                            1124 MHz (1.12 GHz)
  Memory Clock rate:                             900 Mhz
  Memory Bus Width:                              128-bit
  L2 Cache Size:                                 2097152 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 7.5, CUDA Runtime Version = 7.5, NumDevs = 1, Device0 = GeForce GTX 950M
Result = PASS
 

L'instruction Result = PASS à la fin indique que tout fonctionne correctement. Maintenant, exécutez l’autre exemple de bandwidthTest programme suggérée d’une manière similaire. La sortie sera similaire à:

[CUDA Bandwidth Test] - Starting...
Running on...

 Device 0: GeForce GTX 950M
 Quick Mode

 Host to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10604.5

 Device to Host Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            10202.0

 Device to Device Bandwidth, 1 Device(s)
 PINNED Memory Transfers
   Transfer Size (Bytes)    Bandwidth(MB/s)
   33554432            23389.7

Result = PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
 

Encore une fois, l'instruction Result = PASS indique que tout a été exécuté correctement. Tous les autres programmes exemples peuvent être exécutés de la même manière.

Lançons un seul thread CUDA pour dire bonjour

Ce programme CUDA simple montre comment écrire une fonction qui s'exécutera sur le GPU (aka "device"). Le CPU, ou "host", crée des threads CUDA en appelant des fonctions spéciales appelées "noyaux". Les programmes CUDA sont des programmes C ++ avec une syntaxe supplémentaire.

Pour voir comment cela fonctionne, placez le code suivant dans un fichier nommé hello.cu :

#include <stdio.h>

// __global__ functions, or "kernels", execute on the device
__global__ void hello_kernel(void)
{
  printf("Hello, world from the device!\n");
}

int main(void)
{
  // greet from the host
  printf("Hello, world from the host!\n");

  // launch a kernel with a single thread to greet from the device
  hello_kernel<<<1,1>>>();

  // wait for the device to finish so that we see the message
  cudaDeviceSynchronize();

  return 0;
}
 

(Notez que pour utiliser la fonction printf sur le périphérique, vous avez besoin d'un périphérique ayant une capacité de calcul d'au moins 2.0. Voir la présentation des versions pour plus de détails.)

Maintenant, compilons le programme en utilisant le compilateur NVIDIA et exécutons-le:

$ nvcc hello.cu -o hello
$ ./hello
Hello, world from the host!
Hello, world from the device!
 

Quelques informations supplémentaires sur l'exemple ci-dessus:

  • nvcc signifie "NVIDIA CUDA Compiler". Il sépare le code source en composants hôte et périphérique.
  • __global__ est un mot-clé CUDA utilisé dans les déclarations de fonction indiquant que la fonction s'exécute sur le périphérique GPU et est appelée depuis l'hôte.
  • Les crochets à trois angles ( <<< , >>> ) marquent un appel du code hôte vers le code du périphérique (également appelé «lancement du noyau»). Les nombres entre ces crochets indiquent le nombre de fois à exécuter en parallèle et le nombre de threads.

Conditions préalables

Pour démarrer la programmation avec CUDA, téléchargez et installez CUDA Toolkit et le pilote de développement . La boîte à outils comprend nvcc , le compilateur NVIDIA CUDA et d'autres logiciels nécessaires au développement d'applications CUDA. Le pilote garantit que les programmes GPU fonctionnent correctement sur le matériel compatible CUDA , dont vous aurez également besoin.

Vous pouvez confirmer que CUDA Toolkit est correctement installé sur votre machine en exécutant nvcc --version partir d'une ligne de commande. Par exemple, sur une machine Linux,

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jul_12_18:28:38_CDT_2016
Cuda compilation tools, release 8.0, V8.0.32
 

génère les informations du compilateur. Si la commande précédente a nvcc , le CUDA Toolkit n'est probablement pas installé ou le chemin d'accès à nvcc ( C:\CUDA\bin sur les machines Windows, /usr/local/cuda/bin sur les systèmes d'exploitation POSIX) ne fait pas partie de votre Variable d'environnement PATH .

De plus, vous aurez également besoin d'un compilateur hôte qui fonctionne avec nvcc pour compiler et construire des programmes CUDA. Sous Windows, il s'agit de cl.exe , le compilateur Microsoft, cl.exe avec Microsoft Visual Studio. Sur les systèmes d'exploitation POSIX, d'autres compilateurs sont disponibles, y compris gcc ou g++ . Le Guide de démarrage rapide CUDA officiel peut vous indiquer quelles versions du compilateur sont prises en charge sur votre plate-forme particulière.

Pour vous assurer que tout est configuré correctement, compilons et exécutons un programme CUDA trivial pour nous assurer que tous les outils fonctionnent correctement ensemble.

__global__ void foo() {}

int main()
{
  foo<<<1,1>>>();

  cudaDeviceSynchronize();
  printf("CUDA error: %s\n", cudaGetErrorString(cudaGetLastError()));

  return 0;
}
 

Pour compiler ce programme, copiez-le dans un fichier appelé test.cu et compilez-le à partir de la ligne de commande. Par exemple, sur un système Linux, les éléments suivants devraient fonctionner:

$ nvcc test.cu -o test
$ ./test
CUDA error: no error
 

Si le programme réussit sans erreur, alors commençons à coder!

Sommez deux tableaux avec CUDA

Cet exemple montre comment créer un programme simple qui additionnera deux tableaux int à CUDA.

Un programme CUDA est hétérogène et comprend des parties exécutées à la fois sur le processeur et sur le GPU.

Les parties principales d’un programme utilisant CUDA sont similaires aux programmes CPU et se composent de

  • Allocation de mémoire pour les données qui seront utilisées sur le GPU
  • Copie de données de la mémoire hôte vers la mémoire GPU
  • Invoquer la fonction du noyau pour traiter des données
  • Copier le résultat dans la mémoire de la CPU

Pour allouer la mémoire des périphériques, nous utilisons la fonction cudaMalloc . Pour copier des données entre le périphérique et l'hôte, la fonction cudaMemcpy peut être utilisée. Le dernier argument de cudaMemcpy spécifie la direction de l'opération de copie. Il y a 5 types possibles:

  • cudaMemcpyHostToHost - Hôte -> Hôte
  • cudaMemcpyHostToDevice - Hôte -> Périphérique
  • cudaMemcpyDeviceToHost - Périphérique -> Hôte
  • cudaMemcpyDeviceToDevice - Périphérique -> Périphérique
  • cudaMemcpyDefault - Espace d'adressage virtuel unifié par défaut

Ensuite, la fonction du noyau est appelée. L'information entre les chevrons triples est la configuration d'exécution, qui dicte combien de threads de périphérique exécutent le noyau en parallèle. Le premier nombre ( 2 dans l'exemple) spécifie le nombre de blocs et le second ( (size + 1) / 2 dans l'exemple) - nombre de threads dans un bloc. Notez que dans cet exemple nous ajoutons 1 à la taille, de sorte que nous demandons un thread supplémentaire plutôt que d'avoir un thread responsable de deux éléments.

Comme l'invocation du noyau est une fonction asynchrone, cudaDeviceSynchronize est appelée pour attendre que l'exécution soit terminée. Les tableaux de résultats sont copiés dans la mémoire hôte et toute la mémoire allouée sur le périphérique est libérée avec cudaFree .

Pour définir la fonction comme noyau, le __global__ déclaration __global__ est utilisé. Cette fonction sera appelée par chaque thread. Si nous voulons que chaque thread traite un élément du tableau résultant, nous avons besoin d'un moyen de distinguer et d'identifier chaque thread. CUDA définit les variables blockDim , blockIdx et threadIdx . La variable prédéfinie blockDim contient les dimensions de chaque bloc de thread comme spécifié dans le deuxième paramètre de configuration d'exécution pour le lancement du noyau. Les variables prédéfinies threadIdx et blockIdx contiennent respectivement l'index du thread dans son bloc de thread et le bloc de thread dans la grille. Notez que puisque nous demandons potentiellement un thread de plus que des éléments dans les tableaux, nous devons passer en size pour nous assurer de ne pas accéder au-delà de la fin du tableau.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__global__ void addKernel(int* c, const int* a, const int* b, int size) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < size) {
        c[i] = a[i] + b[i];
    }
}

// Helper function for using CUDA to add vectors in parallel.
void addWithCuda(int* c, const int* a, const int* b, int size) {
    int* dev_a = nullptr;
    int* dev_b = nullptr;
    int* dev_c = nullptr;

    // Allocate GPU buffers for three vectors (two input, one output)
    cudaMalloc((void**)&dev_c, size * sizeof(int));
    cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaMalloc((void**)&dev_b, size * sizeof(int));

    // Copy input vectors from host memory to GPU buffers.
    cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
    // 2 is number of computational blocks and (size + 1) / 2 is a number of threads in a block
    addKernel<<<2, (size + 1) / 2>>>(dev_c, dev_a, dev_b, size);
    
    // cudaDeviceSynchronize waits for the kernel to finish, and returns
    // any errors encountered during the launch.
    cudaDeviceSynchronize();

    // Copy output vector from GPU buffer to host memory.
    cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);
}

int main(int argc, char** argv) {
    const int arraySize = 5;
    const int a[arraySize] = {  1,  2,  3,  4,  5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    addWithCuda(c, a, b, arraySize);

    printf("{1, 2, 3, 4, 5} + {10, 20, 30, 40, 50} = {%d, %d, %d, %d, %d}\n", c[0], c[1], c[2], c[3], c[4]);

    cudaDeviceReset();

    return 0;
}