OpenCL - barrier

[Früchtemüsli]

Hi :)

Ich verstehe diese barrier-Befehle in OpenCL noch nicht.

Angenommen, ich habe folgendes:

kernel void x(local float4* a)
{
// irgendein Befehl ...
float4 b = a[0];

// blockieren:
barrier(CLK_GLOBAL_MEM_FENCE);

float4 c = a[1];
}

Wird jetzt blockiert? Ich habe zwar dieses barrier-Befehl, jedoch den GLOBAL. Und danach arbeite ich nur noch mit local Daten.

[Früchtemüsli]

Dann habe ich auch noch so eine ähnliche Frage (zahlt sich nicht aus, deswegen ein neues Thema aufzumachen :) ).

Wieder ein Beispiel:

kernel void x
(
global float4* b,
local float4* c
)
{
c[get_local_id(0)] = b[get_global_id(0)];
}

Was passiert, wenn diese Kernel-Instanz abgelaufen ist? Wird dann der lokale Speicher sofort gelöscht? Oder wird er erst gelöscht, wenn ALLE Kernel-Instanzen innerhalb ihrer Work Group abgelaufen sind?

[Gast]

Hi :)

Ich verstehe diese barrier-Befehle in OpenCL noch nicht.

Angenommen, ich habe folgendes:

kernel void x(local float4* a)
{
// irgendein Befehl ...
float4 b = a[0];

// blockieren:
barrier(CLK_GLOBAL_MEM_FENCE);

float4 c = a[1];
}

Wird jetzt blockiert? Ich habe zwar dieses barrier-Befehl, jedoch den GLOBAL. Und danach arbeite ich nur noch mit local Daten.
Der Befehl stellt sicher dass alle global memory accesses für alle Threads in der work group sichtbar sind. Was gibt es da nicht zu verstehen? Wenn du vor diesem Befehl global memory accesses hast wird blockiert. Wenn nicht, nicht.
Dann habe ich auch noch so eine ähnliche Frage (zahlt sich nicht aus, deswegen ein neues Thema aufzumachen :) ).

Wieder ein Beispiel:

kernel void x
(
global float4* b,
local float4* c
)
{
c[get_local_id(0)] = b[get_global_id(0)];
}

Was passiert, wenn diese Kernel-Instanz abgelaufen ist? Wird dann der lokale Speicher sofort gelöscht? Oder wird er erst gelöscht, wenn ALLE Kernel-Instanzen innerhalb ihrer Work Group abgelaufen sind?
Local memory ist, wie du bestimmt weißt, innerhalb einer work group gültig (=steht allen Threads innerhalb einer workgroup zur Verfügung). Damit ist klar dass der Speicher erst "gelöscht" werden kann wenn alle Kernel der workgroup fertig sind...

[Früchtemüsli]

Local memory ist, wie du bestimmt weißt, innerhalb einer work group gültig (=steht allen Threads innerhalb einer workgroup zur Verfügung). Damit ist klar dass der Speicher erst "gelöscht" werden kann wenn alle Kernel der workgroup fertig sind...
OK, diese Vermutung hatte ich eh schon. Danke :)

Die Antwort zur anderen Frage verstehe ich aber leider von vorne bis hinten nicht :(

Der Befehl stellt sicher dass alle global memory accesses für alle Threads in der work group sichtbar sind.
Was bedeutet dieses "sichtbar"?

Wenn du vor diesem Befehl global memory accesses hast wird blockiert. Wenn nicht, nicht.
Meinst du wirklich "vor"? Oder meinst du "danach"?

Bsp 1:

kernel void x
(
local float* l,
global float* g
)
{
// ein global-Befehl
float4 g0 = g[0];

barrier(CLK_GLOBAL_MEM_FENCE); // blockiert, weil davor ein global-Befehl?

// ein local-Befehl:
float4 l0 = l[0];
}

Bsp 2:

kernel void x
(
local float* l,
global float* g
)
{
// ein local-Befehl
float4 l0 = l[0];

barrier(CLK_GLOBAL_MEM_FENCE); // blockiert nicht, weil davor keine global-Befehle?

// ein global-Befehl:
float4 g0 = g[0]; // der spielt keine Rolle für den barrier-Befehl?
}

[Gast]

Wenn du vor einem barrier(CLK_GLOBAL_MEM_FENCE) Befehl in global memory schreibst, dann wird bei der barrier so lange gewartet, bis alle Threads in der work-group beim Lesen vom global memory die gleichen Werte zurückbekommen würden, die vorher dort hin geschrieben wurden. Ohne die barrier könnte es vorkommen dass einige Threads schon beim Lesen sind, aber andere Threads noch garnicht geschrieben haben.

[Früchtemüsli]

D. h. für die work items einer work group gilt: beim barrier-Befehl wird so lange blockiert, bis alle work items bei diesem Befehl angekommen sind. :)

Sagt ja auch der Prof in der Vorlesung.

Stimmen die Kommentare zu Bsp 1 und Bsp 2?

[Gast]

D. h. für die work items einer work group gilt: beim barrier-Befehl wird so lange blockiert, bis alle work items bei diesem Befehl angekommen sind. :)

Sagt ja auch der Prof in der Vorlesung.

Ja.
Du haettest uebrigens auch einfach in die OpenCL Reference schauen koennen, dort steht es fast genau so drin:
All work-items in a work-group executing the kernel on a processor must execute this function before any are allowed to continue execution beyond the work_group_barrier.
http://www.khronos.org/registry/cl/sdk/2.0/docs/man/xhtml/

Stimmen die Kommentare zu Bsp 1 und Bsp 2?
Konzeptionell ja, in wirklichkeit wird hier nicht blockiert weil deine work items nur lesend auf den global mem zugreifen. Eine Problem ergibt sich erst wenn auf global mem schreibend zugegriffen wird, dann kann die situation eintreten die der Gast zuvor beschrieben hat.

[Früchtemüsli]

Ich glaube, ich müsste mal ein typisches, praktisches, kleines Beispiel sehen, um das zu kapieren :D

[Gast]

Hier ein ziemlich einfaches Beispiel:


__kernel void do_it(__global float* input, __global float* output)
{
__local float cache[N];
int lid = get_local_id(0);

// Jeder Thread liest einen Teil vom input (global) zum cache (local)
for (int i = lid; i < N; i += <local_work_size>)
cache[i] = input[i];

// Warten bis der cache auch wirklich komplett geschrieben wurde
barrier(CLK_LOCAL_MEM_FENCE);

// Hier kann der cache aus lokalem memory gelesen werden. Je nachdem wie oft man das macht kann es schneller sein als wenn man hier immer wieder vom global memory lesen würde
}


Anstatt Daten in einen Cache zu lesen könnte man auch irgendwas initialisieren oder Vorberechnungen machen, die von allen Threads hinter der barrier benötigt werden.

Man kann es sich vielleicht so Denken, dass die barrier dazu dient "Vorbereitungen" von der "Hauptarbeit" zu trennen. Bei der Vorarbeit berechnet jeder Thread in der work group irgendwas und speichert es in local/global memory. Hier ist zu beachten, dass nicht jeder Thread die ganze Vorarbeit leistet, sondern nur einen Teil. Nach der barrier können dann alle Threads die Vorarbeiten von allen anderen Threads (inklusive der eigenen) benutzen. Ohne die barrier wäre nicht sicher dass auch alle Threads die Vorarbeiten abgeschlossen haben - und so würde bei der Hauptarbeit auf invalide/uninitialisierte Daten zugreifen.

Das ist nur ein Beispiel. Wenn man es Verstanden hat sollte man sich selbst Use Cases einfallen lassen können.