Démarrer avec cuda

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 (alias “périphérique”). Le CPU, ou « hôte », 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 doté d’une capacité de calcul d’au moins 2.0. Voir la [vue d’ensemble des versions] (https://www.wikiod.com/fr/cuda/demarrer-avec-cuda /introduction-to-cuda#t=20160918151348344166&a=versions) pour plus de détails.)

Compilons maintenant le programme à l’aide du 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 triples crochets (<<<,>>>) marquent un appel du code hôte au code périphérique (également appelé “lancement du noyau”). Les nombres entre ces triples parenthèses indiquent le nombre de fois à exécuter en parallèle et le nombre de threads.

Conditions préalables

Pour commencer à programmer avec CUDA, téléchargez et installez [CUDA Toolkit and developer driver][1]. 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 s’exécutent correctement sur [matériel compatible CUDA][2], dont vous aurez également besoin.

Vous pouvez vérifier que CUDA Toolkit est correctement installé sur votre ordinateur 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

affiche les informations du compilateur. Si la commande précédente n’a pas réussi, le CUDA Toolkit n’est probablement pas installé, ou le chemin vers 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, fourni avec Microsoft Visual Studio. Sur les systèmes d’exploitation POSIX, d’autres compilateurs sont disponibles, notamment gcc ou g++. Le [Quick Start Guide][1] officiel de CUDA peut vous indiquer quelles versions du compilateur sont prises en charge sur votre plate-forme particulière.

Pour vous assurer que tout est correctement configuré, 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 !

[1] : https://developer.nvidia.com/cuda-downloads [2] : http://www.nvidia.com/object/cuda_gpus.htm

Compilation et exécution des 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. Vérifiez d’abord tous les prérequis. 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 semblable à :

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 basés sur UNIX construira tous les exemples de programmes. Sinon, naviguez vers un sous-répertoire où un autre Makefile est présent et exécutez la commande make à partir de là pour ne construire que cet exemple.

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

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery

La sortie sera similaire à celle illustrée 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 programme suggéré “bandwidthTest” de la même manière. 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 exemples de programmes peuvent être exécutés de la même manière.

Somme de deux tableaux avec CUDA

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

Un programme CUDA est hétérogène et se compose de parties s’exécutant à la fois sur CPU et sur GPU.

Les parties principales d’un programme qui utilise CUDA sont similaires aux programmes CPU et consistent en

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

Pour allouer de la mémoire aux appareils, nous utilisons la fonction cudaMalloc. Pour copier des données entre l’appareil 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 existe 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 noyau est invoquée. L’information entre les triples chevrons 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 ((taille + 1) / 2 dans l’exemple) - le 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.

Étant donné que l’invocation du noyau est une fonction asynchrone, cudaDeviceSynchronize est appelé pour attendre la fin de l’exécution. Les tableaux de résultats sont copiés dans la mémoire hôte et toute la mémoire allouée sur l’appareil est libérée avec cudaFree.

Pour définir la fonction comme noyau, le spécificateur de 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 d’éléments dans les tableaux, nous devons passer size pour nous assurer que nous n’accédons pas 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;
}
Réduction parallèle (par exemple, comment additionner un tableau) →