DE102020127704A1

DE102020127704A1 - Techniken zum effizienten transferieren von daten an einem prozessor

Info

Publication number: DE102020127704A1
Application number: DE102020127704.0A
Authority: DE
Inventors: Andrew Kerr; Jack Choquette; Xiaogang Qiu; Omkar PARANJAPE; Poornachandra Rao; Shirish Gadre; Heinrich Steven J; Manan Patel; Olivier Giroux; Alan Kaatz
Original assignee: Nvidia Corp
Current assignee: Nvidia Corp
Priority date: 2019-10-29
Filing date: 2020-10-21
Publication date: 2021-04-29
Also published as: US11907717B2; US20230185570A1

Abstract

Eine Technik für Blockdatentransfer wird offenbart, die Datentransfer- und Speicherzugriff-Overheads verringert und Multiprozessor-Aktivität und Multiprozessor-Energieverbrauch erheblich verringert. Auf einem Multiprozessor ausgeführte Threads, die im globalen Speicher gespeicherte Daten benötigen, können die benötigten Daten anfordern und in einem chipinternen gemeinsam genutzten Speicher speichern, auf die von den Threads mehrere Male zugegriffen werden kann. Die Daten können aus dem globalen Speicher geladen und in dem gemeinsam genutzten Speicher unter Verwendung einer Anweisung gespeichert werden, welche die Daten in den gemeinsam genutzten Speicher ohne Speichern der Daten in Registern und/oder dem Cache-Speicher des Multiprozessors während des Datentransfers leitet.

Description

GEBIET
Diese Technologie betrifft die effiziente Nutzung von Prozessorspeicher und Anweisungssatzarchitekturen (Instruction Set Architectures; ISAs) von Prozessoren, die Anweisungen umfassen, die eine derartige effiziente Nutzung ermöglichen. Genauer gesagt betrifft die Technologie hier die effiziente Speicherung von Daten in einem Speicher, der von mehreren Kerne gemeinsam genutzt wird, und spezialisierte oder zusätzliche Speicherzugriffsanweisungen, die den Registerdruck während Speichertransfer und/oder Datentransferlatenz verringern.
HINTERGRUND
Massiv parallele Hochleistung-Multithread-Mehrkernverarbeitungssysteme - Systeme, die viele parallel arbeitende Verarbeitungskerne enthalten -verarbeiten Daten viel schneller, als es in der Vergangenheit möglich war. Diese Verarbeitungssysteme können komplexe Berechnungen in kleinere Aufgaben zerlegen, die dann von mehreren parallelen Verarbeitungskernen nebenläufig durchgeführt werden können. Diese „Teile-und-Herrsche“-Vorgehensweise macht es möglich, komplexe Berechnungen in einem kleinen Bruchteil der Zeit durchzuführen, die erforderlich sein würde, wenn lediglich ein oder einige wenige Prozessoren an den gleichen Berechnungen der Reihe nach arbeiten.
Da die Anzahl von nebenläufig verarbeiteten Aufgaben zugenommen hat, hat die Datenmenge, die benötigt wird, um diese vielen parallelen Berechnungen zu unterstützen, ebenfalls zugenommen, was einen Datenzugriffengpass in „nahen“ chipinternen Speicher erzeugt.
Moderne Verarbeitungschips, wie beispielsweise GPUs, enthalten typischerweise erhebliche Mengen an Speicher in der Nähe der parallelen Prozessoren - was Speicherzugriffslatenz verringert. Beispielsweise enthalten, ab dieser Einreichung, einige NVIDIA-GPUs in der Größenordnung von 12GB oder mehr eines lokalen chipinternen Speichers hoher Bandbreite (umfassend z.B. 4GB von Cache/gemeinsam genutzten Speicher), um über 5000 parallel arbeitenden Kernen zu dienen.
Um Datenzugriffslatenz weiter zu verringern, organisieren moderne Verarbeitungssysteme typischerweise Speicher in einer Hierarchie (z.B. Ebene 1 (L1) Cache, Ebene 2 (L2) Cache, Ebene 3 (L3) Cache, Hauptspeicher, usw.). Derartige Speicherhierarchien speichern Daten, an denen die Verarbeitungskerne aktuell arbeiten, näher zu diesen Verarbeitungskerne, so dass sie leichter verfügbar sind. Ein Cache-Speicher, der am nächsten zu den Verarbeitungskernen ist, z.B. ein L1-Cache, kann partitioniert, verteilt oder anderweitig organisiert sein, so dass jeder Verarbeitungskern oder Satz von Verarbeitungskernen exklusiven Zugriff auf seinen eigenen Cache aufweist, um Wartezeiten aufgrund von Speicherkonkurrenz mit anderen Kernen zu vermeiden. Ein derartiger Cache-Speicher wird häufig von Hardwareschaltungen unterstützt, die Tags aufrechterhalten und sich um das automatische Schreiben von „schmutzigen“ Cache-Zeilen in den Hauptspeicher kümmern, bevor die Zeilen geräumt werden - wobei dem Softwareprogrammierer die Notwendigkeit erspart wird, den Cache explizit zu verwalten.
Es gibt ebenfalls eine Notwendigkeit für die gemeinsame Nutzung lokaler Speicher zwischen Verarbeitungskernen. Manchmal benötigt ein Verarbeitungskern(e) Zugriff auf einen oder mehreren Werte, die von einem unterschiedlichen Verarbeitungskern(en) berechnet werden, oder verwenden Speicher als einen Kommunikationskanal, um einem anderen Verarbeitungskern zu etwas signalisieren. Während ein Kern Daten in den Hauptspeicher für den Zugriff durch einen anderen Kern zurückschreiben könnte, wird ein Zugriff auf den Hauptspeicher häufig Tausende von Zyklen benötigen. Die hierarchische Cache-Speicherarchitektur kann manchmal verwendet werden, um derartige Daten zwischen Prozessoren gemeinsam zu nutzen, ohne auf einen Hauptspeicherzugriff zu warten. Um die gemeinsame Nutzung von Daten zwischen Multiprozessoren zweckmäßiger zu machen, werden moderne GPUs jedoch häufig mit einem lokalen chipinternen Speicher versehen, der zwischen Verarbeitungskerne gemeinsam genutzt wird. Des Weiteren wird ein derartiger gemeinsam genutzter Speicher häufig durch Kopie (direkten Speicherzugriff oder DMA) Hardware unterstützt/beschleunigt.
Die Verwendung eines gemeinsam genutzten Speichers zusätzlich zu oder anstelle eines Cache-Speichers kann bestimmte Vorteile bereitstellen. Beispielsweise ermöglichen Anwendungen, die einen gemeinsam genutzten Speicher benutzen, koaleszierendere Zugriffe und können höhere Grade von Parallelität auf Speicherebene erreichen. Siehe beispielsweise den folgenden technischen Artikel, der hier durch Bezugnahme aufgenommen ist, als ob er ausdrücklich dargelegt ist: C. Li, Y. Yang, H. Dai, S. Yan, F. Mueller und H. Zhou, „Understanding the tradeoffs between Software-managed vs. Hardware-managed Caches in GPUs,“ 2014 IEEE International Symposium on Poweranalysis Analysis of Systems und Software (ISPASS), Monterey, CA, 2014, S. 231-242.
In einigen Systemen kann ein Block eines physikalischen Speichers zwischen unterschiedlichen lokalen Speicherfunktionen, wie beispielsweise zwischen gemeinsam genutzten Speicher und Cache-Speicher, flexibel zugeteilt werden. Als Beispiel bietet eine Ausführungsform der NVIDIA-Turing/Volta GPU Architektur eine Software-Programmieroption, um (zur Laufzeit) lokale Speicherzuteilung dynamisch abzustimmen. Im Einzelnen kombiniert die NVIDIA-Volta-Architektur die Funktionalität der L1- und Textur-Caches in einem vereinigten L1/Textur-Cache, der als ein koaleszierender Puffer für Speicherzugriffe wirkt, wobei die Daten, durch die Threads eines Warp angefordert wurden, vor der Lieferung dieser Daten an den Warp eingesammelt werden. In der Volta-Architektur werden der L1-Cache, der Textur-Cache und der gemeinsam genutzte Speicher durch einen kombinierten Daten-Cache untermauert. Der Abschnitt des kombinierten Daten-Cache, der dem gemeinsam genutzten Speicher fest zugeordnet ist (bekannt in CUDA als der „Carveout“), kann zur Laufzeit unter Verwendung einer CUDA-Anweisung cudaFuncSetAttribute() mit dem Attribut cudaFuncAttributePreferredSharedMemoryCarveout ausgewählt werden, das wirksam auswählt, wie viel gemeinsam genutzter Speicher jedem Streaming-Multiprozessor (SM) zugeteilt werden sollte.
Ein Nutzen der Vereinigung des L1-Cache mit dem gemeinsam genutzten Speicher ist verbesserte Latenz und Bandbreite. Für viele Anwendungen engt diese Vorgehensweise die Leistungslücke zwischen explizit verwalteten gemeinsam genutzten Speicher und direkten Zugriff auf Vorrichtungsspeicher ein. Ein derartiger gemeinsam genutzter Speicher ermöglicht parallelen Threads, die auf gleichen oder unterschiedlichen Verarbeitungskerne ausgeführt werden, Daten vorübergehend zu speichern und auszutauschen. Der gemeinsam genutzte Speicher stellt somit einen Kommunikationsmechanismus bereit, der den unabhängigen Threads ermöglicht, miteinander zu kommunizieren. Die Kosten von Registerüberläufen sind ebenfalls niedriger und der Ausgleich von Belegung versus Überlaufen kann neu bewertet werden, um die beste Leistung bereitzustellen.
Herkömmlicherweise erfordert das Speichern von Daten aus dem Hauptspeicher in den gemeinsam genutzten Speicher einen mehrstufigen Prozess. Zuerst führt der Prozessor eine Speicherladeanweisung von dem Hauptspeicher aus, die dazu führt, dass adressierte Daten aus dem Hauptspeicher gelesen, in eine Cache-Zeile(n) des Cache-Speichers gespeichert und dann aus dem Cache-Speicher in ein oder mehreren Register des Prozessors geschrieben werden. Derartige Register können innerhalb einer Registerdatei (die ein anderer Block des lokalen Speichers sein kann) zugeteilt werden - wobei unterschiedliche Register innerhalb der Registerdatei unterschiedlichen Prozessoren oder Prozessorkernen zugeteilt werden. Der Prozessor kann dann eine Speicheranweisung ausführen, um die Daten nun innerhalb seines Registers/seiner in dem gemeinsam genutzten Speicher zu speichern.
Eine derartige herkömmliche Vorgehensweise zum Laden von Daten in den gemeinsam genutzten Speicher kann im Fall von großen Datentransfers, die für bestimmte gemeinsame Transaktionen benötigt werden, wie beispielsweise Matrixmultiplikationen, eine große Anzahl von Registern für einen längeren und häufig unbestimmten Zeitraum verbrauchen. Während dieser Zeit (die in einigen Fällen Tausende von Zyklen aufgrund der langen Latenz des Hauptspeichers oder anderen Abhängigkeiten andauern kann), können die Register gebunden und zur Verwendung für irgendeinen anderen Zweck nicht verfügbar sein. Eine derartige Registerbindung kann die zugeordneten Prozessoren daran hindern, nützliche Arbeit auszuführen, bis die Register freigegeben werden.
Demgemäß verbrauchen, obwohl die Nutzung des gemeinsam genutzten Speichers für bestimmte Berechnungen einen Datenverfügbarkeits-Engpass verringern kann, existierende Verfahren zum Bewegen von Daten in und aus dem gemeinsam genutzten Speicher häufig Ressourcen (d.h. Register), die anderweitig zur Datenverarbeitung verwendet werden könnten - welche die Gesamtverarbeitungsrate der Verarbeitungssystems verlangsamt, wenn bestimmte datenintensive Berechnungen durchgeführt werden. Somit gibt es eine Notwendigkeit, Speicherbandbreitenanforderungen effizienter zu verwalten, während immer noch ein erhöhter Mathematikdurchsatz in Bereichen erreicht wird, wie beispielsweise bei künstlicher Intelligenz (Artificial Intelligence; AI) und tiefem Lernen (Deep Learning; DL).
Figurenliste
Die folgende ausführliche Beschreibung von beispielhaften nicht beschränkenden veranschaulichenden Ausführungsformen ist in Verbindung mit den Zeichnungen zu lesen, von denen:

1A-1B eine beispielhafte nicht beschränkende Parallelverarbeitungsarchitektur veranschaulichen, bei der Daten in den gemeinsam genutzten Speicher über einen Cache und Register geleitet werden;
2A-2B eine beispielhafte nicht beschränkenden Parallelverarbeitungsarchitektur veranschaulichen, bei der Daten in den gemeinsam genutzten Speicher unter Umgehung von Registern geschrieben werden;
3A-3B eine beispielhafte nicht beschränkende Parallelverarbeitungsarchitektur veranschaulichen, bei der Daten in den gemeinsam genutzten Speicher unter Umgehung von Registern und einem Cache geschrieben werden;
4 ein beispielhaftes nicht beschränkendes Blockdiagramm eines Speichersteuersystems zeigt;
5 ein Beispiel eines gemeinsam genutzten Speichers zeigt, der mehrere Bänke umfasst, die mit einer Interconnect-Schaltung gekoppelt sind;
6 eine beispielhafte Darstellung zeigt, wie Daten, die aus einem Hauptspeicher oder globalen Speicher abgerufen werden, in dem gemeinsam genutzten Speicher angelegt werden können, um widersprechende Anforderungen zu verringern;
7 beispielhafte Sektormuster zeigt, die mit einem gemeinsam genutzten Speicher, der vier Sektoren umfasst, angewandt werden können;
8 eine beispielhafte Parallelverarbeitungseinheit veranschaulicht;
9A einen beispielhaften allgemeinen Verarbeitungscluster innerhalb der Parallelverarbeitungseinheit von 8 veranschaulicht;
9B eine beispielhafte Speicherpartitionseinheit der Parallelverarbeitungseinheit von 8 veranschaulicht;
10A einen beispielhaften Streaming-Multiprozessor von 9A veranschaulicht;
10B ein beispielhaftes Konzeptdiagramm eines Verarbeitungssystems ist, das unter Verwendung der Parallelverarbeitungseinheit (PPU) von 8 implementiert ist;
10C ein Blockdiagramm eines beispielhaften Systems ist, bei dem die verschiedene Architektur und/oder Funktionalität der verschiedenen vorherigen Ausführungsformen implementiert werden können; und
11 ein Konzeptdiagramm einer beispielhaften Graphikverarbeitungspipeline ist, die durch die PPU von 8 implementiert wird, gemäß einer Ausführungsform.

AUSFÜHRLICHE BESCHREIBUNG VON BEISPIELHAFTEN NICHT
BESCHRÄNKENDEN AUSFÜHRUNGSFORMEN
Die beispielhafte nicht beschränkende Technologie stellt hier Lösungen für eine wirksame Datenzugriff-Bandbreite bereit, um Streaming-Multiprozessor-Kern-Mathematikeinheiten bei ihrer „Lichtgeschwindigkeit“-Rate zu speisen. Innerhalb des Streaming-Multiprozessors (SM) gibt es eine Notwendigkeit, um die Bandbreite einer Schnittstelle einer Speicher-Eingabe-Ausgabe (Memory Input Output; MIO) zu verringern und den Registerdruck zum Bewegen von Datenoperanden zu verringern, um einen 2X erhöhten Mathematikdurchsatz gegenüber vorherigen Architekturen zu ermöglichen.
Beispielhafte nicht beschränkende Ausführungsformen stellen Techniken des Blockdatentransfers als eine Möglichkeit zum Verringern von Datentransfer- und Speicherzugriff-Overheads bereit, um Multiprozessor (z.B. SM-Ebenen) Aktivität und Energieverbrauch erheblich zu verringern. In einem typischen Rechenprogramm laden die Threads kooperative Daten aus dem globalen Speicher, speichern sie dann im gemeinsam genutzten Speicher, so dass auf sie anschließend mehrere Male zugegriffen werden kann und sie verarbeitet werden können. Beispielhafte Ausführungsformen stellen eine neue Anweisung bereit, welche die aus dem globalen Speicher geladenen Daten leitet, um sie direkt in dem gemeinsam genutzten Speicher zu speichern. Dies beseitigt das Bewegen der Daten durch die SM-Register und minimiert die Datenbewegung durch die MIO. Beispielhafte nicht beschränkende Ausführungsformen stellen die ausführende Anwendung bereit, um die angeforderten Daten aus dem globalen Speicher in den gemeinsam genutzten Speicher (1) unter Umgehung der Register und des LI-Cache, (2) über den LI-Cache unter Umgehung der Register oder (3) über den L1-Cache und den Registern selektiv zu leiten. In beispielhaften nicht beschränkenden Ausführungsformen ist der gemeinsam genutzte Speicher Software-verwaltet und der L1-Cache-Speicher Hardware-verwaltet.
Ein Aspekt der beispielhaften nicht beschränkenden Technologie stellt somit hier eine Anweisungssatzarchitektur (ISA) bereit, der (mindestens) drei unterschiedlichen Typen von Anweisungen des Ladens-aus-Speicher und Speichern-in-Speicher oder Anweisungssequenzen bereitstellt:

(a) eine typische herkömmliche Ladeanweisung, bei der Daten aus dem Hauptspeicher oder dem globalen Speicher abgerufen und in eine Cache-Zeile eines chipinternen Cache-Speicher (der durch Hardware verwaltet werden kann) geschrieben werden, dann aus dem Cache-Speicher in ein oder mehrere Register eines Verarbeitungskerns geschrieben werden (sobald derartige Daten in dem einen oder mehreren Registern des Verarbeitungskerns gespeichert sind, dann können sie aus dem Register(s) in chipinterne gemeinsam genutzte Speicher unter Verwendung einer herkömmliche Speichern-in-Speicheranweisung gespeichert werden); oder ein erster Typ oder Spezies einer Umgehungsladeanweisung aus dem Speicher, die das(die) Prozessorkernregister umgeht, d.h. die in die Cache-Zeile geschriebenen Daten werden in den gemeinsam genutzten Speicher kopiert, ohne zuerst in das(die) Verarbeitungskernregister geschriebenen zu werden; oder ein zweiter Typ oder Spezies einer Umgehungsladeanweisung, die sowohl das(die) Prozessorkernregister als auch den Cache-Speicher umgeht, d.h. die aus dem Hauptspeicher oder globalen Speicher abgerufenen Daten werden direkt in den gemeinsam genutzten Speicher geschrieben, ohne zuerst in die Prozessorkernregister oder dem Cache-Speicher gespeichert zu werden.

Der Entwickler kann unterschiedliche dieser Anweisungen auswählen, um eine Lade-ausdem Speicher-Operation abhängig von der Notwendigkeit und den Leistungsanforderungen durchzuführen.
Beispielsweise wird es in vielen Fällen bedeutsam sein, Speicherkohärenz unter Verwendung der Hardwareunterstützung des chipinternen Cache beizubehalten, und der Entwickler wird wollen, das Speichern von einzelnen Werten insbesondere gemeinsam genutzte Speicherorte unter Verwendung von Anwendung-definierten Adressen explizit zu steuern. In diesen Fällen wird der Entwickler veranlassen, dass die herkömmlichen Lade- und Speicheranweisungen von GLOBAL > CACHE > REGISTER > SHARED ausgeführt werden.
Andererseits wird es Fälle geben, in denen der Entwickler die Vorteile eines Hardwareaufrechterhaltenen Cache-Speicher will, wobei die geladenen Daten jedoch nicht von der Anwendung als durch einen spezifischen Thread geladenen Wert benötigt werden. Sie können nur als zusammen geladene und explizit synchronisierte Daten in dem gemeinsam genutzten Speicher benötigt werden. In derartigen Fällen kann der Entwickler veranlassen, dass der erste Typ oder Spezies der Umgehungs-Ladeanweisung ausgeführt wird, um den Speicherladepfad GLOBAL > CACHE > SHARED bereitzustellen.
In noch anderen Instanzen wird es Fälle geben, in denen der Entwickler wünscht, Daten in den gemeinsam genutzten Speicher so effizient wie möglich zu laden, und wünscht, ein Minimum von anderen Ressourcen dabei zu binden, wobei in diesem Fall der Entwickler veranlassen kann, dass der zweite Typ oder Spezies der Umgehungs-Ladeanweisung ausgeführt wird, um den Speicherladepfad GLOBAL > SHARED bereitzustellen.
Die typische herkömmliche Ladeanweisung, bei der Daten aus dem Hauptspeicher oder globalen Speicher abgerufen und in eine Cache-Zeile eines chipinternen Cache-Speicher geschrieben werden, wird durch zwei separate Anweisungen implementiert. Eine Laden-aus-globalen-Speicheranweisung gefolgt von ein separate Speichern-in-gemeinsam genutzte-Speicheranweisung. Die Laden-aus-globalen Speicheranweisungen laden Daten aus dem globalen Speicher in das Datenregister des anfordernden Thread. Der Speichern-in gemeinsam genutzten Speicher speichert Daten in den Registern in dem gemeinsam genutzten Speicher.
Beispielhafte nicht beschränkende Ausführungsformen stellen eine fusionierte Lade- und Speicheranweisung (LDGSTS) bereit, die Daten aus dem globalen Speicher (LDG) laden und die Daten in den gemeinsam genutzten Speicher (STS) unter Umgehung des Prozessorkernregisters speichern. Die fusionierte Lade- und Speicheranweisung (LDGSTS) kann Daten aus dem globalen Speicher laden und die Daten in den gemeinsam genutzten Speicher mit Optionen laden, um das Prozessorkernregister(s) zu umgehen oder sowohl das Prozessorkernregister(s) als auch den Cache-Speicher zu umgehen. Der erste Typ oder Spezies der Lade/Speicher-Anweisung ist LDGSTS.ZUGRIFF, welche das Prozessorkernregister(s) umgeht. Der zweite Typ oder Spezies der Lade/SpeicherAnweisung ist LDGSTS.BYPASS, welche sowohl das Prozessorkernregister(s) als auch den Cache-Speicher umgeht. Die LDGSTS.ZUGRIFF und LDGSTS.BYPASS Anweisungen können zwei Adressenoperanden, ein Zieladresse des gemeinsam genutzten Speichers und ein globale Quelladresse umfassen. Der Anweisungscode für den herkömmlichen Lade- und Speichervorgang, LDGSTS.ZUGRIFF und/oder LDGSTS.BYPASS, kann in ein Anweisungsregister geladen werden und ein Anweisungsdecodierer kann die Anweisungen in Einstellungen für die Hardware konvertieren. Basierend auf dem Typ oder Spezies der ausgegebenen Anweisung, kann das Laden von Daten aus dem globalen Speicher und das Speichern in der gemeinsam genutzten Speicheranweisung die Hardware steuern, um das Prozessorkernregister(s) und/oder den Cache-Speicher zu umgehen.
In einigen beispielhaften nicht beschränkenden Ausführungsformen kann die Entscheidung, welche der unterschiedlichen Typen oder Spezien des Ladens aus Speicheranweisungen zu verwenden ist, durch einen optimierenden Compiler basierend z.B. auf Schaltern getroffen werden, die der Entwickler in dem Quellcode miteinbezieht, ein Spur von dem, was mit den Daten passiert (z.B. werden sie immer geändert und in den Hauptspeicher oder globalen Speicher zurückgeschrieben) und/oder ob Kerne/Threads/Warps anders als der Kern/Thread/Warp, der die Ladeanweisung aus dem Speicher ausführt, spätere Zugriffe der Daten und/oder anderen Optimierungsfaktoren.
In einigen Beispielen kann die Entscheidung, welche der unterschiedlichen Typen oder Spezien der Laden-aus-Speicheranweisungen zu verwenden ist, dynamische bei Laufzeit durch Bewerten von Bedingungen in dem Quellcode durch den Entwickler durchgeführt werden. Beispielsweise kann die Entscheidung durch die Verwendung eines Speicherdeskriptors getroffen werden, der als ein Anweisungsparameter enthalten ist. Dieser kann ein Hinweis sein, um LDGSTS.ZUGRIFF in LDGSTS.BYPASS zu transformieren, wenn möglich. Beispielsweise kann die LDGSTS.ZUGRIFF in Varianten mit einem Speicherdeskriptor außer Kraft gesetzt werden. Wenn der Speicherdeskriptor einen 11InvDontAllocate Satz aufweist und .sz als 128 spezifiziert ist, dann wird der globale Speicherzugriff behandelt, als ob .BYPASS spezifiziert wäre. Es kann Anweisungseigenschaften, wie beispielsweise Operandengröße für LDGSTS.ZUGRIFF, jedoch nicht für LDGSTS.BYPASS unterstützt, an einigen Ausführungsformen geben, die diese Transformation verhindern.
Es versteht sich für einen Fachmann, dass der gleiche Ausführungs-Thread beliebige oder alle dieser unterschiedlichen Speicherladeanweisungen abhängig von dem Kontext verwenden kann. Somit kann der gleiche Ausführungs-Thread einige Werte aus dem globalen Speicher in den gemeinsam genutzten Speicher unter Verwendung der herkömmliche Ladeanweisung laden, er kann andere Werte aus dem globalen Speicher in den gemeinsam genutzten Speicher unter Verwendung des ersten Typs der Umgehungs-Ladeanweisung laden und er kann noch andere Werte aus dem globalen Speicher in den gemeinsam genutzten Speicher unter Verwendung des zweiten Typs der Umgehungs-Ladeanweisung laden.
Beispielhafte nicht beschränkende Ausführungsformen stellen eines oder mehreres des Folgenden bereit:

Multithread-Kopie mit asynchronem DMA-ähnlichen Verhalten. Die Kopie kann geleitet werden, um durch den LI-Cache zuzugreifen oder ihn zu umgehen.

Erfassung und Codieren von Füllmustern. Die Daten werden umgeordnet („swizzled“), wenn sie in dem gemeinsam genutzten Speicher gespeichert werden, um einen anschließenden Konflikt-freien Zugriff zu ermöglichen. Rather als die Daten durch SM-Register bereitzustellen und sie umzuordnen, wenn sie in dem gemeinsam genutzten Speicher gespeichert werden, erfassen beispielhafte Ausführungsformen das unter kooperativen Threads erzeugte Umordnungsmuster (Swizzle-Muster) und codieren es zur Verwendung, wenn die Daten gefüllt sind.
Nullsetzen auf Byte-Ebene von Füllergebnissen außerhalb der Grenzen.
Programmiermodelle, die unter anderem effizientes Software-Pipelining von Operationen ermöglichen.
Die beispielhafte nicht beschränkende Technologie stellt hier Blockdatentransfertechniken, um Datentransfer- und Speicherzugriff-Overheads zu verringern, um Streaming-Multiprozessor SM-Level-Aktivität und Energieverbrauch erheblich zu verringern. In einem typischen Rechenprogramm laden die Threads kooperativ Daten aus dem globalen Speicher, speichern sie dann in dem gemeinsam genutzten Speicher, so dass auf sie anschließend mehrere Male zugegriffen werden kann und sie verarbeitet werden können.
1A-3B veranschaulichen eine beispielhafte nicht beschränkende Parallelverarbeitungsarchitektur, die einen Software-gesteuerten Transfer von Daten aus dem globalen Speicher 511 in den gemeinsam genutzten Speicher 574 zeigt, auf den mehrere Funktionseinheiten 512₀-512_N zugreifen können. Die Parallelverarbeitungsarchitektur umfasst mehrere SMs 440 (SM0-SMN) mit Zugriff auf den globalen Speicher 511, der extern zu den SMs 440 ist. Der globale Speicher 511 kann einen hierarchischen Cache-Speicher (z.B. Level-2-Cache und/oder Level-3-Cache) und/oder einen DRAM-Speicher umfassen. In einigen Beispielen kann der globale Speicher 511 eine Speicherverwaltungseinheit (Memory Management Unit; MMU) 490, eine X-Bar 370, eine Speicherpartitionseinheit 380 und/oder einen Speicher 304 umfassen, die mit Bezugnahme auf 8-10C beschrieben sind.
Jeder der SMs 440 umfasst mehrere Kerne, wie beispielsweise Funktionseinheiten 512₀-512_N , die konfiguriert sind, um mehrere Threads parallel zu verarbeiten. Ein Thread (z.B. ein Thread der Ausführung) ist eine Instantiierung eines Satzes von Anweisungen, die konfiguriert sind, um durch die Funktionseinheiten 512₀-512_N an einem bestimmten Datensatz ausgeführt zu werden. Threads eines Thread-Blocks können nebenläufig ausgeführt werden und mehrere Thread-Blöcke können nebenläufig ausgeführt werden. In einigen Ausführungsformen gibt eine single-Anweisung, mehrere-daten (Single Instruction, Multiple Data; SIMD) Anweisung Techniken aus, die verwendet werden, um eine parallele Ausführung einer großen Anzahl von Threads ohne Bereitstellen mehrerer unabhängiger Anweisungseinheiten zu unterstützen. In anderen Ausführungsformen werden einzelne-Anweisung, mehrere-Threads (Single Instruction, Multiple Thread; SIMT) Techniken verwendet, um eine parallele Ausführung einer großen Anzahl von im Allgemeinen synchronisierten Threads, unter Verwendung einer gemeinsamen Anweisungseinheit zu unterstützen, die konfiguriert ist, um Anweisungen an einen Satz von Kernen auszugeben. Anders als bei einem SIMD-Ausführungsregime, bei dem alle Kerne typischerweise identische Anweisungen ausführen, ermöglicht die SIMT-Ausführung unterschiedlichen Threads, divergenten Ausführungspfaden durch ein gegebenes Programm leichter zu folgen. Fachleute werden verstehen, dass ein SIMD-Verarbeitungsregime eine funktionale Untermenge eines SIMT-Verarbeitungsregimes darstellt.
Jede der Funktionseinheiten 512₀-512_N ist mit einem Cache-Speicher 572, einem gemeinsam genutzten Speicher 574 und einer Registerdatei 520 über ein Interconnect-Netzwerk, beispielsweise einer hierarchischen Kreuzschiene mit einer oder mehreren Lese- und/oder Schreibkreuzschienen, gekoppelt. Der Cache-Speicher 572, der ein Level-1(L1)-Cache sein kann, und der gemeinsame genutzte Speicher 574 stellen einen chipinternen Speicher mit niedriger Latenz in der Nähe jeder der Funktionseinheiten 512₀-512_N bereit. Die Registerdatei 520 kann Datenregister umfassen, die mittels Software einer unterschiedlichen Funktionseinheit der mehreren Funktionseinheiten 512₀-512_N und/oder unterschiedlichen Warps zugewiesen werden können, die durch den SM 440 ausgeführt werden. Die Registerdatei 520 stellt eine vorübergehende Speicherung für Funktionseinheiten 512₀-512_N bereit.
Die Parallelverarbeitungsarchitektur kann mehrere Adressenräume unterstützen, die lokale, gemeinsam genutzte und globale Adressenräume umfassen, um Datensichtbarkeit für die Threads zu unterstützen. Zusätzliche Nur-Lese-Adressenräume, einschließlich Konstanten und Texturen, können unterstützt werden. Jeder Thread weist seinen eigenen lokalen oder privaten Speicher pro Thread auf, die durch Zuteilung von Registern gesteuert werden können (siehe z.B. US-Patent Nr. 8,555,035 und US-Patent Nr. 7,634,621 , die hiermit durch Bezugnahme aufgenommen sind, als ob sie ausdrücklich dargelegt werden).
Jeder Thread in dem gleichen Thread-Block oder unterschiedlichen Thread-Blöcken kann auf den globalen Speicher 511 unter Verwendung der hierarchischen Cache-Speicher zugreifen. Jeder Thread in dem gleichen Thread-Block kann auf einen zugewiesenen Abschnitt des gemeinsam genutzten Speicher 574 zugreifen, der pro Block gemeinsam genutzten Speichers berücksichtigt werden kann. Jeder ausführende Block von Threads kann einen zugeteilten Abschnitt des gemeinsam genutzten Speichers 574 aufweisen. Der gemeinsam genutzte Speicher 574 ist ein Software-verwalteter Cache, der verwendet wird, um Daten aus dem globalen Speicher zu laden, so dass die Anzahl von chipexternen Speicherzugriffen durch die ausführenden Threads verringert wird. Die Software teilt den gemeinsam genutzten Speicher 574 explizit zu und greift auf diesen zu. Threads in einem Thread-Block werden synchronisiert (z.B. nach kooperativem Laden von Daten aus dem globalen Speicher in den gemeinsam genutzten Speicher), um kritische Ressourcennutzungskonflikte zu vermeiden.
Wenn erwartet wird, dass mehrere Threads in einem Block die gleichen Daten aus dem globalen Speicher 511 verwenden, kann der gemeinsam genutzte Speicher 574 verwendet werden, um diese Daten zu speichern, so dass die Anzahl von Anforderungen an den globalen Speicher 511 verringert wird. Der gemeinsam genutzte Speicher 14 kann ebenfalls verwendet werden, um nicht koaleszierte Speicherzugriffe durch Laden und Speichern von Daten in einem koaleszierten Muster aus dem globalen Speicher 11 zu vermeiden und sie dann in dem gemeinsam genutzten Speicher 14 neu zu ordnen, um Zugriff auf die Daten durch die Threads zu verbessern.
In einigen Ausführungsformen kann ein vereinigtes Cache-System 570 einen Datenspeicher umfassen, der sowohl als ein gemeinsam genutzter Speicher 574 als auch ein lokaler Cache-Speicher 572 konfigurierbar ist. Das vereinigte Cache-System kann in dem gleichen chipinternen Speicher (z.B. SRAM) bereitgestellt werden, der sowohl für den L1-Cache als auch den gemeinsam genutzten Speicher verwendet wird, und umfasst einen Mechanismus, um zuzuteilen, wie viel des vereinigten Speichers dem LI-Cache versus dem gemeinsam genutzten Speicher für jeden Kernel-Aufruf fest zugeordnet ist. Das Kombinieren des L1-Daten-Cache mit dem gemeinsam genutzten Speicher in einem vereinigten Speicher verringert Latenz und stellt höhere Bandbreite bereit. In einigen Beispielen kann das vereinigte Cache-System ebenfalls eine dynamisches konfigurierbare Registerdatei (z.B. Registerdatei 520) umfassen. Für mehr Informationen über ein vereinigtes Cache-System und wie es konfiguriert sein kann, siehe beispielsweise die folgenden Referenzen, die hier durch Bezugnahme aufgenommen sind, als ob sie ausdrücklich dargelegt sind: U.S. Patentanmeldung Veröffentlichung Nr. 2018/0322078 ; und CUDA C Programming Guide, PG-02829-00 I_v 10.1 | Mai 2019 https://docs.nvidia.com/cuda/cuda-c-programmingguide/index.html#shared-memory.
Wenn Software, die auf einer oder mehreren der Funktionseinheiten 512₀-512_N läuft, Daten benötigt, die in dem globalen Speicher 511 gespeichert sind, initiiert die Software einen Thread mit einem „Laden“-aus-Speicher-Befehl. Der Laden-aus-Speicher-Befehl kann Daten aus dem globalen Speicher 511 laden und die Daten in dem gemeinsam genutzten Speicher 574 speichern, was es allen Threads (z.B. allen Threads in einem Block) sichtbar macht. Sobald die Daten in dem gemeinsam genutzten Speicher gespeichert sind, können die Threads auf die Daten mehrere Male zugreifen.
1A und 1B veranschaulichen herkömmliches Bereitstellen von Daten, die aus dem globalen Speicher (z.B. L2) an den gemeinsam genutzten Speicher 574 des SM durch eine hierarchische Kreuzschiene geliefert wurden. Jeder SM 440 kann einen oder mehrere derartige Lese-Ports jeweils mit einer damit verbundenen 32-Byte-Füllbandbreite Die Daten werden unter Verwendung einer Globalen Speicher-Lade(LDG)-Operation gefolgt von einer Speichern-in-Gemeinsam-Genutzten-Speicher(STS)-Operation geladenen. In einigen Beispielen weist der LI-Cache des SM einen vereinigten physikalischen Daten-RAM sowohl für gekennzeichnete Daten als auch den gemeinsam genutzten Speicher auf. Die Organisierung dieses Daten-RAM kann beispielsweise 32 Bänke von jeweils 32 Bit aufweisen. Sogar mit einem vereinigten Speicher bestand die herkömmliche Vorgehensweise jedoch darin, die Daten durch mehrere Speicherorte (d.h. L1-Cache und Register) zu stufen, bevor sie in dem gemeinsam genutzten Speicher gespeichert werden.
In früheren GPUs werden die eingehenden Lese-Port-Daten (32-Byte-Sektorfüllungen) zuerst in den Cache-Speicher 572 (z.B. L1-Cache) zurückgeschrieben. Die LI-Cache-Zeile kann 128 Byte mit 4 32-Byte-Sektoren sein. Ein eingehender Füllsektor wird zu einem der 4 Cache-Zeilen-Sektoren gelenkt. Wenn alle die einem Ladevorgang zugeordneten Sektoren empfangen sind, ist der Ladevorgang abgeschlossen und die Daten werden aus den gefüllten Sektoren in L1 zurückgelesen, mit dem anfordernden Thread durch die Lesekreuzschiene ausgerichtet und dann in die Zielregister zurückgeschrieben. Anschließend wird eine STS-Operation ausgegeben, wobei die Daten aus der SM-Registerdatei gelesen und dann in gemeinsam genutzte Speicher Datenbänke nach Durchlaufen durch die Schreibkreuzschiene zurückgeschrieben werden, um sie mit der gemeinsam genutzten Speicheradresse auszurichten. Diese gesamte Sequenz weist somit 2 breite RAM Lese- und Schreibvorgänge und 2 Durchgänge durch eine lokale Kreuzschiene auf.
Wie in 1A und 1B gezeigt, werden sowohl der Cache-Speicher 572 als auch die Registerdatei 520 während des Datentransfers benutzt. Diese Vorgehensweise kann aus einer Anzahl von Gründen für die Leistung nachteilig sein. Wenn ein Ladevorgang aus dem globalen Speicher 511 ausgegeben wird, werden dem Thread Register zum Darstellen der Ergebnisse zugewiesen. Weil der Ladevorgang Daten aus einhundert unterschiedlichen Orten anfordern kann und die Latenz soweit wie möglich überlappt werden sollte, werden einhundert unterschiedliche Register gebunden. Außerdem werden diese Register solange gebunden, wie der Ladevorgang unerledigt ist, was Tausende von Zyklen sein könnte. Obwohl weniger Register zugeteilt werden können, wird das Zuteilen von weniger Registern für derartige Ladevorgänge die Speicher-Lesevorgänge und -Schreibvorgänge serialisieren, was die Transferzeit erhöht.
Ein beispielhafter Code, der Daten in den gemeinsam genutzten Speicher über den Globalen-Speicher-Ladevorgang(LDG) gefolgt durch die Speichern-in-Gemeinsam Genutzten-Speicher(STS)-Anweisung kann umfassen:

// kooperatives Laden
Rx <- LDG [];
Rz <- LDG [];
//Scoreboard-Abhängigkeit
// gefolgt von gemeinsam genutzten Speichern
STS [], Rx;
STS [], Rz;
<barrier>
// an der Barriere warten
LDS [];
FFMA;

Wie in dem obigen Code gezeigt, lädt die LDG-Anweisung Daten aus dem globalen Speicher in Register und die STS-Anweisungen speichern die Daten aus den Registern in dem gemeinsam genutzten Speicher. Nachfolgend den LDG- und STS-Anweisungen wird eine Laden-innerhalb-des-Gemeinsam-Genutzten-Speicherfenster(LDS)-Anweisung ausgeführt, um Daten in Register zum Durchführen einer fusionierten Multiplizier-Addier-Operation FFMA zu laden. Die Matrixmultiplikation ist eine Schlüsselberechnung innerhalb vieler wissenschaftlichen Anwendungen, insbesondere jene bei tiefem Lernen. Für mehr Informationen über effiziente Matrixmultiplikation auf einer GPU siehe beispielsweise CUTLASS: Fast Linear Algebra in CUDA C++, 5. Dezember 2017 (aktualisiert 21. Mai 2018), https://devblogs.nvidia.com/cutlass-linear-algebra-cuda/, die hier durch Bezugnahme aufgenommen ist. Die vorliegende Technologie ist nicht auf die FFMA-Anweisung beschränkt, sondern kann mit anderen Anweisungen implementiert werden, die in den gemeinsam genutzten Speicher geladene Daten verwenden.
Um das Laden von Daten in den gemeinsam genutzten Speicher zu optimieren, eliminieren beispielhafte nicht beschränkenden Ausführungsformen die Notwendigkeit, Daten durch die Registerdatei des SM bereitzustellen. In beispielhaften nicht beschränkenden Ausführungsformen wird dies als eine einzelne LDGSTS.ZUGRIFF Anweisung implementiert, mit zwei Adressenoperanden, eine Zieladresse des gemeinsam genutzten Speichers und eine globale Quelladresse.
2A und 2B veranschaulichen das Laden von Daten in den gemeinsam genutzten Speicher 574 ohne zuerst die Daten in das Register 520 während des Datentransfers zu laden. Die Daten werden unter Verwendung einer fusionierten Ladeanweisung aus dem globalen Speicher 511 und Speichern der Daten in den gemeinsam genutzten Speicher 574 über den Cache-Speicher 572 transferiert. In dieser Ausführungsform wird ein normaler Ladevorgang mit globalen Adressen durchgeführt und eingehende Fülldatensektoren werden in den Daten-Cache zurückgeschrieben. Dann werden Daten aus zwischengespeicherten Datenreihen gezogen, durch die Lesedatenkreuzschiene ausgerichtet und in den Zielort des gemeinsam genutzten Speichers zurückgeschrieben. Die Schreibkreuzschiene, die normalerweise Speicherdaten ausrichtet, kann verwendet werden, um die Schreibadresse auszurichten. Unter Energieeffizienzgesichtspunkten eliminiert diese Anweisung Lese- und Schreibenergie einer Registerdatei im Vergleich mit separaten LDG- und STS-Anweisungen. Aus Leistungsgesichtspunkten verringert diese Anweisung ebenfalls den Registerdruck, indem Register nicht für Laderückführungen langer Latenz reserviert werden müssen. Weil die Register nicht länger benötigt werden, um transiente Lastwerte zu halten, können sie für andere Aufgaben verwendet werden. Die einzigen Register, die für den fusionierten Lade- und Speichervorgang des Speichers benötigt werden, können Adressen für die SM-Arithmetik umfassen.
Wie nachstehend ausführlicher erläutert wird, ermöglichen beispielhafte nicht beschränkenden Ausführungsformen, dass die Beendigung der Anweisung als eine asynchrone Kopie/DMA unabhängig von anderen Speicheroperationen verfolgt wird.
Ein beispielhafter Code, der ein fusionierten LDGSTS über L1-Cache verwendet, während die Register umgangen werden, kann umfassen:

// kooperatives Laden
LDGSTS.ZUGRIFF [] [];
LDGSTS.ZUGRIFF [] [];
//Scoreboard- Abhängigkeit
// gefolgt von Barriere
<barrier>
// warten an der Barriere
LDS [];
FFMA;

Um das Laden von Daten in den gemeinsam genutzten Speicher weiter zu optimieren, eliminieren beispielhafte nicht beschränkenden Ausführungsformen die Notwendigkeit, Daten durch die Registerdatei des SM und den LI-Cache zu stufen. 3A und 3B veranschaulichen das Laden von Daten in den gemeinsam genutzten Speicher 574, ohne die Daten zuerst in die Registerdatei 520 und den Cache-Speicher 572 während des Datentransfers zu laden. Diese Anweisung wird mit einer einzelnen LDGST.BYPASS-Anweisung bereitgestellt.
In dieser beispielhaften nicht beschränkenden Ausführungsform ist eine weitere Energieverringerung möglich, in dem die eingehenden Fülldaten direkt an dem endgültigen gemeinsam genutzten Speicherziel zurückgeschrieben werden. Dies kann zusätzlich einen L1-Daten-RAM-Lesevorgang und eine Schreibenergie über die in 2A und 2B gezeigte Ausführungsform beseitigen. Unter Leistungsgesichtspunkten übersetzt sich diese beispielhafte Ausführungsform ebenfalls in eine verringerte Nachfrage nach Kreuzschienen- und LIDATA-Bandbreite des gemeinsam genutzten Speichers. In dichten Math-Kernels, wie allgemeine Matrix-zu-Matrix-Multiplikation (General Matrix-to-Matrix-Multiplication; GEMM) mit den warpweiten Matrix-Multiplikation-Akkumulation(MMA)-Anweisungen, können L1 DATEN ein Engpass werden. Das Implementieren eines Datenpfads ohne Bereitstellen der Daten in den Registern und L1 kann eine ausgewachsene Datenkreuzschiene von dem Lese-Port benutzen (als ein Artefakt der beispielhaften Ausführungsform kann der Lese-Port nicht blockiert werden). In 2A umfasst die Schnittstelle zwischen dem globalen Speicherzugriff 511 und dem Cache-Speicher 572 den Lese-Port. In 3B umfasst die Schnittstelle zwischen Fill Steer und L1 den Lese-Port.
Beispielhafter Code unter Verwendung einer fusionierten LDGSTS.BYPASS Anweisung durch Umgehen des L1-Cache und Registern kann umfassen:

// kooperatives Laden
LDGSTS.BYPASS [] [];
LDGSTS.BYPASS [] [];
//Scoreboard-Abhängigkeit
// gefolgt von Barriere
<barrier>
// warten an der Barriere
LDS [];
FFMA;

In einigen Beispielen kann ein Verarbeitungssystem, wie beispielsweise eine GPU, Decodierlogik zum Empfangen und Decodieren der Anweisungen umfassen. Die decodierte Anweisungen können verwendet werden, um Hardware-Laden und -Speichern von Daten basierend auf den Anweisungen zu steuern. Die GPU kann das Decodieren eines fusionierten Lade/Speicher-Anweisungsformat umfassen, das spezifiziert, ob die GPU (a) das Prozessorregister umgehen oder (a) sowohl das Prozessorregister als auch einen Cache-Speicher umgehen soll. Basierend auf dem Decodieren, können die angeforderten Daten geladen und von der Hardware gespeichert werden. In einem Beispiel kann als Reaktion auf das Decodieren die GPU ein Datenwort aus einem ersten Speicher (z.B. dem globalen Speicher) abrufen und das abgerufene Datenwort in einem zweiten Speicher (z.B. dem gemeinsam genutzten Speicher) speichern, ohne die abgerufenen Daten zuerst in der Decodierlogik zu speichern oder ohne zuerst die abgerufenen Daten in der Decodierlogik und dem Cache-Speicher zu speichern, wie durch die decodierte fusionierten Lade/Speicher-Anweisungsformat-Spezifikation ausgewählt. Der zweite Speicher kann von mehreren Threads (z.B. mehreren Threads in einem Thread-Block) gemeinsam genutzt werden und/oder in einem gemeinsamen chipinternen physikalischen Speicheradressenraum mit dem Cache-Speicher angeordnet sein.
Das fusionierte Lade/Speicher-Anweisungsformat kann eine Quelladresse (z.B. globale Speicheradresse), eine Zieladresse (z.B. Adresse im zweiten Speicher) und einen Transferpfad der Daten umfassen. In einem Beispiel kann die fusionierten Lade/SpeicherAnweisung umfassen (i) eine Quellwortadresse, (ii) eine Zielwortadresse und (iii) mindestens ein Bit, das spezifiziert, ob die GPU (a) ein Prozessorregister umgehen oder (b) sowohl das Prozessorregister als auch einen Cache-Speicher umgehen soll, um ein Datenwort aus der Quellwortadresse abzurufen und das abgerufene Datenwort in der Zielwortadresse zu speichern, und das Decodieren umfasst Decodieren des mindestens einen Bits.
Gekennzeichnetes und gemeinsam genutztes Speichersteuersystem
4 zeigt ein beispielhaftes Blockdiagramm eines Speichersteuersystems, das durch Anweisungen geleitet werden kann, die in dieser Anmeldung offenbart werden (z.B. Laden, Speichern und/oder fusionierte Lade- und Speicheranweisungen). Beispiele der vorliegenden Technologie sind nicht auf die in 4 gezeigten spezifischen Speichersteuersystem beschränkt, sondern können auf andere Speichersteuersysteme angewandt werden. Das in 4 gezeigte Speichersteuersystem umfasst Adressenlogik 22, eine Tag-Pipeline 24, eine Tag-zu-Daten-Warteschlange 26, eine Datenspeicherspeicherung 30, ein Interconnect-Netzwerk 580 und eine Registerdatei 520. Eine oder mehrere mit Bezugnahme auf 4 beschriebenen Operationen können durch eine oder mehrere Lade/Speicher-Einheiten (LSUs) 554 (in 10A gezeigt) implementiert werden. Die Datenspeicherspeicherung 30 kann die gemeinsam genutzten Datenspeicherspeicherung 32 und die LI-Datenspeicherspeicherung 34 umfassen, die in einer konfigurierbaren vereinigten Cache-Architektur bereitgestellt werden.
Die Adressenlogik 22 ist konfiguriert, um Speichertransaktionen von ausführenden Threads zu empfangen, Adressen zum Versorgen jener Transaktionen zu erzeugen und die Speichertransaktion basierend auf dem Typ von Speichertransaktion weiterzuleiten. In einigen Beispielen kann die Adressenlogik 22 eine Adressenerzeugungseinheit (Address Generation Unit; AGU) umfassen. Die AGU kann konfiguriert sein, um globale Adressen zum Zugreifen auf den globalen Speicher und lokale Adressen zum Zugreifen auf den gemeinsam genutzten Speicher zu berechnen.
Die durch die Adressenlogik 22 empfangenen Speichertransaktionen können Ladetransaktionen, Speichertransaktionen und fusionierte Lade/Speicher-Transaktionen umfassen. Die Ladetransaktionen können Laden aus dem gemeinsam genutzten Speicher und Laden aus dem globalen Raum umfassen. Die als Reaktion auf die Ladetransaktionen zurückgegeben Daten können den Registerdateien bereitgestellt werden, die den anfordernden Threads zugeordnet sind.
Die Speichertransaktionen können Speichern den in gemeinsam genutzten Speicher und Speichern in den globalen Speicher umfassen. Daten in den Registern, die spezifischen Funktionseinheiten zugeordnet sind, können in dem gemeinsam genutzten Speicher und/oder dem globalen Speicher gespeichert werden.
Abhängig von der Speicheranweisung und der Art der Systemarchitektur können die Daten für die Lade- und/oder Speicheranweisung bestimmte Pfade nehmen. Als Beispiel berechnet, wenn die Adressenlogik 22 eine Anweisung empfängt, um Daten aus dem gemeinsam genutzten Speicher in das Register zu laden, die einem spezifischen Thread zugeordnet sind, die Adressenlogik 22 die gemeinsam genutzten Speicheradresse und leitet die Anweisung direkt entlang des Datenpfads 22a zu dem gemeinsam genutzten Datenspeicher 32, wobei eine Speichertransaktion niedriger Latenz bereitgestellt wird. Anders als bei einem externen Speicher, der Hunderte oder Tausende von Zyklen entfernt von dem gemeinsam genutzten Speicher sein kann, kann der gemeinsam genutzten Speicher 32 einen Zyklus entfernt von den anfordernden Threads sein. Die aus dem gemeinsam genutzten Datenspeicher 32 angeforderten Daten werden der Registerdatei 520, die dem anfordernden Thread zugeordnet ist, über das Interconnect-Netzwerk 580 bereitgestellt.
Wie oben erläutert laufen, wenn Daten in den gemeinsam genutzten Speicher unter Verwendung einer herkömmlichen Ladeanweisung aus dem globalen Speicher geladen werden, die Daten aus dem globalen Speicher in den LI-Cache, aus dem L1-Cache in die Register und aus den Registern in den gemeinsam genutzten Speicher. Die in dieser Anmeldung offenbarten Anweisungen stellen eine Option zum Laden in die gemeinsam genutzte Speicheranweisung bereit, um die Register und/oder den L1-Cache zu umgehen, da die angeforderten Daten in den gemeinsam genutzten Speicher aus dem globalen Speicher gespeichert werden.
Wenn die Adressenlogik 22 eine Anweisung empfängt, um Speicher in den gemeinsam genutzten Speicher zu laden, während die Register (z.B. LDGSTS.ZUGRIFF Anweisung) umgangen werden, wird die Anweisung in zwei separate Unteranweisungen durch eine Adressenerzeugungseinheit (Address Generation Unit; AGU) aufgebrochen, die eine globale Quelladresse und eine Zieladresse des gemeinsam genutzten Speichers berechnet. Die globale Quelladresse wird bei der Ladeoperation verwendet und die Zieladresse des gemeinsam genutzten Speichers wird bei der Speicheroperation in den gemeinsam genutzten Speicher verwendet. Die LDG-Unteranweisung kann als normaler Ladevorgang durch eine Tag-Stufe behandelt werden und führt eine Anforderungszusammenführung von 32 Adressen in 4 unabhängige Tags aus. Die STS-Unteranweisung kann als eine Dummy-Anweisung behandelt werden, um bloß durch die Adressen des gemeinsam genutzten Speichers zu der Pipeline zu gehen.
Die Adressenlogik 22 leitet die Ladeanweisung entlang des Datenpfads 22b zu der Tag-Pipeline 24. Die Tag-Pipeline 24 bestimmt, basierend auf Tag-Information, ob die angeforderten Daten in dem L1-Datenspeicher 34 verfügbar sind. Wenn ein Treffer auftritt, weil die angeforderten Daten in dem L1-Datenspeicher 34 verfügbar sind, leitet die Tag-Pipeline 24 die Speichertransaktion zu dem L1-Datenspeicher 34 über den Datenpfad 24a. Die Daten werden aus den zwischengespeicherten Datenreihen des L1-Datenspeichers 34 gezogen, durch das Interconnect-Netzwerk 580 (z.B. eine Schreibkreuzschiene) ausgerichtet und in die Zieladresse des gemeinsam genutzten Speichers des gemeinsam genutzten Datenspeichers 32 zurückgeschrieben. Dieser Datenpfad unterscheidet sich von einer herkömmlichen Ladeanweisung, bei der angeforderte Daten in dem L1-Datenspeicher 34 zuerst in der Registerdatei 520 über das Interconnect-Netzwerk 580 gespeichert und dann aus der Registerdatei 520 in den gemeinsam genutzten Datenspeicher 32 bewegt werden. Im Fall des Treffers, wird der Speicherpfad (z.B. umfassend einen Seitenkollektor) verwendet, um eine atomähnliche Lese-/Änderung-/Schreib-Operation zu synthetisieren, die aus gekennzeichnete Zeilen lädt und in gemeinsam genutzten Speicherzeilen in Daten-Caches speichert.
Wenn ein Fehltreffer auftritt, weil die angeforderten Daten in dem L1-Datenspeicher 34 nicht verfügbar sind, teilt die Tag-Pipeline 24 einen Tag zu und fordert verpasste Daten aus einer Schnittstelle des globalen Speichers über den Datenpfad 24b an. Die Schnittstelle des globalen Speichers kann Hardware (z.B. direkten Speicherzugriff (Direct Memory Access; DMA) E/A) zum Abrufen von Daten aus anderen Caches und dem Hauptspeicher umfassen. Die Speichertransaktion wird ebenfalls zu der Warteschlange 26 geschoben, die ein Tag-zu-Daten (t2d) First-in First-out (FIFO) umfassen kann. Die Speichertransaktion wird in der Warteschlange 26 blockiert, bis die Fehltrefferdaten zu dem L1-Datenspeicher 34 zurückgegeben werden und die zugeordnete Speichertransaktion den Anfang der Warteschlange 26 erreicht.
Die Tag-Pipeline 24 kann einen Koaleszer 40, einen Tag-Speicher 42 und einen Tag-Prozessor 44 umfassen, um die Speichertransaktionen zu bedienen, die von der Adressenlogik 22 empfangen werden. Der Koaleszer 40 kann einzelne Schreibanforderungen (z.B. Sektoranforderungen) nachbestellen und/oder in einer große Cache-Zeilen-Anforderung zusammenlegen.
Der Tag-Prozessor 44 bestimmt, ob ein Tag für eine gegebene Speichertransaktion in dem Tag-Speicher 42 vorhanden ist. Wenn der Tag-Prozessor 44 bestimmt, dass der Tag vorhanden ist und die dem Tag entsprechenden Daten in dem L1-Datenspeicher 34 vorhanden sind, findet ein Cache-Treffer statt und die Speichertransaktion wird zu dem L1-Datenspeicher 34 geleitet. Wenn der Tag-Prozessor 44 bestimmt, dass der Tag nicht vorhanden ist, teilt der Tag-Prozessor 44 einen Tag in dem Tag-Speicher 42 und zugeordneten Ort innerhalb des L1-Datenspeichers 34 für Daten zu, wenn sie aus dem globalen Speicher zurückgegeben werden.
Wenn die Adressenlogik 22 eine Anweisung (z.B. LDGSTS.BYPASS-Anweisung) empfängt, um Speicher in dem gemeinsam genutzten Speicher zu laden, während das Register und der LI-Cache umgangen werden, wird die Anweisung in zwei separate Unteranweisungen durch die AGU aufgebrochen, die eine globale Quelladresse und eine Zieladresse des gemeinsam genutzten Speichers berechnet. Die globale Quelladresse wird in dem Ladevorgang verwendet und die Zieladresse des gemeinsam genutzten Speichers wird verwendet, um bei der Ladeoperation des gemeinsam genutzten Speichers verwendet zu werden. Darstellungen der globalen Adresse und der gemeinsam genutzten Speicheradresse in der Anweisung können in Registern des anfordernden Threads spezifiziert sein.
Die Adressenlogik 22 leitet die Ladeanweisung entlang des Datenpfads 22b zu der Tag-Pipeline 24. Die Tag-Pipeline 24 bestimmt, basierend auf Tag-Information, ob die angeforderten Daten in dem L1-Datenspeicher 34 gespeichert sind. Wenn ein Treffer auftritt, weil die angeforderten Daten in dem L1-Datenspeicher 34 verfügbar sind, leitet die Tag-Pipeline 24 die Speichertransaktion zu dem L1-Datenspeicher 34 über den Datenpfad 24a. Die Daten werden aus zwischengespeicherten Datenreihen des L1-Datenspeichers 34 gezogen, durch das Interconnect-Netzwerk 580 (z.B. Schreibkreuzschiene) ausgerichtet, in die Zieladresse des gemeinsam genutzten Speichers des gemeinsam genutzten Datenspeichers 32 zurückgeschrieben und der Tag für die geschriebenen Daten wird ungültig gemacht. Der Tag kann ungültig gemacht werden, weil die Daten nun in dem gemeinsam genutzten Speicher gespeichert sind und es nicht erwartet wird, dass die Anwendung die gleichen Daten aus dem globalen Speicher kurzfristig anfordern wird.
Wenn ein Fehltreffer auftritt, teilt die Tag-Pipeline 24 einen Tag zu und fordert verfehlte Daten von der Schnittstelle des globalen Speichers über den Datenpfad 24b an. Die Speichertransaktion wird ebenfalls in die Warteschlange 26 geschoben. Die Speichertransaktion wird blockiert, bis die Fehltrefferdaten zurückgegeben sind. Wenn die Daten über den Datenpfad 511b zurückgegeben sind, werden die Fülldaten direkt in das Endziel des gemeinsam genutzten Speichers in den gemeinsam genutzten Datenspeicher 32 zurückgeschrieben, ohne die Daten in dem L1-Datenspeicher 34 zu speichern, und der Tag für die geschriebenen Daten wird ungültig gemacht. Die Daten werden direkt in den gemeinsam genutzten Datenspeicher 32 über das Interconnect-Netzwerk 580 aus einem Lese-Port geschrieben (siehe das mit Bezugnahme auf 2A und 2B erläuterte Lese-Port).
Der direkte Pfad von dem globalen zu dem gemeinsam genutzten Speicher (ohne durch Cache oder Register zu laufen) kann Kohärenz auf Zeilenebene nicht aufrechterhalten, wobei ein nicht mehr gültiger Wert im Cache gelassen wird, der mit dem Wert inkonsistent sein kann, der in dem gemeinsam genutzten Speicher in einigen Implementierungen sichtbar ist. Um die Inkonsistenz zu vermeiden, sehen oben beschriebene beispielhafte Ausführungsformen eine Implementierung vor, bei der das System snoops und die L1-Zeile ungültig macht.
In einigen Ausführungsformen kann das Interconnect-Netzwerk 580 mehr als eine Kreuzschiene umfassen. Beispielsweise umfasst in einer Ausführungsform das Interconnect-Netzwerk 580 eine Kreuzschiene für allgemeine Lade- und Speichervorgänge und eine separate Kreuzschiene zum Schreiben von Daten, die aus der Speicherschnittstelle empfangen wurden, über den Datenpfad 511b in den gemeinsam genutzten Datenspeicher 32. In dieser Ausführungsform kann die separate Kreuzschiene verwendet werden, da die Kreuzschiene funktionsmäßig benötigt wird, ist dieser Pfad ist einfacher als was für allgemeine Lade- und Speichervorgänge benötigt wird. In diesem Beispiel kann der Datenpfad den Datenpfad 511b von der Speicherschnittstelle durch eine separate Füllkreuzschiene zu dem Inhalt der Datenspeicherspeicherung 30 umfassen. Dies ermöglicht, gleichzeitig aus dem Speicher zu füllen, während aus dem gemeinsam genutzten Datenspeicher 32 oder dem L1-Datenspeicher 34 gelesen wird.
Das Laden von Daten aus dem gemeinsam genutzten Speicher sieht Flexibilität in der Granularität von zurückgegebenen Daten vor. Wenn beispielsweise 32 Wörter aus dem gemeinsam genutzten Speicher gelesen werden, können die 32 Wörter aus unterschiedlichen Zeilen empfangen werden, wenn die angeforderten Wörter über unterschiedliche Bänke des gemeinsam genutzten Speichers gespeichert sind. In L1 würden die 32 Wörter von der gleichen Cache-Zeile kommen. Obwohl die physikalische Struktur des L1 dem gemeinsam genutzten Speicher ähnlich ist, ist es nicht möglich, diese Granularitätsstufe mit L1 zu bekommen, weil L1 darauf beschränkt ist, wie viele Tags (z.B. vier Tags) verwendet werden können. Die L1-Zugriffe können bis zu der Anzahl von unterschiedlichen nichtwidersprechende Zeilen als die Anzahl von Tag-Bänken beschränkt werden. In dem Beispiel von LI, das vier Tag-Bänke verwendet, könnte auf bis zu vier Cache-Zeilen zugegriffen werden.
Um die Bandbreite im gemeinsam genutzten Datenspeicher 32 zu vergrößern, ist der chipinterne Speicher in mehreren Bänken organisiert. 5 zeigt ein Beispiel eines gemeinsam genutzten Speichers, der mehrere mit einem Interconnect-Netzwerk 580 gekoppelte Bänken umfasst. Weil der gemeinsam genutzten Speicher nahe zu den Verarbeitungskernen ist, stellt der gemeinsam genutzte Speicher eine höhere Bandbreite und niedrigere Latenz als ein chipexterner Speicher bereit. Wie die Daten über die mehreren gemeinsam genutzten Speicherbänken verteilt sind, kann jedoch die Bandbreite beeinflussen.
Wie in 5 gezeigt, ist der gemeinsam genutzte Speicher in mehrere Bänken geteilt (z.B. 32 Bänke mit 32 oder 64 Bit pro Taktzyklus-Bandbreite), die jeweils einen Lese-Port und einen Schreib-Port umfassen, aus denen Lese- und Schreibvorgänge gleichzeitig in mehreren Bänken verarbeitet werden können. Speicheranforderungen, die von N Adressen gemacht werden, die in N distinkte Bänke fallen, können gleichzeitig verarbeitet werden. Wenn eine Leseanforderung beispielsweise für Daten bei A0, A1, A3 gemacht wird, dann können die Daten aus Bank 0, Bank 1 und Bank 2 gleichzeitig durch das Interconnect-Netzwerk 580 über die jeweiligen Lese-Ports gelesen werden. Wenn mehrere Adressen einer Speicheranforderung in die gleiche Speicherbank fallen, wird die Anforderung die einzelnen Adressenanforderungen serialisieren. Wenn die Leseanforderung beispielsweise für Daten bei B0, A1 und C0 gemacht wird, dann gibt es einen Bankkonflikt und die Leseanforderung für B0, A1 und C0 wird in drei Lesevorgängen serialisiert.
Daten können in den gemeinsam genutzten Speicherbänken derart gespeichert werden, so dass aufeinanderfolgende Wörter einer zurückgegebenen Cache-Zeile in aufeinanderfolgenden Bänken abgebildet werden. 6 zeigt eine beispielhafte Darstellung, wie Daten aus dem globalen Speicher in dem gemeinsam genutzten Speicher ausgelegt werden können, um widersprechende Anforderungen zu verringern. Das obere Diagramm stellt dar, wie Daten in dem globalen Speicher gespeichert werden können, und das untere Diagramm stellt dar, wie Daten in dem gemeinsam genutzten Speicher nach einer fusionierten LDGSTS-Anweisung gespeichert werden können. Jede Reihe in dem oberen Diagramm entspricht Daten, wie sie in einer Cache-Zeile (LI-Cache oder chipexternen Speicher) ausgelegt werden würden. Jeder Thread würde auf den globalen Speicher in dieser linearen Art und Weise zugreifen. Ein Block von A's in dem globalen Speicher kann 8 zwei-Byte Elemente breit mal 8 Elemente hoch darstellen. Der Block von A's kann einen Block einer Gesamtmatrix darstellen.
Wenn die Threads auf die Cache-Zeile zugreifen, werden die Daten in den gemeinsam genutzten Speicheradressen neu angeordnet, die „swizzled“ (umgeordnet) sind. Beispielsweise wird eine Reihe abgerufen, wie in dem globalen Speicher ausgelegt. Durch die gemeinsam genutzten Adressen, die durch jeden Thread spezifiziert und in einem Token codiert werden, der mit der Anforderung ausgesendet und mit den Daten zurückgegeben wird, wird jedes Element der Zeile zurückgegeben und in dem gemeinsam genutzten Speicher entlang unterschiedlicher Bänke gespeichert.
Wie in 6 gezeigt, wird eine Reihe von A's über die Breite des gemeinsam genutzten Speichers ausgebreitet. Dieser macht es möglich, dass auf jeden der Unterblöcke von einem gleichen Block in einer Konflikt-freien Weise zugegriffen werden kann. Wenn sämtliche der Daten aus Block A in dem gemeinsam genutzten Speicher vertikal gestapelt wären, könnten sie nur seriell gelesen werden. Indem sie ausgebreitet werden, kann ein Zugriff auf sämtliche der A-Blöcke ohne Konflikt durchgeführt werden. 6 zeigt ebenfalls, wie aufeinanderfolgende Elemente von Block C in unterschiedlichen spezifischen Bänken des gemeinsam genutzten Speichers verteilt sein können.
Mit Datenverteilung über die Bänke können, wenn 32 Wörter aus dem gemeinsam genutzten Speicher gelesen werden, der 32 Bänke umfasst, die 32 Wörter aus unterschiedlichen Zeilen gelesen werden. Eine ähnliche Anforderung an den L1 -Cache, würde aus einer gleichen Zeile gelesen werden. Wie oben erläutert, sind in einigen Implementierungen die L1-Zugriffe bis zu der Anzahl von unterschiedlichen, nicht in Konflikt stehenden Zeilen als die Anzahl von Tag-Bänken eingeschränkt.
7 zeigt beispielhafte Sektormuster, die mit dem gemeinsam genutzten Speicher, der vier Sektoren umfasst, angewandt werden können. Ein koaleszierender Puffer kann konfiguriert sein, um Muster zwischen Füllsektoren und Zieladressen des gemeinsam genutzten Speichers zu bestimmen. Die erfassten Muster werden für gemeinsame Muster optimiert, die in existierenden Arbeitslasten beobachtet werden.
Der koaleszierender Puffer kann konfiguriert sein, um ein Permutationsmuster und/oder eine Intra-Sektor-Umordnung zu erfassen und anzuwenden, wenn die Speichertransaktion verarbeitet wird. In einem Beispiel sind die 128B globalen Cache-Zeilen ferner in vier 32B Sektoren angeordnet. Wenn ein globaler Cache-Zeilensektor (z.B. Sektor 0) als Reaktion auf die LDGSTS-Anweisung zurückgegeben wird, kann der Cache-Zeilensektor direkt über dem Interconnect-Netzwerk in Sektor 0 des gemeinsam genutzten Speichers gespeichert werden, oder gedreht und in Sektoren 1, 2, oder 3 des gemeinsam genutzten Speichers gespeichert werden.
Ferner kann die Granularität, wie der zurückgegebene globalen Cache-Zeilensektor in dem gemeinsam genutzten Speicher gespeichert ist, gesteuert werden. Beispielsweise können die Anweisungen mit einer Option versehen werden, um die Hälften des zurückgegebenen globalen Cache-Zeilensektors auszutauschen (z.B. untere 16B in dem hohen Teil des gemeinsam genutzten Speichersektors und die unteren 16B in dem unteren Teil des gemeinsam genutzten Speichersektors speichern). Außerdem kann die Anweisung nur die unteren 16B oder bloß die hohen 16B in den gemeinsam genutzten Speichersektor schreiben.
Wenn mehrere Threads Daten anfordern, schaut der koaleszierende Puffer auf die Adressen der angeforderten Daten und bestimmt, ob die angeforderten Daten zu den gleichen Cache-Zeilen gehören. Wenn die Adressen bestimmt werden, zu der gleichen Cache-Zeile zu gehören, können die Anforderungen für die Daten kombiniert werden und ein Sektormuster wird derart eingestellt, dass, wenn die Daten zurückgegeben werden, die Daten in unterschiedlichen Bänken des gemeinsam genutzten Speichers gespeichert werden. Als Beispiel wird ein erster Sektor in einen oberen Abschnitt der ersten Bank geschrieben und Verschiebung und Drehung wird angewandt, so dass der nächste Sektor in einen unteren Abschnitt der zweiten Bank geschrieben wird (siehe z.B. 6). Wenn das Muster nicht erkannt wird, dann kann die Hardware die Anforderungen für Daten einzeln aussenden.
Das Interconnect-Netzwerk 580 kann eine oder mehrere Kreuzschienen umfassen, die Lesekreuzschiene und/oder Schreibkreuzschienen umfassen. Das Interconnect-Netzwerk 580 unterstützt mehrere parallele Lese- und Schreibvorgänge und unterstützt einen Zugriff auf Daten, die in dem gemeinsam genutzten Datenspeicher 32 gespeichert sind, unter Verwendung direkten Adressierens, und einen Zugriff auf Daten, die in dem L1-Datenspeicher 34 gespeichert sind, unter Verwendung von Tag-Nachschlagen. Das Interconnect-Netzwerk 580 kann eine Anzahl von gleichzeitigen Lese- und Schreibvorgängen unterstützen, die der Anzahl von Bänken in dem gemeinsam genutzten Speicher, dem L1-Cache und/oder den Registern entsprechen. In einem Beispiel kann das Interconnect-Netzwerk 580 eine Anzahl von gleichzeitigen Lese- und Schreibvorgängen unterstützen, die gleich der Anzahl von Speicherbänken in einer Datenspeicherspeicherung 30 (z.B. einem vereinigten Speicher) ist.
Aufspaltung von Speicheranforderungen
Wenn die durch einen Thread angeforderten Daten eine vorbestimmte Größe überschreiten, kann die Speicheranforderung in unabhängige Anforderungen aufgespaltet werden, um eine bestimmte Granularitätsstufe für jede Anforderung bereitzustellen. Wenn ein Thread beispielsweise Daten anfordert, die größer als eine vorbestimmte Größe (z.B. 16 oder 32 Byte) sind, kann der Thread, der die Lade/Speicher-Anforderung ausgibt, in einen oder mehreren Helper-Threads gegabelt werden, die jeweils 16-Byte oder 32-Byte Speicheranforderungen unabhängig ausgeben. Jede Speicheranforderung wird in Cache-Zeilen-Anforderungen zerlegt, die unabhängig ausgegeben werden. Die einzelnen Threads können nach Beendigung der einzelnen Anforderungen verbunden werden.
Setzen von Füllbvtes auf null
LDGSTS kann eine Option (.ZFILL) unterstützen, um Bytes der SMEM-Zieladresse auf null zu setzen, die gefüllt wird. Dies ist nützlich, wenn die von einem Thread aus dem globalen Speicher gelesene Grundfläche eine maximale Randbedingung überspannt. Es ist nicht, dass die globale Adresse ungültig ist oder die Daten selbst ungültig sind, sondern dass die auf null gesetzten Daten nicht auf das aktuelle Arbeitsprodukt eines gegebenen Thread Anwendung findet. Ein Vorteil ist, dass ein Kernel eine konsistente Zugriffbitgröße (z.B. sz: {.32, .64, .128}) verwenden kann, ohne angepassten Code oder Verzweigung unter Verwendung von Lasten fester Größe aufzuweisen.
Mehrere Alternativen können dafür verwendet werden, wie das Nullsetzen spezifiziert werden sollte, z.B. als eine explizite Maske oder Bytezahl in den durch den Thread zugeführten Anweisungsparametern. Stattdessen ist die .ZFILL Option innerhalb der LSBs der globalen Adresse eingebettet, die angibt, wie viele Bytes auf null vom Ende des Bereichs gespeichert werden. Dies hält die gleiche Parameterbewegung und BW mit nicht-.ZFILL-Operationen aufrecht, die eine minimale Störung der Ausgestaltung und die gleiche Ausgabeleistung ermöglichte.
Wenn .ZFILL spezifiziert ist, werden globale Adressenausrichtungsprüfungen suspendiert und die Ausrichtung wird zu der .sz Grenze gezwungen. Dies setzt log₂(.sz Zugriff in Byte) Bits der niederstwertigen Adressierung frei und ermöglicht den LSBs der globalen Adresse als eine Zahl interpretiert zu werden, wie viele Bytes auf null vom Ende des Bereichs gespeichert werden. Wenn .ZFILL in der Anweisung spezifiziert ist, kann der .ZFILL Parameterwert pro Thread spezifiziert werden und kann den pro Thread gespeicherten Bereich eindeutig beeinflussen. Die globale Adresse wird als die Endsumme von (RF + UR + Imm) definiert.
Angesichts der Zahl von Bytes auf null, die in den LSBs der globalen Adresse eingebettet ist, und einer Zugriffgröße kann eine Maske erzeugt werden, die angibt, welche Bytes von den gefüllten Daten (Maskenbit = 1) oder null (Maskenbit = 0) fernzuhalten sind. Wenn .ZFILL nicht verwendet wird, wird die Maske als alles 1's impliziert.
Speichervorgang mit Byte-Schreibfreigabe
Ein komplementäres Merkmal zu .ZFILL kann für globale Speichervorgänge bereitgestellt werden, um zu ermöglichen, dass Bytes übersprungen (.WSKIP) und nicht aus den pro Thread gesendeten Daten gespeichert werden. Die gleiche Technik des Einbettens der Bytezahl in den LSBs der globalen Adresse kann verwendet werden, wobei die Zahl angab, wie viele Byte beim Schreiben von dem Ende des .sz Zugriffs zu überspringen sind, die durch einen gegebenen Thread gespeichert werden. Dies ermöglicht Kerneln ebenfalls, einen einheitlichen .sz Wert (z.B.STG.128.WSKIP) zu verwenden, um maximale Grenzbedingungen zu handhaben und anstatt Speichervorgänge von unterschiedlichen .sz bereitzustellen.
Programmiermodell und Software-Exposition
Einfaches LDG/STS/Barrierenmuster: ein übliches und relativ unkompliziertes Muster ist, wo Threads globale Daten kooperativ laden, sie sofort in einem gemeinsam genutzten Speicher zu speichern und eine Barrierensynchronisation zu verwenden, um diese Daten zu verbrauchen. Obwohl Speicherlatenz hier exponiert ist, stützt sich die Anwendung auf multikooperative Thread-Arrays (CTA) Belegung in SM, um diese exponierte Latenz zu verbergen. Dies ist eine relativ ineffiziente Verwendung von Ressourcen, da sämtliche der Ressourcen, wie beispielsweise Threads, Registerdatei, gemeinsam genutzter Speicher, dupliziert werden. Der Compiler kann jedoch die LDG- und STS-Operationen fusionieren. Der Compiler ist mit einem Mechanismus versehen, um die Barrierensynchronisation abhängig von der STS-Beendigung zu machen. Eine Scoreboard-Freigabefähigkeit wird der LDGSTS-Anweisung bereitgestellt, so dass die Barrierenanweisung abhängig von diesem Scoreboard-Register sein kann.
Stapellademuster: ein fortschrittliche Verwendungsmuster ist, wenn die Threads mehrere LDG-Operationen gefolgt von mehreren STS-Operationen vor dem Ausführen der Barrierensynchronisationen ausgeben. Die fusionierten LDGSTS-Operationen können außerhalb der Reihenfolge in Bezug zueinander beendet werden. Um eine korrekte Abhängigkeitsverfolgung für die Barriere-Operation bereitzustellen, kann ein neue Anweisung, genannt LDGDEPBAR, hinzugefügt werden, welche die Beendigung sämtlicher vorheriger LDGSTS-Operationen einzäunt.
Software-Pipeline-Ladevorgänge: eine effizientere Vorgehensweise, um Speicherlatenz zu verbergen, besteht darin, Daten vorzeitig unter Verwendung von Software-Pipelines und Schleifen-Entrollen (loop unrolling) vorabzurufen. Dies wird durch Zuweisen mehrere Sätze von Registern für mehrere Ladestapel und explizites Verzögern der entsprechenden STS-Operationen erreicht. Für Compiler kann es schwierig sein, diese Technik anzuvisieren, insbesondere wenn sie das Heben von Speicheroperationen durch eine Barriere beinhaltet. Die Kombination von LDGSTS und LDGDEPBAR macht diesen Prozess möglich.
Da es keinen benannten Registerstapel gibt, die zuzuweisen/zu verwalten sind, wird der Prozess des Entrollens leichter. Die LDGDEPBAR-Anweisung ist definiert, In-Order-Semantiken mit früheren LDGDEPBAR-Anweisungen aufzuweisen. Diese borgt eine Technik, die für TEXTURE-Stapelverwaltung mit DEPBAR-Abhängigkeitsverfolgung verwendet wird.
Dieser Prozess ermöglichte ebenfalls einen noch leistungsfähigeren Primitiven des asynchronen DMA-Transfers und Synchronisation von Objekten basierend auf dem gemeinsam benutzten Speicher. In diesem Beispiel:

Das Barrierenzustandsobjekt wird im gemeinsam genutzten Speicher anstelle von dedizierten Zählern implementiert.
Threads können den Zustand mit einer angepassten atomaren Operation ATOMS.ARRIVE.64 aktualisieren.
Ähnlich LDGDEPBAR, wird eine weitere Fence-Operation ARRIVES.LDGSTSBAR.64 definiert, die nach Beendigung des Stapels von LDGSTS-Operationen das Barrierenobjekt des gemeinsam genutzten Speicher aktualisiert.

Diese Konstrukte werden LDGSTS ermöglichen, als eine kooperative asynchrone DMA-Thread-Operation exponiert zu werden, deren Beendigungsstatus durch Abfragen eines Barrierenzustandsobjekts des gemeinsam genutzten Speicher abgefragt werden kann. [Siehe vorläufige US-Patentanmeldung Nr. 62/927,511 , eingereicht am 29. Oktober 2019 und Patentanmeldung Nr. 16/712,236 , eingereicht am 12. Dezember 2019, hier durch Bezugnahme aufgenommen.] Dies ermöglicht dem asynchronen DMA-Transfer, durch Anwendungsprogrammierer ausgenutzt und durch Treiber/Compiler ohne Weiteres unterstützt zu werden.
Wenn eine LDGSTS-Anweisung ausgegeben wird, kann der ausführenden Thread einen logischen Helper-Thread abzweigen, der seinerseits Daten aus dem globalen Speicher in den gemeinsam genutzten Speicher kopiert (ohne SM-Ressourcen, wie Warp-Slots und Registerdatei, zu verwenden). Der Helper-Thread konkurriert jedoch für Speicherbandbreite mit den anderen Threads. Der ausführenden Thread markiert anschließend die Beendigung des Helper-Thread mittels einer Join-Operation, die eine von drei Formen annimmt:

1. ARRIVES.LDGSTSBAR mit einer Arrive-Wait-Barriere verwendet
2. Stapel-zählende Scoreboards unter Verwendung von LDGDEPBAR und DEPBAR.LE implementiert
3. &req auf einem Schreib-Scoreboard (das mittels der LDGSTS-Anweisung aktualisiert wird)

Anweisungsausgabereihenfolge
Alle Anweisungen vor LDGSTS können ausgegeben werden, bevor der Helper-Thread abgezweigt wird. Es kann keine garantierte Ausführungsreihenfolge zwischen a) dem Helper-Thread und b) Anweisungen in dem ausführenden Thread zwischen der LDGSTS-Anweisung und der Join-Anweisung geben. Es wird garantiert, dass der Helper-Thread abschließt, bevor der ausführende Thread Anweisungen ausgibt, welche der Join-Operation in der Programmreihenfolge nachfolgen.
BAR-Anweisungen in dem ausführenden Thread (vor der LDGSTS-Anweisung, zwischen der LDGSTS-Anweisung und der Join-Operation oder nach der Join-Operation) gewährleisten nicht irgendetwas anderes über die Ausführungsreihenfolge des Helper-Thread. Mit anderen Worten nehmen die Helper-Threads, die nicht zu der CTA des ausführenden Thread gehören, nicht in CTA-weiten Synchronisationsbarrieren wie BAR.SYNC und BAR.SYNCALL teil.
Beispielhafte nicht einschränkende Parallelverarbeitung-GPU-Architektur zum Durchführen der oben beschriebenen Operationen und Verarbeitung
Ein Beispiel veranschaulichender Architektur, das auf den Lade- und Speichervorgang gemeinsam genutzter Speicheranweisungen gerichtet werden kann, die in dieser Anmeldung offenbart werden, wird nun beschrieben. Die folgenden Informationen werden für erläuternde Zwecke dargelegt und sollten nicht als in irgendeiner Weise beschränkend ausgelegt werden. Jedes der folgenden Merkmale kann optional mit oder ohne den Ausschluss von anderen beschriebenen Merkmalen aufgenommen werden.
8 veranschaulicht eine Parallelverarbeitungseinheit (PPU) 300 gemäß einer Ausführungsform. In einer Ausführungsform ist die PPU 300 ein Multithread-Prozessor, der auf einer oder mehreren integrierten Schaltungsvorrichtungen implementiert ist. Die PPU 300 ist eine Latenz-verbergende Architektur, die ausgestaltet ist, um eine große Anzahl von Threads bzw. Strängen parallel zu verarbeiten. Ein Thread bzw. Strang (d.h. ein Ausführungsthread) ist eine Instanziierung eines Satzes von Anweisungen, die konfiguriert sind, um von der PPU 300 ausgeführt zu werden. In einer Ausführungsform ist die PPU 300 eine Graphikverarbeitungseinheit (GPU), die konfiguriert ist, um eine Graphik-Rendering-Pipeline zur Verarbeitung von dreidimensionalen (3D) Graphikdaten zu implementieren, um zweidimensionale (2D) Bilddaten zur Anzeige auf einer Anzeigevorrichtung, wie beispielsweise einer Flüssigkristallanzeige(LCD)-Vorrichtung, zu erzeugen. In anderen Ausführungsformen kann die PPU 300 zum Durchführen von Allzweckberechnungen benutzt werden. Während ein beispielhafter paralleler Prozessor hier für veranschaulichende Zwecke bereitgestellt wird, sei nachdrücklich bemerkt, dass ein derartiger Prozessor lediglich für veranschaulichende Zwecke dargelegt wird und dass ein beliebiger Prozessor benutzt werden kann, um dasselbe zu ergänzen und/oder zu ersetzen.
Eine oder mehrere PPUs 300 können konfiguriert sein, um Tausende von HPC(High Performing Computing)-, Datenzentrum- und Maschinenlern-Anwendungen zu beschleunigen. Die PPU 300 kann konfiguriert sein, um zahlreiche Systeme und Anwendungen für tiefgehendes Lernen zu beschleunigen, die autonome Fahrzeugplattformen, tiefgehendes Lernen, hochgenaue Sprache, Bild- und Texterkennungssysteme, intelligente Videoanalyse, molekulare Simulationen, Wirkstoffentdeckung, Krankheitsdiagnose, Wettervorhersage, Analyse großer Datenmengen, Astronomie, Molekulardynamiksimulation, Finanzmodellierung, Robotertechnik, Fabrikautomation, Sprachübersetzung in Echtzeit, Online-Suchoptimierungen und personalisierte Benutzerempfehlungen und dergleichen umfassen.
Wie in 8 gezeigt, umfasst die PPU 300 eine Eingabe/Ausgabe(E/A)-Einheit 305, eine Frontend-Einheit 315, eine Planer-Einheit 320, eine Arbeitsverteilungs-Einheit 325, einen Hub 330, eine Kreuzschiene (XBar) 370, einen oder mehrere allgemeine Verarbeitungscluster (GPCs) 350 und eine oder mehrere Partitions-Einheiten 380. Die PPU 300 kann mit einem Host-Prozessor oder anderen PPUs 300 über einen Interconnect des Hochgeschwindigkeits-NVLink 310 verbunden sein. Die PPU 300 kann ebenfalls mit einem Host-Prozessor oder anderen peripheren Vorrichtungen über einen Interconnect 302 verbunden sein. In einer Ausführungsform kann der lokale Speicher eine Anzahl von Direktzugriffsspeicher(DRAM)-Vorrichtungen umfassen. Die DRAM-Vorrichtungen können als ein HBM(Speicher mit hoher Bandbreite)-Subsystem konfiguriert sein, wobei mehrere DRAM-Dies innerhalb jeder Vorrichtung gestapelt sind.
Der Interconnect des NVLink 310 ermöglicht Systemen, eine oder mehrere PPUs 300 zu skalieren und zu umfassen, die mit einer oder mehreren CPUs kombiniert sind, unterstützt Cache-Kohärenz zwischen den PPUs 300 und CPUs sowie CPU-Mastering. Daten und/oder Befehle können mittels des NVLink 310 durch den Hub 330 an/von anderen Einheiten der PPU 300, wie beispielsweise eine oder mehrere Kopiermaschinen, einen Videocodierer, einen Videodecodierer, eine Leistungsverwaltungseinheit usw. (nicht explizit gezeigt) übertragen werden. Das NVLink 310 wird ausführlicher in Verbindung mit 5B beschrieben.
Die E/A-Einheit 305 ist konfiguriert, um Kommunikationen (d.h. Befehle, Daten usw.) von einem Host-Prozessor (nicht gezeigt) über den Interconnect 302 zu übertragen und zu empfangen. Die E/A-Einheit 305 kann mit dem Host-Prozessor direkt über den Interconnect 302 oder durch eine oder mehrere Zwischenvorrichtungen, wie beispielsweise eine Speicherbrücke, kommunizieren. In einer Ausführungsform kann die E/A-Einheit 305 mit einem oder mehreren anderen Prozessoren, wie beispielsweise eine oder mehrere der PPUs, über den Interconnect 302 kommunizieren. In einer Ausführungsformen implementiert die E/A-Einheit 305 eine PCIe(Peripheral Component Interconnect Express)-Schnittstelle für Kommunikationen über einen PCIe-Bus und der Interconnect 302 ist ein PCIe-Bus. In alternativen Ausführungsformen kann die E/A-Einheit 305 andere Typen von wohlbekannten Schnittstellen zum Kommunizieren mit externen Vorrichtungen umfassen.
Die E/A-Einheit 305 decodiert Pakete, die über den Interconnect 302 empfangen wurden. In einer Ausführungsform stellen die Pakete Befehle dar, die konfiguriert sind, um die PPU 300 zu veranlassen, verschiedene Operationen durchzuführen. Die E/A-Einheit 305 überträgt die decodierten Befehle an verschiedene andere Einheiten der PPU 300, wie es die Befehle spezifizieren können. Beispielsweise können einige Befehle an die Frontend-Einheit 315 übertragen werden. Andere Befehle können an den Hub 330 oder andere Einheiten der PPU 300, wie beispielsweise eine oder mehrere Kopiermaschinen, einen Video-Codierer, einen Video-Decodierer, eine Leistungsverwaltungseinheit usw. (nicht explizit gezeigt) übertragen werden. Mit anderen Worten ist die E/A-Einheit 305 konfiguriert, um Kommunikationen zwischen und unter den verschiedenen logischen Einheiten der PPU 300 weiterzuleiten.
In einer Ausführungsform codiert ein Programm, das von dem Host-Prozessor ausgeführt wird, einen Befehlsstrom in einem Puffer, welcher der PPU 300 Arbeitslasten zur Verarbeitung bereitstellt. Eine Arbeitslast kann mehrere Anweisungen und Daten umfassen, die von diesen Anweisungen zu verarbeiten sind. Der Puffer ist eine Region in einem Speicher, der von sowohl dem Host-Prozessor als auch der PPU 300 zugänglich ist (d.h. Lesen/Schreiben). Beispielsweise kann die Host-Schnittstelleneinheit 305 konfiguriert sein, um auf den Puffer in einem Systemspeicher, der mit dem Interconnect 302 verbunden ist, über Speicheranfragen, die über den Interconnect 302 übertragen werden, zuzugreifen. In einer Ausführungsform schreibt der Host-Prozessor den Befehlsstrom in den Puffer und überträgt dann einen Zeiger zu dem Start des Befehlsstroms an die PPU 300. Die Frontend-Einheit 315 empfängt Zeiger auf einen oder mehrere Befehlsströme. Die Frontend-Einheit 315 verwaltet den einen oder mehrere Ströme, liest Befehle aus den Strömen und leitet Befehle an die verschiedenen Einheiten der PPU 300 weiter.
Die Frontend-Einheit 315 ist mit einer Planer-Einheit 320 gekoppelt, welche die verschiedenen GPCs 350 konfiguriert, um Aufgaben zu verarbeiten, die durch den einen oder mehrere Ströme definiert sind. Die Planer-Einheit 320 ist konfiguriert, um Zustandsinformation zu verfolgen, die verschiedene Aufgaben betrifft, die von der Planer-Einheit 320 verwaltet werden. Der Zustand kann angeben, welchem GPC 350 eine Aufgabe zugewiesen ist, ob die Aufgabe aktiv oder inaktiv ist, ob ein Prioritätsniveau der Aufgabe zugeordnet ist und so weiter. Die Planer-Einheit 320 verwaltet die Ausführung einer Mehrzahl von Aufgaben auf dem einen oder mehreren GPCs 350.
Die Planer-Einheit 320 ist mit einer Arbeitsverteilungs-Einheit 325 gekoppelt, die konfiguriert ist, um Aufgaben zur Ausführung auf den GPCs 350 zu versenden. Die Arbeitsverteilungs-Einheit 325 kann eine Anzahl von eingeplanten Aufgaben verfolgen, die von der Planer-Einheit 320 empfangen werden. In einer Ausführungsform verwaltet die Arbeitsverteilungs-Einheit 325 einen Pool für anstehende Aufgaben und einen Pool für aktive Aufgaben für jeden der GPCs 350. Der Pool für anstehende Aufgaben kann eine Anzahl von Schlitzen (z.B. 32 Schlitze) umfassen, die Aufgaben enthalten, die zugewiesen sind, um von einem bestimmten GPC 350 verarbeitet zu werden. Der Pool für aktive Aufgaben kann eine Anzahl von Schlitzen (z.B. 4 Schlitze) für Aufgaben umfassen, die von den GPCs 350 aktiv verarbeitet werden. Wenn ein GPC 350 die Ausführung einer Aufgabe abschließt, wird diese Aufgabe aus dem Pool für aktive Aufgaben für den GPC 350 geräumt und eine der anderen Aufgaben wird aus dem Pool für anstehende Aufgaben ausgewählt und zur Ausführung auf dem GPC 350 eingeplant. Wenn eine aktive Aufgabe auf dem GPC 350 inaktiv war, wie beispielsweise während darauf gewartet wird, dass eine Datenabhängigkeit behoben wird, dann kann die aktive Aufgabe aus dem GPC 350 geräumt und zu dem Pool für anstehende Aufgaben zurückgeführt werden, während eine weitere Aufgabe in dem Pool für anstehende Aufgaben ausgewählt und zur Ausführung auf dem GPC 350 eingeplant wird.
Die Arbeitsverteilungs-Einheit 325 kommuniziert mit dem einen oder mehreren GPCs 350 über eine Kreuzschiene bzw. XBar 370. Die XBar 370 ist ein Interconnect-Netzwerk, das viele der Einheiten der PPU 300 mit anderen Einheiten der PPU 300 koppelt. Beispielsweise kann die XBar 370 konfiguriert sein, um die Arbeitsverteilungs-Einheit 325 mit einem bestimmten GPC 350 zu koppeln. Obwohl nicht explizit gezeigt, kann eine oder mehrere andere Einheiten der PPU 300 ebenfalls mit der XBar 370 über den Hub 330 verbunden sein.
Die Aufgaben werden von der Planer-Einheit 320 verwaltet und an einen GPC 350 durch die Arbeitsverteilungs-Einheit 325 versandt. Der GPC 350 ist konfiguriert, um die Aufgabe zu verarbeiten und Ergebnisse zu erzeugen. Die Ergebnisse können von anderen Aufgaben innerhalb des GPC 350 verbraucht werden, an einen unterschiedlichen GPC 350 über die XBar 370 weitergeleitet oder im Speicher 304 gespeichert werden. Die Ergebnisse können in den Speicher 304 über die Partitions-Einheiten 380 geschrieben werden, die eine Speicherschnittstelle zum Lesen und Schreiben von Daten in/von dem Speicher 304 implementieren. In einer Ausführungsform umfasst die PPU 300 eine Anzahl U von Speicherpartitions-Einheiten 380, die gleich der Anzahl von getrennten und unterschiedlichen Speichervorrichtungen 304 ist, die mit der PPU 300 gekoppelt sind. Eine Speicherpartitions-Einheit 380 wird nachstehend ausführlicher in Verbindung mit 9B beschrieben.
In einer Ausführungsform führt ein Host-Prozessor einen Treiber-Kernel aus, der eine Anwendungsprogrammmier-Schnittstelle (API) implementiert, die einer oder mehreren Anwendungen ermöglicht, die auf dem Host-Prozessor ausgeführt werden, Operationen zur Ausführung auf der PPU 300 einzuplanen. In einer Ausführungsform werden mehrere Rechenanwendungen gleichzeitig von der PPU 300 ausgeführt und die PPU 300 stellt Isolierung, Dienstqualität (QoS) und unabhängige Adressräume für die mehreren Rechenanwendungen bereit. Eine Anwendung kann Anweisungen (d.h. API-Aufrufe) erzeugen, welche den Treiber-Kernel veranlassen, eine oder mehrere Aufgaben zur Ausführung durch die PPU 300 zu erzeugen. Der Treiberkernel gibt Aufgaben an einen oder mehrere Streams aus, die von der PPU 300 verarbeitet werden. Jede Aufgabe kann eine oder mehrere Gruppen von in Beziehung stehender Threads umfassen, die hier als ein Warp bezeichnet werden. In einer Ausführungsform umfasst ein Warp 32 in Beziehung stehende Threads, die parallel ausgeführt werden können. Kooperierende Threads können sich auf eine Mehrzahl von Threads beziehen, die Anweisungen umfassen, um die Aufgabe durchzuführen, und die Daten durch einen gemeinsam genutzten Speicher austauschen können. Threads und kooperierende Threads werden ausführlicher in Verbindung mit 5A beschrieben.
9A veranschaulicht einen GPC 350 der PPU 300 von 8 gemäß einer Ausführungsform. Wie in 9A gezeigt, umfasst jeder GPC 350 eine Anzahl von Hardwareeinheiten zur Verarbeitung von Aufgaben. In einer Ausführungsform umfasst jeder GPC 350 einen Pipeline-Manager 410, eine Vor-Raster-Operationen-Einheit (PROP) 415, eine Raster-Engine 425, eine Arbeitsverteilungs-Kreuzschiene (WDX) 480, eine Speicherverwaltungseinheit (MMU) 490 und einen oder mehrere Datenverarbeitungscluster (DPCs) 420. Es wird anerkannt, dass der GPC 350 von 9A andere Hardwareeinheiten anstelle von oder zusätzlich zu den in 9A gezeigten Einheiten umfassen kann.
In einer Ausführungsform wird der Betrieb des GPC 350 durch den Pipeline-Manager 410 gesteuert. Der Pipeline-Manager 410 verwaltet die Konfiguration des einen oder mehrerer DPCs 420 zur Verarbeitung von Aufgaben, die dem GPC 350 zugeteilt sind. In einer Ausführungsform kann der Pipeline-Manager 410 mindestens einen des einen oder mehrerer DPCs 420 konfigurieren, um mindestens einen Abschnitt einer Graphik-Rendering-Pipeline zu implementieren. Beispielsweise kann ein DPC 420 konfiguriert sein, um ein Vertex-Shader-Programm auf dem programmierbaren Streaming-Multiprozessor (SM) 440 auszuführen. Der Pipeline-Manager 410 kann ebenfalls konfiguriert sein, um Pakete, die von der Arbeitsverteilungs-Einheit 325 empfangen werden, an die geeigneten logischen Einheiten innerhalb des GPC 350 weiterzuleiten. Beispielsweise können einige Pakete an Festfunktions-Hardwareeinheiten in der PROP 415 und/oder der Raster-Engine 425 weitergeleitet werden, während andere Pakete an die DPCs 420 zur Verarbeitung durch die Primitiven-Engine 435 oder den SM 440 weitergeleitet werden können. In einer Ausführungsform kann der Pipeline-Manager 410 mindestens einen oder mehrere DPCs konfigurieren, um ein neuronales Netzwerkmodell und/oder eine Rechen-Pipeline zu implementieren.
Die PROP-Einheit 415 ist konfiguriert, um Daten, die von der Raster-Engine 425 und den DPCs 420 erzeugt werden, an eine Raster-Operationen-Einheit (ROP-Einheit) weiterzuleiten, die nachstehend ausführlicher in Verbindung mit 9B beschrieben wird. Die PROP-Einheit 415 kann ebenfalls konfiguriert sein, um Optimierungen zur Farbenmischung durchzuführen, Pixeldaten zu organisieren, Adressenübersetzungen und dergleichen durchzuführen.
Graphikverarbeitungs-Pipeline
In einer Ausführungsform umfasst die PPU 300 eine Graphikverarbeitungseinheit (GPU). Die PPU 300 ist konfiguriert, um Befehle zu empfangen, die Shader-Programme zur Verarbeitung von Graphikdaten spezifizieren. Graphikdaten können als ein Satz von Primitiven, wie beispielsweise Punkte, Linien, Dreiecke, Vierecke, Dreieckstreifen und dergleichen, definiert sein. Typischerweise umfasst eine Primitive Daten, die eine Anzahl von Vertices für die Primitive (z.B., in einem Modellraum-Koordinatensystem) sowie auch Attribute spezifizieren, die jedem Vertex der Primitive zugeordnet sind. Die PPU 300 kann konfiguriert sein, um die Graphikprimitive zu verarbeiten, um ein Frame-Puffer zu erzeugen (d.h. Pixeldaten für jedes der Pixel der Anzeige).
Eine Anwendung schreibt Modelldaten für eine Szene (d.h. eine Sammlung von Vertices und Attributen) in einen Speicher, wie beispielsweise einen Systemspeicher oder Speicher 304. Die Modelldaten definieren jedes der Objekte, die auf einer Anzeige sichtbar sein können. Die Anwendung führt dann einen API-Aufruf an dem Treiber-Kernel aus, der die Modelldaten anfragt, die zu rendern und anzuzeigen sind. Der Treiber-Kernel liest die Modelldaten und schreibt Befehle an den einen oder mehrere Ströme, um Operationen durchzuführen, um die Modelldaten zu verarbeiten. Die Befehle können unterschiedliche Shader-Programme referenzieren, die auf den SMs 440 der PPU 300 zu implementieren sind, die einen oder mehrere eines Vertex-Shader, Hull-Shader, Domain-Shader, Geometrie-Shader und eines Pixel-Shader umfassen können. Beispielsweise können eine oder mehrere der SMs 440 konfiguriert sein, um ein Vertex-Shader-Programm auszuführen, das eine Anzahl von Vertices verarbeitet, die durch die Modelldaten definiert sind. In einer Ausführungsform können die unterschiedlichen SMs 440 konfiguriert sein, um unterschiedliche Shader-Programme nebenläufig auszuführen. Beispielsweise kann eine erste Untermenge von SMs 440 konfiguriert sein, um ein Vertex-Shader-Programm auszuführen, während eine zweite Untermenge von SMs 440 konfiguriert sein kann, um ein Pixel-Shader-Programm auszuführen. Die erste Untermenge von SMs 440 verarbeitet Vertexdaten, um verarbeitete Vertexdaten zu erzeugen, und schreibt die verarbeiteten Vertexdaten in den L2-Cache-Speicher 460 und/oder den Speicher 304. Nachdem die verarbeiteten Vertexdaten gerastert sind (d.h. von dreidimensionalen Daten in zweidimensionale Daten im Bildschirmraum transformiert sind), um Fragmentdaten zu erzeugen, führt die zweite Untermenge von SMs 440 einen Pixel-Shader aus, um verarbeitete Fragmentdaten zu erzeugen, die dann mit anderen verarbeiteten Fragmentdaten gemischt werden und in den Frame-Puffer im Speicher 304 geschrieben werden. Das Vertex-Shader-Programm und das Pixel-Shader-Programm können nebenläufig ausgeführt werden, wobei unterschiedliche Daten von der gleichen Szene in einem Pipeline-Verfahren verarbeitet werden, bis alle der Modelldaten für die Szene gerenderten zu dem Frame-Puffer gerendert wurden. Dann wird der Inhalt des Frame-Puffers an einen Anzeigencontroller zur Anzeige auf einer Anzeigevorrichtung übertragen.
11 ist ein Konzeptdiagramm einer von der PPU 300 von 8 implementierten Graphikverarbeitungs-Pipeline 600 gemäß einer Ausführungsform. Die Graphikverarbeitungs-Pipeline 600 ist ein abstraktes Ablaufdiagramm der Verarbeitungsschritte, die implementiert werden, um 2D-Computer-erzeugte Bilder aus 3D-Geometriedaten zu erzeugen. Wie bekannt ist, können Pipeline-Architekturen Operationen mit langer Latenzzeit effizienter durch Aufteilen der Operation in eine Mehrzahl von Stufen durchführen, wobei die Ausgabe jeder Stufe mit dem Eingang der nächsten aufeinanderfolgenden Stufe gekoppelt ist. Somit empfängt die Graphikverarbeitungs-Pipeline 600 Eingabedaten 601, die von einer Stufe an die nächste Stufe der Graphikverarbeitungs-Pipeline 600 übertragen werden, um Ausgabedaten 602 zu erzeugen. In einer Ausführungsform kann die Graphikverarbeitungs-Pipeline 600 eine Graphikverarbeitungs-Pipeline darstellen, die durch die OpenGL^® API definiert ist. Als eine Option kann die Graphikverarbeitungs-Pipeline 600 im Kontext der Funktionalität und Architektur der vorherigen Figuren und/oder etwaiger nachfolgenden Figur(en) implementiert werden.
Wie in 11 gezeigt, umfasst die Graphikverarbeitungs-Pipeline 600 eine Pipeline-Architektur, die eine Anzahl von Stufen umfasst. Die Stufen umfassen, sind jedoch nicht beschränkt auf, eine Datenanordnungs-Stufe 610, eine Vertex-Shading-Stufe 620, eine Primitivenanordnungs-Stufe 630, eine Geometrie-Shading-Stufe 640, eine Darstellungsfeld-Skalierungs-, Aussonderungs- und Abschneide-(VSCC)-Stufe 650, eine Rasterungs-Stufe 660, eine Fragment-Shading-Stufe 670 und eine Raster-Operationen-Stufe 680. Wie oben beschrieben, können die Software-Shading-Algorithmen, die in Verbindung mit derartiger Shading-Software arbeiten, optimiert werden, um Rechenzeit zu verringern.
In einer Ausführungsform umfassen die Eingabedaten 601 Befehle, welche die Verarbeitungseinheiten konfigurieren, um die Stufen der Graphikverarbeitungs-Pipeline 600 zu implementieren, und geometrische Primitive (z.B., Punkte, Linien, Dreiecke, Vierecke, Dreieckstreifen oder Fächer, usw.), die mittels der Stufen zu verarbeiten sind. Die Ausgabedaten 602 können Pixeldaten (d.h. Farbdaten) umfassen, die in einen Frame-Puffer oder einen anderen Typ von Oberflächen-Datenstruktur in einem Speicher kopiert werden.
Die Datenanordnungs-Stufe 610 empfängt die Eingabedaten 601, die Vertexdaten für Oberflächen höherer Ordnung, Primitive oder dergleichen spezifizieren. Die Datenanordnungs-Stufe 610 sammelt die Vertexdaten in einem temporären Speicher oder Warteschlange, wie beispielsweise durch Empfangen eines Befehls von dem Host-Prozessor, der einen Zeiger auf einen Puffer im Speicher umfasst und die Vertexdaten aus dem Puffer liest. Die Vertexdaten werden dann an die Vertex-Shading-Stufe 620 zur Verarbeitung übertragen.
Die Vertex-Shading-Stufe 620 verarbeitet Vertexdaten durch Durchführen eines Satzes von Operationen (d.h. eines Vertex-Shader oder eines Programms) einmal für jede der Vertices. Vertices können beispielsweise als ein 4-Koordinaten-Vektor (d.h. <x, y, z, w>) spezifiziert sein, der einem oder mehreren Vertexattributen (z.B., Farbe, Texturkoordinaten, Oberflächennormalen, usw.) zugeordnet ist. Die Vertex-Shading-Stufe 620 kann einzelne Vertexattribute, wie beispielsweise Position, Farbe, Texturkoordinaten und dergleichen, manipulieren. Mit anderen Worten führt die Vertex-Shading-Stufe 620 Operationen an den Vertex-Koordinaten oder anderen Vertexattributen durch, welche einem Vertex zugeordnet sind. Derartige Operationen umfassen gewöhnlicherweise Beleuchtungs-Operationen (d.h. Modifizieren von Farbattributen für einen Vertex) und Transformations-Operationen (d.h. Modifizieren des Koordinatenraums für einen Vertex). Beispielsweise können Vertices unter Verwendung von Koordinaten in einem Objekt-Koordinatenraum spezifiziert sein, die durch Multiplizieren der Koordinaten mittels einer Matrix transformiert werden, welche die Koordinaten aus dem Objekt-Koordinatenraum in einen Welt-Raum oder einen normierten Vorrichtungskoordinaten(NCD = Normalized Device Coordinates)-Raum übersetzt. Die Vertex-Shading-Stufe 620 erzeugt transformierte Vertexdaten, die an die Primitivenanordnungs-Stufe 630 übertragen werden.
Die Primitivenanordnungs-Stufe 630 sammelt Vertices, die mittels der Vertex-Shading-Stufe 620 ausgegeben werden, und gruppiert die Vertices in geometrische Primitive zur Verarbeitung durch die Geometrie-Shading-Stufe 640. Beispielsweise kann die Primitivenanordnungs-Stufe 630 konfiguriert sein, um alle drei aufeinanderfolgenden Vertices als eine geometrische Primitive (d.h. ein Dreieck) zur Übertragung an die Geometrie-Shading-Stufe 640 zu gruppieren. In einigen Ausführungsformen können spezifische Vertices für aufeinanderfolgende geometrische Primitive erneut verwendet werden (z.B. können zwei aufeinanderfolgende Dreiecke in einem Dreieckstreifen zwei Vertices gemeinsam benutzen). Die Primitivenanordnungs-Stufe 630 überträgt geometrische Primitive (d.h. eine Sammlung von zugeordneten Vertices) an die Geometrie-Shading-Stufe 640.
Die Geometrie-Shading-Stufe 640 verarbeitet geometrische Primitive durch Durchführen eines Satzes von Operationen (d.h. eines Geometrie-Shader oder Programms) an den geometrischen Primitiven. Tessellations-Operationen können eine oder mehrere geometrische Primitive aus jeder geometrischen Primitive erzeugen. Mit anderen Worten kann die Geometrie-Shading-Stufe 640 jede geometrische Primitive in ein feineres Netz von zwei oder mehreren geometrischen Primitiven zur Verarbeitung durch den Rest der Graphikverarbeitungs-Pipeline 600 unterteilen. Die Geometrie-Shading-Stufe 640 überträgt geometrische Primitive an die Darstellungsfeld-SCC-Stufe 650.
In einer Ausführungsform kann die Graphikverarbeitungs-Pipeline 600 innerhalb eines Streaming-Multiprozessors arbeiten und die Vertex-Shading-Stufe 620, die Primitivenanordnungs-Stufe 630, die Geometrie-Shading-Stufe 640, die Fragment-Shading-Stufe 670 und/oder die damit zugeordnete Hardware/Software kann sequenziell Verarbeitungsoperationen durchführen. Sobald die sequenziellen Verarbeitungsoperationen in einer Ausführungsform abgeschlossen sind, kann die Darstellungsfeld-SCC-Stufe 650 die Daten benutzen. In einer Ausführungsform können Primitivendaten, die durch eine oder mehrere der Stufen in der Graphikverarbeitungs-Pipeline 600 verarbeitet wurden, in einen Cache-Speicher (z.B. einen LI-Cache-Speicher, einen Vertex-Cache-Speicher usw.) geschrieben werden. In diesem Fall kann in einer Ausführungsform die Darstellungsfeld-SCC-Stufe 650 auf die Daten in dem Cache-Speicher zugreifen. In einer Ausführungsform sind die Darstellungsfeld-SCC-Stufe 650 und die Rasterungs-Stufe 660 als Festfunktionsschaltungen implementiert.
Die Darstellungsfeld-SCC-Stufe 650 führt Darstellungsfeld-Skalierung, Aussonderung und Abschneidung der geometrischen Primitive durch. Jede Oberfläche, die gerendert wird, wird einer abstrakten Kameraposition zugeordnet. Die Kameraposition stellt einen Ort eines Betrachters dar, der die Szene betrachtet, und definiert einen Betrachtungsstumpf, der die Objekte der Szene einschließt. Der Betrachtungsstumpf kann eine Betrachtungsebene, eine hintere Ebene und vier Abschneideebenen umfassen. Jede geometrische Primitive, die vollständig außerhalb des Betrachtungsstumpfes ist, kann ausgesondert (d.h. verworfen) werden, weil die geometrische Primitive nicht zu der endgültigen gerenderten Szene beitragen wird. Jede geometrische Primitive, die innerhalb des Betrachtungsstumpfs und teilweise außerhalb des Betrachtungsstumpf ist, kann abgeschnitten werden (d.h. in eine neue geometrische Primitive transformiert werden, die innerhalb des Betrachtungsstumpf eingeschlossen ist). Des Weiteren können geometrische Primitive jeweils basierend auf einer Tiefe des Betrachtungsstumpfes skaliert werden. Alle potenziell sichtbaren geometrischen Primitive werden dann an die Rasterungs-Stufe 660 übertragen.
Die Rasterungs-Stufe 660 wandelt die 3D-geometrischen Primitive in 2D-Fragmente um (die z.B. in der Lage sind, zur Anzeige benutzt zu werden, usw.). Die Rasterungs-Stufe 660 kann konfiguriert sein, um die Vertices der geometrischen Primitive zu benutzen, um einen Satz von Ebenengleichungen aufzustellen, von denen verschiedene Attribute interpoliert werden können. Die Rasterungs-Stufe 660 kann ebenfalls eine Abdeckungsmaske für eine Mehrzahl von Pixeln berechnen, die angibt, ob eine oder mehrere Abtastorte für das Pixel die geometrische Primitive schneiden. In einer Ausführungsform kann auch z-Testen durchgeführt werden, um zu bestimmen, ob die geometrische Primitive von anderen geometrischen Primitiven verdeckt wird, die bereits gerastert wurden. Die Rasterungs-Stufe 660 erzeugt Fragmentdaten (d.h. interpolierte Vertexattribute, die einem bestimmten Abtastort für jedes abgedeckte Pixel zugeordnet sind), die an die Fragment-Shading-Stufe 670 übertragen werden.
Die Fragment-Shading-Stufe 670 verarbeitet Fragmentdaten durch Durchführen eines Satzes von Operationen (d.h. eines Fragment-Shader oder eines Programms) an jedem der Fragmente. Die Fragment-Shading-Stufe 670 kann Pixeldaten (d.h. Farbenwerte) für das Fragment erzeugen, wie beispielsweise durch Durchführen von Beleuchtungs-Operationen oder Abtasten von Texturabbildungen unter Verwendung von interpolierten Texturkoordinaten für das Fragment. Die Fragment-Shading-Stufe 670 erzeugt Pixeldaten, die an die Rasteroperationen-Stufe 680 übertragen werden.
Die Rasteroperationen-Stufe 680 kann verschiedene Operationen an den Pixeldaten durchführen, wie beispielsweise Alpha-Tests, Stencil-Tests und Mischung der Pixeldaten mit anderen Pixeldaten, die anderen Fragmente entsprechen, die dem Pixel zugeordnet sind. Wenn die Raster-Operationen-Stufe 680 die Verarbeitung der Pixeldaten (d.h. der Ausgabedaten 602) beendet hat, können die Pixeldaten in ein Render-Ziel, wie beispielsweise einen Frame-Puffer, einen Farbenpuffer oder dergleichen, geschrieben werden.
Die Raster-Engine 425 umfasst eine Anzahl von Festfunktions-Hardwareeinheiten, die konfiguriert sind, um verschiedene Raster-Operationen durchzuführen. In einer Ausführungsform umfasst die Raster-Engine 425 eine Setup-Engine, eine Grobraster-Engine, eine Aussonderungs-Engine, eine Abschneide-Engine, eine Feinraster-Engine und eine Kachel-verschmelzende Engine. Die Setup-Engine empfängt transformierte Vertices und erzeugt Ebenengleichungen, die den geometrischen Primitiven zugeordnet sind, die durch die Vertices definiert werden. Die Ebenengleichungen werden an die Grobraster-Engine übertragen, um Abdeckungsinformation (z.B. eine (x,y)-Abdeckungsmaske für eine Kachel) für die Primitive zu erzeugen. Die Ausgabe der Grobraster-Engine wird an die Aussonderungs-Engine übertragen, wo Fragmente, die der Primitiven zugeordnet sind, die einen z-Test nicht bestehen, ausgesondert werden, und nicht ausgesonderte Fragmente werden an eine Abschneide-Engine übertragen werden, wo Fragmente, die außerhalb eines Betrachtungsstumpfes liegen, abgeschnitten werden. Diejenigen Fragmente, welche die Abschneidung und Aussonderung überleben, können an eine Feinraster-Engine weitergeben werden, um Attribute für die Pixelfragmente basierend auf den Ebenengleichungen zu erzeugen, die durch die Setup-Engine erzeugt werden. Die Ausgabe der Raster-Engine 425 umfasst Fragmente, die beispielsweise von einem Fragment-Shader zu verarbeiten sind, der innerhalb eines DPC 420 implementiert ist.
Es wird anerkannt, dass eine oder mehrere zusätzliche Stufen in der Graphikverarbeitungs-Pipeline 600 zusätzlich zu oder anstatt einer oder mehrerer der oben beschriebenen Stufen enthalten sein können. Verschiedene Implementierungen der abstrakten Graphikverarbeitungs-Pipeline können unterschiedliche Stufen implementieren. Des Weiteren können eine oder mehrere der oben beschriebenen Stufen der Graphikverarbeitungs-Pipeline in einigen Ausführungsformen (wie beispielsweise der Geometrie-Shading-Stufe 640) ausgeschlossen sein. Andere Arten von Graphikverarbeitungs-Pipelines werden betrachtet, als innerhalb des Schutzumfangs der vorliegenden Offenbarung zu liegen. Des Weiteren können beliebige der Stufen der Graphikverarbeitungs-Pipeline 600 von einer oder mehreren dedizierten Hardwareeinheiten innerhalb eines Graphikprozessors, wie beispielsweise der PPU 300, implementiert werden. Andere Stufen der Graphikverarbeitungs-Pipeline 600 können durch programmierbare Hardwareeinheiten, wie beispielsweise dem SM 440 der PPU 300, implementiert werden.
Die Graphikverarbeitungs-Pipeline 600 kann über eine Anwendung implementiert werden, die von einem Host-Prozessor, wie beispielsweise einer CPU, ausgeführt wird. In einer Ausführungsform kann ein Vorrichtungstreiber eine Anwendungsprogrammierschnittstelle (API) implementieren, die verschiedene Funktionen definiert, die von einer Anwendung benutzt werden können, um graphische Daten zur Anzeige zu erzeugen. Der Vorrichtungstreiber ist ein Softwareprogramm, das eine Mehrzahl von Befehlen umfasst, die den Betrieb der PPU 300 steuern. Die API stellt eine Abstraktion für einen Programmierer bereit, die einem Programmierer erlaubt, spezialisierte Graphikhardware zu benutzen, wie beispielsweise die PPU 300, um die graphischen Daten zu erzeugen, ohne zu verlangen, dass der Programmierer den spezifischen Befehlssatz für die PPU 300 benutzen muss. Die Anwendung kann einen API-Aufruf umfassen, der an den Vorrichtungstreiber für die PPU 300 weitergeleitet wird. Der Vorrichtungstreiber interpretiert den API-Aufruf und führt verschiedene Operationen durch, um auf den API-Aufruf zu antworten. In einigen Fällen kann der Vorrichtungstreiber Operationen durch Ausführen von Befehlen auf der CPU 550 durchführen. In anderen Fällen kann der Vorrichtungstreiber Operationen, zumindest teilweise, durch Starten von Operationen auf der PPU 300 durchführen, wobei eine Eingabe/Ausgabe-Schnittstelle zwischen der CPU und der PPU 300 benutzt wird. In einer Ausführungsform ist der Vorrichtungstreiber konfiguriert, um die Graphikverarbeitungs-Pipeline 600 unter Benutzung der Hardware der PPU 300 zu implementieren.
Verschiedene Programme können innerhalb der PPU 300 ausgeführt werden, um die verschiedenen Stufen der Graphikverarbeitungs-Pipeline 600 zu implementieren. Beispielsweise kann der Vorrichtungstreiber ein Kernel auf der PPU 300 starten, um die Vertex-Shading-Stufe 620 auf einem SM 440 (oder mehreren SMs 440) durchzuführen. Der Vorrichtungstreiber (oder den von der PPU 300 ausgeführten Anfangskernel) kann ebenfalls andere Kernels auf der PPU 300 starten, um andere Stufen der Graphikverarbeitungs-Pipeline 600 durchzuführen, wie beispielsweise die Geometrie-Shading-Stufe 640 und die Fragment-Shading-Stufe 670. Außerdem können einige der Stufen der Graphikverarbeitungs-Pipeline 600 auf einer festen Hardwareeinheit implementiert werden, wie beispielsweise einem Rasterer oder einem Daten-Assembler, der innerhalb der PPU 300 implementiert ist. Es wird anerkannt, dass Ergebnisse von einem Kernel durch eine oder mehrere intervenierende Festfunktions-Hardwareeinheiten verarbeitet werden können, bevor sie von einem nachfolgenden Kernel auf einem SM 440 verarbeitet werden.
Jeder DPC 420, der in dem GPC 350 umfasst ist, umfasst einen M-Pipe-Controller (MPC) 430, eine Primitiven-Engine 435 und einen oder mehrere SMs 440. Der MPC 430 steuert den Betrieb des DPC 420, wobei von dem Pipeline-Manager 410 empfangene Pakete an die geeigneten Einheiten im DPC 420 weiterleitet werden. Beispielsweise können einem Vertex zugeordnete Pakete an die Primitiven-Engine 435 weitergeleitet werden, die konfiguriert ist, um der Vertex zugeordnete Vertexattribute von dem Speicher 304 abzurufen. Im Gegensatz dazu können einem Shader-Programm zugeordnete Pakete an den SM 440 übertragen werden.
Der SM 440 umfasst einen programmierbaren Streaming-Prozessor, der konfiguriert ist, um Aufgaben zu verarbeiten, die durch eine Anzahl von Threads dargestellt werden. Jeder SM 440 umfasst mehrere Threads (ist multi-threaded) und ist konfiguriert, um eine Mehrzahl von Threads (z.B. 32 Threads) von einer bestimmten Gruppe von Threads nebenläufig auszuführen. In einer Ausführungsform implementiert der SM 440 eine SIMD(Einzelne-Anweisung, Mehrere-Daten)-Architektur, wobei jeder Thread in einer Gruppe von Threads (d.h. einem Warp) konfiguriert ist, um einen unterschiedlichen Satz von Daten basierend auf dem gleichen Satz von Anweisungen zu verarbeiten. Alle Threads in der Gruppe von Threads führen die gleichen Anweisungen aus. In einer anderen Ausführungsform implementiert der SM 440 eine SIMT(Einzelne Anweisung, Mehrere Threads)-Architektur, wobei jeder Thread in einer Gruppe von Threads konfiguriert ist, um einen unterschiedlichen Satz von Daten basierend auf dem gleichen Satz von Anweisungen zu verarbeiten, wobei jedoch einzelnen Threads in der Gruppe von Threads ermöglicht wird, während der Ausführung zu divergieren. In einer Ausführungsform wird ein Programmzähler, ein Aufrufstapel und ein Ausführungszustand für jeden Warp beibehalten, was eine Nebenläufigkeit zwischen Warps und eine serielle Ausführung innerhalb Warps ermöglicht, wenn Threads innerhalb des Warp divergieren. In einer weiteren Ausführungsform wird ein Programmzähler, ein Aufrufstapel und ein Ausführungszustand für jeden einzelnen Thread beibehalten, was eine gleiche Nebenläufigkeit zwischen allen Threads, innerhalb und zwischen Warps ermöglicht. Wenn der Ausführungszustand für jeden einzelnen Thread beibehalten wird, können Threads, welche die gleichen Anweisungen ausführen, konvergiert und für maximale Effizienz parallel ausgeführt werden. Der SM 440 wird ausführlicher nachstehend in Verbindung mit 10A beschrieben.
Die MMU 490 stellt eine Schnittstelle zwischen dem GPC 350 und der Partitions-Einheit 380 bereit. Die MMU 490 kann eine Übersetzung von virtuellen Adressen in physische Adressen, einen Speicherschutz und eine Arbitrierung von Speicheranfragen bereitstellen. In einer Ausführungsform stellt die MMU 490 einen oder mehrere Adressenübersetzungspuffer (TLBs = translation lookaside buffers) zum Durchführen der Übersetzung von virtuellen Adressen in physische Adressen in dem Speicher 304 bereit.
9B veranschaulicht eine Speicherpartitions-Einheit 380 der PPU 300 von 8 gemäß einer Ausführungsform. Wie in 9B gezeigt, umfasst die Partitions-Einheit 380 eine Raster-Operationen(ROP)-Einheit 450, einen L2-Cache-Speicher 460 und eine Speicherschnittstelle 470. Die Speicherschnittstelle 470 ist mit dem Speicher 304 gekoppelt. Die Speicherschnittstelle 470 kann 32-, 64-, 38-, 1024-Bit-Datenbusse oder dergleichen für Hochgeschwindigkeits-Datentransfer implementieren. In einer Ausführungsform umfasst die PPU 300 U Speicherschnittstellen 470, eine Speicherschnittstelle 470 pro Paar von Speicherpartitions-Einheiten 380, wobei jedes Paar von Speicherpartitions-Einheiten 380 mit einer entsprechenden Speichervorrichtung 304 verbunden ist. Beispielsweise kann die PPU 300 mit bis zu Y Speichervorrichtungen 304, wie beispielsweise Speicherstapel mit hoher Bandbreite oder Graphikdoppeldatenraten, Version 5 SDRAM oder andere Typen von persistenter Speicherung verbunden sein.
In einer Ausführungsform implementiert die Speicherschnittstelle 470 eine HBM2-Speicherschnittstelle und Y ist gleich einem halben U. In einer Ausführungsform sind die HBM2-Speicherstapel auf der gleichen physischen Packung wie die PPU 300 lokalisiert, die wesentliche Leistungs- und Flächeneinsparungen verglichen mit herkömmlichen GDDR5 SDRAM Systemen bereitstellt. In einer Ausführungsform umfasst jeder HBM2-Stapel vier Speicher-Dies und Y ist gleich 4, wobei der HBM2-Stapel zwei 38-Bit Kanäle pro Die für eine Gesamtzahl von 8 Kanälen und eine Datenbusbreite von 1024 Bit umfasst.
In einer Ausführungsform unterstützt der Speicher 304 Fehlerkorrekturcode (ECC) mit Einzelfehlerkorrektur und Doppelfehlerdetektion (SECDED), um Daten zu schützen. Der ECC stellt eine höhere Zuverlässigkeit für Rechenanwendungen bereit, die gegen Datenverfälschung empfindlich sind. Zuverlässigkeit ist besonders wichtig in großen Cluster-Rechenumgebungen, wo PPUs 300 sehr große Datensätze verarbeiten und/oder Anwendungen für längere Zeiträume ausführen.
In einer Ausführungsform implementiert die PPU 300 eine Mehrebenen-Speicherhierarchie. In einer Ausführungsform unterstützt die Speicherpartitions-Einheit 380 einen vereinheitlichten Speicher, um einen einzigen vereinheitlichten virtuellen Adressraum für den Speicher der CPU und den Speicher der PPU 300 bereitzustellen, der Datenteilung zwischen virtuellen Speichersystemen ermöglicht. In einer Ausführungsform wird die Häufigkeit von Zugriffen durch eine PPU 300 auf einen Speicher verfolgt, der auf anderen Prozessoren lokalisiert ist, um sicherzustellen, dass Speicherseiten in den physischen Speicher der PPU 300 bewegt werden, die häufiger auf die Seiten zugreift. In einer Ausführungsform unterstützt das NVLink 310 Adressenübersetzungsdienste, die der PPU 300 erlauben, auf Seitentabellen einer CPU direkt zuzugreifen und einen vollen Zugriff auf den CPU-Speicher durch die PPU 300 bereitstellen.
In einer Ausführungsform transferieren Kopiermaschinen Daten zwischen mehreren PPUs 300 oder zwischen PPUs 300 und CPUs. Die Kopiermaschinen können Seitenfehler für Adressen erzeugen, die nicht in den Seitentabellen abgebildet werden. Die Speicherpartitions-Einheit 380 kann dann die Seitenfehler bedienen, wobei die Adressen in der Seitentabelle abgebildet werden, nachdem die Kopiermaschine den Transfer durchführen kann. In einem herkömmlichen System ist der Speicher für mehrere Kopiermaschinenoperationen zwischen mehreren Prozessoren gesperrt (d.h. nicht auslagerbar), was den verfügbaren Speicher wesentlich verringert. Mit Hardware-Seiten-Faulting können Adressen an die Kopiermaschinen weitergeleitet werden, ohne sich Sorgen zu machen, ob die Speicherseiten resident sind und das Kopierverfahren transparent ist.
Daten von dem Speicher 304 oder einem anderen Systemspeicher können von der Speicherpartitions-Einheit 380 abgerufen und in dem L2-Cache-Speicher 460 gespeichert werden, der Auf-Chip lokalisiert ist und zwischen den verschiedenen GPCs 350 gemeinsam benutzt wird. Wie gezeigt, umfasst jede Speicherpartitions-Einheit 380 einen Bereich des L2-Cache-Speichers 460, der einer entsprechenden Speichervorrichtung 304 zugeordnet ist. Cache-Speicher niedrigerer Ebene können dann in verschiedenen Einheiten innerhalb der GPCs 350 implementiert werden. Beispielsweise kann jeder der SMs 440 einen LI-Cache-Speicher implementieren. Der L1-Cache-Speicher ist ein privater Speicher, der einem bestimmten SM 440 fest zugeordnet ist. Daten von dem L2-Cache-Speicher 460 können abgerufen und in jedem der L1-Cache-Speicher zur Verarbeitung in den Funktionseinheiten der SMs 440 gespeichert werden. Der L2-Cache-Speicher 460 ist mit der Speicherschnittstelle 470 und der XBar 370 gekoppelt.
Die ROP-Einheit 450 führt Graphik-Raster-Operationen durch, welche die Pixelfarbe betreffen, wie beispielsweise Farbenkomprimierung, Pixelmischung und dergleichen. Die ROP-Einheit 450 implementiert ebenfalls Tiefentesten in Verbindung mit der Raster-Engine 425, wobei eine Tiefe für einen Abtastort, der einem Pixelfragment zugeordnet ist, von der Aussonderungs-Engine der Raster-Engine 425 empfangen wird. Die Tiefe wird gegen eine entsprechende Tiefe in einem Tiefenpuffer für einen Abtastort geprüft, der dem Fragment zugeordnet ist. Wenn das Fragment den Tiefentest für den Abtastort besteht, dann aktualisiert die ROP-Einheit 450 den Tiefenpuffer und überträgt ein Ergebnis des Tiefentests an die Raster-Engine 425. Es wird anerkannt, dass sich die Anzahl von Speicherpartitions-Einheiten 380 von der Anzahl von GPCs 350 unterscheiden kann, und daher kann jede ROP-Einheit 450 mit jedem der GPCs 350 gekoppelt werden. Die ROP-Einheit 450 verfolgt Pakete, die von den unterschiedlichen GPCs 350 empfangen werden, und bestimmt, zu welchem GPC 350 ein durch die ROP-Einheit 450 erzeugtes Ergebnis durch die Xbar 470 weitergeleitet wird. Obwohl die ROP-Einheit 450 innerhalb der Speicherpartitions-Einheit 380 in 9B umfasst ist, kann die ROP-Einheit 450 außerhalb der Speicherpartitions- Einheit 380 sein. Beispielsweise kann die ROP-Einheit 450 in dem GPC 350 oder einer anderen Einheit liegen.
10A veranschaulicht den Streaming-Multiprozessor 440 von 9A gemäß einer Ausführungsform. Wie in 10A gezeigt, umfasst der SM 440 einen Anweisungs-Cache-Speicher 505, eine oder mehrere Planer-Einheiten 510, eine Registerdatei 520, einen oder mehrere Verarbeitungskerne 550, eine oder mehrere Spezialfunktionseinheiten (SFUs) 552, eine oder mehrere Lade/Speicher-Einheiten (LSUs) 554, ein Interconnect-Netzwerk 580 und einen gemeinsam genutzten Speicher/L1-Cache-Speicher 570.
Wie oben beschrieben, versendet die Arbeitsverteilungs-Einheit 325 Aufgaben zur Ausführung auf den GPCs 350 der PPU 300. Die Aufgaben werden einem bestimmten DPC 420 innerhalb eines GPC 350 zugeteilt, und wenn die Aufgabe einem Shader-Programm zugeordnet ist, kann die Aufgabe einem SM 440 zugeteilt werden. Die Planer-Einheit 510 empfängt die Aufgaben von der Arbeitsverteilungs-Einheit 325 und verwaltet die Anweisungs-Planung (instruction scheduling) für einen oder mehrere Thread-Blöcke, die dem SM 440 zugewiesen sind. Die Planer-Einheit 510 plant Thread-Blöcke zur Ausführung als Warps von parallelen Threads, wobei jeder Thread-Block mindestens einem Warp zugeordnet ist. In einer Ausführungsform führt jeder Warp 32 Threads aus. Die Planer-Einheit 510 kann eine Mehrzahl von unterschiedlichen Thread-Blöcken verwalten, welche die Warps den unterschiedlichen Thread-Blöcken zuordnet und dann Anweisungen von der Mehrzahl von unterschiedlichen kooperativen Gruppen an die verschiedenen Funktionseinheiten (d.h. Kerne 550, SFUs 552 und LSUs 554) während jedes Taktzyklus versendet.
Cooperative Groups ist ein Programmiermodell zum Organisieren von Gruppen von kommunizierenden Threads, die Entwicklern ermöglicht, die Granularität auszudrücken, bei der Threads kommunizieren, wobei der Ausdruck von reicheren, effizienten Parallelzerlegungen ermöglicht wird. Cooperative-Start-APIs unterstützen die Synchronisierung unter Thread-Blöcken für die Ausführung von parallelen Algorithmen. Herkömmliche Programmiermodelle stellen einen einzigen, einfachen Aufbau zum Synchronisieren von kooperierenden Threads bereit: eine Barriere über alle Threads eines Threadblocks (d.h. die Funktion syncthreads()). Programmierer würden jedoch häufig gerne Gruppen von Threads bei kleineren als Thread-Block-Granularitäten definieren und innerhalb der definierten Gruppen synchronisieren, um größere Leistung, Gestaltungsflexibilität und Software-Wiederverwendung in der Form von kollektiven gruppenweiten Funktionsschnittstellen zu ermöglichen.
Cooperative Groups ermöglicht Programmierern, Gruppen von Threads explizit bei Sub-Block(d.h. so klein wie ein einziger Thread)- und Mehr-Block-Granularitäten zu definieren und kollektive Operationen, wie beispielsweise Synchronisierung, an den Threads in einer kooperativen Gruppe durchzuführen. Das Programmiermodell unterstützt eine saubere Zusammensetzung über Softwaregrenzen, so dass Bibliotheken und Hilfsfunktionen innerhalb ihres lokalen Kontexts sicher synchronisieren können, ohne Annahmen über Konvergenz machen zu müssen. Cooperative-Groups-Primitive ermöglichen neue Muster von kooperativer Parallelität, die Erzeuger-Verbraucher Parallelität, opportunistische Parallelität und globale Synchronisierung über ein gesamtes Gitter von Threadblöcken umfassen.
Eine Versandeinheit 515 ist konfiguriert, um Anweisungen an eine oder mehrere der Funktionseinheiten zu übertragen. In der Ausführungsform umfasst die Planer-Einheit 510 zwei Versandeinheiten 515, die ermöglichen, dass zwei unterschiedliche Anweisungen von dem gleichen Warp während jedes Taktzyklus versandt werden. In alternativen Ausführungsformen kann jede Planer-Einheit 510 eine einzige Versandeinheit 515 oder zusätzliche Versandeinheiten 515 umfassen.
Jeder SM 440 umfasst eine Registerdatei 520, die einen Satz von Registern für die Funktionseinheiten des SM 440 bereitstellt. In einer Ausführungsform ist die Registerdatei 520 zwischen jeder der Funktionseinheiten aufgeteilt, so dass jeder Funktionseinheit ein zugehöriger Abschnitt der Registerdatei 520 zugeteilt ist. In einer anderen Ausführungsform ist die Registerdatei 520 zwischen den unterschiedlichen Warps aufgeteilt, die von dem SM 440 ausgeführt werden. Die Registerdatei 520 stellt temporäre Speicherung für Operanden bereit, die mit den Datenpfaden der Funktionseinheiten verbunden sind. 18. Jeder SM 440 umfasst L Verarbeitungskerne 550. In einer Ausführungsform umfasst der SM 440 eine große Anzahl (z.B., 128 usw.) von unterschiedlichen Verarbeitungskernen 550. Jeder Kern 550 kann eine vollständig in einer Pipeline angeordnete (fully-pipelined) Einfach-Präzisions- Verarbeitungseinheit umfassen, die eine Gleitkommaarithmetik-Logikeinheit und eine Ganzzahlarithmetik-Logikeinheit umfasst. In einer Ausführungsform implementieren die Gleitkommaarithmetik-Logikeinheiten den IEEE 754-3008 Standard für Gleitkommaarithmetik. In einer Ausführungsform umfassen die Kerne 550 64 Einfach-Präzisions-(32-Bit)-Gleitkommakerne, 64 Integer-Kerne, 32 Doppel-Präzisions-(64-Bit)-Gleitkommakerne und 8 Tensorkerne.
Tensorkerne sind konfiguriert, um Matrix-Operationen durchzuführen, und in einer Ausführungsform sind ein oder mehrere Tensorkerne in den Kernen 550 enthalten. Insbesondere sind die Tensorkerne konfiguriert, um tiefgehendes Lernen-Matrix-Arithmetik, wie beispielsweise Faltungsoperationen für neuronales Netzwerktraining und Inferenzieren, durchzuführen. In einer Ausführungsform arbeitet jeder Tensorkern an einer 4x4 Matrix und führt eine Matrix-Multiplikations- und Akkumulations-Operation D=A×B+C durch, wobei A, B, C und D 4x4 Matrizen sind.
In einer Ausführungsform sind die Matrix-Multiplikations-Eingaben A und B 16-Bit-Gleitkomma-Matrizen, während die Akkumulationsmatrizen C und D 16-Bit-Gleitkomma- oder 32-Bit-Gleitkomma-Matrizen sein können. Tensorkerne arbeiten an 16-Bit-Gleitkomma-Eingabedaten mit 32-Bit-Gleitkomma-Akkumulation. Die 16-Bit-Gleitkomma-Multiplikation erfordert 64 Operationen und ergibt ein Produkt voller Präzision, das dann unter Verwendung einer 32-Bit-Gleitkomma-Addition mit den anderen Zwischenprodukten für eine 4x4x4-Matrix-Multiplikation akkumuliert wird. In der Praxis werden Tensorkerne verwendet, um viel größere zweidimensionale oder höherdimensionale Matrixoperationen durchzuführen, die von diesen kleineren Elementen aufgebaut werden. Eine API, wie beispielsweise die CUDA 9 C++ API, stellt spezialisierte Matrix-Lade-, Matrix-Multiplikations- und Matrix-Akkumulations- und Matrix-Speicher-Operationen bereit, um Tensorkerne von einem CUDA-C++ Programm effizient zu verwenden. Bei der CUDA-Ebene nimmt das Warp-Schnittstellenniveau 16×16 große Matrizen an, die alle 32 Threads des Warp überspannen.
In einigen Ausführungsformen ist Transpositionshardware in den Verarbeitungskernen 550 oder einer anderen Funktionseinheit (e.g., SFUs 552 oder LSUs 554) enthalten und ist konfiguriert, um durch Diagonale gespeicherte Matrixdaten zu erzeugen und/oder die ursprüngliche Matrix und/oder die transponierten Matrix aus den durch Diagonale gespeicherten Matrixdaten zu erzeugen. Die Transpositionshardware kann innerhalb des gemeinsam genutzten Speichers 570 dem Ladepfad der Registerdatei 520 des SM 440 bereitgestellt werden.
In einem Beispiel können die durch Diagonale gespeicherten Matrixdaten von dem DRAM abgeholt und in dem gemeinsam genutzten Speicher 570 gespeichert werden. Wenn die Anweisung, die Verarbeitung unter Verwendung der durch Diagonale gespeicherten Matrixdaten durchzuführen, verarbeitet wird, kann die Transpositionshardware, die in dem Pfad des gemeinsam genutzten Speichers 570 angeordnet ist, und die Registerdatei 520, die ursprüngliche Matrix, die transponierte Matrix, die kompaktierten ursprüngliche Matrix und/oder die kompaktierten transponierte Matrix bereitstellen. Bis zur allerletzten Speicherung vor der Anweisung können die durch Diagonale gespeicherten einzelnen Matrixdaten beibehalten werden und die durch die Anweisung gekennzeichnete Matrixart wird nach Bedarf in der Registerdatei erzeugt.
Jeder SM 440 umfasst ebenfalls M SFUs 552, die Sonderfunktionen durchführen (z.B. Attributauswertung, reziproke Quadratwurzel und dergleichen). In einer Ausführungsform können die SFUs 552 eine Baumtraversierungseinheit umfassen, die konfiguriert ist, um eine hierarchische Baumdatenstruktur zu durchlaufen. In einer Ausführungsform können die SFUs 552 eine Textureinheit umfassen, die konfiguriert ist, um Texturkarten-Filteroperationen durchzuführen. In einer Ausführungsform sind die Textureinheiten konfiguriert, um Texturkarten (z.B. eine 2D-Anordnung von Texeln) von dem Speicher 304 zu laden und die Texturkarten abzutasten, um abgetastete Texturwerte zum Gebrauch in Shader-Programmen zu erzeugen, die durch den SM 440 ausgeführt werden. In einer Ausführungsform werden die Texturkarten in dem gemeinsam genutzten Speicher/L1-Cache-Speicher 470 gespeichert. Die Textureinheiten implementieren Texturoperationen, wie beispielsweise Filteroperationen, unter Verwendung von Mip-Maps (d.h. Texturkarten von veränderlichem Detaillierungsgrad). In einer Ausführungsform umfasst jeder SM 340 zwei Textureinheiten.
Jeder SM 440 umfasst ebenfalls N LSUs 554, die Lade- und Speicheroperationen zwischen dem gemeinsam genutzten Speicher/L1-Cache-Speicher 570 und der Registerdatei 520 implementieren. Jeder SM 440 umfasst ein Interconnect-Netzwerk 580, das jede der Funktionseinheiten mit der Registerdatei 520 und die LSU 554 mit der Registerdatei 520, dem gemeinsam genutzten Speicher/L1-Cache-Speicher 570 verbindet. In einer Ausführungsform ist das Interconnect-Netzwerk 580 eine Kreuzschiene, die konfiguriert sein kann, um irgendeine der Funktionseinheiten mit irgendeinem der Register in der Registerdatei 520 zu verbinden und die LSUs 554 mit der Registerdatei und Speicherorten im gemeinsam genutzten Speicher/L1-Cache-Speicher 570 zu verbinden.
Der gemeinsam genutzte Speicher/L1 -Cache-Speicher 570 ist eine Auf-Chip-Speicheranordnung, die Datenspeicherung und Kommunikation zwischen dem SM 440 und der Primitiven-Engine 435 und zwischen Threads in dem SM 440 ermöglicht. In einer Ausführungsform umfasst der gemeinsam genutzte Speicher/L1-Cache-Speicher 570 38KB von Speicherkapazität und ist in dem Pfad von dem SM 440 zu der Partitions-Einheit 380. Der gemeinsam genutzte Speicher/L1 -Cache-Speicher 570 kann verwendet werden, um Lese- und Schreibvorgänge zwischenzuspeichern. Einer oder mehrere des gemeinsam genutzten Speicher/L1-Cache-Speichers 570, L2-Cache-Speichers 460 und des Speichers 304 sind Hintergrundspeicher.
Das Kombinieren von Daten-Cache und gemeinsam genutzter Speicherfunktionalität in einem einzigen Speicherblock stellt die beste Gesamtleistung für beide Arten von Speicherzugriffen bereit. Die Kapazität ist als ein Cache-Speicher von Programmen benutzbar, die den gemeinsam genutzten Speicher nicht verwenden. Wenn der gemeinsam genutzte Speicher beispielsweise konfiguriert ist, um die Hälfte der Kapazität zu verwenden, können Textur- und Lade/Speicher-Operationen die verbleibende Kapazität verwenden. Integration innerhalb des gemeinsam genutzten Speichers/L1-Cache-Speicher 570 ermöglicht dem gemeinsam genutzten Speicher/L1-Cache-Speicher 570 als eine Leitung mit hohem Durchsatz zum Streamen von Daten zu arbeiten, während gleichzeitig eine höhere Bandbreite und ein latenzarmer Zugriff auf häufig wiederverwendete Daten bereitgestellt werden.
Wenn für Allzweck-Parallelberechnung konfiguriert wird, kann im Vergleich mit der Graphikverarbeitung eine einfachere Konfiguration verwendet werden. Im Einzelnen werden die in 3 gezeigten Festfunktions-Graphikverarbeitungseinheiten umgangen, wobei ein viel einfacheres Programmiermodell erzeugt wird. In der Allzweck-Parallelberechnungs-Konfiguration werden Blöcke von Threads von der Arbeitsverteilungs-Einheit 325 direkt den DPCs 420 zugewiesen und verteilt. Die Threads in einem Block führen das gleiche Programm unter Verwendung einer eindeutigen Thread-ID in der Berechnung, um sicherzustellen, dass jeder Thread eindeutige Ergebnisse erzeugt, unter Verwendung des SM 440, um das Programm auszuführen und Berechnungen durchzuführen, eines gemeinsam genutzten Speicher/L1-Cache-Speichers 570, um zwischen Threads zu kommunizieren, und der LSU 554 aus, um globalen Speicher durch den gemeinsam genutzten Speicher/L1-Cache-Speicher 570 und die Speicherpartitions-Einheit 380 zu lesen und zu beschreiben. Wenn für Allzweck-Parallelberechnung konfiguriert, kann der SM 440 ebenfalls Befehle schreiben, welche die Planer-Einheit 320 verwenden kann, um neue Arbeit auf den DPCs 420 zu starten.
Die PPU 300 kann in einem Tischcomputer, einem Laptop-Computer, einem Tablet-Computer, einem Smartphone (z.B. einer drahtlosen handgehaltenen Vorrichtung), einem persönlichen digitalen Assistenten (PDA), einer Digitalkamera, einer handgehaltenen elektronischen Vorrichtung und dergleichen enthalten sein. In einer Ausführungsform ist die PPU 300 auf einem einzelnen Halbleitersubstrat verkörpert. In einer anderen Ausführungsform ist die PPU 300 in einem System-auf-Chip (SoC) zusammen mit einer oder mehreren anderen Vorrichtungen, wie beispielsweise zusätzlichen PPUs 300, dem Speicher 304, einem Rechner-mit-reduziertem-Befehlssatz(RISC)-CPU, einer Speicherverwaltungseinheit (MMU), einem Digital/Analog-Wandler (DAC) und dergleichen
In einer Ausführungsform kann die PPU 300 auf einer Graphikkarte enthalten sein, die eine oder mehrere Speichervorrichtungen 304 umfasst. Die Graphikkarte kann konfiguriert sein, um sich mit einem PCIe-Schlitz auf einer Hauptplatine eines Tischcomputers schnittstellenmäßig zu verbinden. In noch einer anderen Ausführungsform kann die PPU 300 eine integrierte Graphikverarbeitungseinheit (iGPU) oder ein Parallelprozessor sein, der in dem Chipsatz der Hauptplatine umfasst ist.
Beispielhaftes Rechensystem
Systeme mit mehrere GPUs und CPUs werden in einer Vielfalt von Industrien verwendet, da Entwickler mehr Parallelität in Anwendungen, wie beispielsweise Rechnen für künstliches Intelligenz, freilegen und wirksam einsetzen. Hochleistungs-GPU-beschleunigte Systeme mit zehn bis vielen Tausenden von Rechenknoten werden in Datenzentren, Forschungsanlagen und Supercomputern eingesetzt, um immer größere Probleme zu lösen. Da die Anzahl von Verarbeitungsvorrichtungen innerhalb der Hochleistungssysteme zunimmt, müssen die Kommunikations- und Datentransfermechanismen angepasst werden, um die erhöhte Bandbreite zu unterstützen.
10B ist ein Konzeptdiagramm eines Verarbeitungssystems 500, das unter Verwendung der PPU 300 von 8 implementiert wird, gemäß einer Ausführungsform. Das beispielhafte System 500 kann konfiguriert sein, um die in dieser Anmeldung offenbarten Verfahren zu implementieren (z.B. in 5, 6 oder 8 gezeigte Verfahren). Das Verarbeitungssystem 500 umfasst eine CPU 530, einen Schalter 555 und jeweils mehrere PPUs 300 und jeweilige Speicher 304. Das NVLink 310 stellt Hochgeschwindigkeits-Kommunikationslinks zwischen jeder der PPUs 300 bereit. Obwohl eine bestimmte Anzahl von Verbindungen des NVLink 310 und des Interconnect 302 in 10B veranschaulicht werden, kann die Anzahl von Verbindungen zu jeder PPU und der CPU 530 variieren. Der Schalter 555 verbindet sich schnittstellenmäßig zwischen dem Interconnect 302 und der CPU 530. Die PPUs 300, die Speicher 304 und die NVLinks 310 können auf einer einzigen Halbleiterplattform gelegen sein, um ein Parallelverarbeitungsmodul 525 zu bilden. In einer Ausführungsform unterstützt der Schalter 555 zwei oder mehrere Protokolle, um sich schnittstellenmäßig zwischen verschiedenen unterschiedlichen Verbindungen und/oder Links zu verbinden.
In einer anderen Ausführungsform (nicht gezeigt) stellt das NVLink 310 ein oder mehrere Hochgeschwindigkeits-Kommunikationslinks zwischen jeder der PPUs 300 und der CPU 530 bereit und der Schalter 555 ist schnittstellenmäßig zwischen dem Interconnect 302 und jeder der PPUs 300 verbunden. Die PPUs 300, die Speicher 304 und der Interconnect 302 können auf einer einzigen Halbleiterplattform situiert sein, um ein Parallelverarbeitungsmodul 525 zu bilden. In noch einer anderen Ausführungsform (nicht gezeigt) stellt der Interconnect 302 ein oder mehrere Kommunikationslinks zwischen jeder der PPUs 300 und der CPU 530 bereit und der Schalter 555 ist schnittstellenmäßig zwischen jeder der PPUs 300 unter Verwendung des NVLink 310 verbunden, um ein oder mehrere Hochgeschwindigkeits-Kommunikationslinks zwischen den PPUs 300 bereitzustellen. In einer anderen Ausführungsform (nicht gezeigt) stellt das NVLink 310 ein oder mehrere Hochgeschwindigkeits-Kommunikationslinks zwischen den PPUs 300 und der CPU 530 durch den Schalter 555 bereit. In noch einer anderen Ausführungsform (nicht gezeigt) stellt der Interconnect 302 ein oder mehrere Hochgeschwindigkeits-Kommunikationslinks zwischen jeder der PPUs 300 direkt bereit. Ein oder mehrere der NVLink 310 Hochgeschwindigkeits-Kommunikationslinks können als ein physischer NVLink-Interconnect oder entweder als ein Auf-Chip- oder Auf-Die-Interconnect unter Verwendung des gleichen Protokolls wie das NVLink 310 implementiert werden.
Im Kontext der vorliegenden Beschreibung kann sich eine einzige Halbleiterplattform auf eine einzelne unitäre Halbleiter-basierte integrierte Schaltung beziehen, die auf einem Die oder Chip angefertigt ist. Es sei bemerkt, dass sich der Begriff einzige Halbleiterplattform ebenfalls auf Mehr-Chip-Module mit erhöhter Konnektivität beziehen kann, die eine Auf-Chip-Operation simulieren und die wesentlichen Verbesserungen gegenüber der Benutzung einer herkömmlichen Busimplementierung vornehmen. Natürlich können die verschiedenen Schaltungen oder Vorrichtungen ebenfalls getrennt oder in verschiedenen Kombinationen von Halbleiterplattformen gemäß den Wünschen des Benutzers situiert sein. Alternativ kann das Parallelverarbeitungsmodul 525 als ein Leiterplattensubstrat implementiert sein und jede der PPUs 300 und/oder Speicher 304 können verpackte Vorrichtungen sein. In einer Ausführungsform sind die CPU 530, der Schalter 555 und das Parallelverarbeitungsmodul 525 auf einer einzigen Halbleiterplattform situiert.
In einer Ausführungsform ist die Signalrate von jedem NVLink 310 gleich 20 bis 25 Gigabit/s und jede PPU 300 umfasst sechs NVLink 310 Schnittstellen (wie in 10B gezeigt, fünf NVLink 310 Schnittstellen sind für jede PPU 300 umfasst). Jedes NVLink 310 stellt eine Datentransferrate von 25 Gigabyte/s in jeder Richtung bereit, wobei sechs Verknüpfungen 300 Gigabyte/s bereitstellen. Die NVLinks 310 können exklusiv für PPU-zu-PPU Kommunikation, wie in 10B gezeigt, oder einer Kombination von PPU-zu-PPU und PPU-zu-CPU verwendet werden, wenn die CPU 530 ebenfalls eine oder mehrere NVLink 310 Schnittstellen umfasst.
In einer Ausführungsform ermöglicht das NVLink 310 einen direkten Lade/Speicher/atomaren Zugriff von der CPU 530 auf jeden Speicher 304 der PPU 300. In einer Ausführungsform unterstützt das NVLink 310 Kohärenzoperationen, die ermöglichen, das von dem Speicher 304 gelesene Daten in der Cache-Hierarchie der CPU 530 gespeichert werden können, was die Cachezugriffslatenz für die CPU 530 verringert. In einer Ausführungsform umfasst das NVLink 310 Unterstützung für Adressübersetzungsdienste (Address Translation Services; ATS), was der PPU 300 ermöglicht, auf Seitentabellen innerhalb der CPU 530 direkt zuzugreifen. Eines oder mehrere der NVLinks 310 können ebenfalls konfiguriert sein, um in einem Niedrigleistungsmodus zu arbeiten.
10C veranschaulicht ein beispielhaftes System 565, bei dem die verschiedene Architektur und/oder Funktionalität der verschiedenen vorherigen Ausführungsformen implementiert werden kann. Das beispielhafte System 565 kann konfiguriert sein, um die in dieser Anmeldung offenbarten Verfahren zu implementieren (z.B. in 5, 6,oder 8 gezeigte Verfahren).
Wie gezeigt, wird ein System 565 bereitgestellt, das mindestens eine zentrale Verarbeitungseinheit 530 umfasst, die mit einem Kommunikationsbus 575 verbunden ist. Der Kommunikationsbus 575 kann unter Verwendung jedes geeigneten Protokolls, wie beispielsweise PCI (Peripheral Component Interconnect), PCI-Express, AGP (Accelerated Graphics Port), Hyper-Transport oder irgendeinem anderen Bus- oder Punkt-zu-Punkt-Kommunikationsprotokoll(en), implementiert sein. Das System 565 umfasst ebenfalls einen Hauptspeicher 540. Steuerlogik (Software) und Daten werden in dem Hauptspeicher 540 gespeichert, der die Form eines Direktzugriffspeichers (RAM) annehmen kann.
Das System 565 umfasst ebenfalls Eingabevorrichtungen 560, das Parallelverarbeitungssystem 525 und Anzeigevorrichtungen 545, z.B. eine herkömmliche CRT (Kathodenstrahlröhre), eine LCD (Flüssigkristallanzeige), eine LED (lichtemittierende Diode), eine Plasmaanzeige oder dergleichen. Eine Benutzereingabe kann von den Eingabevorrichtungen 560, z.B. Tastatur, Maus, Berührfeld, Mikrophon und dergleichen empfangen werden. Jedes der vorhergehenden Module und/oder Vorrichtungen kann sogar auf einer einzigen Halbleiterplattform situiert sein, um das System 565 zu bilden. Alternativ können die verschiedenen Module auch getrennt oder in verschiedenen Kombinationen von Halbleiterplattformen gemäß den Wünschen des Benutzers situiert sein.
Ferner kann das System 565 mit einem Netzwerk (z.B. einem Telekommunikationsnetzwerk, Lokalbereichsnetzwerk (LAN), drahtlosen Netzwerk, Weitbereichsnetzwerk (WAN), wie beispielsweise dem Internet, Peer-zu-Peer-Netzwerk, Kabelnetzwerk oder dergleichen) durch eine Netzwerkschnittstelle 535 für Kommunikationszwecke gekoppelt sein.
Das System 565 kann ebenfalls einen Sekundärspeicher umfassen (nicht gezeigt). Der Sekundärspeicher umfasst beispielsweise ein Festplattenlaufwerk und/oder ein entfernbares Speicherlaufwerk, das ein Diskettenlaufwerk darstellt, ein Magnetbandlaufwerk, ein Kompaktdisklaufwerk, ein DVD(digitale versatile disk)-Laufwerk, eine Aufzeichnungsvorrichtung und einen Universal-Serial-Bus-(USB)-Flash-Speicher. Das entfernbare Speicherlaufwerk liest von und/oder schreibt auf eine entfernbare Speichereinheit auf eine wohlbekannte Art und Weise.
Computerprogramme oder Computersteuerlogik-Algorithmen können in dem Hauptspeicher 540 und/oder dem Sekundärspeicher gespeichert sein. Derartige Computerprogramme, wenn ausgeführt, ermöglichen dem System 565, verschiedene Funktionen durchzuführen. Der Speicher 540, die Speicherung, und/oder jede andere Speicherung sind mögliche Beispiele von computerlesbaren Medien.
Die Architektur und/oder Funktionalität der verschiedener vorherigen Figuren kann im Kontext eines allgemeinen Computersystems, eines Schaltungsplatinensystems, eines Unterhaltungszwecken fest zugeordneten Spielkonsolensystems, eines anwendungsspezifischen Systems und/oder jedem anderen gewünschten System implementiert sein. Beispielsweise kann das System 565 die Form eines Tischcomputers, eines Laptop-Computers, eines Tablet-Computers, von Servern, von Supercomputern, eines Smartphones (z.B. einer drahtlosen handgehaltenen Vorrichtung), eines persönlichen digitalen Assistenten (PDA), einer Digitalkamera, eines Fahrzeugs, einer am Kopf angebrachten Anzeige, einer handgehaltenen elektronischen Vorrichtung, eines mobilen Telefongeräts, eines Fernsehers, einer Arbeitsstation, einer Spielkonsole, eines eingebetteten Systems und/oder von jedem anderen Typ von Logik annehmen.
Alle Patente und gedruckte Veröffentlichungen, auf die vorstehend Bezug genommen wurde, werden durch Verweis hier einbezogen, als ob sie ausdrücklich dargelegt wären.
Während die Erfindung in Verbindung mit derzeit als am praktischsten und bevorzugtesten angesehenen Ausführungsformen beschrieben wurde, versteht sich, dass die Erfindung nicht auf die offenbarten Ausführungsformen zu beschränken ist, sondern im Gegenteil darauf abzielt, verschiedene Modifikationen und äquivalente Anordnungen abzudecken, die im Rahmen und Umfang der beigefügten Ansprüche enthalten sind.
ZITATE ENTHALTEN IN DER BESCHREIBUNG
Diese Liste der vom Anmelder aufgeführten Dokumente wurde automatisiert erzeugt und ist ausschließlich zur besseren Information des Lesers aufgenommen. Die Liste ist nicht Bestandteil der deutschen Patent- bzw. Gebrauchsmusteranmeldung. Das DPMA übernimmt keinerlei Haftung für etwaige Fehler oder Auslassungen.
Zitierte Patentliteratur

US 8555035 [0034]
US 7634621 [0034]
US 2018/0322078 [0037]
US 62/927511 [0092]
US 16/712236 [0092]

Claims

Verfahren, das von einem Verarbeitungssystem durchgeführt wird, das einen Multithread-Prozessor und einen chipinternen Speicher umfasst, der einen Cache-Speicher und einen durch Software-verwalteten gemeinsam genutzten Speicher umfasst, wobei das Verfahren umfasst: nebenläufiges Ausführen mehrerer Threads, die einen Thread umfassen, der eine fusionierte Lade- und Speicheranweisung ausführt, um Daten zu laden, die in einem Speicher extern zu dem Verarbeitungssystem gespeichert sind, und die Daten in dem gemeinsam genutzten Speicher zu speichern, wobei die fusionierte Lade- und Speicheranweisung konfigurierbar ist zwischen Umgehen (1) von Prozessorregistern, die dem Thread zugeordnet sind, der eine fusionierten Lade- und Speicheranweisung ausführt, oder (2) den Prozessorregistern und dem Cache-Speicher, wobei die Anweisung das System anleitet, um: die Daten aus dem externen Speicher abzurufen; und die abgerufenen Daten in dem gemeinsam genutzten Speicher zu speichern, ohne zuerst die abgerufenen Daten in den Prozessorregistern zu speichern oder ohne zuerst die empfangenen Daten in Prozessorregistern und dem Cache-Speicher zu speichern.
Verfahren gemäß Anspruch 1, wobei mehrere Threads die fusionierten Lade- und Speicheranweisungen ausführen, um in dem externen Speicher gespeicherte Daten in den gemeinsam genutzter Speicher zu laden, und die Anweisungen das System anleiten, gemeinsam genutzten Speicherzieladressen für Daten zu bestimmen, die von den Anweisungen angefordert wurden, ein Muster zwischen Füllsektoren der abgerufenen Daten und der Zieladressen des gemeinsam genutzten Speichers zu bestimmen, und die Füllsektoren der abgerufenen Daten in gemeinsam genutzten Speichersektoren zu speichern, die durch das bestimmte Muster identifiziert werden.
Verfahren gemäß Anspruch 1 oder 2, wobei der gemeinsam genutzte Speicher mehrere in Bänken angeordnete Sektoren umfasst, und Daten, die in mindestens einem der Sektoren gespeichert sind, auf ein Intra-Sektor-Umordnung (swizzle) angewandt werden.
Verfahren gemäß einem der vorangehenden Ansprüche, wobei die fusionierten Lade- und Speicheranweisungen eine einzelne Anweisung ist, die eine Zieladresse des gemeinsam genutzten Speichers und eine globale Quelladresse kennzeichnet.
Verfahren gemäß einem der vorangehenden Ansprüche, wobei, wenn (1) die fusionierte Lade- und Speicheranweisung konfiguriert ist, um die Prozessorregister und den Cache-Speicher zu umgehen, und (2) die von der Anweisung angeforderten Daten in dem Cache-Speicher gespeichert sind, die fusionierte Lade- und Speicheranweisung das System anleitet, um: die in dem Cache-Speicher angeforderten Daten in dem gemeinsam genutzten Speicher zu speichern und den Tag in dem Cache-Speicher für die angeforderten Daten ungültig zu machen.
Verfahren gemäß einem der vorangehenden Ansprüche, wobei die fusionierte Lade- und Speicheranweisung während der Laufzeit dynamisch konfigurierbar ist.
Verarbeitungssystem, umfassend: einen Multithread-Prozessor, der konfiguriert ist, um mehrere Threads nebenläufig auszuführen; mehrere Datenregister, wobei jedes der Datenregister zugewiesen ist, um Threads auszuführen; und einen chipinternen Speicher, der einen Cache-Speicher und einen gemeinsam genutzten Speicher umfasst, wobei der Cache-Speicher konfiguriert ist, um den mehreren ausführenden Threads zu ermöglichen, auf in dem Cache-Speicher gespeicherte gekennzeichneten Daten zuzugreifen, und der gemeinsam genutzten Speicher konfiguriert ist, um den mehreren ausführenden Threads zu ermöglichen, auf nicht gekennzeichnete Daten zuzugreifen, die in dem gemeinsam genutzten Speicher gespeichert sind, wobei das System konfiguriert ist, um als Reaktion auf mindestens einen der Threads, der eine Anweisung ausführt, um in einem Speicher extern zu dem Verarbeitungssystem gespeicherte Daten zu laden und die Daten in dem gemeinsam genutzten Speicher selektiv zu speichern, ohne zuerst die Daten in den mehreren Datenregistern zu speichern, oder ohne zuerst die Daten in den mehreren Datenregistern und dem Cache-Speicher zu speichern, die Daten aus dem externen Speicher abzurufen und die abgerufenen Daten in dem gemeinsam genutzten Speicher (1) selektiv zu speichern, ohne zuerst die abgerufenen Daten in den mehreren Datenregistern zu speichern oder (2) ohne zuerst die abgerufenen Daten in den mehreren Datenregistern und dem Cache-Speicher zu speichern.
Verarbeitungssystem gemäß Anspruch 7, wobei die abgerufenen Daten in dem gemeinsam genutzten Speicher (1) gespeichert werden, ohne zuerst die abgerufenen Daten in den mehreren Datenregistern zu speichern, oder (2) ohne zuerst die abgerufenen Daten in den mehreren Datenregistern und dem Cache-Speicher zu speichern, basierend auf einer dynamischen Auswahl während der Laufzeit.
Verarbeitungssystem gemäß Anspruch 7 oder 8, wobei der Cache-Speicher und der gemeinsam genutzte Speicher ein vereinigter physikalischer Direktzugriff-Speicher ist.
Verarbeitungssystem gemäß Anspruch 9, wobei der vereinigte physikalische Direktzugriff-Speicher eine Registerdatei umfasst, welche die mehreren Datenregister umfasst, die den ausführenden Threads dynamisch zugewiesen werden können.
Verarbeitungssystem gemäß einem der Ansprüche 7 bis 10, wobei das Verarbeitungssystem Hardwareschaltungen umfasst, die konfiguriert sind, um: mehrere Anweisungen zu empfangen, um in dem externen Speicher gespeicherte Daten zu laden und die Daten in dem gemeinsam genutzten Speicher zu speichern, und basierend auf den mehreren Anweisungen ein Sektorfüllmuster des gemeinsam genutzten Speichers zum Speichern der abgerufenen Daten in dem gemeinsam genutzten Speicher zu bestimmen.
Verarbeitungssystem gemäß einem der Ansprüche 7 bis 11, wobei der gemeinsam genutzten Speicher mehrere in Bänken angeordnete Sektoren umfasst und die abgerufenen Daten in mindestens einem der Sektoren mit einer Intra-Sektor-Umordnung gespeichert werden.
Verarbeitungssystem gemäß einem der Ansprüche 7 bis 12, wobei der gemeinsam genutzte Speicher mehrere in Bänken angeordnete Sektoren umfasst und Abschnitte der abgerufenen Daten in den Sektoren in unterschiedlichen Bänken gespeichert werden.
Verarbeitungssystem gemäß einem der Ansprüche 7 bis 13, wobei die Anweisung Daten aus einem Sektor des externen Speichers lädt und die geladenen Daten in Teilen der mehreren Sektoren in dem gemeinsam genutzten Speicher speichert.
Verarbeitungssystem gemäß einem der Ansprüche 7 bis 14, wobei der gemeinsam genutzte Speicher ein Software-verwalteter Cache ist.
Verarbeitungssystem gemäß einem der Ansprüche 7 bis 15, wobei die abgerufenen Daten in einer oder mehreren Cache-Zeilen des Cache-Speichers gespeichert werden und das System ferner eine Interconnect-Schaltung umfasst, die konfiguriert ist, um Daten, die in der einen oder mehreren Cache-Zeilen gespeichert sind, in Sektoren des gemeinsam genutzten Speichers zu transferieren, ohne zuerst die abgerufenen Daten in den mehreren Datenregistern zu speichern.
Verfahren, das von einem Verarbeitungssystem durchgeführt wird, umfassend: einen Multithread-Prozessor, der konfiguriert ist, um mehrere Threads nebenläufig auszuführen; mehrere Register, die den ausführenden Threads zugewiesen sind; einen Hardware-verwalteten Cache-Speicher, der konfiguriert ist, um den mehreren ausführenden Threads zu ermöglichen, auf in dem Cache-Speicher gespeicherte gekennzeichnete Daten zuzugreifen; und einen Software-verwalteten gemeinsam genutzten Speicher, der konfiguriert ist, um den mehreren ausführenden Threads zu ermöglichen, auf in dem gemeinsam genutzten Speicher gespeicherte nicht gekennzeichnete Daten zuzugreifen, wobei das Verfahren umfasst: nebenläufiges Ausführen mehrerer Threads, wobei mindestens einer der Threads eine Anweisung ausführt, um Daten zu laden, die in einem Speicher extern zu dem Verarbeitungssystem gespeichert sind, und die Daten in dem gemeinsam genutzten Speicher zu speichern, wobei die Anweisung konfigurierbar ist, um (1) die Daten in dem gemeinsam genutzten Speicher zu speichern, ohne die Daten zuerst in Registern zu speichern, die den ausführenden Threads zugeordnet sind, oder (2) die Daten in dem gemeinsam genutzten Speicher zu speichern, ohne zuerst die Daten in den Registern und dem Cache-Speicher zu speichern; als Reaktion auf die Anweisungen, die konfiguriert sind, um die Daten in dem gemeinsam genutzten Speicher zu speichern, ohne die Daten zuerst in den Registern zu speichern, Abrufen der Daten aus dem externen Speicher und Speichern der abgerufenen Daten in dem gemeinsam genutzten Speicher, ohne zuerst die abgerufenen Daten in den Registern zu speichern; und als Reaktion auf die Anweisungen, die konfiguriert sind, die Daten in dem gemeinsam genutzten Speicher zu speichern, ohne die Daten zuerst in den Registern und dem Cache-Speicher zu speichern, Abrufen der Daten aus dem externen Speicher und Speichern der abgerufenen Daten in dem gemeinsam genutzten Speicher, ohne zuerst die abgerufenen Daten in den Registern und dem gemeinsam genutzten Speicher zu speichern.
Verfahren gemäß Anspruch 17, ferner umfassend zwei oder mehrere der ausführenden Threads, die Anweisungen ausführen, um in dem gemeinsam genutzten Speicher gespeicherte Daten in Register zu laden, die den beiden oder mehreren ausführenden Threads zugewiesen sind, und eine Berechnung unter Verwendung der in den Registern gespeicherten Daten durchzuführen.
Verfahren gemäß Anspruch 18, wobei die Berechnung eine Matrix-Multiplikations- und Akkumulations-Operation ist.
Verfahren gemäß einem der Ansprüche 17 bis 19, wobei die Anweisung, in dem Speicher extern zu dem Verarbeitungssystem gespeicherte Daten zu laden und die Daten in dem gemeinsam genutzten Speicher zu speichern, eine einzelne Anweisung ist, die ein Zieladresse des gemeinsam genutzten Speichers und eine globale Quelladresse umfasst.
Verfahren gemäß einem der Ansprüche 17 bis 20, wobei der gemeinsam genutzte Speicher mehrere in Bänken angeordnete Sektoren umfasst und die abgerufenen Daten in mindestens einem der Sektoren gespeichert und eine Intra-Sektor-Umordnung angewandt wird.
Verfahren gemäß einem der Ansprüche 17 bis 21, wobei der gemeinsam genutzte Speicher mehrere in Bänken angeordnete Sektoren umfasst und die abgerufenen Daten in den Sektoren in unterschiedlichen Bänken gespeichert werden.
GPU-Ausführungsverfahren, umfassend: Decodieren eines fusionierten Lade/Speicher-Anweisungsformats, das spezifiziert, ob die GPU (a) ein Prozessorregister umgehen oder (b) sowohl das Prozessorregister als auch einen Cache-Speicher umgehen soll; und als Reaktion auf das Decodieren, Abrufen eines Datenworts aus einem ersten Speicher und Speichern des abgerufenen Datenworts in einem zweiten Speicher ohne zuerst die abgerufenen Daten in dem Prozessorregister zu speichern oder ohne zuerst die abgerufenen Daten in dem Prozessorregister und dem Cache-Speicher zu speichern, wie durch die decodierte fusionierte Lade/Speicher-Anweisungsformatspezifikation ausgewählt.
GPU-Ausführungsverfahren gemäß Anspruch 23, ferner umfassend eine gemeinsame Nutzung des zweiten Speichers zwischen mehreren Threads.
GPU-Ausführungsverfahren gemäß Anspruch 23 oder 24, wobei der Cache-Speicher und der zweite Speicher in einem gemeinsamen chipinternen physikalischen Speicheradressenraum angeordnet sind.
GPU-Ausführungsverfahren gemäß einem der Ansprüche 23 bis 25, wobei das fusionierte Lade/Speicher-Anweisungsformat umfasst: (i) eine Quellwortadresse, (ii) eine Zielwortadresse und (iii) mindestens ein Bit, das spezifiziert, ob die GPU (a) ein Prozessorregister umgehen oder (b) sowohl das Prozessorregister als auch einen Cache-Speicher umgehen soll, um ein Datenwort aus der Quellwortadresse abzurufen und das abgerufene Datenwort in die Zielwortadresse speichern, und das Decodieren umfasst das Decodieren des mindestens einen Bits.
GPU-Ausführungsverfahren, umfassend: Decodieren eines Lade/Speicher-Speicherzugriffsanweisungsformats; als Reaktion auf das Decodieren, Initiieren eines Speicherzugriffs auf den gemeinsam genutzten Speicher der GPU, der ein selektives Umgehen der Zwischenspeicherung in mindestens eines von einem Prozessorregister und einem Cache-Speicher umfasst; und Verfolgungsbeendigung des Speicherzugriffs als eine asynchrone Kopier-/ Direktspeicherzugriffsoperation.
Verfahren gemäß Anspruch 27, ferner umfassend Schreiben in den gemeinsam genutzten Speicher während dynamisch ausgewählt wird, ob jede Zwischenspeicherung zur Laufzeit zu umgehen ist.
Verfahren gemäß Anspruch 27 oder 28, wobei der Speicherzugriff Lesen eines Sektors des globalen Speichers und Schreiben desselben in Teilen von mehreren Sektoren in dem gemeinsam genutzten Speicher umfasst.