ZFX + Developia

Okay, es ist offensichtlich essentiell, dass wesentlich mehr Threads als Cores da sind, damit die GPU überhaupt die Chance hat, Latenzen zu verstecken. Im Moment erhalte ich optimale Performance ab 64 Gruppen mit jeweils 64 Threads, was die Anzahl meiner Cores definitiv weit übersteigt. Langsam wird die Sache interessant. :)

Ich habe hier 336 CUDA Cores. Die neue Kepler nV GTX 690 hat 3072 CUDA Cores. Die effizienteste Art der Optimierung wäre wohl Zeitreise. ;)

Und ich brauche wohl doch eine API zur Ermittlung der Corezahl, sonst laufe ich selbst mit großzügiger Überschätzung der Gruppenzahl in 3 Jahren vermutlich nur noch bei 1/128 GPU-Auslastung. :P

Bin gerade ein wenig schockiert wieviel über CUDA ich in dem letzten Jahr schon wieder vergessen habe, aber ich bin mir sehr sicher dass transparente Skalierung auf neuere Hardware eines der angepriesenen Features ist.

Die best practice war (damals und IIRC) wenn du wie in deinem Beispiel 64 Threads pro Gruppe hast, nicht auch 64 Gruppen zu erstellen, sondern Problemgröße/64 viele und den Rest der Hardware zu überlassen. Das mag andere performance drawbacks haben, aber zumindest lastest du die Hardware immer aus, solange dein Problem groß genug ist.

Disclaimer: Ich finde das Single-Instruction-Multiple-Data Modell aus dem nVidia Marketing relativ überzeugend. Ich finde auch die Logik, dass diese schiere Anzahl von Cores und Threads nur durch strunzdämliches brute-force scheduling möglich ist überzeugend. Klar ist dass für ein konkretes Problem auf einer konkreten Karte eine handverlesene Anzahl von Threads und Blocks vermutlich besser ist als der brute-force Ansatz "einfach mal irre viele Threads erzeugen". Ich habe auch (durchaus mit gewissem Erstaunen) das ein oder andere Abstract zur Kenntnis genommen in dem es darum ging auf der CPU "bessere" scheduling Logik für die GPU laufen zu lassen. Ich selbst würde einfach nicht versuchen "zu schlau" zu sein, gerade weil Moores Gesetz dich dabei sowieso überholen wird, sondern stattdessen das Problem auf die Grundprinzipien zurechtschneidern die nVidia angibt und die hoffentlich auch in 10 Jahren noch gelten werden. Vielleicht bin ich aber auch zu herstellergläubig.

Das ist grundsätzlich richtig, und solange du pro Problemeinheit kohärenten Programmfluss in etwa gleicher Länge hast, ziemlich sicher auch optimal. In meinem konkreten Fall habe ich leider trotz weitestgehend kohärenten Programmflusses je nach Problemeinheit extrem unterschiedliche Ablauflängen. In diesem Fall laste ich mit der von dir beschriebenen Form des datenorientierten Parallelismus die GPU tatsächlich nicht voll aus, weil neue Problemeinheiten erst angegangen werden, wenn alle vorangegangenen Problemeinheiten in der jeweiligen Gruppe vollständig abgeschlossen worden sind. Dieses Problem dürfte sich mit der Anzahl der Cores sogar noch vergrößern, weil damit noch mehr Cores für einige Zeit unbeschäftigt wären.

Persistente Threads bieten gewissermaßen die Möglichkeit, der GPU zusätzliche Information über die Beschaffenheit des Problems zu übermitteln, nämlich die, dass einzelne Problemeinheiten unabhängig ihrer Länge über weite Strecken immer denselben Schleifenrumpf durchlaufen.

Ein Spin Lock in HLSL scheint gar keine gute Idee zu sein, weil das Scheduling kohärente Threads offenbar so lange laufen lässt wie nur möglich. :-(

Jetzt bin ich ziemlich aufgeschmissen, wie warte ich so lange, bis in einem anderen Thread eine Transaktion (mehrere Writes) abgeschlossen ist? Gibt es eine Möglichkeit, Threads auf Eis zu legen (sowas wie sleep())?

in CUDA innerhalb von Blocks __syncthreads(), blockübegreifend gar nicht.

Okay, ich habs geschafft, indem ich das Spin Lock bis in die äußerste Schleife aufgebogen habe, sodass sich die Verarbeitung von locked wie unlocked Datenelementen regelmäßig an gemeinsamen Kohärenzpunkten trifft.

So sagt das natürlich wenig aus, deshalb hier ein wenig Pseudocode, falls andere mal in ähnliche Situation kommen: Korrigierte Fassung: Tatsächlich ist auch die korrigierte Fassung wohl nicht 100% sicher, weil sie sich ungesicherter Annahmen über das Scheduling bedient, funktioniert aber zumindest auf meiner GPU hervorragend. :mrgreen:

Mit Spinlocks musst du wahnsinnig vorsichtig sein, denn da kann dich das Scheduling auf gewissen Karten deadlocken. NVIDIA scheduled Threads in so genannten Warps, die als SIMD Einheit im Lockstep ausgeführt werden. Branches funktionieren so, dass alle Threads im einen Zweig deaktiviert werden, während der andere Zweig ausgeführt wird und umgekehrt. Ich denk das ist dir soweit bewusst. Das kann aber nun problematisch werden, wenn ein Thread in einem Warp ein Lock hält und die restlichen Threads des Warp auf dieses Lock warten. Da die wartenden Threads nun in einer Endlosschleife laufen, kann es passieren, dass die Branch des einen Thread der das Lock hält nie mehr aktiv wird -> Deadlock.

Ja, genau diese Situation hatte ich ja vor der Korrektur. Mit dem von dir beschriebenen nVidia-Scheduling sollte die Lösung robust laufen, die Frage ist, wie das auf anderen Karten ausschaut. ATI werde ich morgen testen. ;)

Ah, sry, sollte wohl den ganzen Thread lesen... ;)

Der letzte Post hätte auch gelangt. ;) Egal, dafür ist es jetzt nochmal im Detail klar. :)

So, die Lösung mit dem aufgebogenen Spin Lock funktioniert auch auf ATI einwandfrei, wies ausschaut.

Sehe ich das richtig, dass ich die Anzahl der Elemente in einer Compute Shader Gruppe nicht nur hartcodieren muss, sondern obendrein noch nicht mal abfragen kann, und somit auch auf Anwendungsseite von vorneherein festsetzen muss?

Nachtrag: Doch nicht, es ist nur erstaunlicherweise eine separate Funktion in der Shader Reflection: ID3D11ShaderReflection::GetThreadGroupSize()