OpenCL auf NVidia - Unterschied zwischen Fermi und Kepler?!

[Nakai]

Ich bin gerade dabei eine Applikation in der medizinischen Bildverarbeitung auf mehrere Rechner zu verteilen. Es handelt es sich um einen EM-Clustering-Algorithmus im Bereich der Neurochirurgie zur Segmentierung des gesamten Hirnstamms. Das nur zur Info.

Zur Beschleunigung wird OpenCL für GPUs verwendet. Ich entwickle OpenCL grundsätzlich auf AMD-Karten(FirePRO M4100), weil es einfach leichter ist. :)
Ich muss die Software auch auf NVidia-Maschinen verteilen. Die Applikation lauft super auf einem Rechner mit einer Tesla C2075(Fermi), aber auf einem Rechner mit einer Quadro K4000(GK106) gibt es Probleme. Die Ausführung ist arschlangsam, die Benutzeroberfläche, also Desktop, während der Ausführung reagiert träge, es zuckelt alles. Der Rechner hat nur eine K4000 als GPU enthalten, keine sonstige GPU. Ich vermute es liegt an fehlendem Hyper-Q auf GK106 oder einfach OpenCL auf Nvidia. GK106 hat CC3.0, aber ich finde hierzu wenig zu Hyper-Q. Weißt jemand hierzu mehr, oder liegt es an den Kernels? Auf Fermi hat es wunderbar funktioniert(zwei GPUs hier im System: einmal Fermi und einmal eine AMD). Liegt es am fehlendem Hyper-Q? Fühlt sich jedenfalls so an...:freak:

[del_4901]

Windows/Linux? Unter Windows schau einfach mal mit WPR/GPUView wo es klemmt.

[Nakai]

Bei diesem Rechner ist es Windows, ansonsten für Mac, Linux und Windows.
Werd ich mal überprüfen. Danke für das Feedback.

Also für Hyper-Q, also mehrere work queues, wird CC3.5 benötigt. GK104, GK106, GK107/8 haben wohl nur CC3.0 und daher kein Hyper-Q. Scheint jedenfalls daran zu liegen und auch die Symptomatik passt darauf.

[Locuza]

Hyper-Q kam auch erst mit dem GK110.
Und wieso sollte es dann auf einer Tesla C2075 gut laufen, die darüber auch nicht verfügt?

Ich kann es mir als Laie schwer vorstellen, dass Mars, welcher meines Wissens auf GCN 1.0 basiert, mit 2 bis maximal 4 Queues maßgeblich bei 384 ALUs davon profitieren könnte.
Ganz davon ab, ob OpenCL 1.2 die ACEs unterstützt?
Gibt es Erweiterungen von AMD diese zu unterstützen?

Und falls man eine GPU mit Hyper-Q erwischt, besteht mit OpenCL überhaupt die Möglichkeit das zu verwenden, da ich ich schlimme Sachen ahne, aka CUDA exklusiv und schleifender OCL-Support.

[Nakai]

Hyper-Q kam auch erst mit dem GK110.
Und wieso sollte es dann auf einer Tesla C2075 gut laufen, die darüber auch nicht verfügt?

Ich kann es mir als Laie schwer vorstellen, dass Mars, welcher meines Wissens auf GCN 1.0 basiert, mit 2 bis maximal 4 Queues maßgeblich bei 384 ALUs davon profitieren könnte.
Ganz davon ab, ob OpenCL 1.2 die ACEs unterstützt?
Gibt es Erweiterungen von AMD diese zu unterstützen?

Und falls man eine GPU mit Hyper-Q erwischt, besteht mit OpenCL überhaupt die Möglichkeit das zu verwenden, da ich ich schlimme Sachen ahne, aka CUDA exklusiv und schleifender OCL-Support.

Der Tesla-Rechner hat zwei GPUs, irgendeine AMD und eben die C2075 für Compute. Mein eigener Rechner ist ein Notebook mit einer FirePro M4100(Mars) und einer integrated Intel. Es ist ein HP ZBook14.
Hyper-Q und ACEs sollte doch vollkommen transparent für den Entwickler aussehen, oder liege ich da falsch? Es geht hierbei nicht darum, dass mehrere OpenCL-Kernel parallel auf die GPU geschmissen werden, nein, diese werden sequentiell ausgeführt. Ich vermute, es kollidiert einfach mit dem Grafik-Task, welcher auch auf diese GPU ausgeführt wird.

[Locuza]

Ich hoffe für dich, dass man im Profiler schnell etwas erkennt.
Wenn es um die Theorie Rendering + Computing geht, dürftest du bei dem Tesla-Rechner die AMD ausbauen?

Ich habe einfach meine Zweifel daran, dass eine Radeon mit halber ALU-Anzahl deswegen flüssig läuft und eine Quadro K4000 so einbricht, dass der Desktop zuckelt.

[Nakai]

Ausbauen bräuchte ich gar nicht, es würde reichen, mal die Grafikdarstellung auf die Tesla zu berechnen und gleichzeitig einen Compute-Kernel auszuführen(C2075 hat sogar einen Grafikausgang).
Wenns dann auch ruckelt und zuckelt, dann liegt es einfach daran, dass NV-Karten(Pre Hyper-Q) eben nicht Compute und Grafik gleichzeitig berechnen können.

Windows/Linux? Unter Windows schau einfach mal mit WPR/GPUView wo es klemmt.

Das ist auf meinem Notebook schonmal installiert. Ich vergleich es einfach mal mit dem anderen Rechner, mal gucken, was dabei rauskommt.

@Locuza:
Ich habe einfach meine Zweifel daran, dass eine Radeon mit halber ALU-Anzahl deswegen flüssig läuft und eine Quadro K4000 so einbricht, dass der Desktop zuckelt.

Das wundert mich auch, vor allem, da es auf einer C2075(Fermi) hervorragend läuft. OpenCL auf NV ist einfach ein Graus...ich hatte schon ganze andere Probleme damit.

[Nakai]

So ich hab mal mit GPUView die Queues angeguckt. Hier der Dropbox-Link für die Merges:
https://www.dropbox.com/sh/fl2v24iu8gs5ipz/AADpmQYlkOK-tSEdPX7NkTIPa?dl=0

Bei meiner FirePro M4100 sieht es gut aus, bei Quadro K4000 nicht. Bei fast jedem DMA Packet wird ein Preemption Packet hinterhergeschmissen. Dadurch wird die Ausführungzeit dramatisch nach oben getrieben. Die Software ist FAST-CL.exe bzw. FAST-CL_GPU.exe.

[del_4901]

Ich hab noch nicht reingeschaut, aber DMA Paket ist verdeachtig. Ist dein System unter Memory pressure? und WDDM versucht die ganze Zeit zu pagen? Wie sieht es mit page faults aus?

Preemption Packet: http://systemscenter.ru/gpuview.en/gpuview/selections_in_the_gpu_hardware_queue.htm
Windows versucht wohl die seine UI zu rendern. :)

[Nakai]

Also ich habe mit log.cmd nur mit min erstellt, da es sonst zu groß wird. Nicht wundern, wenn einiges fehlt. Auf den Merge der K4000 nicht soviel geben, der ist nicht vollständig. Ich wollte eher auf die Pre-emption Packete abzielen.

Memory Pressure kann ich mir schwer vorstellen. Beide Systeme haben genug Speicher. Die Quadro K4000 hat 3GB die FirePRO M4100 nur 1GB. Auf der AMD läuft es ohne Probleme.

Windows versucht wohl die seine UI zu rendern.

Ja, das denke ich mir auch. Interessant finde ich die Größe der DMA Packete. Auf der AMD sind diese immer 288Byte groß, während auf der NV nur 4 Byte große DMA-Packete versendet werden. Außerdem wird pro 4 Byte DMA-Packet ein Pre-emption Packet dazugeschmissen.

€: Wie sieht es mit page faults aus?

Sollte da nicht Paging Packete verschickt werden?

[Nakai]

Windows versucht wohl die seine UI zu rendern.

Ich hab es nochmal etwas genauer angeguckt. Ich habe ehrlich gesagt nicht viel Ahnung, was passiert. Mittlerweile habe ich auf einer Tesla C2075 auch den GPU-Ausgang nutzen können, und es führt zu keinen solchen Problemen. Bei gewissen Vsync Interrupt Events wird dann ein Preemption Paket geschmissen und gleichzeitig immer eine Warning.

Ja, sieht so aus, dass Windows seine GUI rendern will, aber diese Probleme habe ich woanders nicht. Gibt es irgendwelche WDDM-Einstellungen, Treiber-Settings, o.Ä mit denen ich das beheben kann? Würde eine zweite GPU das Problem beheben(nur für Compute?).


mfg

[del_4901]

Ich hab es nochmal etwas genauer angeguckt. Ich habe ehrlich gesagt nicht viel Ahnung, was passiert. Mittlerweile habe ich auf einer Tesla C2075 auch den GPU-Ausgang nutzen können, und es führt zu keinen solchen Problemen. Bei gewissen Vsync Interrupt Events wird dann ein Preemption Paket geschmissen und gleichzeitig immer eine Warning.

Ja, sieht so aus, dass Windows seine GUI rendern will, aber diese Probleme habe ich woanders nicht. Gibt es irgendwelche WDDM-Einstellungen, Treiber-Settings, o.Ä mit denen ich das beheben kann? Würde eine zweite GPU das Problem beheben(nur für Compute?).


mfg
eine 2te GPU wird warscheinlich das Problem beheben. Ansonsten: Willkommen in Club der Leute die mit WDDM zu kaempfen haben. Allerdings bin ich recht stolz auf dich, dass du es geschafft hast, das WDDM keine Arbeit mehr gescheduled bekommt, ich versuche genau das zu erreichen. :)

[Nakai]

https://devtalk.nvidia.com/default/topic/559274/cuda-programming-and-performance/opencl-asynchronous-kernel-launches/

Womöglich liegt es auch daran. NVs Implementierung von OpenCL ist anscheinend buggy. Ich habe die Software etwas besser profilen können und die Ausführung ist hauptsächlich da langsam, wo zwei Kernels sequentiell in einer Schleife sehr häufig ausgeführt werden(ähnlich dem Link). Dabei handelt es sich um drei Lowpass-Kernels(Z,loop{X,Y}; In dieser Reihenfolge). LowpassFilterung in X und Y sind zwei Kernels, welche in einer Schleife sequentiell öfter ausgeführt werden(X, Y, X, Y, X, Y, ...), bis die Iterationstiefe erreicht ist.

Anscheinend ist NVs clEnqueueNDRangeKernel-Aufruf immer blockierend. Ich benutze eine OpenCL-Wrapper-Bibliothek(ITK; medizinische Bildverarbeitung), welche nach jedem clEnqueueNDRangeKernel einmal clFinish aufruft, um auf die GPU zu warten(es gibt auch andere Methoden, welche nicht warten).

Sonderbar ist es, dass es auf einer Nvidia Tesla C2075(Fermi) funktioniert, aber auf einer K4000 NICHT!

€: Auf einer Quadro 4000 läufts auch. Ich teste es demnächst auch auf einem Mac Book Pro. Hat NV bei den Keplers was mit OpenCl geändert?

[Nakai]

Ein kurzes Feedback.

Ich bin mir nicht sicher was genau passiert. Die GUI ist wohl nicht schuld. Sogenannte Preemption Packets erscheinen nur auf dem Trace, wenn ich wirklich versuche etwas an dem Desktop zu ändern bzw. weitere Programme parallel laufen lasse.

Es gibt jedoch eine Neuigkeit. Ich habe meine Kernels umstrukturiert, sodass diese deutlich weniger VGPRs verwenden(hier mit CodeXL gecheckt). Hierzu einige Private Variablen zusammengestrichen, lokale Speicher erstellt zum auslagern, Code umstrukturiert, etc. Der Effekt war sofort spürbar, denn die Windows GUI wurde deutlich reaktiver. Auch die Ausführungsgeschwindigkeit ging dramatisch hoch, natürlich bei weitem immer noch viel zu langsam, aber immerhin besser.
Ich muss mir das nochmal genauer anschauen...aber es sieht irgendwie so aus, dass Kepler hier in eine dramatische Register/Memory Pressure gerutscht ist.

Nur zum Vergleich:
GCN 256KB Registerfile 64KB Local Memory, 64SPs
Fermi 32KB Registerfile, 16/48KB, 32SPs
Kepler 64KB Registerfile, 64KB, 192SPs

Interessant ist eher, dass dies hauptsächlich in nur einer bestimmten Schleife passiert, welche sequentiell einen Kernel für Lowpass in X und dann einen Kernel für Lowpass in Y ausführt. Schleife wird 100mal ausgeführt. Ich befürchte, dass diese auf wechselnden Aufrufe ein Problem verursachen.

[Nakai]

Erstmal ein Nachtrag zum vorherigen Post.

Fermi 128KB Registerfile, 16/48KB, 32SPs
Kepler 256KB Registerfile, 64KB, 192SPs

Ich könnte mich in den Arsch beissen. Der Fehler ist gefunden...

Ich habe in meinem Tiefpassfilter ein paar Barriers gehabt. Die sahen so aus:

barrier(CLK_GLOBAL_MEM_FENCE|CLK_LOCAL_MEM_FENCE);

wurde zu

barrier(CLK_LOCAL_MEM_FENCE);

Mir ist sehr wohl bewusst, was der Unterschied ist. Also einmal Flag für globalen Speicher und lokalen Speicher -> nur lokalen Speicher(ich mach ja nur lokale Speicher-Operationen). Ich hätte so einen Performancezugewinn, nur dadurch, NICHT erwartet.

Geschwindigkeitsboost für den Tiefpassfilter Faktor ~30. Auf Fermi war das egal. GCN sowieso.