ZFX + Developia

Ich kann gerade irgendwie nicht glauben, dass DirectX 11 Capability Flags wie TileBasedDeferredRenderer anbietet, ich aber keine Möglichkeit finde, die Anzahl der verfügbaren Compute Cores zu erfragen. Bin ich blind?

Also ich kann auf die Schnelle auch keinen Weg finden um diese Information abzufragen. Ehrlich gesagt wundert mich das aber auch absolut ganz und gar nicht, denn es lässt sich doch nichtmal architekturunabhängig definieren, was genau so ein "Compute Core" eigentlich sein sollte... ;)

Wofür genau brauchst du diese Info? Mir fällt atm kaum ein Anwendungsfall ein wo mir das Wissen über die Anzahl der "Compute Cores" was bringen würde, selbst wenn ich es hätte. Für die Performance wichtig ist in der Regel die richtige Aufteilung des ganzen in Thread Groups. Und D3D11 garantiert dir da einfach gewisse Limits für die Anzahl der Groups die du dispatchen kannst, die Anzahl der Threads pro Group und das verfügbare Groupshared Memory. Wenn man einfach schaut, dass die Anzahl der Threads in einer Group ein Vielfaches von 64 ist und genug Groups da sind, dann sollte man zumindest auf ATI und NVIDIA relativ optimale Performance erreichen (zumindest was das angeht).

Es geht mir um optimiertes Scheduling bei repetetiven Abarbeitungsprozessen mit sehr inkohärenten Wiederholungszahlen. In diesem Fall führt klassisches datenorientiertes Scheduling durch die GPU, wie du es beschreibst, dazu, dass das Arbeitspaket mit der größten Wiederholungszahl alle anderen Cores, deren Arbeitspaket kleiner ausgefallen ist, so lange gefangen hält, bis wirklich alle aktuell durch die entsprechende Gruppe bearbeiteten Pakete fertiggestellt sind. Erledigt man das Scheduling dagegen selbst, indem man einfach alle Cores mit genügend endlos laufenden Gruppen abdeckt, lassen sich die Arbeitspakete in Teilpakete kohärenter länge unterteilen, wobei arbeitslose Threads schon nach jedem Teilpaket die Gelegenheit bekommen, sich neue Arbeit zu beschaffen.

Für den Moment starte ich also wohl einfach genügend Threads für meine GPU. Auf GPUs mit weniger Cores werden damit eben einige Gruppen zu viel abgeschickt, was aber auch nicht weiter schlimm ist. Am Ende stellen diese einfach fest, dass keine Arbeit mehr übrig ist, und beenden sich im Anschluss sofort wieder.

Rein aus Neugier: was hast Du denn vor?

AMD: OpenCL ;)

dot hat geschrieben:AMD: OpenCL ;)

Ja gut, das wusste ich auch schon. ;) Keine Ahnung heiß noch immer eXile'isch-keine-Ahnung, nicht dass ich dumm wie Brot bin. :) Schnelle Googlierung hat noch CL_DEVICE_MAX_COMPUTE_UNITS hervorgebracht, in dessen Umgebung ich mal weiter suchen würde.

Obligatorische Jammerei: Der aktuelle Zustand zur Optimierung von CUDA/OpenCL/DirectCompute ist doch beschissen. Zwar gibt es mittlerweile hinreichend gute Tools von nVidia und AMD, aber im Endeffekt optimiert man nur auf eine Graphikgeneration hin (wenn's schlimm kommt: Auf nur eine Graphikkarte); und das ist ja wohl bescheuert. Schöne neue Graphikwelt.

Oh und in fast jedem Optimierungs-„Guide“ steht sinngemäß drin: „Wenn alles bis hier hin nicht geklappt hat, sollten sie mal versuchen, mit den Parametern wild herum zu experimentieren!“ No shit Sherlock! Da ist es fast schon verständlich, dass Microsoft auf eine so rigide Hardware-Abstraktion gesetzt hat, dass so etwas wie Warp-Size oder Anzahl SMs vollkommen unbekannt hinter der HAL ist.

Also über DirectCompute kann ich nicht urteilen, da ich es noch nie wirklich verwendet hab. Es ist wohl von allen genannten die API mit den größten Einschränkungen. Ich würde aber mal davon ausgehen, dass die verfügbaren DirectCompute Implementierungen dafür ziemlich gut funktionieren. Von OpenCL bin ich alles andere als begeistert. CUDA ist wohl die mit sehr großem Abstand am weitesten entwickelte Plattform, aber auch noch lange nicht wirklich toll...

Danke für die Hinweise, für präzise Informationen komme ich dann wohl kaum um proprietäre Hilfsbibliotheken drumrum.

Ich bin insgesamt erschüttert, wie wenig Information sich im Internet zum Thema Compute Shaders findet. Am hilfreichsten ist da noch der etwas ausführlichere SM-5.0-ASM-Teil der HLSL-Referenz, aber im Augenblick kann auch der mir nicht bei der Frage helfen, warum ein Aufruf von InterlockedAdd() auf einem RWTexture2DArray<uint> sofort zu undefiniertem Inhalt führt. :evil:

Bist du dir sicher dass Atomic Operations für ein RWTexture2DArray überhaupt verfügbar sind? Ich dachte bisher die funktionieren nur auf UAVs...

Ein RWTexture2DArray ist prinzipiell *ein* UAV. Das ist es eben, die Nuancen stehen nirgends.

Laut ASM-Abteil arbeitet atomic_iadd einfach auf u#-Registern, welche wiederum auch Texture2DArrays schlucken. Allerdings verlässt auch mich langsam der Glaube, dass das wirklich geht.

WTF, ich muss Unordered Access Views in der Compute Stage nach Gebrauch von Hand entbinden, und zwar nicht mit CSSetUnorderedAccessViews(0, 0, nullptr, nullptr), sondern mit einem extra nullptr-Array, welches alle CS-Register überdeckt.

Dann sind atomare Operationen aber tatsächlich auch auf mehrdimensionalen Texturen definiert.

CodingCat hat geschrieben:WTF, ich muss Unordered Access Views in der Compute Stage nach Gebrauch von Hand entbinden, und zwar nicht mit CSSetUnorderedAccessViews(0, 0, nullptr, nullptr), sondern mit einem extra nullptr-Array, welches alle CS-Register überdeckt.

Ist das nicht sowieso mit allen derartigen Methoden so, also auch bei den ganzen *SetShaderResource() Methoden etc.?
Ich mach's zumindest immer so, da ich in der Doku noch nichts finden konnte was eine einfachere Variante nahelegen würde...

Ja, ich kümmere mich nur sonst nie darum, weil an anderer Stelle in der Pipeline eh keine Konflikte auftreten können. In der Regel werden Input Resource Views von Direct3D automatisch entbunden, sobald konfligierende Output Views gesetzt werden. Der umgekehrte Fall tritt in einer normalen Rendering Pipeline praktisch nie ein. Wie wir nun gelernt haben, gilt dieser Automatismus umgekehrt aber offenbar nicht, d.h. Output Views dominieren Input Views.

Und mit Vollgas gegen die Wand:
Ich sehe schon, alles mit Strahlen wird verdammt haarig (rote Pixel wurden mehr als 3000x überschrieben, dass da noch eine 4 steht, grenzt schon an ein GPU-Wunder). :|

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()