Warum hat Nvidia Pascal sowohl FP32- als auch FP64-Kerne? Warum kann ich sie nicht gleichzeitig verwenden?

3135
AstrOne

Ich versuche, die GPU-Architektur von Nvidia zu verstehen, stecke aber etwas fest, das recht simpel erscheint. Jeder Streaming-Multiprozessor im Pascal besteht aus 64xFP32- und 32xFP64-Kernen. Und hier sind meine zwei Fragen:

  • Warum hat Nvidia sowohl FP32- als auch FP64-Einheiten in den Chip eingebaut? Warum nicht einfach FP64-Einheiten einsetzen, die 2xFP32-Operationen pro Befehl ausführen können (wie die SIMD-Befehlssätze in CPUs).
  • Warum kann ich nicht alle FP32- und FP64-Einheiten gleichzeitig verwenden?

Ich denke, beides sind Hardwareauslegungsentscheidungen, aber ich würde gerne mehr Details zu diesem Thema erfahren. Alle Informationen dazu sind mehr als willkommen!

EDIT1:

  • Wenn es möglich ist, FP32 und FP64 gleichzeitig auszuführen, bedeutet dies, dass eine GPU mit 8TFLOPS SP und 4TFLOPS DP (theoretisch) 12 TFLOPS gemischte TFLOPS bieten kann?
    • Wie wird dies bei CUDA erreicht? Benutze ich einfach Doppel- und Floats gleichzeitig in meinem Kernel? Oder muss ich eine Flagge an NVCC übergeben?
1
FP64-Einheiten verbrauchen wahrscheinlich erheblich mehr Immobilien und (als Ergebnis) eine höhere Leistungsaufnahme als ein FP32-Core. Wenn Sie nur FP64 verwenden, haben Sie einen viel geringeren Energieverbrauch und weniger Kerne. Dies bedeutet weniger Leistung für einfache FP32-Aufgaben in einer bestimmten Chipgröße. In Heimanwender-Situationen wäre dies eine inakzeptable Verschwendung von Leistung und Leistung, da die meisten Spiele kein FP64 benötigen. Wenn Sie mehr Kerne in einen Chip stecken, können Sie dies kompensieren. Größere Chips bedeuten jedoch weniger Chips pro Fertigungswafer und somit teurer pro Chip. Mokubai vor 7 Jahren 1
Ein FP64, das 2xFP32-Anweisungen zusammenführen, berechnen und dann erneut splitten kann, erfordert wahrscheinlich viel Steuerungslogik, entweder in Hardware (mehr Platzverschwendung) oder in Software, die an Leistung einbüßt. Mokubai vor 7 Jahren 1
Wenn überhaupt, verwenden nur wenige Consumer-Anwendungen die FP64-Funktionalität. Ein vollständiges FP64 auf Verbraucherkarten würde die Herstellungskosten und den Stromverbrauch erhöhen, da ein erheblicher Platzbedarf für die Funktionalität für die meisten Verbraucher unbrauchbar ist. Dies würde auch eine geringere Leistung für Spiele bedeuten. Heutige High-End-GPUs sind häufig auf die Strom- und Wärmeabgabe beschränkt. Durch die Verschwendung von Energie bei ungenutzter Funktionalität wird der für die nützliche Arbeit verfügbare Strom reduziert. bwDraco vor 7 Jahren 0
Es ist auch teilweise aus Gründen der Marktsegmentierung. Da es sich in der Regel um einen spezialisierten Markt handelt, für den GP64 mit GPU beschleunigt werden muss, kann die Beschränkung der vollen FP64-Leistung auf Spezialkarten, die für diese Märkte konzipiert wurden, einen wesentlich höheren Preis für diese Karten verlangen. Der höhere Preis deckt die zusätzlichen (* sehr * teuren) Validierungs- und Anbieterzertifizierungen ab, die für kritische Geschäftsanwendungen erforderlich sind, und erhöht die Gewinnmargen. Siehe auch: [Warum kosten Workstation-Grafikkarten weit mehr als vergleichbare Grafikkarten für Endverbraucher?] (Http://superuser.com/q/690388) bwDraco vor 7 Jahren 0

1 Antwort auf die Frage

2
huseyin tugrul buyukisik

Warum hat Nvidia sowohl FP32- als auch FP64-Einheiten in den Chip eingebaut?

Ich denke, es geht um Marktdurchdringung, so viele wie möglich zu verkaufen. Ohne FP64 können Wissenschaftler der wissenschaftlichen Forschung nicht einmal eine Demo mit wissenschaftlich wichtiger gpgpu-Software versuchen, die FP64 verwendet (und sogar Spiele könnten gelegentlich doppelte Präzision verwenden). Ohne FP32 wären Spielphysik und Simulationen sehr langsam oder GPU würde einen Kernreaktor benötigen. Ohne FP16 kein schnelles neuronales Netzwerk. Wenn es nur FP32 gäbe, würde eine neuronale Netzwerksimulation mit halber Geschwindigkeit arbeiten, oder eine FP64-Summierung würde nicht funktionieren.

Wer weiß, vielleicht gibt es in Zukunft FP-Raytrace-dedizierte Kerne, die ultraschnell Raytracing durchführen, so dass DX12 DX11 DX9 keine schmerzhaften Aktualisierungen und bessere Grafiken enthält.

Letztendlich würde ich nicht nein für eine FPGA-basierte GPU sagen, die einige Kerne von FP64 in FP32 oder einige spezielle Funktionskerne für eine Anwendung konvertieren kann, dann alle FP64 für eine andere Anwendung konvertiert und sogar alles in einen einzigen Fettkern konvertiert führt sequentielle Arbeit aus (z. B. das Erstellen von Shadern). Dies würde für Leute von Vorteil sein, die auf einem Computer viele verschiedene Dinge tun. Ich brauche zum Beispiel möglicherweise mehr Multiplikationen als Ergänzungen, und FPGA könnte hier helfen. Aber jetzt spricht Geld und es heißt "Feste Funktion für jetzt" und bestes Einkommen wird mit einer Mischung aus FP64 und FP32 (und FP16 in letzter Zeit) erzielt.

Warum nicht einfach FP64-Einheiten einsetzen, die 2xFP32-Operationen pro Befehl ausführen können (wie die SIMD-Befehlssätze in CPUs).

SIMD erwartet immer dieselbe Operation für mehrere Daten und weniger Spaß für skalare GPGPU-Kernel. Um aus einem FP64 2xFP32 zu machen, würde man mehr Transistoren benötigen als reines FP64, mehr Wärme, vielleicht mehr Latenz.

Mehr Transistoren = höhere Produktionsausfallwahrscheinlichkeit, daher könnte eine 1024 FP32-GPU mit größerer Wahrscheinlichkeit als eine 512-FP64-flexible GPU hergestellt werden.

Warum kann ich nicht alle FP32- und FP64-Einheiten gleichzeitig verwenden?

Mixed-Precision-Computing kann in cuda und opencl ausgeführt werden, sodass Sie mit allen Kernen noch schneller werden können, jedoch nur für Situationen geeignet, in denen es keinen Speicherplatz gibt, was selten und schwer zu codieren ist.

Beantworten, um 1 zu bearbeiten:

Hier ist eine ausführliche Quelle: http://www.nvidia.com/content/PDF/sc_2010/CUDA_Tutorial/SC10_Accelerating_GPU_Computation_Through_Mixed-Precision_Methods.pdf

lange Geschichte kurz, fügen sie nicht hinzu, es gibt "geringere Renditen", die die Skalierung von% 100 wegen der erforderlichen "zusätzlichen Zyklen" zwischen verschiedenen Präzisionsberechnungen irgendwie nicht zulassen. Wenn sie nicht gemischt werden, benötigen sie "zusätzliche Iterationen" zwischen Blöcken, die ebenfalls keine Skalierung von% 100 zulassen. Es scheint sinnvoller zu sein, "FP64" statt "FP32" nach unten zu beschleunigen (aber da viele FP64-Kerne vorteilhaft sein sollten (für FP32), könnten Sie sie mit so etwas wie einem Body-Kernel testen (der nicht über einen Engpass im Speicher verfügt). ). FP64 ist sehr speicherintensiv (und Cachezeilen (und lokaler Speicher)). Aus diesem Grund habe ich vorgeschlagen, einige Daten für N-Zeiten (> 64k) wiederzuverwenden. Meine GPU verfügt über 1/24 FP64-Leistung, daher traue ich meinem Computer nicht. Du hast einen titan Du solltest es versuchen,

Diese Quelle: http://www.nvidia.com/content/nvision2008/tech_presentations/NVIDIA_Research_Summit/NVISION08-Mixed_Precision_Methods_on_GPUs.pdf

sagt "hervorragende Leistung und Genauigkeit", aber ich konnte keinen Physik-Löser für Spiele finden, die FP32 + FP32 (abgeschnittener FP64) verwenden. Vielleicht spricht das Geld noch einmal. Wenn jemand dies macht, wäre dies "herausragende Leistung und Einbruch" beim Spielen. (vielleicht schlimmer als Furmark explodierender gpus)

Menschen verwenden hier sogar ganze Zahlen (Integer-Punkt-Produkt) auf Floats: https://devblogs.nvidia.com/parallelforall/mixed-precision-programming-cuda-8/

Wie wird dies bei CUDA erreicht? Benutze ich einfach Doppel- und Floats gleichzeitig in meinem Kernel? Oder muss ich eine Flagge an NVCC übergeben?

ein Beispiel zur iterativen Verfeinerung mit fp64 + fp32 in derselben Funktion:

https://www.sciencesmaths-paris.fr/upload/Contenu/HM2012/07-dongarra_part2.pdf

Seiten 26-28.


Für den Opencl-Teil ist hier ein Evergreen (HD5000-Serie) in der Lage, in jedem Zyklus 1dp fma + 1 sp (oder 1 sf) auszugeben.

http://www.microway.com/download/whitepaper/gpgpu_architecture_and_performance_comparison_2010.pdf

Ich werde auf meinem R7-240 etwa einen Nbody testen, der morgen 1/24 oder 1/26 Power von FP32 als FP64 ist.

Edit: es funktioniert.

__kernel void sumGPU(__global float * a,__global float * b) { int idx = get_global_id(0); float a0=a[idx]; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; //a0=convert_float(convert_double(a0)+2.0); //a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; a0+=2.0f; b[idx] = a0;  }  

Es wechselt zwischen 13,02 ms und 12,85 ms, wenn nur einer der Kommentare deaktiviert ist.


Hinweis: Kerne sind nicht selbst fp32. Es gibt keine Kerne. Es gibt Scheduler, die Hardwareressourcen (fp32, fp64, special_function, register) an Kernel-Anweisungen von Threads binden. Threads sind auch keine echten Threads. Wenn Sie also fp32, dann fp64, dann fp32 und dann fp64_square_root verwenden, werden die erforderlichen Ressourcen reserviert, wenn sie benötigt werden. Wenn nicht, sind sie Optionen für andere Arbeitselemente (aber ein einzelnes Arbeitselement kann nicht mehr als 1-2 fp32-ALUs verwenden, die ich vermute (idk, ich habe mir das ausgedacht)).


Edit (2018/03): Wird FP_raytrace(der zweite Absatz dieser Antwort oben) Wirklichkeit?

(NVIDIA) https://www.geforce.com/whats-new/articles/nvidia-rtx-real-time-game-ray-tracing

(AMD) https://www.gamingonlinux.com/articles/amd- hat- radeon- rays-an- open- source- ray- tracing- sdk- using- vulkan.11461 angekündigt

Oder ist es ein weiteres Marketing-Gimmick? Wenn es eine Hardware-Seite hat, können Raytracer-Leute schneller arbeiten, aber es wäre keine Hilfe für einen Moba-Gamer oder einen Ray-Tracer-freien Physik-Simulator. Warum sollte ich für diese Raytracer mehr bezahlen, wenn ich Videos bearbeiten möchte? Vielleicht können diese auch als andere segmentiert werden, wahrscheinlich aber für mehr Geld.

Hallo, mein Freund. Vielen Dank für Ihre Antwort. Ich werde in meinem Beitrag eine weitere Frage zum Mixed Precision Computing hinzufügen. Mir war nicht bewusst, dass gemischte Präzision möglich ist. Ich werde Ihre Antwort heute bis Ende des Tages annehmen, unabhängig davon, ob Sie die neue Frage beantworten oder nicht. Natürlich wäre ich dankbar, wenn Sie es würden! :) Danke noch einmal. AstrOne vor 7 Jahren 0
fügte eine Quelle hinzu, jetzt nach Opencl-Teil huseyin tugrul buyukisik vor 7 Jahren 0
Beeindruckend! Viele Sachen! Vielen Dank mein Freund! AstrOne vor 7 Jahren 0
Ich habe gerade mein funktionierendes Opencl-Beispiel zu meinem schlechten 1/16 oder 1/24 FP64 hinzugefügt. Die Umwandlung von Float in Double muss durch ein weiteres FP32 ausgeblendet werden. Ein Titan würde dann mit 3-4 Zeilen gemischtem Code arbeiten huseyin tugrul buyukisik vor 7 Jahren 0
auch gpu 1 ° C heißer gemacht huseyin tugrul buyukisik vor 7 Jahren 0