C++ OpenCL kernel beschleunigen

Furtano · 15. März 2014

Hallo,

ich habe gelsen ,dass wenn man den Adress Space Qualifier in OpenCL Kernel von Global auf Local setzt, der Kernel schneller arbeiten kann.
Wenn ich allerdings die Variablen auf __local setze funktioniert das Programm nicht mehr.

Eigentlich müsste man sie doch auf Local setzen können, oder?
Weil im Kernel wird ja nur das Work Item aus der eigenen Work-Group aufgerufen.

Hier mein Kernel:

Code:

__kernel void stepSizeAnt(
						  __global  float *antX,
						  __global  float *antY,
						  __global  float *newAntX,
						  __global  float *newAntY,
						  __global  int *targetX,
						  __global  int *targetY,
						  __global  float *resultX,
						  __global  float *resultY,
						  __global  int *lifeStadium, 
						  __global  int *speed,
						  __global  int *pheromonMap
						  )
{
	int gid = get_global_id(0);


	
		resultX[gid] = 0;
		resultY[gid] = 0;

	if (lifeStadium[gid] == 1){
		//newAntX[gid] = antX[gid];
		//newAntY[gid] = antY[gid];

		// Neuen Richtungsvektor berechnen (in Welche Richtung soll die Ameise laufen)
		resultX[gid] = (targetX[gid]-newAntX[gid]);
		resultY[gid] = (targetY[gid]-newAntY[gid]);

		// Normvektor ausrechnen normalisieren
		float norm = native_sqrt(pown(resultX[gid],2) + pown(resultY[gid],2));


		// Richtungsvektor normalisieren
		resultX[gid] = (resultX[gid] / norm);
		resultY[gid] = (resultY[gid] / norm);


		// Ameisenschritt machen
		// neue Ameisenposition = Alte Position + (neue Schrittposition * Geschwindigkeit)
		newAntX[gid] +=   (resultX[gid])*speed[gid];
		newAntY[gid] +=   (resultY[gid])*speed[gid];

		
	
	}

	// Auf 2 Stellen runden
		//newAntX[gid] += 0.5;
		newAntX[gid] = (float)((int)(newAntX[gid]*10000))/10000;
		//newAntY[gid] += 0.5;
		newAntY[gid] = (float)((int)(newAntY[gid]*10000))/10000;


		// Wenn die Ameise den Rand betritt

		if (newAntX[gid] < 0 || newAntY[gid] < 0){
			newAntX[gid] = 0;
			newAntY[gid] = 0;

		}
		if (newAntX[gid] > 900 || newAntY[gid] > 900){
			newAntX[gid] = 0;
			newAntY[gid] = 0;

		}

}

Danke

Hancock · 15. März 2014

Zuerst: Verwende float2, wenn es geht, das kann deine Grafikkarte besser...
Und dann gibt es normalize, welches dir deine Wurzel erspart.

Du kannst in dem Fall deine Parameter nicht __local machen, da du den Status über einen __kernel-Aufruf hinaus behalten willst.

Aber was ist denn langsam? Wenn ich mir deinen Code so anschaue, sollte das relativ flüssig laufen (10 GAnt/s oder so). Ein Problem könnte sein, dass du sehr viele Speicherzugriffe auf den globalen Speicher hast, da könntest du etwas optimieren, das macht's immer langsam.

BTW: Soll die Ameise am rechten Rand wirklich zum linken teleportiert werden?

Furtano · 15. März 2014

Hi,

ok, welchen Status möchte ich über einen Kernel Aufruf hinaus behalten?

Wie benutze ich die normalized Methode?

Achso und der Kernel läuft flix durch (nicht mal 1ms pro Aufruf). Trotzdem möchte ich ihn optimieren

Mit dem Rand ist toter Code

Hancock · 15. März 2014

So wie dein Kernel zur Zeit aufgebaut ist, integrierst du nur die Position in der Zeit.
Du musst wissen, dass ein Kernel eine fixe Setup-Zeit pro Aufruf von ein paar Duzend µs hat, der bei solch kurzen Kernels die meiste Zeit frisst.

Ich hab mal dein Code genommen und so umgebaut, wie ich es machen würde:

Code:

typedef struct ant_state{
    float2 pos;
    float2 target;
    float speed;
    int lifeStadium;
}ant_state;
__kernel void stepSizeAnt(
						  __global  ant_state *ant_states,
						  __global  int *pheromonMap
						  )
{
	int gid = get_global_id(0);
    ant_state me=ant_states[gid];
 
	if (me.lifeStadium == 1){
 
		// Richtungsvektor normalisieren
		float2 result = normalize(me.target-me.pos);
 
 
		// Ameisenschritt machen
		// neue Ameisenposition = Alte Position + (neue Schrittposition * Geschwindigkeit)
		me.pos +=   (result)*me.speed;
	}
 
	// Auf 2 Stellen runden
		//newAntX[gid] += 0.5;
		//newAnt[gid] = (float2)((int2)(newAnt[gid]*10000.0f))/10000.0f;
        me.pos.x = (float)((int)(me.pos.x*10000.0f))/10000.0f;
        me.pos.y = (float)((int)(me.pos.y*10000.0f))/10000.0f;
 
		// Wenn die Ameise den Rand betritt
 
		if (me.pos.x < 0 || me.pos.y < 0){
			me.pos = 0;
		}
		if (me.pos.x > 900 || me.pos.y > 900){
			me.pos = 0;
		}
    ant_states[gid]=me;
}

Laut Code-Analyzer ist das zwar nur minimal schneller, aber mehr als nichts...

.

Falls du wissen willst, was auf Maschinencodeebene wirklich lange braucht, such dir einen Code-Analyzer, vom AMD gibt's da den KernelAnalyzer2 kostenlos zum DL. NVIDIA weiß ich nicht.

Für mehr Durchsatz kannst du dir eine Schleife drumrumbauen, die dann gleich ~100 Schritte macht, damit wird die GPU erst richtig "warm". (Falls du eine feste Schrittweite nimmst, greift da die Loop-Unroll-Optimization.)

__local loht sich erst, wenn du eine Workgroup hast, die die Daten auch untereinander teilt.

Nai · 16. März 2014

Ich würde das Kernel ersteinmal weitestgehend so umgestalten wie Hancock, abgesehen davon dass ich bei den Structures of Arrays bleiben würde und davon abraten würde ein Array of Structures wie folgt zu verwenden:

Code:

typedef struct ant_state{
float2 pos;
float2 target;
float speed;
int lifeStadium;
}ant_state;
....
__global  ant_state *ant_states 
....

Das macht das Coalescing bei sequentiellen Zugriff auf einer GPU komplett kaputt. Denn der NVIDIA Compiler macht beim Laden einer solchen Struct aus dem global Memory für jedes dieser float2s zuerst zwei 8 Byte Ladeoperationen, dann für den einzelnen float und den einzelnen int jeweils noch eine 4 Byte Ladeoperation. Dadurch haben 2 benachbarte Warpthreads jeweils bei einem Speicherzugriff einen Stride von der größe der Struct, welche 24 Byte betragen sollte. Beim Store verhält es sich ähnlich, abgesehen davon, dass zusätzlich kostbare Speicherbandbreite verschwendet wird, da man target, speed und Lifestadium nicht ändert.

Zuerst: Verwende float2, wenn es geht, das kann deine Grafikkarte besser...
Und dann gibt es normalize, welches dir deine Wurzel erspart.

Float2 oder Float ist mehr oder weniger egal, da moderne GPUs Skalararchitekturen sind. Außerdem ist das reine normalize teuer, da der Compiler es durch viele Operationen berechnen lassen muss. fast_normalize wäre da besser. Das existiert aber auch nicht als Operation sondern der Compiler lässt das (zumindest bei NVIDIA) berechnen per:
vec_norm = vec * __rsqrt(vec.x*vec.x+vex.y*vec-y)
Ursache hierfür ist, dass die GPUs nicht direkt die Wurzel sondern nur den Kehrwert der Wurzel (rsqrt) berechnen können, und danach in einer weiteren Operation den Kehrwert bilden müssen.

Ansonsten, dein Kernel ist komplett Speicherbandbreiten limitiert. Da kannst du das versuchen:
http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf
Allerdings wenn du dir die Ergebnisse immer noch über den PCI-E zurückziehst nach jeder Ausführung, bist du durch die PCI-E Bandbreite limitiert; und zwar vermutlich so sehr, dass es sinnvoller wäre das ganze nicht auf der GPU laufen zu lassen. Wenn ja kannst du versuchen, den PCI-Transfer eventuell mit Zero Copy oder DMA zu beschleunigen.
Alternativ könntest du gegen beides auch noch mehr andere Berechnungen deiner Simulation in dein Kernel einbauen, so dass du von diesem Speicherbandbreiten-Limit weg bist.

Des Weiteren ist deine Workgroup Size denn immer noch eins ? Das zerhaut dir die Performance auf jeder GPU. Nimm mal lieber einen Wert von 128, 256 oder so.

Falls du wissen willst, was auf Maschinencodeebene wirklich lange braucht, such dir einen Code-Analyzer, vom AMD gibt's da den KernelAnalyzer2 kostenlos zum DL. NVIDIA weiß ich nicht.

NVIDIA hat den Visual Profiler. Da braucht man allerdings eine sehr alte Version (CUDA Toolkit 4.0 oder so), weil NVIDIA den OpenCL support rausgehauen hat.

Eigentlich müsste man sie doch auf Local setzen können, oder?
Weil im Kernel wird ja nur das Work Item aus der eigenen Work-Group aufgerufen.

Bei dem local Memory handelt es sich um einen Scratchpadspeicher in jedem (Streaming)-Prozessorkern deiner GPU, während es sich bei dem global Memory um den DRAM deiner Graphikkarte handelt. Deshalb kannst du den Local Memory nur dafür verwenden um zuerst sachen manuell im Kernel reinzuschreiben,
und später wieder auszulesen.

Des Weiteren würde ich dir Raten mal die Programming Guides von den Herstellern (NVIDIA und AMD) durchzulesen. Sind zwar ~ 100 Seiten lang, aber da lernst du sehr schnell wie GPUs ticken, und woran du schrauben musst, um die Performance zu verbessern.

Furtano · 16. März 2014

globalWorkSize == Anzahl der Ants

PHP:

clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL,
		globalWorkSize, NULL,
		0, NULL, NULL);

So besser?

Das mit dem Struct verstehe ich nicht, wie soll ich das ins Hauptprogramm integrieren? Muss ja auch die Argumente vom Kernel setzen.

Nai · 16. März 2014

Würde das tendentiell so machen mit der Workgroupgröße:

Code:

 size_t LocalSize = 256;
 size_t GlobalSize = (AntCount / LocalSize)*LocalSize;

 if(AntCount % LocalSize!= 0 )
           GlobalSize +=LocalSize;

 clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, GlobalSize , LocalSize ,0, NULL, NULL);
//oder alternativ stattdessen nun wieder:
// clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, GlobalSize , NULL,0, NULL, NULL);

Dann noch im Kernel gleich zu Beginn:

Code:

if(get_global_id(0) >= AntCount)
    return;

Dadurch vermeidet man, dass die Implementierung die Workgroupgröße nach Gutdünken selbst wählen muss. Ausserdem vermeidet man, dass, da die Globalsize immer ein Vielfaches der Workgroupgröße sein muss, die Implementierung bei manchen Globalsizes nur sehr schlechte Workgroupgrößen wählen kann.
(z.B. wenn die Globalsize Prim ist, dann kann die Implementierung nur eine Workgroupgröße von 1 wählen). Hier suckt OpenCL m.E. wieder, da man de facto nicht Workitems sondern Workgroups startet. Wenn man wie bei CUDA die Workgroupgröße und die Zahl der gestarteten Workgroups angeben müsste und nicht die Anzahl der gestarteten Workitems, dann käme es zu diesem Problem nicht.

Das mit dem Struct verstehe ich nicht, wie soll ich das ins Hauptprogramm integrieren? Muss ja auch die Argumente vom Kernel setzen.

Falls du meine Erklärung nicht verstanden hast: Mach das mit den Structs *nicht*; das ist schlecht für die Performance.

P.S. Ich wiederhole mich nocheinmal: Ich würde dir wirklich raten die Programming Guides durchzulesen, denn da würden sich solche Fragen mehr oder weniger von selbst beantworten. Denn es kann m.E. nicht wirklich der Sinn deiner Bachelor Thesis sein, dass du dir die Optimierung, welche ja eben ein Teil deiner Arbeit sein sollte, dir von Grundauf im Forum hier machen lässt. Etwas Eigeninitiative wäre da angebracht.

Furtano · 16. März 2014

Ok danke

Ich erwarte auch nicht dass ihr meine ganze Arbeit macht, ich möchte nur ein paar Unsicherheiten erst mal aus dem Weg Räumen

Ich lese mir jetzt mal diesen Guide durch:
http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_ProgrammingGuide.pdf

Suche

C++ OpenCL kernel beschleunigen

Furtano

Gast

Hancock

Captain

Furtano

Gast

Hancock

Captain

Nai

Lt. Commander

Furtano

Gast

Nai

Lt. Commander

Furtano

Gast

Ähnliche Themen