News DirectX 12: Nvidia-Treiber soll Async Compute nachreichen

[ampre]

Entscheidend ist wie gesagt der Anteil an zu kleinen Polygonen weil das den Verschnitt erhöht. Bei einer klassischen Pipeline ohne Deferred Rendering und einem entsprechend größeren Pixelshader, kann dieser Verschnitt AMD mal eben bis zu 50% langsamer machen als Nvidia, im absoluten Worstcase, da der Verschnitt mit dem Anteil der Pixelshader an der Gesamtlaufzeit auf multipliziert wird.

Da reichen bei einer ungünstig geschriebenen Pipeline bereits ein paar Objekte im Hintergrund bei denen LOD nicht korrekt funktioniert um die Performance bereits spürbar runter zu ziehen. Das lässt sich bei AMD bis zu dem Punkt treiben, wo die GPU zwar 100% ausgelastet ist, aber effektiv nur 1/64tel der Rechenleistung tatsächlich genutzt wird. Bei Nvidia "nur" bis runter auf 1/32tel. Oder bei AMD runter auf 1/2, während Nvidia noch überhaupt keine Verluste hin nimmt, wenn man die Situation böswillig herbeiführen will.

Und warum ist das kein Problem für Nvidia was machen die da anderster? Kannst du deine Aussagen etwas vereinfachen. Das ist für einen GPU Anfänger doch reichlich Komplex. Trozdem Danke für die Ausführungen

 

[Ext3h]

Das ist für Nvidia genauso ein Problem, nur um eine Größenordnung kleiner. AMD kann aufgrund Limitierungen der Hardware Threads immer nur in Gruppen von 64 bearbeiten, Nvidia in Gruppen von 32. Dabei können allerdings nur Pixel die zum gleichen Polygon gehörten auch in der gleichen Gruppe behandelt werden.

Bei Polygonen die z.B. exakt 32 Pixel erzeugen kann Nvidia exakt noch eine Gruppe voll bekommen, sprich volle Effizienz, bei AMD hingegen muss die Hälfte der Threads in einer Gruppe das Ergebnis wieder verwerfen, sprich nur noch 50% Effizienz.

 

[ampre]

Aber ist das nicht der Vorteil von Async Compute diese Abhängikgkeit aufzulösen?

 

[Nai]

Defakto "kann" die Hardware von GK110, GK20X, GM10X und GM20X "Async Compute" nicht - bzw. nicht vollständig. Theoretisch "könnte" man dafür die HyperQ-Funktionalität verwenden, allerdings fehlt der ein wichtiges Feature: Barrieren zwischen Queues.

Die Idee bei den Async Queues bei DX12, ist dass diese über simple Barrieren miteinander synchronisiert werden können, und diese Barrieren in Hardware, oder zu mindestens Hardware-nah aufgelöst werden so dass die GPU nahezu verzögerungsfrei zwischen zwischen verschiedenen Aufgaben wechseln kann, abhängig davon was gerade bereits fertig ist.

Bist du dir da sicher? Ich habe die Tage einmal das ganze in CUDA getestet. Da besitzen eine Synchronisation zwischen zwei Streams per Events circa 10 µs Latenz und die Synchronisation über Software und die CPU circa 40 µs. Der große Unterschied deutet m.E. stark darauf hin, dass die Synchronization auf der GPU selbst stattfindet.

Die Anzahl der Polygone an sich sagt erst mal nichts - immerhin können die sich ja immer noch überlappen. Problematisch sind wirklich nur übertrieben kleine Polygone. Sowohl Nvidia als auch AMD können AFAIK aus dem Rasterizer immer nur mindestens volle Wavefronts rausschicken, und eine Wavefront kann nur Pixel aus einem Polygon beinhalten.

Ich nehme stark an, dass das falsch ist. Das erkennt man wahrscheinlich zum Beispiel an den Occupancy-Regeln für den Fragmentshader beim GCN. Dort darf der VS dem Rasterisierer für maximale Occupancy maximal 64 Byte zur Interpolation übergeben, weil ansonsten die 64 kibyte Scratchpad-Speicher der CU ausgehen. Das ergibt wiederum 192 Byte pro Dreieck. Jede Wavefront besitzt jedoch bei einer maximaler Occupancy von 40 circa 1638 Byte. Ergo werden die Interpolationsattribute nicht für Wavefronts sondern für kleinere Einheiten (1638 / 192 ~ 8 Pixel) alloziert. Diese feinere Allozierung wäre wiederum sinnfrei, wenn sowieso alle Fragmente einer Wavefront das selbe Polygon bearbeiten müssten. M.E. aber insgesamt sehr schade, dass kein GPU Hersteller seine Rasterpipeline-Implementierung ausführlich dokumentiert, da so das alles starkes stochern des Nebels ist.

 

[Ext3h]

Aber ist das nicht der Vorteil von Async Compute diese Abhängikgkeit aufzulösen?

Nein, Async Compute arbeitet auf einem viel, viel höheren Level. Nur um hier mal die Größenordnungen zu verdeutlichen:

So eine Fiji-GPU hat 4096 Prozessoren, die jeweils (stark vereinfacht gesagt) jeweils in Gruppe von 64 zusammen gefasst sind, so dass jeweils 64 Prozessoren immer exakt das gleiche Programm ausführen müssen. Dabei beherrscht jede dieser Prozessorgruppe noch mal eine Art von "Hyperthreading" durch die sie 4-10 Threadgruppen gleichzeitig ausführt. Das sind 16-40.000 Threads die da parallel laufen.

Gestartet werden diese Threads entweder über die Grafik-Pipeline, wobei erst mal jeder Vertex eines Polygons einen Thread bedeutet, und dann noch mal jeder Pixel eines Polygons einen weiteren. (Ebenfalls stark vereinfacht, kann wegen Tesselation noch mal mehr werden.) Oder über ein Compute Command, das effektiv sagt "Ich will x*y*z Threads haben, mit der Möglichkeit jeweils X Stück davon zu synchronisieren".

Async Compute bedeutet jetzt, dass die Compute-Programme nicht nur abwechselnd mit Drawcalls ausgeführt werden können, sondern völlig unabhängig davon gleichzeitig.

So kann die GPU z.B. gerade noch damit beschäftigt sein über Drawcalls die Shadowmap aufzubauen, und während dessen beginnt ein Compute-Shader schon mal damit die (bereits vorher in einen G-Puffer gerenderte) Szene fertig zu texturieren, während wieder ein anderer Compute-Shader asynchron die Beleuchtung berechnet, und wieder ein anderer Compute-Shader schon mal damit beginnt Motion Blur vorzubereiten. Wieder nur stark vereinfachte, und nicht unbedingt realistische Beispiele.

Das hat jetzt allerdings nichts damit zu tun wie gut die Prozessorgruppen intern ausgelastet sind, sondern sorgt lediglich dafür, dass die GPU praktisch immer etwas zu tun hat.

Ergänzung (24. September 2015)

Bist du dir da sicher? Ich habe die Tage einmal das ganze in CUDA getestet. Da besitzen eine Synchronisation zwischen zwei Streams per Events circa 10 µs Latenz und die Synchronisation über Software und die CPU circa 40 µs. Der große Unterschied deutet m.E. stark darauf hin, dass die Synchronization auf der GPU selbst stattfindet.

Nein, ich bin mir nicht völlig sicher. Allerdings muss man da jetzt unterscheiden zwischen Streams die bereits laufen, sprich bereits in den Shadern sitzen und dann dort auch Register etc. blockieren während sie warten, und Streams die noch nicht mal im Work Distributor angekommen sind. Bei ersteren geht das garantiert direkt über die GPU, da sieht auch die CPU nichts von.

Bei den zweiten, ich befürchte dass ausgerechnet die eben nicht funktionieren. Solange die Warps noch nicht dispatched sind, ist Synchronisation auf der GPU selbst nicht vorgesehen.

Allerdings sind auch so 10 µs eigentlich eine ganze Menge, das ist immer noch genug Zeit für einen CPU-Roundtrip. Nicht genug für einen Roundtrip in den Userspace, was du mit "Synchronisation per Software" wahrscheinlich impliziert hast, aber genug für den Treiber.

Ich nehme stark an, dass das falsch ist. Das erkennt man wahrscheinlich zum Beispiel an den Occupancy-Regeln für den Fragmentshader beim GCN. Dort darf der VS dem Rasterisierer für maximale Occupancy maximal 64 Byte zur Interpolation übergeben, weil ansonsten die 64 kibyte Scratchpad-Speicher der CU ausgehen. Das ergibt wiederum 192 Byte pro Dreieck. Jede Wavefront besitzt jedoch bei einer maximaler Occupancy von 40 circa 1638 Byte. Ergo werden die Interpolationsattribute nicht für Wavefronts sondern für kleinere Einheiten (1638 / 192 ~ 8 Pixel) alloziert. Diese feinere Allozierung wäre wiederum sinnfrei, wenn sowieso alle Fragmente einer Wavefront das selbe Polygon bearbeiten müssten. M.E. aber insgesamt sehr schade, dass kein GPU Hersteller seine Rasterpipeline-Implementierung ausführlich dokumentiert, da so das alles starkes stochern des Nebels ist.

Herzlichen Dank für den Hinweis.

8 Pixel? Klingt immer noch plausibel, ich hatte so einen Verdacht dass etwas in der Art möglich wäre, aber hatte selber keine Indizien gefunden. OK, das verschiebt die Grenzen in deutlich realistischere Größenordnungen. Bedeutet aber auch dass unterhalb der 8 Pixel immer noch Probleme mit der Auslastung auftreten. Und würde indirekte Adressierung, bzw. zu mindestens Adressierung per Offset implizieren. Aber auf jeden Fall macht AMD die Interpolation im Shader, nicht im Rasterizer, das ist bekannt.

Gut, das mit den 64 FS Instanzen war auch schon eigentlich unrealistisch hoch.

Bedeutet auch, dass das bei Nvidia tatsächlich noch mal feiner aufgeteilt wird, wenn die da weniger empfindlich reagieren. Eventuell ist an den Gerüchten was dran, dass die Interpolation bei denen im Rasterizer durchgeführt wird, würde den Platzbedarf pro Fragment auf 1/3 reduzieren.

 

[ampre]

Ich glaube AMD ist auch noch mal feiner aufgelöst. Jede CU hat 4 SIMDs die unterschiedliche Wavefronts abarteiben können.

The SIMDs hide memory latency by having several wavefronts in flight at the same time, allowing the compute unit scheduler to switch between different wavefronts. For example, while one wavefront is waiting for results from memory, other wavefronts can issue memory requests. Each SIMD supports a maximum of 10 simultaneous wavefronts in flight. However, whether a particular kernel (i.e. shader) can achieve the maximum depends on several factors. For HLSL pixel shaders, the limiting factor is often VGPR usage.

http://developer.amd.com/community/blog/2014/05/16/codexl-game-developers-analyze-hlsl-gcn/

 

[Ext3h]

Das mit CU/SIMD habe ich erst mal bewusst unterschlagen, genauso wie Nvidia analog zu den CU-Einheiten auch noch mal SMX/SMM-Einheiten hat die jeweils eine Gruppe von SIMD-Prozessoren zusammen fasst. Die striktesten Limits sind aber jeweils pro SIMD-Einheit, nicht pro CU/SMM. Die Größe der SIMD-Einheiten sind 64 bei GCN, und 32 Threads bei Nvidia.

 

[ampre]

So eine Fiji-GPU hat 4096 Prozessoren, die jeweils (stark vereinfacht gesagt) jeweils in Gruppe von 64 zusammen gefasst sind, so dass jeweils 64 Prozessoren immer exakt das gleiche Programm ausführen müssen. Dabei beherrscht jede dieser Prozessorgruppe noch mal eine Art von "Hyperthreading" durch die sie 4-10 Threadgruppen gleichzeitig ausführt. Das sind 16-40.000 Threads die da parallel laufen.

Was meinst du hier mit gleichem Programm? Führt die CU ein Porgramm aus? Ich dachte die CU managet nur wie die Pixelshader ausgelastet werden sollen.

 

[Ext3h]

Was meinst du hier mit gleichem Programm? Führt die CU ein Porgramm aus? Ich dachte die CU managet nur wie die Pixelshader ausgelastet werden sollen.

Nein, die CU ist eine Ebene über den SIMD-Einheiten, und verwaltet jeweils 4 Stück davon. Die SIMD-Einheiten führen die Shader aus, und haben eine Breite von 64 Threads. Sobald die Threads erst mal in einer SIMD-Einheit drinnen sind kommen die auch nicht mehr raus bis fertig. Die SIMD-Einheit muss dabei für alle 64 Threads immer jeweils exakt die gleiche Instruktion ausführen, geht nicht anders

 

[ampre]

Ah jetzt verstehe ich. Nvidia hat also den Vorteil das sie zwar kleinere Simds davon aber mehrere haben. Dadurch können unterschiedliche Threads besser gehandhabt werden. Bei AMD wären also Threads von gleicher Natur vorteilhaft.

 

[Ext3h]

Ah jetzt verstehe ich. Nvidia hat also den Vorteil das sie zwar kleinere Simds davon aber mehrere haben. Dadurch können unterschiedliche Threads besser gehandhabt werden. Bei AMD wären also Threads von gleicher Natur vorteilhaft.

Schön wärs, so einfach ist das schon wieder nicht.

AMD hat da noch die "eigenartige" Eigenschaft, dass pro SIMD eigentlich immer mindestens 4 und bis zu 10 verschiedene Prozesse mit jeweils 64 Threads laufen. Und bei weniger als 4 verschwendet man auch schon wieder Rechenzeit.

Nvidia ist da bereits mit einem einzigen "Prozess" pro SIMD zufrieden.

Also so viel zu der Annahme "mehrere SIMDs". Tatsächlich erfordert AMD gleichzeitig größere Thread-Gruppen, und oben drein auch noch wesentlich mehr davon.

Das mit "von gleicher Natur" gilt jeweils immer nur für eine einzige Threadgruppe - eine von mindestens 256 und bis zu 640 die gerade aktiv ist.

 

[Nai]

Bei den zweiten, ich befürchte dass ausgerechnet die eben nicht funktionieren. Solange die Warps noch nicht dispatched sind, ist Synchronisation auf der GPU selbst nicht vorgesehen.

Das ist doch afaik nirgendwo möglich, und wäre mit dem Programmiermodel auch relativ komisch. Das einzige ähnliche was in dieser Hinsicht geht ist die Synchronisation per Dynamic Parallelism auf gestartete Kind-Gitter.

Allerdings sind auch so 10 µs eigentlich eine ganze Menge, das ist immer noch genug Zeit für einen CPU-Roundtrip. Nicht genug für einen Roundtrip in den Userspace, was du mit "Synchronisation per Software" wahrscheinlich impliziert hast, aber genug für den Treiber.

Deswegen war ich mir bei meiner Deutung auch nicht 100 \% sicher. Aber 30 µs kamen mir für einen Roundtrip etwas groß vor.

8 Pixel? Klingt immer noch plausibel, ich hatte so einen Verdacht dass etwas in der Art möglich wäre, aber hatte selber keine Indizien gefunden. OK, das verschiebt die Grenzen in deutlich realistischere Größenordnungen.

Die 8 Pixel sind auch nur eine etwas grobere Milchmädchenschätzung meinerseits. Es kann sein, dass noch ein paar Byte Scratchpad pro Fragment benötigt werden und es zusätzliche Allozierungsrestriktionen für den Scratchpad gibt, so dass sich der Wert nocheinmal auf zum Beispiel 16 Pixel erhöht. Meine Verumutung basierte auch eher darauf, dass es algorithmisch gesehen leicht möglich wäre, weshalb ich nach "Indizen" gesucht habe, dass es die GPUs diese Optimierung auch tatsächlich nutzen.

Interessanterweise sind es immer mindestens 4 Pixel oder ein Vielfaches davon, da GPUs beim Zeichnen das Bild in Quadrate von 2 auf 2 Pixel unterteilen. Das Unterteilen ist wiederum nötig, weil sie dann dadurch im Fragmentshader innerhalb des Quadrats per Forward und Backward Difference leicht die Screen-Space-Gradienten von beliebigen Werten berechnen kann. Diese Gradienten werden zum Beispiel wiederum für die Texturierung benötigt.

 

[ampre]

Das wird hier ganz schön komplex kann das einer noch mal vereinfacht zusammenfassen?

Verstehe ich das jetzt Richtig das AMD also kein Problem hat Polygonen in Pixel umzusetzen, sondern die anfallenden Pixel alle in den Cu's abzuarbeiten?

 

[Krautmaster]

http://www.anandtech.com/show/9659/fable-legends-directx-12-benchmark-analysis/2

schon bessere Benchmarks zu DX12...

 

[Nai]

So habe die Synchronisation von zwei Streams per Events und ohne Roundtrip nocheinmal etwas getestet: Die Latenzen skalieren nicht mit dem CPU-Takt, aber ist skalieren fast linear mit dem GPU-Takt. Ergo scheint es komplett auf der GPU abzulaufen.

 

[Ext3h]

Das ist doch afaik nirgendwo möglich, und wäre mit dem Programmiermodel auch relativ komisch. Das einzige ähnliche was in dieser Hinsicht geht ist die Synchronisation per Dynamic Parallelism auf gestartete Kind-Gitter.

Das Lustige ist: Bei AMD geht es. Die ACEs sind speziell darauf ausgelegt Barrieren im Global Data Share zu überwachen und entsprechend dann einen der sich bereits in einer der 8 Queues befindlichen Compute Commands zur Ausführung zu bringen, noch BEVOR der aktualisierte Wert der Barriere an die CPU übermittelt wurde. Und das dann wiederum dann noch mal unabhängig von den 64 Slots pro ACE die für bereits laufende Commands reserviert sind. Und das dann noch mal x8 weil 8 ACEs. Das ist einfach pervers was die da für einen Durchsatz vorgesehen haben.

Interessanterweise sind es immer mindestens 4 Pixel oder ein Vielfaches davon, da GPUs beim Zeichnen das Bild in Quadrate von 2 auf 2 Pixel unterteilen. Das Unterteilen ist wiederum nötig, weil sie dann dadurch im Fragmentshader innerhalb des Quadrats per Forward und Backward Difference leicht die Screen-Space-Gradienten von beliebigen Werten berechnen kann. Diese Gradienten werden zum Beispiel wiederum für die Texturierung benötigt.

Selbst 4 Pixel wären immer noch plausibel, wenn man "Full Occupancy" nicht wörtlich nimmt, sondern das untere Limit von 4 Wavefronts pro SIMD annimmt. Das wird aber schwer das wirklich nachzuvollziehen. Ansonsten wären 16 Pixel aber wirklich plausibler. Wenn da jetzt an der Stelle nicht auch noch das Abtastmuster des jeweiligen AA-Verfahren mit rein spielen würde was das ganze noch mal komplexer macht...

Verstehe ich das jetzt Richtig das AMD also kein Problem hat Polygonen in Pixel umzusetzen, sondern die anfallenden Pixel alle in den Cu's abzuarbeiten?

Ja, es scheint da Randfälle zu geben in denen es nicht mehr optimal funktioniert.

Ergänzung (25. September 2015)

So habe die Synchronisation von zwei Streams per Events und ohne Roundtrip nocheinmal etwas getestet: Die Latenzen skalieren nicht mit dem CPU-Takt, aber ist skalieren fast linear mit dem GPU-Takt. Ergo scheint es komplett auf der GPU abzulaufen.

10 µs sind dafür aber eigentlich trotzdem noch zu lang. Es sei denn...

Das wird dann wahrscheinlich doch von der Grid Managment Unit unterstützt. Die damit mit ziemlicher Sicherheit ein µC ist, so langsam wie die arbeitet. Und an den GPU-Takt gekoppelt ist.

 

[ampre]

Und wie hängt nun die Wavefront mit der Auslastung der SIMDs zusammen?

 

[Nai]

Das Lustige ist: Bei AMD geht es. Die ACEs sind speziell darauf ausgelegt Barrieren im Global Data Share zu überwachen und entsprechend dann einen der sich bereits in einer der 8 Queues befindlichen Compute Commands zur Ausführung zu bringen, noch BEVOR der aktualisierte Wert der Barriere an die CPU übermittelt wurde

Die Synchronisation von Schlangen mit Barrieren ist wieder ein anderes Bier, was ja auch bei NVIDIA gemäß meines Benchmarks sehr wahrscheinlich ohne auf die CPU zu warten geht. Ich dachte du meintest hier die Synchronsiation von Threads mit denjenigen Threads, die noch nicht In-Flight sind. Das wird eben nicht unterstützt.

Und das dann wiederum dann noch mal unabhängig von den 64 Slots pro ACE die für bereits laufende Commands reserviert sind. Und das dann noch mal x8 weil 8 ACEs. Das ist einfach pervers was die da für einen Durchsatz vorgesehen haben.

Gibt es für die genaue Funktionsweise der ACEs irgendwelche Quellen von AMD? Vieles was ich momentan dazu lese kommt mir immer ein bisschen wie "Spekulation" vor.

10 µs sind dafür aber eigentlich trotzdem noch zu lang. Es sei denn...

Das wird dann wahrscheinlich doch von der Grid Managment Unit unterstützt. Die damit mit ziemlicher Sicherheit ein µC ist, so langsam wie die arbeitet. Und an den GPU-Takt gekoppelt is

Für "globale" Synchronisationen oder "globale" Befehle sind solche Zeiten bei GPUs allgemein nicht unüblich. Zum Beispiel gibt es bereits eine Latenz von 2 µs zwischen zwei Befehle in einer einzigen Schlange. Alternativ kosten Befehle mit Pipelineänderungen der GPU auch immer ein paar 10 µs.

Btw: Aus Transparenzgründen noch der Benchmarkcode:

Spoiler

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "helper_math.h"
#include <stdio.h>
#include <iostream>

__global__ void TestKernel(float* In)
{
	int Index = threadIdx.x;

	float Temp = threadIdx.x;
	
	for (int i = 0; i < 100; i++)
		Temp += Temp;

	if (Temp == -1)
		In[0] = -1;
}



int main()
{
	float* Test;
	cudaEvent_t Start;
	cudaEvent_t End;
	cudaEventCreate(&Start);
	cudaEventCreate(&End);
	cudaStream_t StreamA;
	cudaStream_t StreamB;

	cudaStreamCreate(&StreamA);
	cudaStreamCreate(&StreamB);

	cudaMalloc(&Test, sizeof(float));
	static const int IterCount = 1024;

	//Sequentiell
	cudaEventRecord(Start, StreamA);
	for (int i = 0; i < 2*IterCount; i++)
	{
		TestKernel << <1, 32, 0, StreamA >> >(Test);
	}
	cudaEventRecord(End, StreamA);
	cudaStreamSynchronize(StreamA);
	float TimeOneStream;
	cudaEventElapsedTime(&TimeOneStream, Start, End);


	cudaEvent_t SyncAtoB[IterCount];
	cudaEvent_t SyncBtoA[IterCount];


	for (int i = 0; i < IterCount; i++)
	{
		cudaEventCreate(&SyncAtoB[i]);
		cudaEventCreate(&SyncBtoA[i]);

	}
	//Roundtrip
	/*
	cudaEventRecord(Start, StreamA);
	for (int i = 0; i < 2 * IterCount; i++)
	{
		TestKernel << <1, 32, 0, StreamA >> >(Test);
		cudaStreamSynchronize(StreamA);
	}
	cudaEventRecord(End, StreamA);
	cudaStreamSynchronize(StreamA);*/

	//Events
	cudaEventRecord(Start,StreamA);
	for (int i = 0; i < IterCount; i++)
	{
		if (i != 0)
			cudaStreamWaitEvent(StreamA, SyncAtoB[i - 1],0);

		TestKernel << <1, 32, 0, StreamA >>>(Test);
		cudaEventRecord(SyncBtoA[i],StreamA);
		cudaStreamWaitEvent(StreamB, SyncBtoA[i], 0);
		TestKernel << <1, 32, 0, StreamB >>>(Test);
		cudaEventRecord(SyncAtoB[i], StreamB);
	}

	cudaEventRecord(End, StreamB);
	cudaStreamSynchronize(StreamB);
	float TimeTwoStream;
	cudaEventElapsedTime(&TimeTwoStream, Start, End);
	printf("SyncCosts: %f ms \n", (TimeTwoStream - TimeOneStream) / (2.f*(float)IterCount));// 
	system("pause");
}

Edit: Eben in der NVIDIA-Doku gesehen, dass sie es auch schreiben, dass die Synchronisation auf dem Device statt findet: http://developer.download.nvidia.co...STREAM_gfe68d207dc965685d92d3f03d77b0876.html

Makes all future work submitted to stream wait until event reports completion before beginning execution. This synchronization will be performed efficiently on the device. The event event may be from a different context than stream, in which case this function will perform cross-device synchronization.

 

[Ext3h]

Gibt es für die genaue Funktionsweise der ACEs irgendwelche Quellen von AMD? Vieles was ich momentan dazu lese kommt mir immer ein bisschen wie "Spekulation" vor.

Leider nein. Nichts außer den Produktvorstellungen die man öffentlich findet, darüber hinaus nur synthetische Benchmarks mit zum Teil wirklich seltsamen Ergebnissen.

Ein paar unzusammehängende Rahmendaten von AMD:

• Dispatchrate von 1 Wavefront pro Taktzyklus
• 8 Queues pro ACE, nur Compute
• Zugriff auf Global Data Share, darüber Barrieren-Synchronisation mit Graphics Command Processor und anderen ACE
• Zugriff auf Hauptspeicher über L2 Cache
• Integrierter Work Distributor, unabhängig vom Work Distributer des Graphics Command Processor

Ansonsten experimentell bestätigt:

• Integrierter Work Distributor schafft 64 aktive Grids pro ACE, 128 Grids bei den größeren HWS in Fiji

Was die Teile jetzt tatsächlich sind, ob ASIC oder Microcontroller, das weiß niemand. Und was die großen HWS auf Fiji jetzt noch an Zusatzfunktionalität dazu bekommen haben, dazu gibt es noch weniger Infos.

Eben in der NVIDIA-Doku gesehen, dass sie es auch schreiben, dass die Synchronisation auf dem Device statt findet: http://developer.download.nvidia.co...STREAM_gfe68d207dc965685d92d3f03d77b0876.html

Dann ergibt das überhaupt keinen Sinn mehr, dass die Compute Queues in DX12 von Nvidia nicht auf die Hyper-Qs gemappt wurden. Es wäre überhaupt kein Problem gewesen, mal eben in Software schnell die als Counter implementierten Barrieren auf reine Events zu mappen, dann käme man sogar ohne ein Firmwareupdate aus.

 

[Nai]

Ansonsten experimentell bestätigt:

Integrierter Work Distributor schafft 64 aktive Grids pro ACE, 128 Grids bei den größeren HWS in Fiji

Der Punkt kommt mir massiv komisch vor, da Schlangen ja zunächst sequentiell sind und die GPU für die Korrektheit des Programmes die Schlangen auch sequentiall abarbeiten muss. Deshalb könnte die GPU maximal pro Schlange nur ein aktives Grid berechnen. Diese Sequentialität liesse sich zwar nocheinmal durch Dependency-Bars in der Schlange vermeiden; alles zwischen zwei Dependency-Bars könnte parallel berechnet werden. Da moderne APIs Bindless sind, kann weder die GPU noch der Treiber diese Abhängigkeiten feststellen und automatisch Dependency-Bars einfügen. Deshalb wäre es die Aufgabe der Anwendung für das Einfügen der Dependency-Bars zu sorgen; aber weder in Mantle noch in DX12 gibt es afaik solche Dependency-Bars.