CUDA Programmierung

Aus Thomas-Krenn-Wiki
Wechseln zu: Navigation, Suche

Für den Einstieg in die CUDA Programmierung sind Grundkenntnisse der Programmiersprache C sehr hilfreich. Auch gilt es zu verstehen, welche Vorarbeiten die CPU durchführen muss und welche Teile die GPU berechnet. Der wichtigste Punkt aber ist die Parallelisierung des vorhandenen Algorithmus, sodass dieser von den Threads optimal ausgeführt werden kann.

Die Parallelisierung eines Algorithmus erfordert einiges an Umdenken zur herkömmlichen seriellen Programmierung. Denn im Gegensatz zur herkömmlichen Programmierung führen alle gestarteten Threads die Kernel-Funktion parallel aus! Das heißt Datenzugriffe müssen unter Umständen auf Race Conditions analysiert werden bzw. muss der optimale parallel Weg zur Berechnung der Lösung gefunden werden.

Von der Programmiersprache her genügt es die CUDA spezifischen Erweiterungen zu kennen und anwenden zu können. Ansonsten können durchgehend C bzw. Teile von C++ verwendet werden. Es folgen nun zwei kleinere Beispiele die demonstrieren wie ein einfacher Kernel auf der GPU ausgeführt werden kann. Auch wird das Zusammenspiel zwischen nvcc und gcc kurz vorgezeigt. Das verwendete Betriebssystem ist ein Ubuntu 10.04 64bit mit CUDA Version 3.2 und aktuellem Nvidia Developer Treiber.

Ein erstes Beispiel

In diesem Beispiel, ähnlich dem Beispiel "vectorAdd" aus dem CUDA-SDK [1], werden die Elemente zweier Vektoren miteinander addiert und in einem dritten Vektor gespeichert. Jeder Thread auf der GPU addiert hierbei 2 Elemente und schreibt das Ergebnis in den Vektor C. Zum Beispiel führt der Thread mit dem threadIdx.x 10 folgende Addition aus:

C[10] = A[10] + B[10]

Parallel dazu führen auch alle anderen gestarteten Threads, die natürlich einen anderen threadIdx.x besitzen und somit mit anderen Feld-Elementen operieren, dieselbe Operation aus. Auf der CPU würden diese Operationen typischerweise in einer Schleife abgearbeitet werden. Auf der GPU entfällt die Schleife, da die Threads die Operationen der Schleife durch Parallelität ersetzen.

#include <cuda.h>
#include <stdio.h>
#include <stdlib.h>

// Variables
//Host Variablen bekommen den Prefix h
int* h_A;
int* h_B;
int* h_C;
//Device Variablen bekommen den Prefix d
int* d_A;
int* d_B;
int* d_C;

//Forward Deklaration der Funktionen
void RandomInit(int* data, int n);
int CheckResults(int* A, int* B, int* C, int n);

// Device code (Kernel, GPU)
__global__ void VecAdd(const int* A, const int* B, int* C, int N)
{
    /*
     *Die Variable i dient für den Zugriff auf das Array. Da jeder Thread die Funktion VecAdd
     *ausführt, muss i für jeden Thread unterschiedlich sein. Ansonsten würden unterschiedliche
     *Threads auf denselben Index im Array schreiben. blockDim.x ist die Anzahl der Threads der x-Komponente
     *des Blocks, blockIdx.x ist die x-Koordinate des aktuellen Blocks und threadIdx.x ist die x-Koordinate des
     *Threads, der die Funktion gerade ausführt.
    */
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    
    /*
     *Jeder Thread addiert nun seine Indices der Arrays. Die Variable i setzt sich aus der Block-Dimension, 
     *dem Block-Index und dem Thread-Index zusammen und ist somit für jeden Thread unterschiedlich. Auf der CPU
     *müsste diese Zeile als Schleife für jedes Vektor-Element abgearbeitet werden
    */
    if (i < N)
        C[i] = A[i] + B[i];
}

// Host Code (CPU)
int main(int argc, char** argv)
{
    printf("Vector addition\n");
    int i;
    int N = 50000;
    size_t size = N * sizeof(int);

    // Speicher am Host wird allokiert
    h_A = (int*)malloc(size);
    h_B = (int*)malloc(size);
    h_C = (int*)malloc(size);

    // Variablen werden mit Zufallszahlen initialisiert
    RandomInit(h_A, N);
    RandomInit(h_B, N);

    // Speicher am Device (GPU) wird allokiert
    cudaMalloc((void**)&d_A, size);
    cudaMalloc((void**)&d_B, size);
    cudaMalloc((void**)&d_C, size);

    // Die Vektoren werden vom Host zum Device kopiert
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // Der Kernelaufruf erfolgt
    //Festlegung der Threads pro Block
    int threadsPerBlock = 256;
    //Es werden soviele Blöcke benötigt, dass alle Elemente der Vektoren abgearbeitet werden können
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    //Der Kernel wird gestartet
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    //Das Ergebnis wird zurück auf den Host kopiert
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    //Überprüfung der Ergebnisse
    if(CheckResults(h_A,h_B,h_C,N)==0)
	printf("\nCPU-Überprüfung wurde erfolgreich durchgeführt!\n");
    else
	printf("\nGPU- und CPU-Addition stimmen nicht überein!\n");
    
    //Freigeben der Speicher
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);

    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

// Vektoren werden mit zufälligen Integer-Werten befüllt.
void RandomInit(int* data, int n)
{
    for (int i = 0; i < n; ++i)
        data[i] = rand() % (int) 100;
}

int CheckResults(int* A, int* B, int* C, int n){

    int i;
    for(i=0;i<n;i++){
    	if((A[i]+B[i]) != C[i])
	    return -1;
    }
    return 0;
}

Wird dieser Code in einer Quelldatei "vector_add.cu" abgespeichert und mit dem Befehl

nvcc -o vector_add -lcudart vector_add.cu

kompiliert, so kann anschließend die Applikation mit

./vector_add

ausgeführt werden.

Allgemein empfiehlt sich zu Beginn vor allem die einfachen Beispiele, die mit dem CUDA-SDK mitgeliefert werden, durchzugehen. In Verbindung mit dem CUDA C Programming Guide wird der Start in die CUDA-Programmierung erheblich erleichtert. Auch im Blog zu Supercomputing for the Masses finden sich einige sehr hilfreiche Informationen und Beispiele zur CUDA Programmierung.

CUDA Fehlerbehandlung

Da es vor allem zu Beginn des häufigeren zu Fehlern kommen kann, ist eine robuste Fehlerbehandlung durchaus notwendig. Alle CUDA Runtime-Funktionen geben einen Error-Code zurück, der ausgelesen werden kann. Wichtig ist es zu beachten, dass im Falle der Verwendung von asynchronen Funktionen die Fehlerbehandlung ohne Synchronisierung nicht richtig durchgeführt werden kann. Bei asynchronen Funktionen wird die Kontrolle sofort nach dem Aufruf wieder an den Host übergeben, auch wenn die aufgerufene Funktion ihre Tätigkeiten noch nicht beendet hat. Asynchrone Funktionen sind z.B.:

  • Kernel Launches
  • Device-Device Speicherkopiervorgänge
  • Host-Device Speicherkopiervorgänge für 64KB Blöcke oder weniger
  • Speicherkopiervorgänge die durch Funktionen mit dem Prefix "Async" durchgeführt werden
  • Memory Set Funktionen

Das heißt tritt z.B. bei einem Kernel Launch ein Fehler auf dem Device so kann dieser Error Code nicht reported werden, da die asynchrone Funktion die Kontrolle wieder an den Host zurückgegeben hat. Die einzige Möglichkeit, um den Fehler nach dem Aufruf der asynchronen Funktion zu behandeln, ist eine explizite Synchronisierung. Diese kann z.B. durch den Aufruf der Funktion "cudaThreadSynchronize()" durchgeführt werden. So könnte die Fehlerbehandlung nach einem Kernel Launch dann aussehen:

    //Der Kernel wird gestartet
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    //Explizite Synchronisierung
    cudaThreadSynchronize();
    //Fehlerbehandlung
    cudaError_t err = cudaGetLastError();
    if( err != cudaSuccess){
        fprintf(stderr, "Cuda launch error: %s\n",cudaGetErrorString(err) );
        exit(EXIT_FAILURE);
    } 
    //Das Ergebnis wird zurück auf den Host kopiert
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

Integration in ein C/C++-Projekt

Wie im ersten Abschnitt erwähnt wird für die Kompilierung der Cuda-Quelldateien der Compiler "nvcc" verwendet. Hervorzuheben ist, dass nvcc Object Files generiert, die mit dem Standard-Linker gelinkt werden können.

Angenommen eine Applikation besitzt folgende Quelldateien:

  • cuda_test.c: Diese Datei enthält die main-Funktion, aus der der Kernel aufgerufen werden soll.
  • vector_add.cu: Diese Datei enthält den Kernel, der die Addition von 2 Vektoren vornimmt. Des weiteren besitzt die Datei eine Funktion, die die benötigten Vorarbeiten durchführt und dann den Kernel aufruft.

vector_add.cu:

#include <cuda.h>

// Device code (Kernel, GPU)
__global__ void VecAdd(const int* A, const int* B, int* C, int N)
{
   int i = blockDim.x * blockIdx.x + threadIdx.x;
    
    //Jeder Thread addiert nun seine Indices der Arrays
    if (i < N)
        C[i] = A[i] + B[i];
}

/* Interface zum Kernel - führt benötigte Vorarbeiten durch und wird
aus der main-Funktion aufgerufen*/
extern "C" int vector_add(int* h_A,int* h_B,int* h_C,int N)
{
    //Device Variablen bekommen den Prefix d
    int* d_A;
    int* d_B;
    int* d_C;
    
    //Berechnung der Größe der Vektoren
    size_t size = N * sizeof(int);

    //Speicher am Device (GPU) wird allokiert
    cudaMalloc((void**)&d_A, size);
    cudaMalloc((void**)&d_B, size);
    cudaMalloc((void**)&d_C, size);

    //Die Vektoren werden vom Host zum Device kopiert
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    //Der Kernelaufruf erfolgt
    //Festlegung der Threads pro Block
    int threadsPerBlock = 256;
    //Es werden soviele Blöcke benötigt, dass alle Elemente der Vektoren abgearbeitet werden können
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    //Der Kernel wird gestartet
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    //Das Ergebnis wird zurück auf den Host kopiert
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    //Freigeben der Speicher
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    
    return 0;
}

cuda_test.c:

#include <stdio.h>
#include <stdlib.h>

// Variables
//Host Variablen bekommen den Prefix h
int* h_A;
int* h_B;
int* h_C;

//Forward deklaration der Funktionen
void RandomInit(int* data, int n);
int CheckResults(int* A, int* B, int* C, int n);
//Forward deklaration des Interfaces zum Kernel
int vector_add(int* h_A,int* h_B,int* h_C,int N);

//Host Code (CPU)
int main(int argc, char** argv)
{
    printf("Vector Addition\n");
    int N = 50000;
    size_t size = N * sizeof(int);

    // Speicher am Host wird allokiert
    h_A = (int*)malloc(size);
    h_B = (int*)malloc(size);
    h_C = (int*)malloc(size);

    //Variablen werden mit Zufallszahlen initialisiert
    RandomInit(h_A, N);
    RandomInit(h_B, N);
    
    printf("vector_add wird aufgerufen...\n");
    //Interface zum Kernel wird aufgerunfen. Die Funktion startet den Kernel auf der GPU
    vector_add(h_A, h_B, h_C, N);

    //Überprüfung der Ergebnisse
    if(CheckResults(h_A,h_B,h_C,N)==0)
	printf("CPU-Überprüfung wurde erfolgreich durchgeführt!\n");
    else
	printf("GPU- und CPU-Addition stimmen nicht überein!\n");
	
    free(h_A);
    free(h_B);
    free(h_C);
    
    return 0;
}

// Vektoren werden mit zufälligen Integer-Werten befüllt.
void RandomInit(int* data, int n)
{
    int i;
    for (i = 0; i < n; ++i)
        data[i] = rand() % (int) 100;
}

int CheckResults(int* A, int* B, int* C, int n){

    int i;
    for(i=0;i<n;i++){
    	if((A[i]+B[i]) != C[i])
	    return -1;
    }
    return 0;
}

Im nächsten Schritt werden beide Daten kompiliert und anschließend gelinkt:

nvcc -c -o vector_add.o vector_add.cu
gcc -c -o cuda_test.o cuda_test.c

Beim Linken ist darauf zu achten, dass man die CUDA-Bibliotheken mit angibt:

gcc -o cuda_test cuda_test.o vector_add.o -L/usr/local/cuda/lib -lcudart

Das dabei entstandene Binary wird mittels

./cuda_test

ausgeführt.

Einzelnachweise

  1. CUDA Webseite der SDK-Beispiele

Das könnte Sie auch interessieren

CUDA
CUDA Installation
PNY NVIDIA Quadro Grafikkarten