cuda Tutoriel
Commencer avec cuda
Recherche…
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.
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;
}
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.
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.