Wie verwendet man async_work_group_copy in OpenCL?

Question

Ich würde gerne verstehen, wie Sie den Aufruf async_work_group_copy() in OpenCL korrekt verwenden. Lassen Sie uns über ein vereinfachtes Beispiel einen Blick:.

__kernel void test(__global float *x) { 
    __local xcopy[GROUP_SIZE]; 

    int globalid = get_global_id(0); 
    int localid = get_local_id(0); 
    event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0); 
    wait_group_events(1, &e); 
} 

Der Referenz http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html sagt: „Führen Sie eine Asynchron-Kopie von num_elements gentype Elemente aus src zu dst Die Asynchron-Kopie von allen Workitems in einer Arbeitsgruppe durchgeführt wird, und Diese integrierte Funktion muss daher von allen Arbeitselementen in einer Arbeitsgruppe angetroffen werden, die den Kernel mit den gleichen Argumentwerten ausführt, andernfalls sind die Ergebnisse undefiniert. "

Aber das ist nicht meine Fragen klären ...

Ich mag würde wissen, wenn die folgenden Annahmen korrekt sind:

1. Der Aufruf(), um async_work_group_copy muss von allen Arbeiten ausgeführt werden -Elemente in der Gruppe.
2. Der Aufruf sollte so erfolgen, dass die Quelladresse für alle Workitems identisch ist und auf das erste Element des zu kopierenden Speicherbereichs verweist.
3. Da meine Quelladresse basierend auf der globalen Arbeitselement-ID des ersten Arbeitselements in der Arbeitsgruppe relativ ist. Also muss ich die lokale ID subtrahieren, damit die Adresse für alle Work-Items identisch ist ...
4. Ist der dritte Parameter wirklich die Anzahl der Elemente (nicht die Größe in Bytes)?

Bonus Fragen:

a. Kann ich einfach die Barriere (CLK_LOCAL_MEM_FENCE) anstelle von wait_group_events() verwenden und den Rückgabewert ignorieren? Wenn ja, wäre das wahrscheinlich schneller?

b. Macht eine lokale Kopie auch Sinn für die Verarbeitung auf CPUs oder ist das Overhead, da sie sich ohnehin einen Cache teilen?

Grüße, Stefan

Answer 1

Einer der Hauptgründe für diese Funktion existiert, ist der Treiber/kernel-Compiler zu ermöglichen, effizient den Speicher ohne den Entwickler kopieren zu müssen Annahmen über die Hardware machen.

Sie beschreiben, welchen Speicher Sie kopieren müssen, als wäre es eine single-threaded-Kopie, und async_work_group_copy erledigt das für Sie, indem Sie die parallele Hardware verwenden.

Für Ihre spezifische Fragen:

1. Ich habe noch nie async_work_group_copy durch nur einen Teil der Arbeitselemente in einer Gruppe verwendet gesehen. Ich habe immer angenommen, das ist, weil es es erfordert. Ich denke, dass die blockierende Natur von wait_group_events alle Arbeitselemente dazu zwingt, Teil der Kopie zu sein.

2. Ja. Quell- (und Ziel-) Adressen müssen für alle Arbeitselemente identisch sein.

3. Sie könnten Ihre lokale ID subtrahieren, um die richtige Adresse zu erhalten, aber ich finde, dass die Lösung der Adresse auf groupId dieses Problem ebenfalls löst. (get_group_id)

4. Ja. Der letzte Parameter ist die Anzahl der Elemente, nicht die Größe in Bytes.

a. Nein. Bei der ereignisbasierten Methode werden Sie feststellen, dass Ihre Barriere fast sofort von den Arbeitselementen getroffen wird und die Daten nicht unbedingt kopiert werden. Dies ist sinnvoll, da einige Opencl-Hardware die Recheneinheiten nicht einmal für die eigentliche Kopieroperation verwendet.

b. Ich denke, dass cpu Opencl-Implementierungen L1-Cache-Nutzung garantieren können, wenn Sie lokalen Speicher verwenden. Die einzige Möglichkeit, um sicher zu sein, ob diese Leistung besser ist, besteht darin, Ihre Anwendung mit verschiedenen Einstellungen zu vergleichen.