cuda Tutorial
Iniziare con cuda

Osservazioni

Versioni

Prerequisiti

Per iniziare a programmare con CUDA, scarica e installa CUDA Toolkit e il driver dello sviluppatore . Il toolkit include nvcc , il compilatore NVIDIA CUDA e altro software necessario per sviluppare applicazioni CUDA. Il driver assicura che i programmi GPU vengano eseguiti correttamente su hardware compatibile CUDA , di cui avrete anche bisogno.

È possibile verificare che CUDA Toolkit sia installato correttamente sulla macchina eseguendo nvcc --version da una riga di comando. Ad esempio, su una macchina 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

restituisce le informazioni del compilatore. Se il comando precedente non ha avuto esito positivo, probabilmente il CUDA Toolkit non è installato, oppure il percorso di nvcc ( C:\CUDA\bin su macchine Windows, /usr/local/cuda/bin su sistemi POSIX) non fa parte del tuo Variabile d'ambiente PATH .

Inoltre, avrai anche bisogno di un compilatore host che lavori con nvcc per compilare e creare programmi CUDA. Su Windows, questo è cl.exe , il compilatore Microsoft, fornito con Microsoft Visual Studio. Nei sistemi operativi POSIX sono disponibili altri compilatori, inclusi gcc o g++ . La Guida rapida ufficiale CUDA può dirti quali versioni del compilatore sono supportate sulla tua piattaforma specifica.

Per assicurarci che tutto sia impostato correttamente, compila ed esegui un banale programma CUDA per assicurarti che tutti gli strumenti funzionino correttamente.

__global__ void foo() {}

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

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

  return 0;
}

Per compilare questo programma, copialo in un file chiamato test.cu e compilarlo dalla riga di comando. Ad esempio, su un sistema Linux, il seguente dovrebbe funzionare:

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

Se il programma ha esito positivo senza errori, allora iniziamo a programmare!

Sommare due array con CUDA

Questo esempio illustra come creare un semplice programma che sommerà due array int con CUDA.

Un programma CUDA è eterogeneo e consiste di parti eseguite sia su CPU che su GPU.

Le parti principali di un programma che utilizzano CUDA sono simili ai programmi della CPU e consistono in

• Allocazione di memoria per i dati che verranno utilizzati sulla GPU
• Copia dei dati dalla memoria dell'host alla memoria delle GPU
• Richiamo della funzione del kernel per elaborare i dati
• Copia il risultato nella memoria della CPU

Per allocare la memoria dei dispositivi cudaMalloc funzione cudaMalloc . Per copiare i dati tra dispositivo e host è possibile utilizzare la funzione cudaMemcpy . L'ultimo argomento di cudaMemcpy specifica la direzione dell'operazione di copia. Ci sono 5 tipi possibili:

• cudaMemcpyHostToHost - Host -> Host
• cudaMemcpyHostToDevice - Host -> Dispositivo
• cudaMemcpyDeviceToHost - Dispositivo -> Host
• cudaMemcpyDeviceToDevice - Dispositivo -> Dispositivo
• cudaMemcpyDefault : spazio di indirizzi virtuali unificato basato su predefinito

Successivamente viene richiamata la funzione del kernel. Le informazioni tra i triple chevrons sono la configurazione di esecuzione, che stabilisce quanti thread di dispositivo eseguono il kernel in parallelo. Il primo numero ( 2 nell'esempio) specifica il numero di blocchi e il secondo ( (size + 1) / 2 nell'esempio): numero di thread in un blocco. Nota che in questo esempio aggiungiamo 1 alla dimensione, in modo da richiedere un thread in più anziché avere un thread responsabile per due elementi.

Poiché la chiamata al kernel è una funzione asincrona, cudaDeviceSynchronize viene chiamato ad attendere fino al completamento dell'esecuzione. Gli array di risultati vengono copiati nella memoria dell'host e tutta la memoria allocata sul dispositivo viene liberata con cudaFree .

Per definire la funzione viene usato lo specificatore della dichiarazione __global__ kernel. Questa funzione sarà invocata da ogni thread. Se vogliamo che ogni thread elabori un elemento dell'array risultante, abbiamo bisogno di un mezzo per distinguere e identificare ciascun thread. CUDA definisce le variabili blockDim , blockIdx e threadIdx . Il blockDim variabile predefinitoDim contiene le dimensioni di ciascun blocco di thread come specificato nel secondo parametro di configurazione dell'esecuzione per l'avvio del kernel. Le variabili predefinite threadIdx e blockIdx contengono l'indice del thread all'interno del suo blocco di thread e il blocco di thread all'interno della griglia, rispettivamente. Nota che, poiché potenzialmente richiediamo un thread in più rispetto agli elementi negli array, dobbiamo passare in size per assicurarci di non accedere oltre la fine dell'array.

#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;
}

Lanciamo un singolo thread CUDA per dire ciao

Questo semplice programma CUDA dimostra come scrivere una funzione che verrà eseguita sulla GPU (ovvero "dispositivo"). La CPU, o "host", crea i thread CUDA chiamando funzioni speciali chiamate "kernel". I programmi CUDA sono programmi C ++ con sintassi aggiuntiva.

Per vedere come funziona, inserisci il seguente codice in un file chiamato 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;
}

(Si noti che per utilizzare la funzione printf sul dispositivo, è necessario un dispositivo con una capacità di calcolo di almeno 2.0. Vedere la panoramica delle versioni per i dettagli.)

Ora compiliamo il programma usando il compilatore NVIDIA ed eseguiamolo:

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

Alcune informazioni aggiuntive sull'esempio precedente:

• nvcc sta per "NVIDIA CUDA Compiler". Separa il codice sorgente in componenti host e dispositivo.
• __global__ è una parola chiave CUDA utilizzata nelle dichiarazioni di funzione che indica che la funzione viene eseguita sul dispositivo GPU e viene chiamata dall'host.
• Le parentesi angolari triple ( <<< , >>> ) contrassegnano una chiamata dal codice host al codice dispositivo (chiamato anche "avvio kernel"). I numeri all'interno di queste parentesi quadre indicano il numero di volte da eseguire in parallelo e il numero di thread.

Compilazione ed esecuzione dei programmi di esempio

La guida all'installazione di NVIDIA termina con l'esecuzione dei programmi di esempio per verificare l'installazione di CUDA Toolkit, ma non indica esplicitamente come. Innanzitutto controlla tutti i prerequisiti. Controllare la directory CUDA predefinita per i programmi di esempio. Se non è presente, può essere scaricato dal sito web ufficiale di CUDA. Passare alla directory in cui sono presenti gli esempi.

$ cd /path/to/samples/
$ ls

Dovresti vedere un risultato simile a:

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

Assicurarsi che il Makefile sia presente in questa directory. Il comando make nei sistemi basati su UNIX costruirà tutti i programmi di esempio. In alternativa, accedere a una sottodirectory in cui è presente un altro Makefile ed eseguire il comando make da lì per creare solo quel campione.

Esegui i due programmi di esempio suggeriti: deviceQuery e bandwidthTest :

$ cd 1_Utilities/deviceQuery/
$ ./deviceQuery 

L'output sarà simile a quello mostrato di seguito:

./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'istruzione Result = PASS alla fine indica che tutto funziona correttamente. Ora, esegui l'altro programma di esempio sample bandwidthTest in modo simile. L'output sarà simile a:

[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.

Anche in questo caso, la dichiarazione Result = PASS indica che tutto è stato eseguito correttamente. Tutti gli altri programmi di esempio possono essere eseguiti in modo simile.