Seite 1 von 1

barrier(???)

Verfasst: 27. Jan 2012 17:26
von Lukg
Hi,

kann mir einer "genau" erklären was barrier() macht und wie es funktioniert? Wann muss CLK_LOCAL_MEM_FENCE und wann CLK_GLOBAL_MEM_FENCE übergeben werden? Wie hängt das mit __local-Variablen zusammen? Die kryptischen einzeiligen Erklärungen aus dem Internet helfen nicht sonderlich gut.

Mein Problem:
Soweit ich es verstanden habe, erarbeitet ja jede Workgroup, mit je 25 Threads, einen Eintrag in der Result-Matrix. Einfachheitshalber habe ich also die Berechnung aus dem CPU-Teil übernommen, jedoch nur auf dem ersten der 25 Threads einer Workgroup ausführen lassen. Das OpenCL-Programm sieht dann in etwa so aus:

Code: Alles auswählen

__kernel void oclMIcomputation(__global uchar * sequences, __global float * onePointProbs, __global float * result, uint sequenceLength, uint numSequences)
{
	const double epsilon=1e-6;
	int i = get_group_id(0);
	int j = get_group_id(1);
	__local int twoPointOccs[NUMCHARS][NUMCHARS];

	if (get_local_id(0) == 0 && get_local_id(1) == 0 ) {
		//absolute number of occurrences of character pairs x,y: N_ij(x,y)
		for (int x = 0; x < NUMCHARS; x++) for (int y = 0; y < NUMCHARS; y++) twoPointOccs[x][y] = 0;
		...
		//iterate through all sequences and compute two-point occurrences
		...
		//sum over all x and y
		...
	}
}
Ich habe hier das Array twoPointOccs als __local definiert, da dieses ja innerhalb der Workgroup "gleich" bleibt.
Außerdem habe ich "memset(twoPointOccs, 0, sizeof(twoPointOccs));" ersetzt durch "for (int x = 0; x < NUMCHARS; x++) for (int y = 0; y < NUMCHARS; y++) twoPointOccs[x][y] = 0;".

Und wow!, es funktioniert. Ich erhalte exakt die gleichen Matrix-Werte wie aus dem CPU-Teil.

Jetzt kommt barrier() dazu. Ich ersetze also obigen Code durch folgenden (die Zeilen mit //*** haben sich geändert):

Code: Alles auswählen

__kernel void oclMIcomputation(__global uchar * sequences, __global float * onePointProbs, __global float * result, uint sequenceLength, uint numSequences)
{
	const double epsilon=1e-6;
	int i = get_group_id(0);
	int j = get_group_id(1);
	__local int twoPointOccs[NUMCHARS][NUMCHARS];

	if (get_local_id(0) == 0 && get_local_id(1) == 0 ) {
		//absolute number of occurrences of character pairs x,y: N_ij(x,y)
		for (int x = 0; x < NUMCHARS; x++) for (int y = 0; y < NUMCHARS; y++) twoPointOccs[x][y] = 0;
		barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //***
		...
		//iterate through all sequences and compute two-point occurrences
		...
		//sum over all x and y
		...
	} else barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE); //***
} 
Das heißt, hier wird darauf gewartet bis der erste Thread mit dem Leeren von twoPointOccs fertig ist. Nach meinem Verständnis dürfte sich also NICHTS geändert haben, da die ganze Berechnung sowieso nur durch einen Thread ausgeführt wird.
Die restlichen 24 Threads machen also weiterhin gar nichts.

... Diesmal besteht die Result-Matrix aber nur noch aus Nullen.

Gibt es dafür eine Erklärung?

Re: barrier(???)

Verfasst: 27. Jan 2012 17:48
von Narida
Soweit ich das verstanden habe:

Alle threads in einer work-group müssen diesen Aufruf ereichen, bevor die folgenden Befehle ausgeführt werden dürfen.
So kannst du zB erreichen, das alle Threads einen bestimmt Code-Block ausführen bevor alle mit dem nächsten weitermachen dürfen.

Wenn man dann noch CLK_LOCAL_MEM_FENCE übergibt, werden alle Threads evtl. gecachte Werte in den Speicher schreiben. So kann man sicherstellen, dass der Speicherinhalt für alle Threads der Workgroup konsistent ist.

http://www.khronos.org/registry/cl/sdk/ ... rrier.html

Wo dein Fehler mit den Nullen herkommt ist mir nicht wirklich ersichtlich...

Re: barrier(???)

Verfasst: 27. Jan 2012 18:12
von Drno
If barrier is inside a conditional statement, then all work-items must enter the conditional if any work-item enters the conditional statement and executes the barrier.
Laut deinem Code landen nicht alle Threads an der selben barrier. Einer im if-Zweig der Rest im Else-Zweig. Terminiert dein Code überhaupt?

Re: barrier(???)

Verfasst: 28. Jan 2012 11:52
von mw1039
Grundsaetzlich sollte man es vermeiden barriers in ifs oder Schleifen reinzulegen. Nur wenn man sich sicher ist, dass alle Threads einer Workgroup sich im if oder der Schleife identisch verhalten, sollte man das machen.

Ausserdem brauchst du glaube ich das
Lukg hat geschrieben:| CLK_GLOBAL_MEM_FENCE
nicht. CLK_LOCAL_MEM_FENCE sollte reichen, wenn du nur innerhalb einer Workgroup synchronisieren moechtest.

Warum arbeitet bei dir eigentlich nur der erste Thread daran twoPointOccs mit 0en zu fuellen? Solche Aufgaben sollte man immer ueber alle Threads verteilen.

Re: barrier(???)

Verfasst: 28. Jan 2012 15:28
von Dennis Albrecht
mw1039 hat geschrieben:Warum arbeitet bei dir eigentlich nur der erste Thread daran twoPointOccs mit 0en zu fuellen? Solche Aufgaben sollte man immer ueber alle Threads verteilen.
Denke mal, dass er es einfach Stück für Stück parallelisieren wollte. :D

Gruß

Re: barrier(???)

Verfasst: 30. Jan 2012 21:17
von Lukg
Danke Euch,
das Problem lag wirklich bei der If-Condition. Ich dachte, ich hätte das durch das barrier() im else-Zweig umgehen können.
Das Programm ist aber trotzdem terminiert. Naja, man muss nicht alles verstehen.
Denke mal, dass er es einfach Stück für Stück parallelisieren wollte. :D
So ist es ;)