kramann.info
© Guido Kramann

GPU parallel verwenden mit CUDA

(EN google-translate)

(PL google-translate)

Die GPU des Jetson Nano besteht aus 128 Einheiten, die parallel rechnen können. Als Beispiel soll die Berechnung der Ausgänge einer Neuronenschicht damit umgesetzt werden. Hinweise zu diesem Thema finden sich hier:

01_NeuronaleNetze/01_Neuron

Die N=128 Neuronen sind vollständig vernetzt und es gibt 128 Eingänge und Ausgänge. Das bedeutet, dass N*N=16384 Gewichtswerte w vorhanden sein müssen, die jeden Eingang mit jedem Ausgang verbinden.

Eine C-Funktion, die auf der CPU läuft und die geforderte Aufgabe erledigt sieht so aus:

void neuroSchicht(float* out, float* w, float* in)
{
    for(unsigned int i=0;i<N;i++)
    {
        out[i]=0.0;
        for(unsigned int k=0;k<N;k++)
        {
             out[i] += in[k] * w[k+i*N];
        }
        out[i]  = 1.0 / (1.0 + exp(-d*out[i]));
    }
}

Code 0-1: C-Funktion,

Dazu im Vergleich sieht die Funktion, die das gleiche macht, aber die Verarbeitung verteilt auf 128 Prozessoren ausführt so aus:

__global__ void neuroSchichtGPU(float* out, float* w, float* in) 
{
    int i = blockIdx.x;
    int k=0;
    if(i<N)
    {
        out[i]=0.0;
        for(k=0;k<N;k++)
        {
             out[i] += in[k] * w[k+i*N];
        }
        out[i]  = 1.0 / (1.0 + exp(-d*out[i]));
    }
}

Code 0-2: C-Funktion für GPU.

__global__ legt fest, dass es sich um eine GPU-Funktion handelt.
Die äußere Schleife, die alle Ausgangsneuronen durchgeht fehlt.
Stattdessen wird mit int i = blockIdx.x; ausgelesen, auf welchem Block der Code ausgeführt wird und der korrespondierende Output berechnet.

Prozesse können Blöcken und darin Threads zugewiesen werden.
Threads im gleichen Block können Daten untereinander austauschen, solche in verschiedenen Blöcken nicht.
Das Produkt aus Threads und Blöcken darf höchstens die maximale Anzahl paralleler Prozesse ergeben, hier 128.
Im nachfolgenden Programm wird die gleiche Berechnung aller Ausgänge der Neuronenschicht
sowohl auf der CPU, als auch mit 128 parallen Prozessen auf der GPU durchgeführt.
Schließlich wird sie auch noch mit nur einem Prozeß testweise auf der GPU ausgeführt.
Mit clock() wird die Dauer der Berechnungen jeweils gemessen.

#include <iostream>
#include <cuda_runtime.h>
#include <math.h>
#include <time.h>
#include <stdlib.h>

#define N 127
#define d 1.0

/*
    CPU-Funktion
    repräsentiert Neuronenschicht mit 
    N Neuronen also:
    N Eingängen in
    N Ausgängen out
    N*N Gewichten w

    
*/
void neuroSchicht(float* out, float* w, float* in)
{
    for(unsigned int i=0;i<N;i++)
    {
        out[i]=0.0;
        for(unsigned int k=0;k<N;k++)
        {
             out[i] += in[k] * w[k+i*N];
        }
        out[i]  = 1.0 / (1.0 + exp(-d*out[i]));
    }
}

//##### gleiche Funktion parallelisiert für N Blöcke, bei der die Block-ID bestimmt, welcher Output berechnet wird

__global__ void neuroSchichtGPU(float* out, float* w, float* in) 
{
    int i = blockIdx.x;
    int k=0;
    if(i<N)
    {
        out[i]=0.0;
        for(k=0;k<N;k++)
        {
             out[i] += in[k] * w[k+i*N];
        }
        out[i]  = 1.0 / (1.0 + exp(-d*out[i]));
    }
}

//##### Variante für nur EINE nicht parallelisierte GPU:
__global__ void neuroSchichtEinBlock(float* out, float* w, float* in)
{
    for(unsigned int i=0;i<N;i++)
    {
        out[i]=0.0;
        for(unsigned int k=0;k<N;k++)
        {
             out[i] += in[k] * w[k+i*N];
        }
        out[i]  = 1.0 / (1.0 + exp(-d*out[i]));
    }
}


int main() 
{
//##### Declare the variables
float *inCPU;
float *outCPU;
float *wCPU;
float *in;
float *out;
float *w;

//##### Speicher auf CPU allokieren:

inCPU = (float*)calloc(N,sizeof(float));
outCPU = (float*)calloc(N,sizeof(float));
wCPU = (float*)calloc(N*N,sizeof(float));

cudaMallocManaged(&in, N*sizeof(float));
cudaMallocManaged(&out, N*sizeof(float));
cudaMallocManaged(&w, N*N*sizeof(float));

// Seeding the random number generator
srand(1000);

//##### Gewichte und Inputs zufällig initialisieren
for (unsigned int i = 0; i < N; i++) 
{
    outCPU[i] = 0.0;
    inCPU[i] = ((float)rand()/(float)(RAND_MAX)); // ]0..1[
    for(unsigned int k = 0;k<N;k++)
    {
        wCPU[k+i*N] = 0.01*((float)rand()/(float)(RAND_MAX)); // ]0..0.01[
    }
}


//##### Berechnung 1000000 mal auf CPU ausführen

clock_t startzeit = clock();
//for (unsigned int i = 0; i < 1000000; i++)
    neuroSchicht(outCPU,wCPU,inCPU);
clock_t endzeit = clock();
double dt_ms = 1000.0*((double)(endzeit - startzeit)/ CLOCKS_PER_SEC);
std::cout << "1 Aufruf auf CPU dauerten "<< dt_ms<<" Millisekunden." << std::endl;


//##### Ergebnis auf Konsole ausgeben
/*
std::cout << "Outputwerte:" << std::endl;
for (unsigned int i = 0; i < N; i++) 
{
   std::cout<<"out("<<i<<")="<<outCPU[i]<<std::endl;    
}
*/

//##### Ausführung auf GPU in N Blöcken
//startzeit = clock();
//for (unsigned int i = 0; i < 1000000; i++)
//{
    //##### Variablen der CPU auf GPU-Arrays kopieren:
    cudaMemcpy(in,inCPU, N*sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpy(out,outCPU, N*sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpy(w,wCPU, N*N*sizeof(float),cudaMemcpyHostToDevice);


startzeit = clock();
    neuroSchichtEinBlock<<<1, 1>>>(out,w,in);
//    neuroSchichtGPU<<<N, 1>>>(out,w,in);
    //cudaDeviceSynchronize(); //Warten, bis alle Blöcke fertig sind
endzeit = clock();


    //##### Variablen der GPU auf CPU-Arrays zurück kopieren:
    cudaMemcpy(inCPU,in, N*sizeof(float),cudaMemcpyDeviceToHost);
    cudaMemcpy(outCPU,out, N*sizeof(float),cudaMemcpyDeviceToHost);
    cudaMemcpy(wCPU,w, N*N*sizeof(float),cudaMemcpyDeviceToHost);

//}
//endzeit = clock();
dt_ms = 1000.0*((double)(endzeit - startzeit)/ CLOCKS_PER_SEC);
std::cout << "1 Aufruf auf EINER GPU dauerten "<< dt_ms<<" Millisekunden." << std::endl;


//for (unsigned int i = 0; i < 1000000; i++)
//{
    //##### Variablen der CPU auf GPU-Arrays kopieren:
    cudaMemcpy(in,inCPU, N*sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpy(out,outCPU, N*sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpy(w,wCPU, N*N*sizeof(float),cudaMemcpyHostToDevice);


startzeit = clock();
//    neuroSchichtEinBlock<<<1, 1>>>(out,w,in);
    neuroSchichtGPU<<<N, 1>>>(out,w,in);
    //cudaDeviceSynchronize(); //Warten, bis alle Blöcke fertig sind
endzeit = clock();


    //##### Variablen der GPU auf CPU-Arrays zurück kopieren:
    cudaMemcpy(inCPU,in, N*sizeof(float),cudaMemcpyDeviceToHost);
    cudaMemcpy(outCPU,out, N*sizeof(float),cudaMemcpyDeviceToHost);
    cudaMemcpy(wCPU,w, N*N*sizeof(float),cudaMemcpyDeviceToHost);

//}
//endzeit = clock();
dt_ms = 1000.0*((double)(endzeit - startzeit)/ CLOCKS_PER_SEC);
std::cout << "1 Aufruf auf 127 GPU dauerten "<< dt_ms<<" Millisekunden." << std::endl;



//##### Ergebnis auf Konsole ausgeben
/*
std::cout << "Outputwerte:" << std::endl;
for (unsigned int i = 0; i < N; i++) 
{
   std::cout<<"out("<<i<<")="<<outCPU[i]<<std::endl;    
}
*/


//##### Releasing the memory

cudaFree(in);
cudaFree(out);
cudaFree(w);

free(inCPU);
free(outCPU);
free(wCPU);

std::cout << "Code Execution Completed" << std::endl;
return 0;
}

Code 0-3: neuro003.cu Gesamtprogramm zur Berechnung einer Neuronenschicht.

Wenn der Programmquelltext in einer Datei neuro003.cu steht, dann wird er folgendermaßen kompiliert und gestartet:
nvcc neuro003.cu -o neuro003
./neuro003

Code 0-4: Verwendung

Zunächst wird auf der CPU gestartet, dann Daten für die GPU vorbereitet und dann erst die GPU-Funktion gestartet.
Die ganzen Speicher-Allokierungen und -freigaben müssen so vorgenommen werden, da die Speicherverwltung von CPU und GPU getrennt ist und spezielle Kopierfunktionen nötig sind.
Die GPU-Aufrufe geben in den spitzen Klammern an wieviele Blöcke B und wieviele Threads T benutzt werden <<>>.

Ergebnisse

Nun würde man erwarten, dass die Abarbeitung auf der GPU 100mal so schnell geht. Jedoch ist diese nur circa vier mal so schnell:

1 Aufruf auf CPU dauerten 0.284 Millisekunden.
1 Aufruf auf EINER GPU dauerten 0.065 Millisekunden.
1 Aufruf auf 127 GPU dauerten 0.07 Millisekunden.

Code 0-5: Ergenisse

Mögliche Ursachen (steht in Frage):

Verwendung mathermatischer float-Operationen?
Keine tatsächliche Parallelisierung?
Keine optimale Ausnutzung der Parallelisierungsmöglichkeiten?
Fake-Gerät?

CUDA2.zip -- alle Testdateien in diesem Zusammenhang.

Weitere Schritte: Auch hierfür eine JNI basierte Library für Processing erstellen.