C CUDA Einstieg bzw. Frage

[KillerPinockel]

Hallo,

da ich mich nun etwas mit CUDA Programmierung beschäftigt habe, wollte ich gern mal einen kleinen Test machen.

Hier mein CODE

#define SIZE 216000

__global__ void calc(double *dev_data, double *dev_result) {
	dev_result = dev_data + 1;

}

/**
 * Host function that prepares data array and passes it to the CUDA kernel.
 */
int main(void) {
	//Datei
	FILE *datei;
	datei = fopen(
			"daten.dat",
			"r");

	double data[SIZE], result;
	double *dev_data, *dev_result;

	//Device
	cudaMalloc(&dev_data, SIZE * sizeof(double));
	cudaMalloc(&dev_result, sizeof(double));

	//Host
	fread(data, sizeof(data), 1, datei);
	printf("Test-Ausgabe: %f ", data[0]);

	cudaMemcpy(&dev_data, data, SIZE * sizeof(double), cudaMemcpyHostToDevice);

	calc<<<1, 1>>>(dev_data, dev_result);

	cudaMemcpy(&result, dev_result, sizeof(double), cudaMemcpyDeviceToHost);
	printf("Ergebnis %f: ", result);

	cudaFree(dev_data);
	cudaFree(dev_result);

	return 0;
}

Wirklich nur ein kleines Beispiel, aber so richtig klappen will es nicht.
Eigentlich soll er ein Array lesen (funktioniert), dieses an den Kernel weitergeben und +1 rechnen.

Leider macht er gar nichts im Kernel Hier mal die Ausgabe ...

Test-Ausgabe: -4.162908 Ergebnis 0.000000:

Irgendwas klappt mit der Übergabe nicht, wenn ich das recht sehe.

Freue mich über eure Hilfe

 

[ZuseZ3]

Ohne genauer draufzugucken, du willst ja ein Array bearbeiten, rufst in der Zeile 31 aber nur 1 x 1 thread auf...
Außerdem glaube ich, dass du in der selben Zeile dev_data und dev_result vertauscht hast.
Hier ist mal als bsp einer meiner codes:

Spoiler

    #include <stdio.h>
     
    __global__ void collatz(float * d_out, float * d_in)
    {
    int idx = threadIdx.x;
    float collnumber = d_in[idx];
    float steps = 0;
    while (collnumber > 1)
    {
    if (int(collnumber) % 2 == 0){collnumber /= 2;}
    else {collnumber *= 3; collnumber += 1;}
    steps += 1;
    }
    d_out[idx] = steps;
    }
     
    int main(int argc, char** argv) {
    const int ARRAY_SIZE = 128;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(float);
     
    float h_in[ARRAY_SIZE];
    for(int i = 0; i < ARRAY_SIZE; i++){
    h_in[i] = float(i);
    }
    float h_out[ARRAY_SIZE];
     
     
    float * d_in;
    float * d_out;
     
    cudaMalloc((void**) &d_in, ARRAY_BYTES);
    cudaMalloc((void**) &d_out, ARRAY_BYTES);
     
     
    cudaMemcpy(d_in, h_in, ARRAY_BYTES, cudaMemcpyHostToDevice);
     
    collatz<<<1, ARRAY_SIZE>>>(d_out, d_in);
     
    cudaMemcpy(h_out, d_out, ARRAY_BYTES, cudaMemcpyDeviceToHost);
     
     
    cudaFree(d_in);
    cudaFree(d_out);
     
    for (int i =0; i < ARRAY_SIZE; i++) {
    printf("%f", h_out[i]);
    printf(((i % 4) != 3) ? "\t" : "\n"); }
    }

 

[KillerPinockel]

Hey, danke für die Antwort.

An sich will ich eigentlich mehrere 1D Array übergeben und diese dann im Kernel mit ARRAY[X] benutzen. Ist sowas möglich?

Danke schonmal!

 

[Nai]

dev_result = dev_data + 1;

sollte
*dev_result = *dev_data + 1 sein.

An sich will ich eigentlich mehrere 1D Array übergeben und diese dann im Kernel mit ARRAY[X] benutzen. Ist sowas möglich?

Dafür kannst du eine Structure of Arrays verwenden:
struct SoA
{
Double* Arrays[MAXARRAYCOUNT];
}
Diese kannst du als solche als Kernelargument übergeben. Jeden der Pointer musst du dann davor allerdings einzeln mit cudamalloc allozieren. Die SoA hat aber den Nachteil, dass die gesamte Konstruktion im CUDA Konstantenspeicher landet. Wenn du dann auf die SoA so zugreifst, dass nicht jeder Thread eines Warps auf das selbe Array zugreift, dann werden die Zugriffe auf die Zeiger im Konstantenspeicher sequentailisiert.

Alternativ kannst du das Ganze dadurch verwirklichen, dass du die eindimensionalen Arrays zu einem einzigen eindimensionalen Array kombinierst. Dadurch ergeben sich diese Nachteile nicht.

 

[KillerPinockel]

@ZuseZ3

Ich habe deinen Quellcode mal genommen und wollte die Arraygröße auf 216000 erhöhen. Ja das ist mein größtes Array was ich brauche. Die anderen sind alle kleiner.

const int ARRAY_SIZE = 216000;

Damit läuft das Programm nicht. Es zeigt zwar keinen Fehler, aber macht auch nix. Bei 500 oä geht es.


@Nai
Klingt kompliziert. Da muss ich mich erst mal einlesen. Danke auf jeden Fall!

In C konnte ich den Aufruf so realisieren ...

Aufruf

Funktion(Wert1, Wert2, Array1, Array2, Array3, Array4, Wert 3, Array4, Wert4, Wert5);

Definition

double Funktion(int Wert1, int Wert2, double Array1[], 
double Array2[], double Array3[], double Array4[], double Wert 3, double Array4[], double Wert4, double Wert5)) {

...

}

 

[Nai]

Er startet durch collatz<<<1, ARRAY_SIZE>>>(d_out, d_in) nur einen einzigen ThreadBlock mit der Größe von Array_Size. Die Thread-Block Größe ist bei Kepler jedoch auf 1024 beschränkt. Deshalb musst du die 1 dementsprechend anpassen und die die Berechnung der ID deines Threads im Kernel durch ThreadID = blockIdx.x * blockDim.x + threadIdx.x ersetzen.

Des Weiteren gilt anzumerken, dass die GPU viele Thread-Blocks für die Peak-Performance benötigt. Bei GK 110 sind es bei 1024 Threads per Block in etwa 30 Thread-Blocks. Schau dir mal das an:
https://www.olcf.ornl.gov/tutorials/cuda-vector-addition/

Da steht wie man den Kernelstart richtig macht im Quelltext.

double Funktion(int Wert1, int Wert2, double Array1[],
double Array2[], double Array3[], double Array4[], double Wert 3, double Array4[], double Wert4, double Wert5)) {

...

}

Kannst du doch auch in CUDA so machen, wenn du noch das __global__ Keyword vor die Funktion setzt. Sogar wenn ich so viel Parameter exterm unellegant finde.

 

[KillerPinockel]

Ich nochmal. Danke für eure Hilfe, ich konnte meinen Code nun zum laufen bekommen und er funktioniert auch. Nur ist dieser vernichtend ineffizient. Das liegt sicher daran, dass ich den Aufruf des Kernels schlecht gemacht habe, aber ich weiß nicht genau, wie ich genau nun optimieren soll.

Der Aufruf des Kernels, sieht so aus

for (i = 0; i < 500 i++) {
			for (j = 0; j < 256 j++) {
				kernel<<<1, 1>>>(i, j, ... );
			}
		}

Ich bin mir sicher, dass das deutlich besser geht, wie zum Beispiel die Schleifen weglassen und anstatt mit 1 Thread und 1 Block zu arbeiten. Ich habe den Code schon in C mit OpenMP parallelisiert, vllt hilft das ja etwas weiter.

#pragma omp parallel
			{
#pragma omp for private(j)
				for (i = 0; i < 500 i++) {
					for (j = 0; j < 256; j++) {
						data_erg[i * 256 + j] = funktion(i, j, ...);
					}
				}
			}

Dieser läuft deutlich schneller und nun möchte ich natürlich den GPU Code optimieren
Könnt ihr mir da Hilfestellung geben, wie ich den Code optimieren kann?

Vielen Dank schon mal!

 

[Nai]

#define JSIZE 256 //wichtig das zu definen für Kerneloptimierung da bei 2er Potenzen div und mod billiger

__global__ void HappyCUDAKernel(int TotalThreadCount  , float* Result,  .......)
{
int HappyThreadID = blockIdx.x*blockDim.x+threadIdx.x;
if( HappyThreadID  >= TotalThreadCount )
   return;

if(HappyThreadID == 0)
 printf("I'm so happy!");

int i = HappyThreadID / JSIZE ;
int j = HappyThreadID % JSIZE ;

Result[HappyThreadID] = Funktion(i,j,wasweissich);
}


void KernelCall()
{
int IterationsOfI = 500;
int TotalThreadCount =  IterationsOfI * JSIZE;

int BlockSize =128;
int BlockCount =  TotalThreadCount  / BlockSize +1;

HappyCUDAKernel
#ifdef __CUDACC__
 <<< BlockCount , BlockSize  >>>
#endif 
(TotalThreadCount  , Result, .....);
}

 

[KillerPinockel]

Hey Danke für deine fixe Antwort.

Ich habe den Inhalt der "Funktion(...)" bereits in den "Kernel" transferiert. Also gibt es die "Funktion" nicht mehr.
Im Prinzip gilt Funktion = Kernel

 

[Nai]

Du kannst __host__ __device__ float MyHappyFunction(.......) verwenden. Dadurch kannst du beliebige Funktionen von nem Kernel genauso wie vom Host aus aufrufen. Notfalls schreibe sie halt an der entsprechenden stelle ins Kernel.

 

[KillerPinockel]

Ah hab es so übernommen und geht soweit ... letzte Frage erst mal Nochmal Danke für deine tolle Hilfe bis jetzt ... was wäre wenn der Quellcode vorher so aussehen müsste. Ich berechne ein Bild und das hätte ich jetzt, aber ich möchte nun X Bilder berechnen (BSP 1000)

omp_set_num_threads(8);

for (x = 0; x < 1000; x++) {

//Daten lesen
fread(data, sizeof(data), 1, datei_data);
#pragma omp parallel
			{
#pragma omp for private(j)
				for (i = 0; i < 500 i++) {
					for (j = 0; j < 256 j++) {
						data_2[i * 256 + j] = Funktion(i, j, ...);
					}
				}
			}
//Daten schreiben
fwrite(data_2, sizeof(data_2), 1, datei_data_2);
}

Aktuell läuft der Code zwar deutlich schneller als vorher (ohne Optimierung) aber immer noch schlechter, als wenn ich Ihn auf 8 CPU Threads zerlege. Könnte es sein das die GPU nicht schneller sein kann?

 

[Nai]

Was tust du überhaupt ? Was ist die Funktion ?