barrier(???)

Lukg
Erstie
Erstie
Beiträge: 11
Registriert: 14. Okt 2010 17:32

barrier(???)

Beitrag von Lukg » 27. Jan 2012 17:26

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?

Narida
Neuling
Neuling
Beiträge: 5
Registriert: 27. Jan 2012 14:40

Re: barrier(???)

Beitrag von Narida » 27. Jan 2012 17:48

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...

Drno
Sonntagsinformatiker
Sonntagsinformatiker
Beiträge: 236
Registriert: 10. Feb 2005 20:16

Re: barrier(???)

Beitrag von Drno » 27. Jan 2012 18:12

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?
Programming today is a race between software engineers striving to build bigger and better idiot proof programs, and the Universe trying to produce bigger and better idiots. So far, the Universe is winning.

mw1039
Computerversteher
Computerversteher
Beiträge: 346
Registriert: 12. Apr 2011 12:18

Re: barrier(???)

Beitrag von mw1039 » 28. Jan 2012 11:52

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.

Dennis Albrecht
Sonntagsinformatiker
Sonntagsinformatiker
Beiträge: 222
Registriert: 4. Okt 2010 18:15

Re: barrier(???)

Beitrag von Dennis Albrecht » 28. Jan 2012 15:28

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ß

Lukg
Erstie
Erstie
Beiträge: 11
Registriert: 14. Okt 2010 17:32

Re: barrier(???)

Beitrag von Lukg » 30. Jan 2012 21:17

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 ;)

Antworten

Zurück zu „Archiv“