Worauf wartet meine GPU?

11

Ich schreibe ein OpenCL-Programm zur Verwendung mit meiner GPU der AMD Radeon HD 7800-Serie. Laut AMDs OpenCL-Programmierhandbuch verfügt diese GPU-Generation über zwei Hardware-Warteschlangen, die asynchron arbeiten können.

5.5.6 Befehlswarteschlange

Für Südinseln und höher unterstützen Geräte mindestens zwei Hardware-Rechenwarteschlangen. Dadurch kann eine Anwendung den Durchsatz kleiner Dispatches mit zwei Befehlswarteschlangen für die asynchrone Übermittlung und möglicherweise Ausführung erhöhen. Die Hardware-Berechnungswarteschlangen werden in der folgenden Reihenfolge ausgewählt: Erste Warteschlange = gerade OCL-Befehlswarteschlangen, zweite Warteschlange = ungerade OCL-Warteschlangen.

Zu diesem Zweck habe ich zwei separate OpenCL-Befehlswarteschlangen erstellt, um Daten an die GPU weiterzuleiten. Das Programm, das auf dem Host-Thread ausgeführt wird, sieht ungefähr so aus:

static const int kNumQueues = 2;
cl_command_queue default_queue;
cl_command_queue work_queue[kNumQueues];

static const int N = 256;
cl_mem gl_buffers[N];
cl_event finish_events[N];

clEnqueueAcquireGLObjects(default_queue, gl_buffers, N);

int queue_idx = 0;
for (int i = 0; i < N; ++i) {
  cl_command_queue queue = work_queue[queue_idx];

  cl_mem src = clCreateBuffer(CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, ...);

  // Enqueue a few kernels
  cl_mem tmp1 = clCreateBuffer(CL_READ_WRITE);
  clEnqueueNDRangeKernel(kernel1, queue, src, tmp1);

  clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp1);

  cl_mem tmp2 = clCreateBuffer(CL_READ_WRITE);
  clEnqueueNDRangeKernel(kernel2, queue, tmp1, tmp2);

  clEnqueueNDRangeKernel(kernel3, queue, tmp2, gl_buffer[i], finish_events + i);

  queue_idx = (queue_idx + 1) % kNumQueues;
}

clEnqueueReleaseGLObjects(default_queue, gl_buffers, N);
clWaitForEvents(N, finish_events);

Mit kNumQueues = 1funktioniert diese Anwendung so ziemlich wie beabsichtigt: Sie sammelt die gesamte Arbeit in einer einzigen Befehlswarteschlange, die dann vollständig ausgeführt wird, wobei die GPU die ganze Zeit über ziemlich beschäftigt ist. Ich kann dies anhand der Ausgabe des CodeXL-Profilers sehen:

Geben Sie hier die Bildbeschreibung ein

Wenn ich jedoch einstelle kNumQueues = 2, erwarte ich, dass dasselbe passiert, aber die Arbeit gleichmäßig auf zwei Warteschlangen verteilt ist. Wenn überhaupt, erwarte ich, dass jede Warteschlange die gleichen Eigenschaften wie die eine Warteschlange hat: Sie beginnt nacheinander zu arbeiten, bis alles erledigt ist. Wenn Sie jedoch zwei Warteschlangen verwenden, kann ich feststellen, dass nicht die gesamte Arbeit auf die beiden Hardware-Warteschlangen aufgeteilt ist:

Geben Sie hier die Bildbeschreibung ein

Zu Beginn der Arbeit der GPU können die Warteschlangen einige Kernel asynchron ausführen, obwohl es so aussieht, als ob keiner der beiden die Hardware-Warteschlangen jemals vollständig belegt (es sei denn, mein Verständnis ist falsch). Gegen Ende der GPU-Arbeit scheinen die Warteschlangen nur einer der Hardware-Warteschlangen nacheinander Arbeit hinzuzufügen, aber es gibt sogar Zeiten, in denen keine Kernel ausgeführt werden. Was gibt? Habe ich ein grundlegendes Missverständnis darüber, wie sich die Laufzeit verhalten soll?

Ich habe einige Theorien, warum dies geschieht:

  1. Die eingestreuten clCreateBufferAufrufe zwingen die GPU, Geräteressourcen aus einem gemeinsam genutzten Speicherpool synchron zuzuweisen, wodurch die Ausführung einzelner Kernel blockiert wird.

  2. Die zugrunde liegende OpenCL-Implementierung ordnet logische Warteschlangen keinen physischen Warteschlangen zu und entscheidet nur, wo Objekte zur Laufzeit platziert werden sollen.

  3. Da ich GL-Objekte verwende, muss die GPU beim Schreiben den Zugriff auf den speziell zugewiesenen Speicher synchronisieren.

Stimmt eine dieser Annahmen? Weiß jemand, was dazu führen könnte, dass die GPU im Szenario mit zwei Warteschlangen wartet? Jeder Einblick wäre dankbar!

gpgpu opencl Mokosha
quelle
Ich kann nicht finden, wo im Programmierhandbuch zwei Hardware-Warteschlangen stehen. Könnten Sie ein Angebot aus dem Dokument veröffentlichen? Erwähnen Sie, in welchem ​​Kapitel zwei Warteschlangen stehen. Ist die Anzahl der Hardware-Warteschlangen zur Laufzeit mit OpenCL abfragbar?
Andreas
Ich habe meinen Beitrag aktualisiert. Es heißt mögliche Ausführung, aber wenn es ein paar kann, warum kann es dann nicht alle? Außerdem hat die OpenCL-Laufzeit keine Vorstellung von der Hardware-Warteschlange, sodass Sie diese nicht abfragen können.
Mokosha

Antworten:

2

Berechnungswarteschlangen im Allgemeinen bedeuten nicht unbedingt, dass Sie jetzt 2x Versendungen parallel ausführen können. Eine einzelne Warteschlange, die die Recheneinheiten vollständig ausfüllt, hat einen besseren Durchsatz. Mehrere Warteschlangen sind nützlich, wenn eine Warteschlange weniger Ressourcen verbraucht (gemeinsam genutzter Speicher oder Register). Dann können sich sekundäre Warteschlangen auf derselben Recheneinheit überlappen.

Für das Echtzeit-Rendering ist dies insbesondere bei Dingen wie dem Schatten-Rendering der Fall, die für Computer / Shader sehr leicht, für Hardware mit festen Funktionen jedoch sehr schwer sind, wodurch der GPU-Scheduler frei wird, die sekundäre Warteschlange asynchron auszuführen.

Fand dies auch in den Versionshinweisen. Ich weiß nicht, ob es das gleiche Problem ist, aber es könnte sein, dass CodeXL nicht großartig ist. Ich würde erwarten, dass es möglicherweise nicht die beste Instrumentierung hat, für die Sendungen im Flug sind.

https://developer.amd.com/wordpress/media/2013/02/AMD_CodeXL_Release_Notes.pdf

Bei einer Anwendung, die gleichzeitig asynchrone Datenübertragung und Kernelausführung ausführt, werden diese Vorgänge in der in der Sitzungsansicht "Application Trace" angezeigten Zeitleiste nicht überlappt. Dies liegt daran, dass der Treiber und die Hardware diese Vorgänge während der Profilerstellung synchronisieren. (333981)

Calvin
quelle