DE102013208558A1

DE102013208558A1 - Verfahren und System zur Verarbeitung verschachtelter Stream-Events

Info

Publication number: DE102013208558A1
Application number: DE102013208558A
Authority: DE
Inventors: Luke Durant
Original assignee: Nvidia Corp
Current assignee: Nvidia Corp
Priority date: 2012-05-09
Filing date: 2013-05-08
Publication date: 2013-11-14
Also published as: US9928109B2; CN103559014B; TW201413456A; CN103559014A; US20130305258A1

Abstract

Ein Ausführungsbeispiel der vorliegenden Erfindung belegt eine Technik für das Erzwingen von Stream-übergreifenden Abhängigkeiten in einem Parallelverarbeitungssubsystem wie etwa einer Grafikverarbeitungseinheit (GPU). Zu der Technik gehört, Warte-Events in eine Warteschlange zu stellen, um Stream-übergreifende Abhängigkeiten zu erzeugen, sowie Signalisierungs-Events, um den Warte-Events den Abschluss anzuzeigen. Ein Scheduler-Kernel untersucht eine Datenstruktur, die aus einem korrespondierenden Stream stammt, für den Status einer Aufgabe und aktualisiert Abhängigkeitszähler für Aufgaben und Events innerhalb des Streams. Sobald jede Aufgabenabhängigkeit für einen Warte-Event erfüllt worden ist, darf eine damit verknüpfte Aufgabe mit ihrer Ausführung beginnen.

Description

Hintergrund der Erfindung
Technisches Anwendungsgebiet
Die vorliegende Erfindung bezieht sich allgemein auf Computerarchitekturen mit Multi-Threading und im Besonderen auf ein Verfahren und System zur Verarbeitung verschachtelter Stream-Events.
Beschreibung der betreffenden Anwendungstechnik
In herkömmlichen Computersystemen, die sowohl über eine Zentralverarbeitungseinheit (CPU) und eine Grafikverarbeitungseinheit (GPU) verfügen, bestimmt die CPU, welche bestimmten Verarbeitungsaufgaben von der GPU ausgeführt werden und in welcher Reihenfolge. Eine Verarbeitungsaufgabe der GPU umfasst stark parallelisierte, sehr ähnliche Operationen über eine parallele Datenmenge hinweg, wie etwa eine Grafik oder eine Reihe von Grafiken. In einem herkömmlichen GPU-Ausführungsmodell stößt die CPU eine bestimmte Verarbeitungsaufgabe an, indem sie ein entsprechendes Thread-Programm auswählt und die GPU anweist, eine Reihe paralleler Instanzen dieses Thread-Programms auszuführen. Im konventionellen GPU-Ausführungsmodell darf nur die CPU die Ausführung eines Thread-Programms auf der GPU initiieren. Nachdem alle Thread-Instanzen ihre Ausführung abgeschlossen haben, muss die GPU die CPU entsprechend benachrichtigen und auf eine weitere Verarbeitungsaufgabe warten, die die CPU ausgeben muss. Das Benachrichtigen der CPU und das Warten auf die nächste Aufgabe ist typischerweise eine blockierende, serialisierte Operation, die bestimmte Ressourcen innerhalb der CPU untätig sein lässt, wodurch sich die Gesamtleistung des Systems verringert.
Die Leistung lässt sich in bestimmten Szenarien steigern, indem sequentielle Berechnungsaufgaben in einen Pushbuffer verschoben werden, aus dem die GPU Arbeit zur Ausführung beziehen kann, ohne auf die CPU zu warten. Berechnungsaufgaben, die feste, Datenfluss verarbeitende Pipelines umfassen, profitieren von diesem Puffermodell, wenn die CPU in der Lage ist, Arbeit für die GPU schnell genug zu erzeugen, so dass im Pushbuffer immer genügend Arbeit vorrätig ist, sobald die GPU in der Lage ist, eine neue Aufgabe zu beginnen. Datenabhängige Berechnungsaufgaben sind allerdings weiterhin mit einer sequentiellen Abhängigkeit zwischen GPU-Ergebnissen, CPU-Aufgabenverwaltung und nachfolgender GPU-Aufgabenausführung, die von der CPU gestartet werden muss, belastet.
Multi-threaded-Berechnungsmodelle organisieren Arbeit üblicherweise in geordneten Streams von Aufgaben, die in einer definierten Reihenfolge fertiggestellt werden müssen. In solchen Berechnungsmodellen schreibt die Ausführungssemantik vor, dass eine Aufgabe beendet worden sein muss, bevor eine davon abhängige Aufgabe ausgeführt werden darf. In einem einfachen Szenario lässt sich eine serielle Abhängigkeit innerhalb einer beliebigen Sequenz von Aufgaben in eine Warteschlange in einem Pushbuffer stellen, um so die effiziente Ausführung durch die GPU zu ermöglichen. Bestimmte Verarbeitungsmodelle erlauben allerdings Abhängigkeiten über verschiedene Streams hinweg, wodurch eine Aufgabe in einem Stream von der Fertigstellung zweier oder mehr verschiedener Aufgaben abhängt, die sich in zwei oder mehr anderen Streams befinden. In solchen Szenarien terminiert die CPU die Aufgaben, um eine Blockade zu verhindern. Der Prozess, auf den Abschluss bestimmter Aufgaben zu warten, bevor andere Aufgaben terminiert werden, um eine Blockade zu vermeiden, erzeugt zusätzliche serielle Abhängigkeiten zwischen Aufgabenabschlüssen in CPU und GPU, was die Gesamteffizienz beeinträchtigt.
Wie das oben Gesagte illustriert, ist im Fachgebiet ein Verfahren vonnöten, um die effizientere und semantisch vollständigere GPU-Ausführung zu ermöglichen.
Zusammenfassung der Erfindung
Ein Ausführungsbeispiel der vorliegenden Erfindung stellt ein auf einem Computer implementiertes Verfahren dar, eine Mehrzahl von Aufgaben über eine Gruppe von Threads hinweg zu verarbeiten, wobei das Verfahren folgendes umfasst: Aufrufen eines ersten Elementes aus einer Warteschlange; Feststellen, dass das erste Element keine Aufgabe aufweist; Feststellen, ob das erste Element einen Warte-Event oder einen Signalisierungs-Event aufweist; in Reaktion darauf Dekrementieren eines Zählers; und Entfernen des ersten Elementes aus der Warteschlange.
Andere Ausführungsbeispiele der vorliegenden Erfindung weisen, ohne einschränkend zu wirken, ein Computer-lesbares Speichermedium inklusive Anweisungen auf, welche, sobald sie von einer Verarbeitungseinheit ausgeführt werden, die Verarbeitungseinheit veranlassen, die hier beschriebenen Techniken auszuführen, sowie ein Berechnungsvorrichtung, das eine Verarbeitungseinheit umfasst, die eingerichtet ist, die hier beschriebenen Techniken auszuführen.
Ein Vorteil des hier offengelegten Verfahrens besteht darin, dass eine GPU ohne Involvieren einer CPU die Ausführungsreihenfolge in solchen Aufgaben korrekt und effizient erzwingen kann, welche Stream-übergreifende Abhängigkeiten aufweisen.
Kurze Beschreibung der Zeichnungen
Damit die Weise, in der die oben aufgezählten Merkmale der vorliegenden Erfindung im Einzelnen wirken, verstanden wird, soll eine eingehendere Beschreibung der Erfindung, wie sie oben kurz zusammengefasst wurde, durch Hinweis auf die Ausführungsbeispiele, von denen einige in den beigefügten Zeichnungen dargestellt sind, geliefert werden. Es sollte jedoch vermerkt werden, dass die beigefügten Zeichnungen lediglich typische Ausführungsbeispiele dieser Erfindung illustrieren und daher nicht als Einschränkungen des Potentials der Erfindung angesehen werden sollten, denn die Erfindung lässt weitere, ebenso wirksame Ausführungsbeispiele zu.
1 ist ein Blockdiagramm, das ein Computersystem illustriert, das konfiguriert ist, einen oder mehrere Aspekte der vorliegenden Erfindung zu implementieren.
2 ist ein Blockdiagramm eines parallelverarbeitenden Subsystems für das Computersystem von 1 gemäß einem Ausführungsbeispiel der vorliegenden Erfindung.
3A ist ein Blockdiagramm des Front-Ends von 2 gemäß einem Ausführungsbeispiel der vorliegenden Erfindung.
3B ist ein Blockdiagramm, das einen allgemein verarbeitenden Cluster innerhalb der Parallelverarbeitungseinheiten in 2 gemäß einem Ausführungsbeispiel der vorliegenden Erfindung darstellt.
3C ist ein Blockdiagramm, das einen Teilbereich des Streaming-Multiprozessors in 3B gemäß einem Ausführungsbeispiel der vorliegenden Erfindung zeigt.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung illustriert 4 die verschachtelte Aufgabenausführung auf einem parallelverarbeitenden Subsystem;
5 stellt eine Event-Abfolge für ein Eltern-Thread-Programm gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar, um einen Kind-Grid zu starten.
6 stellt Systemelemente für ein Subsystem für die verschachtelte Ausführung gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
7 stellt einen exemplarischen Graphen für hierarchische Ausführung einschließlich damit assoziierter Warteschlangen und Aufgaben für Aufgaben-Metadaten-Deskriptoren gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
8 stellt einen damit verknüpften Graphen für hierarchische Ausführung einschließlich damit assoziierter Warteschlangen und Aufgaben für Aufgaben-Metadaten-Deskriptoren gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
9A stellt einen exemplarischen Graphen für die hierarchische Ausführung von Aufgaben, die Stream-übergreifende Abhängigkeiten aufweisen gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
9B stellt einen gleichwertigen Graphen für die hierarchische Ausführung von Aufgaben und Events für das Durchsetzen einer Ausführungsreihenfolge unter den abhängigen Aufgaben in verschiedenen Streams einem Ausführungsbeispiel der vorliegenden Erfindung dar.
10 stellt eine Datenstruktur für einen Thread-Gruppenkontext, einschließlich Parametern und Kontextinformationen, die mit einer Thread-Gruppe verknüpft sind, gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
11A stellt eine Datenstruktur für den Aufgabenstatus einschließlich Parametern, die mit einer Berechnungsaufgabe verknüpft sind, gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
11B stellt die Datenstruktur für einen Signalisierungs-Event-Status einschließlich Parametern, die mit einem Signalisierungs-Event verknüpft sind, gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
11C stellt eine Datenstruktur eines Warte-Event-Status einschließlich der Parameter, die mit einer Berechnungsaufgabe verknüpft sind, gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar.
12 ist ein Flussdiagramm von Verfahrensschritten, um festzustellen, dass die Abhängigkeiten für die Aufgabenausführung berücksichtigt worden sind, gemäß einem Ausführungsbeispiel der vorliegenden Erfindung.
Detaillierte Beschreibung
In der folgenden Beschreibung werden zahlreiche spezifische Details dargelegt, um ein tiefergehendes Verständnis der vorliegenden Erfindung zu ermöglichen. Einer Person mit Kenntnissen auf diesem Fachgebiet dürfte jedoch offensichtlich sein, dass die vorliegende Erfindung ohne diese oder mit mehr dieser spezifischen Details umgesetzt werden kann.
Überblick über das System
1 ist ein Blockdiagramm, das ein Computersystem 100 zeigt, welches konfiguriert ist, einen oder mehr Aspekte der vorliegenden Erfindung zu implementieren. Computersystem 100 weist eine Zentralverarbeitungseinheit (CPU) 102 und einen Systemspeicher 104 auf, die miteinander über einen Verbindungspfad kommunizieren, der eine Speicherbrücke 105 umfassen kann. Die Speicherbrücke 105, die beispielsweise in einem Northbridge-Chip bestehen kann, ist über einen Bus oder einen anderen Kommunikationspfad 106 (z. B. eine HyperTransport-Verknüpfung) mit einer I/O-Brücke (I/O: Input/Output) 107 verbunden. Die I/O-Brücke 107, die beispielsweise in einem Southbridge-Chip bestehen kann, empfängt Nutzereingaben über ein oder mehrere Vorrichtungen für die Nutzereingabe 108 (wie etwa eine Tastatur, eine Maus) und leitet die Eingaben über den Kommunikationspfad 106 und die Speicherbrücke 105 an die CPU 102 weiter. Ein Subsystem zur Parallelverarbeitung 112 ist mit Hilfe eines Busses oder einen zweiten Kommunikationspfad 113 (zum Beispiel ein Peripheral Component Interconnect Express [PCIe], Accelerated Graphics Port [AGP] oder HyperTransport-Verbindung) an die Speicherbrücke 105 gekoppelt; in einem Ausführungsbeispiel ist das Parallelverarbeitungssystem 112 ein Graphiksubsystem, das Bildpunkte an eine Anzeigevorrichtung 110 (zum Beispiel eine konventionelle Kathodenstrahlröhre oder ein Flüssigkristallbildschirm) liefert. Eine System-Festplatte 114 ist ebenfalls mit der I/O-Brücke 107 verbunden. Ein Schalter 116 stellt Verbindungen zwischen I/O-Brücke 107 und anderen Komponenten wie etwa einem Netzwerkadapter 118 und verschiedenen Zusatzkarten 120 und 121 her. Andere Komponenten (die nicht explizit gezeigt werden), darunter ein universeller serieller Bus (USB) oder andere Steckverbindungen, CD-Laufwerke, DVD-Laufwerke, Filmaufzeichnungsgeräte und dergleichen lassen sich ebenfalls mit der I/O-Brücke 107 verbinden. Die verschiedenen Kommunikationswege in 1, darunter die spezifisch genannten Kommunikationswege 106 und 113, lassen sich mit Hilfe geeigneter Protokolle wie etwa PCI Express, AGP (Accelerated Graphics Port), HyperTransport oder einem anderen Bus oder Punkt-zu-Punkt-Kommunikationsprotokollen implementieren, und Verbindungen zwischen verschiedenen Geräten dürfen verschiedene Protokolle nutzen, die in dem Fachgebiet bekannt sind.
In einem Ausführungsbeispiel umfasst das Parallelverarbeitungssubsystem 112 Schaltkreise, die für Graphik- und Videoverarbeitung optimiert sind, darunter beispielsweise auch Schaltkreise für die Videoausgabe, was eine graphische Verarbeitungseinheit (GPU) konstituiert. In einem anderen Ausführungsbeispiel weist das Parallelverarbeitungssubsystem 112 Schaltkreise auf, die für allgemeine Verarbeitungsaufgaben optimiert sind, während sie weiterhin die grundlegende Computerarchitektur beibehalten, die hier eingehend beschrieben wird. In einem weiteren Ausführungsbeispiel kann das Parallelverarbeitungssubsystem 112 mit einem oder anderen Systemelementen in einem einzigen Subsystem integriert sein, indem es beispielsweise mit der Speicherbrücke 105, CPU 102 und der I/O-Brücke 107 verbunden ist, um ein ”System auf einem Chip” (SoC) zu bilden.
Es sei darauf hingewiesen, dass das hier gezeigte System rein illustrativen Zwecken dient und dass Variationen sowie Modifikationen möglich sind. Die Topologie der Verbindungen, darunter die Anzahl und Anordnung der Brücken, die Anzahl der CPUs 102, und die Anzahl der Parallelverarbeitungssubsysteme 112 kann nach Wunsch modifiziert werden. In manchen Ausführungsbeispielen ist beispielsweise der Systemspeicher 104 mit der CPU 102 statt über eine Brücke vielmehr direkt verbunden, und andere Geräte kommunizieren mit dem Systemspeicher 104 über eine Speicherbrücke 105 und die CPU 102. In weiteren alternativen Topologien ist das Parallelverarbeitungssubsystem 112 mit der I/O-Brücke 107 oder direkt mit der CPU 102 verbunden statt mit der Speicherbrücke 105. In weiteren Ausführungsbeispielen können die I/O-Brücke 107 und die Speicherbrücke 105 in einem einzigen Chip integriert sein statt als ein oder mehrere Geräte zu existieren. Umfangreiche Ausführungsbeispiele können zwei oder mehr CPUs 102 und zwei oder mehr Parallelverarbeitungssysteme 112 umfassen. Die einzelnen Komponenten in diesem Dokument sind optional; beispielsweise kann jede beliebige Anzahl von Erweiterungskarten oder Zusatzgeräten unterstützt werden. In manchen Ausführungsbeispielen wird der Switch 116 weggelassen und der Netzwerkadapter 118 sowie die Erweiterungssteckkarten 120 und 121 sind direkt mit der I/O-Brücke 107 verbunden.
2 illustriert ein Parallelverarbeitungssubsystem 112 gemäß der Beispielausführung in der vorliegenden Erfindung. Wie gezeigt wird, umfasst das Parallelverarbeitungssubsystem 112 eine oder mehrere Parallelverarbeitungseinheiten (PPUs) 202, von denen jede an einen lokalen Parallelverarbeitungsspeicher (PP) 204 gekoppelt ist. Allgemein gesagt, umfasst ein Parallelverarbeitungssubsystem eine Anzahl U von PPUs, wobei U ≥ 1 ist. (In dem vorliegenden Dokument werden mehrere Instanzen gleicher Objekte mit Referenznummern bezeichnet, die das Objekt identifizieren, sowie mit Nummern in Parenthese, die erforderlichenfalls die Instanz bezeichnen.) Die PPUs 202 und Parallelverarbeitungsspeicherbausteine 204 lassen sich mit Hilfe eines oder mehrerer Schaltkreise implementieren, wie etwa programmierbaren Prozessoren, anwendungsspezifischen Schaltkreisen (ASICs) oder Speicherbausteinen oder auf jede andere technisch machbare Weise.
Mit Verweis auf die 1 und 2 lässt sich sagen, dass in manchen Ausführungsbeispielen manche oder alle PPUs 202 im Parallelverarbeitungssubsystem 112 Graphikprozessoren sind, die über Rendering-Pipelines verfügen, die sich so konfigurieren lassen, dass sie verschiedene Operationen ausführen, die das Erzeugen von Pixeldaten aus Graphikdaten betreffen, die über eine Speicherbrücke 105 und den zweiten Kommunikationspfad 113 von einer CPU 102 und/oder aus dem Systemspeicher 104 kommen. Dabei interagieren die Speicherbrücke 105 und der zweite Kommunikationspfad 113 mit dem lokalen Parallelverarbeitungsspeicher 204 (der sich als Graphikspeicher nutzen lässt, so etwa als konventioneller Frame-Zwischenspeicher), um Pixeldaten zu speichern und zu aktualisieren, wobei Pixeldaten an die Anzeigevorrichtung 110 oder dergleichen geliefert werden. In manchen Ausführungsbeispielen kann das Parallelverarbeitungssubsystem 112 eine oder mehrere PPUs 202 umfassen, die als Graphikprozessoren fungieren, sowie eine oder mehrere PPUs 202, die für Berechnungen allgemeiner Natur verwendet werden. Die PPUs können identisch oder unterschiedlich sein, und jede PPU kann einen dedizierten Parallelverarbeitungsspeicherbaustein besitzen oder auch keinen. Eine oder mehrere PPUs 202 im Parallelverarbeitungssubsystem 112 können Daten an die Anzeigevorrichtung 110 ausgeben, oder jede PPU 202 im Parallelverarbeitungssubsystem 112 kann Daten an eine oder mehr Anzeigevorrichtungen 110 ausgeben.
Beim Betrieb des Rechners ist die CPU 102 der Hauptprozessor des Computersystems 100, wobei sie Operationen anderer Systemkomponenten steuert und koordiniert. Im Besonderen gibt die CPU 102 Befehle aus, die die Tätigkeit der PPUs 202 steuern. In manchen Ausführungsbeispielen schreibt die CPU 102 einen Stream von Befehlen für jede PPU 202 in eine Datenstruktur (die weder in 1 noch 2 explizit angezeigt wird), die im Systemspeicher 104, im Parallelverarbeitungsspeicher 204 oder in einer anderen Speichereinheit vorliegen kann, die sowohl für CPU 102 als auch PPU 202 zugreifbar ist. Ein Pointer, der auf jede Datenstruktur verweist, wird in einen Pushbuffer geschrieben, um die Verarbeitung des Streams von Befehlen in der Datenstruktur anzustoßen. Die PPU 202 liest Befehls-Streams aus einem oder mehreren Pushbuffer aus und führt anschließend Befehle aus, die asynchron zur Tätigkeit der CPU 102 erfolgen. Die Prioritäten der Ausführung lassen sich für jeden Pushbuffer durch ein Anwendungsprogramm über den Gerätetreiber 103 festlegen, um so die Terminierung der verschiedenen Pushbuffer steuern zu können.
Wie in 1 wie auch in 2 zu erkennen, weist jede PPU 202 eine I/O-(Input/Output bzw. Ein-/Ausgabe)Einheit 205 auf, die mit dem Rest des Computersystems 100 über den Kommunikationspfad 113 kommuniziert, der mit der Speicherbrücke 105 (oder in einem alternativen Ausführungsbeispiel direkt mit der CPU 102) verbunden ist. Die Verbindung der PPU 202 mit dem Rest des Computersystems 100 kann ebenfalls unterschiedlich ausfallen. In manchen Ausführungsbeispielen wird das Parallelverarbeitungssubsystem 112 als eine Erweiterungskarte implementiert, die sich in einen Erweiterungssteckplatz des Computersystems 100 stecken lässt. In anderen Ausführungsbeispielen kann eine PPU 202 mit Hilfe einer Busbrücke, wie etwa Speicherbrücke 105 oder I/O-Brücke 107, auf einem einzigen Rechenbaustein implementiert werden. In weiteren Ausführungsbeispielen können einige oder alle Elemente von PPU 202 mit CPU 102 auf einem einzigen Chip integriert sein.
In einem Ausführungsbeispiel ist der Kommunikationspfad 113 eine PCI-Express-Verbindung, in der jeder PPU 202 dedizierte ”Fahrspuren” zugewiesen sind, wie dem Fachmann bekannt ist. Auch andere Kommunikationspfade lassen sich nutzen. Eine I/O-Einheit 205 erzeugt Datenpakete (oder andere Signale) für eine Übertragung auf Kommunikationspfad 113 und empfängt zudem alle ankommenden Datenpakete (oder andere Signale) von Kommunikationspfad 113, wobei sie die ankommenden Pakete an die entsprechenden Komponenten von PPU 202 weiterleitet. So können beispielsweise Befehle, die Verarbeitungsaufgaben betreffen, an ein Host-Interface 206 geleitet werden, wohingegen Befehle, die Speicheroperationen betreffen (zum Beispiel das Auslesen aus oder Schreiben in Parallelverarbeitungsspeicher 204), an eine Speicherklammereinheit 210 geleitet werden könnten. Das Host-Interface 206 liest jeden Pushbuffer aus und gibt den dort gespeicherten Befehlsstrom an das Frontend 212 aus.
Durch jede jede PPU 202 wird vorteilhafterweise eine hochgradig parallele Verarbeitungsarchitektur implementiert. Wie im Detail dargestellt wird, umfasst PPU 202(0) ein Processing Cluster Array 230, das eine Anzahl C von General Processing Clustern (GPC) 208 aufweist, wobei C ≥ 1 ist. Jeder GPC 208 ist in der Lage, eine große Zahl (das heißt Hunderte oder Tausende) von Threads gleichzeitig auszuführen, wobei ein Thread eine Instanz eines Programms darstellt. In verschiedenen Anwendungen können verschiedene GPCs 208 der Verarbeitung verschiedener Programmtypen zugewiesen sein oder um verschiedene Berechnungstypen auszuführen. Die Zuweisung von GPCs 208 kann von der Arbeitslast abhängig sein, die mit jedem Typ von Programm oder Berechnungsweise verbunden ist.
Die GPCs 208 empfangen Verarbeitungsaufgaben, die auszuführen sind, von einer Aufgabenverteilungseinheit innerhalb der Aufgaben/Arbeit-Einheit 207. Die Aufgabenverteilungseinheit empfängt Pointer, die zu Verarbeitungsaufgaben weisen, die als Aufgaben-Metadaten (Task Meta Data, TMD) enkodiert und im Systemspeicher gespeichert sind. Die Pointer auf TMDs sind im Befehlsstrom enthalten, der als Pushbuffer gespeichert ist, und werden von der Frontend-Einheit 212 vom Host Interface 206 empfangen. Zu den Verarbeitungsaufgaben, die als TMDs kodiert sein können, gehören Indizes der zu verarbeitenden Daten sowie Statusparameter und Befehle, wie die Daten zu verarbeiten sind (zum Beispiel welches Programm ausgeführt werden soll). Die Aufgaben/Arbeit-Einheit 207 empfängt Aufgaben vom Frontend 212 und sorgt dafür, dass GPCs 208 in einem gültigen Zustand konfiguriert sind, bevor die Verarbeitung, die von einer der TMDs festgelegt wird, angestoßen wird. Für jede TMD lässt sich eine Priorität festlegen, um auf diese Weise die Ausführung der Verarbeitungsaufgabe zu terminieren. Verarbeitungsaufgaben können auch vom Processing Cluster Array 230 entgegengenommen werden. Optional kann eine TMD einen Parameter enthalten, der steuert, ob die TMD am Anfang oder am Schluss einer Liste von Verarbeitungsaufgaben (oder eine Liste von Pointern auf Verarbeitungsaufgaben) eingefügt wird, so dass eine weitere Ebene der Kontrolle über die Prioritäten bereitsteht.
Die Speicherschnittstelle 214 enthält eine Anzahl D von Partitionseinheiten 215, die jeweils direkt an einen Abschnitt des Parallelverarbeitungsspeichers 204 gekoppelt sind, wobei D ≥ 1 ist. Wie in der Figur dargestellt, entspricht die Anzahl der Partitionseinheiten 215 im allgemeinen der Anzahl von Chips für Dynamic Random Access Memory (DRAM) 220. In anderen Ausführungsbeispielen kann die Anzahl der Partitionseinheiten 215 nicht ebenso hoch wie die der Speicherbausteine sein. Der Fachmann wird erkennen, dass sich DRAM 220 durch andere geeignete Speichergeräte ersetzen lässt und von allgemein konventionellem Aufbau sein darf. Daher wird auf eine detaillierte Beschreibung verzichtet. Rendering-Ziele wie etwa Frame Buffer oder Texture Maps lassen sich über mehrere DRAMs 220 hinweg speichern, was es Partitionseinheiten 215 erlaubt, Portionen jedes Rendering-Ziels parallel zu schreiben, um so die verfügbare Bandbreite des Parallelverarbeitungsspeichers 204 effizient zu nutzen.
Jede der GPCs 208 kann Daten verarbeiten, die in jeden der DRAMs 220 innerhalb des Parallelverarbeitungsspeicher 204 geschrieben werden. Die Crossbar-Einheit 210 ist so konfiguriert, dass sie die Ausgabe jeder GPC 208 zur Eingabe jeder Partitionseinheit 215 oder zu einer anderen GPC 208 zur weiteren Verarbeitung leitet. Die GPCs 208 kommunizieren über die Crossbar-Einheit 210 mit der Speicherschnittstelle 214, um auf verschiedene externe Speichergeräte zu schreiben oder von dort zu lesen. In einem Ausführungsbeispiel verfügt die Crossbar-Einheit 210 über eine Verbindung zur Speicherschnittstelle 214, um mit der I/O-Einheit 205 zu kommunizieren, sowie über eine Verbindung zum lokalen Parallelverarbeitungsspeicher 204, wodurch es den Prozessorkernen innerhalb der verschiedenen GPC 208 ermöglicht wird, mit dem Systemspeicher 104 oder anderem Speicher zu kommunizieren, der nicht in PPU 202 lokalisiert ist. In dem in 2 gezeigten Ausführungsbeispiel ist die Crossbar-Einheit 210 direkt mit der I/O-Einheit 205 verbunden. Die Crossbar-Einheit 210 kann virtuelle Kanäle nutzen, um Datenströme, die zwischen den GPCs 208 und den Partitionseinheiten 215 fließen, aufzuteilen.
Noch einmal: GPCs 208 lassen sich programmieren, Verarbeitungsaufgaben für eine breite Palette von Anwendungen auszuführen, wozu auch, aber nicht nur, lineare und nichtlineare Datentransformationen, ein Filtern von Video- und/oder Audiodaten, Modellierungsoperationen (z. B. das Anwenden der Gesetze der Physik, um Position, Geschwindigkeit oder andere Eigenschaften von Objekten zu bestimmen), Bildrenderingoperationen (z. B. Tesselations-Schattier-, Winkel-Schattier-, Geometrie-Schattier- und/oder Pixel-Schattierprogramme) und so weiter. Die PPUs 202 können Daten aus dem Systemspeicher 104 und/oder den lokalen Parallelverarbeitungsspeichern 204 in den internen Speicher (auf dem Chip) übertragen, die Daten verarbeiten und die Ergebnisdaten zurück in den Systemspeicher 104 oder die lokalen Parallelverarbeitungsspeicher 204 schreiben, wo solche Daten für andere Systemkomponenten, darunter die CPU 102 oder ein anderes Parallelverarbeitungssubsystem 112, zugänglich sind.
Eine PPU 202 kann mit beliebig viel Parallelverarbeitungsspeicher 204 versehen werden, darunter auch keinem lokalem Speicher, und kann lokalen und Systemspeicher in beliebigen Kombinationen nutzen. Eine PPU 202 kann beispielsweise ein Graphikprozessor in dem Ausführungsbeispiel einer Unified Memory Architecture (UMA) sein. In solchen Ausführungsbeispielen würde wenig oder gar kein dedizierter Graphik- (also parallel verarbeitender) Speicher bereitgestellt werden, und die PPU 202 würde ausschließlich oder fast ausschließlich den Systemspeicher nutzen. In Ausführungsbeispielen von UMA kann eine PPU 202 mit einem Brückenbaustein oder Prozessorbaustein integriert sein oder als separater Baustein bereitgestellt werden, der die PPU 202 über eine Hochgeschwindigkeitsverbindung (z. B. PCI Express) über einen Brückenbaustein oder andere Kommunikationsmittel mit dem Systemspeicher verbindet.
Wie oben erwähnt lässt sich jede beliebige Anzahl von PPUs 202 in einem Parallelverarbeitungssubsystem 112 unterbringen. Mehrere PPUs lassen sich beispielsweise auf einer einzigen Erweiterungskarte bereitstellen, oder mehrere Erweiterungskarten können mit dem Kommunikationspfad 113 verbunden werden, oder eine oder mehrere PPUs 202 können mit einem Brückenbaustein integriert werden. PPUs 202, die sich in einem Multi-PPU-System befinden, können zueinander identisch oder verschieden sein. So können verschiedenartige PPUs 202 auch unterschiedliche Mengen von Rechenkernen aufweisen, unterschiedliche Mengen von lokalem Parallelverarbeitungsspeicher und so weiter. Wo mehrere PPUs 202 vorliegen, lassen sich diese PPUs parallel verwenden, um Daten mit einem höheren Datendurchsatz zu verarbeiten, als mit einer einzelnen PPU 202 möglich ist. Systeme, die eine oder mehrere PPUs 202 enthalten, lassen in vielerlei Konfigurationen und Formfaktoren implementieren, darunter Desktop, Laptop oder Handheld-Rechner, Server, Workstations, Spielekonsolen, Embedded-Systeme und dergleichen.
Terminierung mehrerer gleichzeitig ausgeführter Aufgaben
Mehrere Verarbeitungsaufgaben können auf den GPCs 208 gleichzeitig ausgeführt werden, und eine Verarbeitungsaufgabe kann während der Abarbeitung eine oder mehr ”Tochter”-Verarbeitungsaufgabe(n) („child”) generieren. Die Aufgaben/Arbeit-Einheit 207 nimmt die Aufgaben entgegen und terminiert dynamisch die Verarbeitungsaufgaben und Tochter-Verarbeitungsaufgaben zur Ausführung durch die GPCs 208.
3A ist ein Blockdiagramm der Aufgabe/Arbeit-Einheit 207 aus 2, gemäß einem der Ausführungsbeispiele der vorliegenden Erfindung. Die Aufgabe/Arbeit-Einheit 207 umfasst eine Aufgabenverwaltungseinheit 300 und die Arbeitszuteilungseinheit 340. Die Aufgabenverwaltungseinheit 300 organisiert zu terminierende Aufgaben nach ihrem jeweiligen Ausführungsprioritätsgrad. Für jeden Prioritätsgrad speichert die Aufgabenverwaltungseinheit 300 eine Liste von Pointern auf die TMDs 322, die den Aufgaben in der Terminierungstabelle 321 entsprechen, wobei die Liste als verknüpfte Liste implementiert sein kann. Die TMDs 322 können im PP-Speicher 204 (PP: parallelverarbeitend) oder im Systemspeicher 104 abgelegt sein. Die Schnelligkeit, mit der die Aufgabenverwaltungseinheit 300 Aufgaben entgegennimmt und in der Scheduler-Tabelle 321 ablegt, ist von der Schnelligkeit entkoppelt, mit der die Aufgabenverwaltungseinheit 300 Aufgaben zur Ausführung terminiert. Folglich kann die Aufgabenverwaltungseinheit 300 eine Reihe von Aufgaben sammeln, bevor sie diese terminiert. Die gesammelten Aufgaben lassen sich dann je nach Prioritätsinformation oder mit Hilfe anderer Techniken, wie etwa Reihum-Terminierung (Round-Robin-Scheduling), terminieren.
Die Arbeitszuteilungseinheit 340 umfasst eine Aufgabentabelle 345 mit ”Fächern” (Slot), die jeweils von der TMD 322 für eine Aufgabe (Task), die gerade verarbeitet wird, belegt werden. Die Aufgabenverwaltungseinheit 300 kann Aufgaben zur Ausführung einteilen, sobald es ein freies Fach in der Aufgabentabelle 345 gibt. Gibt es hingegen kein freies Fach, kann eine Aufgabe höherer Priorität, die kein Fach hat, eine Aufgabe geringerer Priorität, die ein Fach besetzt, hinauswerfen. Sobald eine Aufgabe hinausgeworfen wird, wird sie angehalten, und falls ihre Ausführung noch nicht vollständig ist, wird ein Zeiger auf diese Aufgabe der Liste der zu terminierenden Aufgabenzeiger hinzugefügt, so dass die Ausführung der Aufgabe zu einem späteren Zeitpunkt wiederaufgenommen wird. Wenn eine Tochter-Verarbeitungsaufgabe während der Aufgabenausführung erzeugt wird, wird auch ein Zeiger auf die Tochter-Verarbeitungsaufgabe in die Liste der Zeiger auf terminierte Aufgaben aufgenommen. Eine Tochter-Verarbeitungsaufgabe kann von einem TMD 322 erzeugt werden, der im Processing Cluster Array 230 ausgeführt wird.
Anders als eine Aufgabe, die die Aufgaben/Arbeit-Einheit 207 vom Frontend 212 entgegennimmt, werden Tochter-Verarbeitungsaufgaben vom Processing Cluster Array 230 entgegengenommen. Tochter-Verarbeitungsaufgaben werden nicht in Pushbuffer eingefügt oder an das Frontend übertragen. Die CPU 102 wird nicht verständigt, wenn eine Tochter-Verarbeitungsaufgabe erzeugt wird oder wenn Daten für die Tochter-Verarbeitungsaufgabe im Speicher abgelegt werden. Ein weiterer Unterschied, der zwischen Tochter-Verarbeitungsaufgaben und Aufgaben besteht, die von Pushbuffern geliefert werden, besteht darin, dass letztere vom Anwendungsprogramm definiert werden, wohingegen Tochter-Verarbeitungsaufgaben dynamisch während der Aufgabenausführung generiert werden.
Überblick über die Aufgabenverarbeitung
3B ist das Blockdiagramm eines GPC 208 innerhalb einer der PPUs 202 in 2, in Übereinstimmung mit einem Ausführungsbeispiel der vorliegenden Erfindung. Jedes GPC 208 lässt sich so konfigurieren, dass es eine Vielzahl von Threads parallel ausführt, wobei der Begriff ”Thread” eine Instanz eines bestimmten Programms bezeichnet, das auf eine bestimmte Menge von Eingabedaten hin ausgeführt wird. In manchen Ausführungsbeispielen werden Single-Instruction, Multiple-Data(SIMD)-Befehlsausgabetechniken verwendet, um eine parallele Ausführung einer Vielzahl von Threads zu unterstützen, ohne mehrere voneinander unabhängige Befehlssätze bereitzustellen. In anderen Ausführungsbeispielen werden Single-Instruction, Multiple-Thread(SIMT)-Techniken genutzt, um die parallele Ausführung einer Vielzahl von allgemein synchronisierten Threads zu unterstützen, unter Nutzung einer gemeinsamen Befehlseinheit, die für die Befehlsausgabe an eine Reihe von Verarbeitungsmaschinen innerhalb jedes der GPCs 208 konfiguriert ist. Anders als im SIMD-Befehlsausgabeverfahren, bei dem alle Verarbeitungsmaschinen typischerweise identische Anweisungen ausführen, gestattet es die SIMT-Befehlsausführung verschiedenen Threads, leichter divergierenden Ausführungspfaden durch das jeweilige Programm zu folgen. Ein Fachmann wird verstehen, dass ein SIMD-Verarbeitungsverfahren eine funktionale Untermenge des SIMT-Verarbeitungsverfahren darstellt.
Die Arbeitsweise des GPC 208 wird vorteilhaft mit Hilfe eines Pipeline-Managers 305 gesteuert, der Verarbeitungsaufgaben an Streaming-Multiprozessoren (SMs) 310 verteilt. Der Pipeline-Manager 305 lässt sich zudem so konfigurieren, dass er einen Arbeitsverteilungs-Crossbar 330 steuert, indem er Bestimmungsorte für die von SMs 310 ausgegebenen Verarbeitungsdaten festlegt.
In einem Ausführungsbeispiel enthält jedes GPC 208 eine Anzahl M von SMs 310, wobei M ≥ 1 und jede SM 310 so konfiguriert ist, eine oder mehr Thread-Gruppen zu verarbeiten. Darüber hinaus weist jede SM 310 eine identische Menge von Funktionsausführungseinheiten (das heißt, Ausführungseinheiten und Laden-Speichern-Einheiten – sie werden in 3C als Exec-Einheiten 302 und LSUs 303 gezeigt) auf, die sich in einer Reihe anordnen lassen, um es einem neuen Befehl zu gestatten, ausgegeben zu werden, bevor ein vorhergehender Befehl vollständig abgearbeitet worden ist, wie im Fachgebiet bekannt. Es lässt sich jede beliebige Kombination von funktionalen Ausführungseinheiten bereitstellen. In einem Ausführungsbeispiel unterstützen die Funktionseinheiten eine Vielfalt von Operationen, darunter Integer- und Fließkomma-Arithmetik (wie etwa Addition und Multiplikation), Vergleichsoperationen, Boolesche Operationen (AND, OR, XOR), bitweise Verschiebungen, Berechnung verschiedener Algebrafunktionen (z. B. planare Interpolation, trigonometrische, Exponential- und logarithmische Funktionen usw.); die gleiche Funktionseinheiten-Hardware lässt sich nutzen, um verschiedene Rechenoperationen auszuführen.
Die Befehlsfolge, die an ein bestimmtes GPC 208 übertragen wird, konstituiert einen Thread, wie er oben definiert wurde, und die Sammlung einer bestimmten Anzahl gleichzeitig in den Parallelverarbeitungsmaschinen (nicht gezeigt) einer SM 310 ausgeführter Threads wird als ”Warp” oder ”Thread-Gruppe” bezeichnet. Eine Thread-Gruppe, wie der Begriff hier verwendet wird, bezeichnet eine Gruppe von Threads, die gleichzeitig das gleiche Programm mit unterschiedlichen Eingabedaten verarbeiten, wobei jeweils ein Thread der Gruppe einer anderen Verarbeitungsmaschine in einer SM 310 zugewiesen ist. Eine Thread-Gruppe kann weniger Threads enthalten als die Anzahl der Verarbeitungsmaschinen in der SM 310; in diesem Fall werden manche Verarbeitungsmaschinen während der Zyklen, in denen die Thread-Gruppe verarbeitet wird, untätig sein. Eine Thread-Gruppe kann auch mehr Threads als die Anzahl der Verarbeitungsmaschinen in der SM 310 enthalten; in diesem Fall findet die Verarbeitung über aufeinander folgende Clock-Zyklen statt. Da jede SM 310 gleichzeitig bis zu G Thread-Gruppen unterstützt, ergibt sich daraus, dass in GPC 208 G × M Thread-Gruppen auf einmal ausgeführt werden können.
Darüber hinaus kann in einem GPC 208 eine Mehrzahl von verbundenen Thread-Gruppen zur gleichen Zeit aktiv (in unterschiedlichen Phasen der Ausführung) sein. Diese Sammlung von Thread-Gruppen wird hier als ”cooperative thread array” (CTA) oder ”thread array” bezeichnet. Die Größe eines bestimmten CTA entspricht m·k, wobei k die Anzahl der gleichzeitig ausgeführten Threads in einer Thread-Gruppe ist und typischerweise ein Integer-Vielfaches der Anzahl von Parallelverarbeitungsmaschinen in der SM 310 ist, und wobei m der Anzahl von Thread-Gruppen entspricht, die gleichzeitig in der SM 310 aktiv sind. Die Größe eines CTA wird im Allgemeinen vom Programmierer festgelegt und vom Umfang der Hardwareressourcen begrenzt, etwa durch Speicher oder Register, die dem CTA zur Verfügung stehen.
Jede SM 310 umfasst einen Level-eins-(L1)-Cache (der in 3C dargestellt ist) oder nutzt Platz in einem korrespondierenden L1-Cache außerhalb der SM 310, der zum Laden und Speichern von Rechenoperationen verwendet wird. Jede SM 310 hat zudem Zugriff auf Level-zwei-(L2)-Cache-Speicher, den sich alle GPCs 208 teilen und der für die Übertragung von Daten zwischen Threads genutzt werden kann. Schließlich haben SMs 310 Zugriff auf ”globalen” Speicher außerhalb ihrer Bausteine, wozu beispielsweise Parallelverarbeitungsspeicher 204 und/oder Systemspeicher 104 gehören kann. Es sollte verstanden werden, dass jeder beliebige Speicher außerhalb von PPU 202 sich als globaler Speicher nutzen lässt. Darüber hinaus kann ein Level-einskommafünf-Speicher (L1,5) 335 im GPC 208 enthalten sein, der konfiguriert ist, Daten zu empfangen und vorzuhalten, die von der SM 310 angefordert und über die Speicherschnittstelle 214 aus dem Speicher geholt wurden, darunter Befehle, uniforme Daten und konstante Daten, und die angeforderten Daten an die SM 310 zu liefern. Ausführungsbeispiele, die über mehrere SMs 310 in GPC 208 verfügen, teilen sich günstigerweise gemeinsame Befehle und Daten, die im L1,5-Speicher 335 zwischengelagert sind.
Jedes GPC 208 kann eine Speicherverwaltungseinheit (Memory Management Unit, MMU) 328 aufweisen, die konfiguriert ist, virtuelle auf physische Adressen abzubilden (Mapping). In anderen Ausführungsbeispielen kann sich die MMU 328 innerhalb der Speicherschnittstelle 214 befinden. Die MMU 328 weist eine Reihe von Seiten-Tabelleneinträgen (Page Table Entries, PTE) auf, die verwendet werden, um eine virtuelle auf eine physische Adresse einer Kachel abzubilden, und optional zudem einen Cache Line Index. Die MMU 328 kann Adressübersetzungsablesepuffer (Address Translation Lookaside Buffers, TLB) oder Caches umfassen, die sich im Multiprozessor SM 310 oder im L1-Cache oder in GPC 208 befinden können. Die physische Adresse wird verarbeitet, um den Standort eines Oberflächendatenzugriffs zu verteilen, um so effizientes Anforderungs-Interleaving zwischen Partitionseinheiten 215 zu gestatten. Der Cache Line Index lässt sich verwenden, um zu bestimmen, ob die Anforderung einer bestimmten Cache-Zeile ein Treffer oder ein Fehlzugriff (hit or miss) ist.
In Graphik- und Rechen-Anwendungen kann ein GPC 208 so konfiguriert werden, dass jede SM 310 mit einer Textur-Einheit 315 gekoppelt ist, um Textur-Mapping-Operationen auszuführen, z. B. eine Bestimmung von Positionen von Textur-Mustern, Lesen von Texturdaten und Filtern von Texturdaten. Texturdaten werden aus einem internen Textur-L1-Cache (nicht dargestellt) ausgelesen oder, in manchen Ausführungsbeispielen, aus dem L1-Cache in der SM 310 und werden aus einem L2-Cache geholt bzw. aufgerufen, den sich alle GPCs 208 teilen, je nach Bedarf auch aus dem Parallelverarbeitungsspeicher 204 oder aus dem Systemspeicher 104. Jede SM 310 gibt verarbeitete Aufgaben an den Arbeitsverteilungs-Crossbar 330 aus, um die verarbeitete Aufgabe zur weiteren Bearbeitung an andere GPC 208 zu liefern oder um die erledigte Aufgabe mit Hilfe der Crossbar-Einheit 210 in einem L2-Cache, einem Parallelverarbeitungsspeicher 204 oder im Systemspeicher 104 abzulegen. Ein PreROP (pre-raster operations, vorab erledigte Rasteroperationen) 325 ist so konfiguriert, dass er Daten von der SM 310 entgegennimmt, Daten an ROP-Einheiten in den Partitionseinheiten 215 leitet und Optimierungen hinsichtlich Farbübergängen ausführt, Pixelfarbdaten organisiert und Adressübersetzungen ausführt.
Es wird anerkannt werden, dass die Kernarchitektur, die hier beschrieben wird, illustrativ ist und Variationen sowie Modifikationen möglich sind. Jede beliebige Anzahl von Verarbeitungseinheiten, so etwa SMs 310 oder Textureinheiten 315 oder preROPs 325, lassen sich in einer GPC 208 unterbringen. Wie in 2 dargestellt, kann eine PPU 202 eine beliebige Anzahl von GPCs 208 aufweisen, die einander vorteilhafterweise funktional ähnlich sind, so dass ihr Verhalten bei der Befehlsausführung nicht davon abhängt, welche GPC 208 eine bestimmte Verarbeitungsaufgabe empfängt. Jede GPC 208 arbeitet vorteilhafterweise unabhängig von anderen GPCs 208 und nutzt separate und klar definierte Verarbeitungseinheiten sowie L1-Caches, um Aufgaben für eine oder mehrere Applikationen zu erledigen.
Ein Fachmann wird verstehen, dass die in den 1, 2, 3A und 3B beschriebene Architektur keineswegs den Umfang der aktuellen Erfindung begrenzt und die darin gelehrten Techniken sich auf jeder ordnungsgemäß eingerichteten Recheneinheit implementieren lassen, die ohne Begrenzung eine oder mehrere CPUs, eine oder mehrere Mehrkern-CPUs, eine oder mehr PPUs 202 eine oder mehr GPCs 208, eine oder mehr Graphik- oder Sonderzweckverarbeitungseinheiten oder ähnliches umfasst, ohne von dem Umfang der aktuellen Erfindung abzuweichen.
In Ausführungsbeispielen der vorliegenden Erfindung ist es wünschenswert, die PPU 202 oder (einen) andere(n) Prozessor(en) eines Rechnersystems zu verwenden, um Berechnungen allgemeiner Art mit Hilfe von Thread-Arrays auszuführen. Jedem Thread im Thread-Array ist ein eindeutiger Thread-Identifikator (unique thread identifier, ”Thread Kennzeichen”, ID) zugewiesen, auf den der Thread während seiner Ausführung zugreifen kann. Die Thread-Kennzeichen, die als ein- oder mehrdimensionaler numerischer Wert definiert werden kann, steuert verschiedene Aspekte des Verarbeitungsverhaltens des Threads. So lässt sich beispielsweise mit der Thread-Kennzeichen festlegen, welche Teilmenge der Eingabedatenmenge ein Thread verarbeiten soll und/oder welche Teilmenge einer Ausgabedatenmenge ein Thread erzeugen oder schreiben soll.
Eine Abfolge von Befehlen pro Thread kann wenigstens einen Befehl enthalten, der festlegt, wie kooperativ sich der repräsentative Thread gegenüber einem oder mehreren anderen Threads des Thread-Arrays verhält. So kann die Abfolge von Pro-Thread-Befehlen einen Befehl enthalten, der die Ausführung von Operationen des repräsentativen Threads an einem bestimmten Punkt in der Abfolge stoppt, bis einer oder mehrere der anderen Threads diesen bestimmten Punkt erreichen; einen Befehl an den repräsentativen Thread, Daten in einem gemeinsamen Speicher, auf den einer oder mehrere der anderen Threads Zugriff haben, zu speichern; einen Befehl an den repräsentativen Thread, Daten, die in einem gemeinsamen Speicher, auf den einer oder mehrere der anderen Threads gemäß ihrer Thread-IDs Zugriff haben, vorliegen, einzeln zu lesen oder zu aktualisieren; oder dergleichen mehr. Ein CTA-Programm kann ebenfalls einen Befehl enthalten, eine Adresse im gemeinsamen Speicher, aus dem Daten gelesen werden sollen, zu berechnen, wobei die Adresse eine Funktion der Thread-Kennzeichen ist. Indem geeignete Funktionen festgelegt und Synchronisationstechniken bereitgestellt werden, lassen sich von einem Thread eines CTA Daten an eine bestimmte Adresse im gemeinsamen Speicher schreiben und von einem anderen Thread des gleichen CTA auf vorhersagbare Weise lesen. Als Folge daraus kann jedes gewünschte Muster der gemeinsamen Datennutzung unter Threads unterstützt werden, und jeder beliebige Thread in einem CTA kann Daten mit jedem anderen Thread im gleichen CTA austauschen. Das Ausmaß des Datenaustauschs, sofern vorhanden, zwischen Threads eines CTA wird vom CTA-Programm festgelegt; daher soll verstanden werden, dass in einer bestimmten Applikation, die CTAs verwendet, die Threads eines CTA Daten gemeinsam nutzen können oder nicht, je nach CTA-Programm; und dass die Termini ”CTA” und ”Thread-Array” in diesem Dokument synonym verwendet werden.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung ist 3C ein Blockdiagramm der SM 310 aus 3B. Die SM 310 weist einen L1-Cache 370 an Befehlen auf, der so eingerichtet ist, dass er Befehle und Konstanten aus dem Speicher über den L1,5-Cache 335 entgegennimmt. Eine Einheit für Befehle und Terminierung von Warps 312 empfängt Befehle und Konstanten vom L1-Cache für Befehle 370 und steuert die lokale Registerdatei 304 und Funktionseinheiten von SM 310 gemäß den Befehlen und Konstanten. Die Funktionseinheiten von SM 310 enthalten N Exec-(Ausführung oder Verarbeitung)-Einheiten 302 und P Lade-Speicher-Einheiten (load-store units, LSU) 303.
Die SM 310 stellt (internen) Datenspeicher auf ihrem Chip bereit, wobei der Zugriff darauf verschiedene Ebenen aufweist. Spezielle Register (nicht dargestellt) können von LSU 303 gelesen, aber nicht geschrieben werden und werden verwendet, um Parameter zu speichern, die die ”Position” jedes Threads bestimmen. In einem Ausführungsbeispiel enthalten die Spezialregister ein Register pro Thread (oder pro Ausführungseinheit 302 in der SM 310), das ein Thread-Kennzeichen speichert; auf jedes Thread-Kennzeichen-Register kann nur ein entsprechendes Thread-Kennzeichen in Exec-Einheit 302 zugreifen. Spezialregister können also zusätzliche Register umfassen, die für alle Threads lesbar sind, die die gleiche Verarbeitungsaufgabe ausführen, die ihnen von TMD 322 (oder von allen LSUs 303) präsentiert wird, das einen CTA-Identifikator, die CTA-Dimensionen, die Dimensionen eines Grids, zu dem der CTA gehört (oder eine Warteschlangenposition, falls das TMD 322 eine Wartenschlangenaufgabe anstelle einer Grid-Aufgabe kodiert) und einen Identifikator des TMD 322, dem der CTA zugewiesen ist, speichert.
Sollte das TMD 322 ein Grid-TMD sein, verursacht eine Ausführung des TMD 322 einen Start und eine Ausführung einer festgelegten Anzahl von CTAs, um die festgelegte Menge von Daten zu verarbeiten, die in der Warteschlange 525 gespeichert sind. Die Anzahl der CTAs ist festgelegt als das Produkt aus Breite, Höhe und Tiefe des Grid. Die festgelegte Menge von Daten lässt sich in TMD 322 ablegen oder TMD 322 kann einen Pointer auf die Daten speichern, der von den CTAs verarbeitet wird. TMD 322 speichert zudem eine Startadresse des Programms, das von den CTAs ausgeführt wird.
Sollte TMD 322 hingegen ein Warteschlangen-TMD sein, dann wird ein Warteschlangenmerkmal von TMD 322 genutzt, was bedeutet, dass die zu verarbeitende Datenmenge nicht zwangsläufig festgelegt ist. Die Einträge in der Warteschlange speichern von den CTAs zu verarbeitende Daten, wobei die CTAs dem TMD 322 zugewiesen sind. Die Einträge in der Warteschlange können zudem eine Tochter-Aufgabe darstellen, die von einem anderen TMD 322 während der Ausführung eines Thread generiert wird, so dass ein geschachtelter Parallelismus bereitgestellt wird. Typischerweise wird die Ausführung eines Thread beziehungsweise eines CTA, der den Thread enthält, ausgesetzt, bis die Ausführung der Tochter-Aufgabe abgeschlossen ist. Die Warteschlange lässt sich in TMD 322 selbst oder separat realisieren; in letzterem Fall speichert das TMD 322 einen Warteschlangen-Pointer auf die Warteschlange. Daten, die von der Tochter-Aufgabe erzeugt werden, können vorteilhafterweise in die Warteschlange geschrieben werden, während sich das TMD 322, das die Tochter-Aufgabe aufweist, in Ausführung befindet. Die Warteschlange lässt sich als ringförmige Warteschlange implementieren, so dass die Gesamtmenge an Daten nicht auf den Umfang der Warteschlange begrenzt ist.
CTAs, die zu einem Grid gehören, verfügen implizit über die Parameter hinsichtlich Breite, Höhe und Tiefe, die die Position des jeweiligen CTA innerhalb des Grids angeben. Spezielle Register werden während der Initialisierungsphase geschrieben, als Reaktion auf Befehle, die über das Frontend 212 vom Gerätetreiber 103 empfangen werden und sich während der Ausführung einer Verarbeitungsaufgabe nicht ändern. Das Frontend 212 terminiert die Ausführung jeder Verarbeitungsaufgabe. Jedes CTA ist mit einem bestimmten TMD 322 zwecks gleichzeitiger Ausführung von einer oder mehreren Aufgaben verknüpft. Zusätzlich kann ein einzelner GPC 208 mehrere Aufgaben gleichzeitig ausführen.
Ein Parameterspeicher (nicht eingezeichnet) speichert Laufzeitparameter (Konstanten), die von jedem Thread innerhalb des gleichen CTA (oder jedes LSU 303) gelesen, aber nicht geschrieben werden können. In einem Ausführungsbeispiel liefert der Gerätetreiber 103 Parameter an den Parameterspeicher, bevor er die SM 310 anweist, die Ausführung einer Aufgabe zu beginnen, die diese Parameter verwendet. Jeder beliebige Thread in jedem beliebigen CTA (oder jede beliebige Exec-Einheit 302 in der SM 310) hat über die Speicherschnittstelle 214 Zugriff auf den globalen Speicher. Teilmengen des globalen Speichers können im L1-Cache 320 gespeichert sein.
Die Lokale Registerbank 304 wird von jedem Thread als Scratch-Speicherplatz benutzt; jedes Register ist für die exklusive Verwendung durch einen einzigen Thread allokiert, und Daten, die sich in einer lokalen Registerbank 304 befinden, sind nur für denjenigen Thread zugänglich, dem das Register zugewiesen ist. Die lokale Registerbank 304 lässt sich als Registerbank implementieren, die physisch oder logisch in P Steuerleitungen aufgeteilt ist, wobei jede eine Anzahl Einträge aufweist (in denen jeder Eintrag beispielsweise ein 32-Bit-Wort speichern kann). Jeweils eine Steuerleitung ist jeder der N Ausführungseinheiten 302 und P Lade-Speicher-Einheiten (LSU) 303 zugewiesen, und korrespondierende Einträge in anderen Steuerleitungen lassen sich mit Daten für andere Threads füllen, die das gleiche Programm ausführen, um so die SIMD-Ausführung zu erleichtern. Unterschiedliche Abschnitte der Steuerleitungen lassen sich anderen aus den G gleichzeitig ausgeführten Thread-Gruppen zuweisen, so dass ein gegebener Eintrag in einer lokalen Registerbank 304 nur für einen bestimmten Thread zugänglich ist. In einem Ausführungsbeispiel sind bestimmte Einträge in der lokalen Registerbank 304 für das Ablegen von Thread-Identifikatoren reserviert, um so eines der Sonderregister zu implementieren. Darüber hinaus speichert ein uniformer L1-Cache 375 uniforme oder konstante Werte für jede Steuerleitung der N Exec-Einheiten 302 und P Lade-Speicher-Einheiten (LSU) 303.
Gemeinsamer Speicher 306 ist Threads in einer einzigen CTA zugänglich; mit anderen Worten, jeder Ort im gemeinsamen Speicher 306 ist jedem beliebigen Thread zugänglich, der sich in der gleichen CTA (oder für jede Verarbeitungsmaschine in einem SM 310) befindet. Gemeinsamer Speicher lässt sich als gemeinsame Registerbank oder gemeinsamer Cachespeicher auf einem Chip mit einem Interkonnektor implementieren, der es jeder Verarbeitungsmaschine erlaubt, von jedem Ort in einem gemeinsamen Speicher von/an einen beliebigen anderen Ort im Speicher zu lesen bzw. zu schreiben. In anderen Ausführungsbeispielen kann ein gemeinsamer Zustandsraum auf eine, für jedes CTA spezifische Region des Chip-externen Speichers abgebildet und im L1-Cache 320 zwischengespeichert werden. Der Parameterspeicher lässt sich als eine designierte Sektion innerhalb des gleichen gemeinsamen Registerspeichers oder des gleichen gemeinsamen Cache-Speichers implementieren, der den gemeinsamen Speicher 306 implementiert, oder als eine separate gemeinsame Registerbank oder Chip-internen Cache-Speicher, auf den die LSUs 303 einen Nur-lesen-Zugriff haben. In einem weiteren Ausführungsbeispiel wird der Abschnitt, der den Parameterspeicher implementiert, auch verwendet, um die CTA-Kennzeichen und Aufgaben-Kennzeichen sowie Dimensionen von CTA und Grid oder Warteschlangenpositionen zu speichern, so dass Teilmengen des Sonderregisters implementiert werden. Jede LSU 303 in der SM 310 ist an eine Unified Adress Mapping Unit 352 gekoppelt, die eine Adresse, die für Lade- und Speicher-Befehle bereitsteht, die in einem vereinten (uniformen) Speicherraum festgelegt sind, in eine Adresse in jedem abgegrenzten Speicherraum übersetzt. Als Folge davon lässt sich ein Befehl verwenden, um auf jeden der lokalen, gemeinsamen oder globalen Speicherräume zuzugreifen, indem man eine Adresse im vereinten Speicherraum angibt.
Der L1-Cache 320 in jeder SM 310 lässt sich dazu verwenden, private und lokale Prä-Thread-Daten sowie globale Pro-Applikation-Daten zwischen zu speichern. In manchen Ausführungsbeispielen lassen sich die für jedes CTA gemeinsam vorliegenden Daten im L1-Cache zwischenspeichern. Die LSUs 303 sind über einen Speicher- und Cache-Interkonnektor 380 an den gemeinsamen Speicher 306 und den L1-Cache 320 gekoppelt.
Verschachtelte Ausführungs-Streams
Verschachtelter Parallelismus ermöglicht es Threads in einer Thread-Gruppe, wie sie oben beschrieben wurden, unabhängig voneinander einen oder mehr Tochter-Threads zu starten und entweder Thread-Synchronisationsgrenzen in einer Tochter-Thread-Gruppe einzuhalten oder in einer Mehrzahl von Kind-Thread-Gruppen die korrekte Ausführungsreihenfolge beizuhalten. Diese grundlegende Fähigkeit ermöglicht eine breitgestreute Klasse von Algorithmen, die es erfordern, dass eine bedingte Ausführung effizient auf das Parallelverarbeitungssubsystem 112 abgebildet wird. Stream-übergreifende Abhängigkeiten zwischen Aufgaben, die Thread-Gruppen enthalten, werden ermöglicht, indem ein Pfad für hierarchische Ausführung transformiert wird, um auslösende Events und solche wartende Events zu umfassen, die auf einen auslösenden Event oder eine Aufgabe, die vor der Fertigstellung steht, warten. Indem Stream-übergreifende Abhängigkeiten im Parallelverarbeitungssubsystem 112 unterstützt werden, ist es möglich, das ein vollständigeres semantisches Modell für die Berechnungen realisiert wird, ohne dass Ineffizienzen hervorgerufen werden, wie sie mit der CPU-basierten Aufgabenverwaltung verbunden sind.
4 illustriert gemäß einem Ausführungsbeispiel der vorliegenden Erfindung eine verschachtelte Aufgabenausführung im Parallelverarbeitungssubsystem 112. Wie dargestellt initiiert die CPU 102 eine Ausführung exemplarischer Aufgaben 420 im Parallelverarbeitungssubsystem 112. Nachdem Aufgabe 420(0) abgeschlossen ist, wird Aufgabe 420(1) ausgeführt. Nachdem Aufgabe 420(1) abgeschlossen ist, wird Aufgabe 420(2) ausgeführt. Während des Ablaufs der Ausführung ruft Aufgabe 420(1) beispielsweise die Aufgaben 430(0) bis 430(2) auf, um ein Zwischenergebnis zu berechnen, das dann von Aufgabe 420(1) verwendet wird. Um die richtige Reihenfolge der Ausführung aufrechtzuerhalten, sollte die Aufgabe 420(1) warten, bis die Aufgaben 430 abgeschlossen sind, bevor sie fortfährt. Um so zu warten, kann Aufgabe 420(1) an einer Thread-Synchronisationsbarriere für die Aufgaben 430 eine Blockade errichten. Jede Aufgabe 420 oder 430 lässt sich von einem oder mehr Threads, CTAs oder Grids ausführen, wie sie oben definiert wurden.
In diesem Beispiel ist Aufgabe 420(1) ein Elternteil der Aufgaben 430, welche daher seine Töchter sind. Obwohl nur eine Ebene der Eltern-Tochter-Hierarchie (also Verschachtelungstiefe) in 4 gezeigt wird, ist es in der Praxis möglich, eine beliebig tiefe Hierarchie zu etablieren. In einem Ausführungsbeispiel wird die Verschachtelungstiefe durch eine Anzahl von Terminierungsgruppen begrenzt. In diesen Terminierungsgruppen lässt sich die Priorität festlegen, die die Tochter-Ausführung gegenüber der Eltern-Ausführung haben soll. In einem Ausführungsbeispiel führen die Aufgaben 420 und 430 jeweils mindestens eine Thread-Gruppe aus, oder zumindest eine CTA in SM 310 von 3B. Um es Thread-Programmen, die eine Eltern-Tochter-Beziehung aufweisen, zu ermöglichen, auf SM 310 ausgeführt zu werden, sollten drei Systemelemente implementiert werden: Hardware-Funktionalität für Parallelverarbeitungssubsystem 112, Software-Runtime-Funktionalität für Parallelverarbeitungssubsystem 112 und schließlich Konstrukte für die Sprachunterstützung des Parallelverarbeitungssubsystems 112.
Die Hardware-Funktionalität, die zur Unterstützung des Starts eines Tochter-Threads, CTA oder Grids durch einen Eltern-Thread im Parallelverarbeitungssubsystem 112 erforderlich ist, umfasst den Start eines neuen Grids oder CTAs von Arbeit aufgrund einer Anfrage, die von SM 310 erzeugt und zur Ausführung in die Aufgaben/Arbeitseinheit 207 in eine Warteschlange gestellt wurde. Dadurch wird der Ausführungsstatus für SM 310 gespeichert, auf dieser Grundlage die Ausführung in SM 310 weitergeführt und die Speicherkohärenz zwischen einer Elter- und einer Tochter-Aufgabe ermöglicht. Zu den Laufzeitmerkmalen, die für den Start eines Tochter-Threads, eines CTA oder Grids im Parallelverarbeitungssubsystem 112 durch einen Elter-Thread nötig sind, gehören: Starten eines neuen Grids als Reaktion auf eine Anfrage durch einen im SM 310 ausgeführten Thread; einen Eltern-Thread in die Lage versetzen, eine Tochter-Synchronisationsbarriere auf eine Tochter-Threadgruppe anzuwenden; Speicherkohärenz zwischen dem Eltern-Thread und der Tochtergruppe sicherstellen; Terminieren von Arbeit und Weiterführung synchronisierter Threadgruppen im Hinblick auf den sichergestellten voranschreitenden Berechnungsvorgang; und Sicherstellen korrekter Ausführungssemantik für Eltern-Threads und Tochtergruppen.
Zu den Sprachunterstützungskonstrukten gehört ein Mechanismus, um den Start eines Tochter-Thread-Programms durch einen Eltern-Thread zu spezifizieren, sowie die Anwendung einer Synchronisationsbarriere auf den Tochter-Thread.
Das Parallelverarbeitungssubsystem 112 wird mit Hilfe einer Thread-orientierten Programmierumgebung wie zum Beispiel CUDA^TM von NVIDIA^TM programmiert. In einem Ausführungsbeispiel wird die CUDA-Sprachspezifikation erweitert, um ein Konstrukt für den Start einer Tochter aufzuweisen, das die Details für den Start eines Tochter-CUDA-Grids spezifiziert. Das Tochter-Start-Konstrukt, das hier als ”A<<<B>>>C” angezeigt wird, enthält den Namen eines Tochter-Programms (A), die Grid-Startparameter (B) und Eingabeparameter (C). Die CUDA-Laufzeitumgebung ist erweitert worden, damit ein Eltern-Thread eine Synchronisationsbarriere auf dem Tochter-CUDA-Grid ausführen kann. Obwohl die aktuelle Erörterung Ausführungsbeispiele der Erfindung im Kontext der CUDA-Programmierumgebung illustriert, wird ein Fachmann erkennen, dass die hier offengelegten Techniken sich auf jede Parallelprogrammierungsumgebung und jedes parallelverarbeitende System anwenden lassen. Daher dienen Hinweise auf CUDA lediglich illustrativen Zwecken und sind nicht dazu gedacht, den Umfang oder den Sinn der aktuellen Offenlegung einzuschränken.
Die untenstehende Tabelle 1 illustriert die Verwendung des Tochter-Start-Konstrukts und der Synchronisationsbarriere in einem beispielhaften CUDA-Programm:
In diesem Beispiel startet eine Instanz des Thread-Programms ”foo()” einen Tochter-Grid unter Verwendung eines Thread-Programms ”A” mit einem Pointer (*ptr) auf Speicher, den foo() zugewiesen hat. Auf den zugewiesenen Speicher können Threads innerhalb des Tochter-Grids zugreifen. Der Eltern-Thread ”foo()” kann fortgesetzt werden, nachdem Tochter-Grid A beendet wurde. Das wird durch eine Rückkehr von einem Funktionsaufruf für eine blockierende Synchronisationsbarriere, die hier cudaThreadSynchronize() genannt ist, angezeigt.
5 stellt gemäß einem Ausführungsbeispiel der vorliegenden Erfindung eine Sequenz von Ereignissen 500 dar, in deren Verlauf ein Eltern-Thread-Programm einen Tochter-Grid startet. Mit Hinblick auf Tabelle 1 kann der Eltern-Thread eine Instanz des Thread-Programms foo() sein, während der Tochter-Grid den Tochter-Grid (A) aufweisen kann, der gemäß des Grid-Start-Konstrukts gestartet wurde.
Die Ereignisabfolge 500 beginnt zum Zeitpunkt T0, an dem ein Softwaretreiber für das Parallelverarbeitungssubsystem 112, wie etwa der Gerätetreiber 103, im SM 310 ein CTA initialisiert. Beispielsweise kann der Softwaretreiber foo() aus Tabelle 1 als ein CTA in SM 310 initialisieren. Daraufhin erfolgt der Beginn der Ausführung des CTA als mindestens ein Thread. Ein Thread von ”foo()” mit dem Thread-Kennzeichen x==0, y==0 beginnt beispielsweise mit der Ausführung. Der Thread führt sodann einen Tochter-Start aus, indem er das Tochter-Start-Konstrukt ”A<<<1, 1>>>” verwendet. Der Tochter-Startvorgang, der weiter unten eingehender beschrieben wird, veranlasst den Beginn der Ausführung von Tochter-Grid A. An diesem Punkt ist der Thread ein Eltern-Thread des Tochter-Grid A. Der Eltern-Thread führt einen Aufruf an cudaThreadSynchronize() aus, das wiederum die Ausführung des Elter-Thread solange blockiert, bis der Tochter-Grid A abgeschlossen ist. Weil der Aufruf von cudaThreadSynchronize() der Synchronisation eines Tochter-Grids dient, kann das mit dem Eltern-Thread verknüpfte CTA angehalten werden, während das CTA auf den Abschluss des Tochter-Grids wartet. Um das CTA anzuhalten, wird dessen Ausführungsstatus in einem Fortsetzungszustandspuffer gespeichert. Dieser Fortsetzungszustandspuffer kann sich im PP-Speicher 204 befinden, im Systemspeicher 104 oder jedem anderen technisch machbaren Speichersubsystem, das mit dem Parallelverarbeitungssubsystem 112 verbunden ist. Das CTA setzt seine Ausführung aus und wartet auf den Abschluss von Tochter-Grid A. Durch das Aussetzen wird SM 310 frei und kann die Zuweisung anderer Ausführungsaufgaben entgegennehmen, solange Tochter-Grid A ausgeführt wird. In einem Ausführungsbeispiel werden Ressourcen, die mit dem CTA verknüpft und vom SM 310 allokiert worden waren, de-allokiert, also freigegeben. Eine Blockade wird vermieden, weil das CTA foo() Ressourcen freigibt, um das Voranschreiten des Berechnungsvorgangs zu ermöglichen.
Sobald Tochter-Grid A abgeschlossen ist, wird ein Scheduler wie etwa die Aufgaben/Arbeitseinheit 207 aus 2 benachrichtigt. Das CTA, das den Eltern-Thread aufweist, wird dann so terminiert, dass es neu gestartet wird, wobei sein Ausführungszustand aus dem Fortsetzungszustandspuffer wiederhergestellt wird. Das CTA wird benachrichtigt, dass der Tochter-Grid abgeschlossen worden ist, wodurch die Ausführung des Eltern-Thread durch die Synchronisationsbarriere, die zuvor durch einen Aufruf von cudaThreadSynchronize() erzeugt worden ist, hindurch weitergeführt werden kann. Daraufhin wird das CTA zu Ende geführt und jeder damit verknüpfte Grid von CTAs ist in der Lage, abgeschlossen zu werden.
In einem Ausführungsbeispiel weist der Prozess des Starts eines Kind-Grids durch einen Eltern-Thread die Vorbereitung der Speicherelemente im Parallelverarbeitungssubsystem 112 dahingehend auf, dass dem Tochter-Grid eine konsistente Sicht auf, den Speicherraum präsentiert wird, der mit dem Elternteil verknüpft ist. Um den Tochter-Grid zu starten, lassen sich CUDA-Laufzeitaufrufe nutzen, darunter eine Barriere zum Systemspeicher (membar.sys), um so die Konsistenz des Speichers zwischen Elternteil und Tochter sicherzustellen. Das Einrichten der Schranke zum Systemspeicher bewirkt, dass alle Daten, die vom Eltern in den Speicher – beispielsweise PP-Speicher 204 – geschrieben werden sollen, gelöscht werden. Das Löschen aller anstehenden Schreibvorgänge erlaubt es jedem Thread im Tochter-Grid, beliebige Daten aus dem Speicherraum des Elternteils zu lesen, während er, der Thread, auf einem beliebigen SM 310 ausgeführt wird. Ein Fachmann wird verstehen, dass das Löschen von Cachespeichern ein Verfahren ist, um Speicherkonistenz sicherzustellen, aber es stehen auch andere Verfahren zur Verfügung.
Sobald die Speicherbarriere eingerichtet worden ist, kann ein ein CUDA-Laufzeitaufruf, hier als cudaRTLaunch() bezeichnet, ausgeführt werden, um den Tochter-Grid zu starten. In einem Ausführungsbeispiel stellt der CUDA-Laufzeitaufruf an cudaRTLaunch() eine neue Aufgabe zur Ausführung in den Scheduler. Dazu präsentiert der Aufruf einer Speicherverwaltungseinheit eine Posted-Compare-and-Swap-Nachricht (PCAS). Die Speicherverwaltungseinheit, beispielsweise MMU 328 in 3B, leitet die Nachricht zurück an den Scheduler. Eine solche PCAS-Operation stellt ein Verfahren dar, wie SM 310 Arbeit terminiert. Die PCAS-Operation wird als eine blockierende (”posted”) Synchronisationsoperation implementiert, die von der MMU 328 ausgeführt wird. Die PCAS-Operation vergleicht Stück für Stück (”atomisch”) einen aktuellen Speicherwert an einer bestimmten Speicherstelle mit einem ersten spezifizierten Wert und überschreibt die Speicherstelle mit einem zweiten spezifizierten Wert, falls – und nur dann – der aktuelle Speicherwert mit dem ersten spezifizierten Wert übereinstimmt.
Der Aufruf cudaThreadSynchronize() zur Einrichtung einer Threadsynchronisationssperre, die in Tabelle 1 gezeigt wird, initiiert eine Serie von Ereignissen, um den aktuellen Ausführungsstatus des aufrufenden (und baldigen Eltern-) Threads zu speichern. In diesem Beispiel ist der aufrufende Thread foo(). Der Aufruf an cudaThreadSynchronize() kann explizit auf einen Tochter-Grid verweisen, der von dem aufrufenden Thread gestartet wird. In einem Ausführungsbeispiel wird das Parallelverarbeitungssubsystem 112 so konfiguriert, dass es alle relevanten Ausführungszustände eines CTA, der in einem GPC 208 ausgeführt wird, einschließlich aller relevanten Architekturzustände für jede verknüpfte SM 310, in den Fortsetzungsstatuspuffer sichert. In bestimmten Ausführungsbeispielen residiert ein oder mehrere Fortsetzungsstatuspuffer an vorher festgelegten Orten im Speicher.
Sobald der Tochter-Grid abgeschlossen ist, wird ein Terminierungsverfahren aufgerufen, um zu entscheiden, welche Aufgabe(n) als nächste auszuführende terminiert werden soll(en). Ein Thread-Programm wird zudem gemeinhin als ein ”Kernel” bezeichnet. In einem Ausführungsbeispiel wird ein Scheduler-Kernel als eine ”AtExit”-Aufgabe aufgerufen, die ausgeführt wird, sobald ein Grid, wie etwa der Tochter-Grid, abgeschlossen wird. Der Scheduler-Kernel entdeckt, welche damit verbundene Arbeit abgeschlossen wurde und welche Arbeit es noch zu tun gibt. Der Scheduler-Kernel entscheidet dann, was als nächstes zur Ausführung terminiert werden soll. In dem Beispiel der Tabelle 1 startet der Scheduler-Kernel den CTA foo() erneut.
Der Vorgang des erneuten Starts eines Eltern-CTAs nach dem Abschluss eines Tochterprozesses wird in diesem Dokument als 'Fortsetzungsprozess' bezeichnet. In einem Ausführungsbeispiel gehört dazu die Ausführung eines Wiederherstellungs-Kernels, der als ”AtEntry”-Aufgabe aufgerufen wurde, um für den Grid den Ausführungszustand aus einem Fortsetzungszustandspuffer wiederherzustellen. Während sich ein Wiederherstellungs-Kernel für die Wiederherstellung eines Ausführungszustands implementieren lässt, so kann doch jede realisierbare Technik implementiert werden, um einen Ausführungszustand wiederherzustellen, ohne vom Umfang und dem Sinn der aktuellen Erfindung abzuweichen. Um bei der fortgesetzten Ausführung die Speicherkonsistenz sicherzustellen, werden Cachespeicher für SM 310, in dem der Grid ausgeführt wird, für ungültig erklärt. Dadurch werden alle irrelevanten oder überholten Zustandsangaben, die von einem vorhergehenden Grid im gleichen SM 310 übriggeblieben sein könnten, zwangsläufig gelöscht. Ein Fachmann wird verstehen, dass das Löschen von Cachespeichern ein Verfahren darstellt, um für Speicherkonsistenz zu sorgen, dass es aber auch andere Verfahren gibt. Nach der Wiederherstellung des Ausführungszustands und der Sicherstellung der Speicherkonsistenz nimmt der Wiederherstellungs-Kernel die Ausführung des Elter-CTA wieder auf, indem er auf die Instruktion springt, die dem Aufruf von cudaThreadSynchronize() folgt.
Wie in jedem CUDA-Thread-Programm wird auch hier jeder Thread unabhängig ausgeführt. Wenn also ein Thread-Programm gemäß Befehl einen unbedingten Start ausführen soll, wird daher jede ausführende Instanz des Thread-Programms diesen Startbefehl ausführen. In dem exemplarischen CUDA-Programm, das in Tabelle 2 zu sehen ist, führt jede Instanz von foo() einen unbedingten Start des Kernels ”bar()” aus, wodurch sich eine Summe von einhundert ausgeführten bar()-Kerneln ergibt.
Eine abgewandelte Version des CUDA-Programms in Tabelle 2 ist in Tabelle 3 dargestellt. In diesem modifizierten CUDA-Programm wird nur ein Thread (mit x==0) von foo() ausgeführt, so dass tatsächlich nur dieser eine Thread von foo(), der unabhängig ausgeführt wird, einen Tochter-Grid startet. In diesem Beispiel wird insgesamt nur ein Tochter-Grid, der das Thread-Programm bar() enthält, aus allen hundert ausgeführten Instanzen von foo() gestartet.
Obwohl jeder CUDA-Thread unabhängig ausgeführt wird, lassen sich doch CUDA-Befehlskontrukte von allen Threads innerhalb eines Thread-Blocks nutzen. Beispielsweise darf ein Thread einen CUDA-Stream erzeugen, und jeder andere Thread innerhalb des Thread-Blocks darf diesen Stream verwenden. Jeder Thread-Block bleibt allerdings weiterhin unabhängig. Ein exemplarisches CUDA-Thread-Programm wird in Tabelle 4 gezeigt, wo ein Eltern-Thread-Programm foo() einen Stream erzeugt, der gemeinsam von Threads im Tochter-CTA des Thread-Programms bar() genutzt wird.
In einem Ausführungsbeispiel wird eine Hierarchie von ausgeführten Kerneln, die einen Eltern- und beliebig viele Tochter-Kernel aufweisen, zusammensetzbar strukturiert. Das bedeutet, dass von außerhalb der Hierarchie nur der Eltern-Kernel sichtbar ist. Anders ausgedrückt, wenn ein Eltern-Kernel Tochter-Kernel startet, dann erscheinen die Tochter-Kernel als Bestandteile des Eltern-Kernels. Das bedeutet wiederum, dass sowohl der Eltern-Kernel als auch alle Tochter-Kernel ihre Ausführung abschließen müssen, bevor die Ausführung des Elternteils als abgeschlossen betrachtet werden kann. In einem Beispiel ausgedrückt bedeutet dies, dass die gesamte Arbeitslast, die mit der Aufgabe 520(0) in 5 verbunden ist, erst erledigt sein muss, bevor Aufgabe 520(1) gestartet werden kann. In ähnlicher Weise muss die gesamte Arbeit, die mit Aufgabe 520(1) und den Kind-Aufgaben 530(0) bis 530(2) verbunden ist, erst abgeschlossen sein muss, bevor Aufgabe 520(1) als abgeschlossen betrachtet wird. Sobald die Aufgabe 520(1) abgeschlossen ist, darf Aufgabe 520(2) gestartet werden.
Es ist vorteilhaft, dass Ausführungsbeispiele der vorliegenden Erfindung einen Eltern-Thread in die Lage versetzen, aus dem Parallelverarbeitungssubsystem 112 einen oder mehrere Grids von Tochter-Threads unabhängig voneinander zu starten. Die Fähigkeit, Tochter-Threads zu starten, ermöglicht neue Ansätze und Verfahren im Hinblick auf den Entwurf von Thread-Programmen. So kann ein Thread-Programm beispielsweise eine extern definierte Bibliotheksfunktion aufrufen, indem es einen entsprechenden Kernel startet, der diese Bibliotheksfunktion als ein Tochterprozess implementiert. In einem weiteren Beispiel lässt sich die Kontrolle über einen substantiellen Teil des Datenflusses in einer Anwendung ausüben, indem Thread-Programme in Parallelverarbeitungssubsystem 112 ohne Intervention seitens einer damit verbundenen Anwendung, die in CPU 102 läuft, ausgeführt werden. Wenn man zum Beispiel eine übergreifende Kontrollschleife für eine Anwendung mit einem oder mehr Threads, die im Parallelverarbeitungssubsystem 112 ausgeführt wird, implementiert, kann man bewirken, dass die Anwendung auf effiziente Weise datenabhängige Algorithmen, rekursive Algorithmen und Algorithmen mit komplexen Ausführungsbedingungen, wie etwa mehrfach verschachtelte bedingte Schleifen, ausführt.
Eine exemplarische Funktion, main(), die verschachtelte bedingte Schleifen aufweist, ist in Tabelle 5 dargestellt. Diese Funktion weist eine unbedingte äußere Schleife mit dem Schleifenindex ”i” auf, eine eingebettete bedingte Schleife mit dem Index ”j” und eine bedingte Ausführung der Funktion do_stuff() auf. Das statische Entfalten von main() ist ebenso wenig machbar wie der Versuch, zu berechnen, welche bedingten Schleifenbereiche ausgeführt werden sollten. Konventionelle Parallelverarbeitungssysteme sind daher nicht in der Lage, diese und andere Typen bedingter Ausführungsstrukturen zu verarbeiten, die vielen wichtigen Klassen von Algorithmen im Fachgebiet inhärent sind.
Ausführungsbeispiele der vorliegenden Erfindung stellen das Konstrukt für den Kind-Start bereit, das ein unabhängig ausgeführtes Thread-Programm in die Lage versetzt zu berechnen, wann die Funktion do_stuff() aufzurufen ist. In Tabelle 6 wird eine Instanz von cuda_inner_loop() bedingt als ein Tochter-Grid gestartet, der die Anzahl kmax Instanzen nur in jenen Iterationen aufweist, in denen die Bedingung ”condition2” zutrifft. Die Funktion do_stuff() wird bedingt in cuda_inner_loop() ausgeführt, wenn die Bedingung ”condition3” zutrifft. Man beachte, dass jeder Tochterprozess des gestarteten cuda_inner_loop() vorteilhafterweise asynchron und im Gleichklang mit anderen Instanzen von cuda_inner_loop() ausgeführt wird. Das Implementieren dieser Schleifenstruktur unter Verwendung des konventionellen GPU-Modells wäre kompliziert und ineffizient gewesen, denn die CPU hätte jeden Tochter-Grid von cuda_inner-loop bedingt dann starten müssen, wenn die Bedingung condition2 zugetroffen hätte.
Ein weiteres Ausführungsbeispiel der eingebetteten Schleife, die in den Tabellen 5 und 6 gezeigt wird, ist in Tabelle 7 zu sehen. Hier kann die Bedingung für die äußere Schleife parallel dazu durch separate Threads bewerten, und jede Schleifenebene lässt sich nun bedingt als Kinder starten, sofern die passende(n) Bedingung oder Bedingungen zutreffen.
6 stellt Systemelemente für ein verschachteltes Ausführungssubsystem 600 gemäß einem Ausführungsbeispiel der vorliegenden Erfindung dar. Das verschachtelte Kernel-Ausführungssubsystem 600 umfasst Hardware- und Software-Strukturen, die mit dem Parallelverarbeitungssubsystem 112 aus 1 implementiert werden. Dazu gehört eine Warteschlange 650 von Metadatendeskriptoren von Gridaufgaben (GTMD: grid task metadata descriptor), welche die Anwendungsarbeit 612 empfängt und speichert, beispielsweise von der CPU 102 in 1. Die Anwendungsarbeit 612 enthält eine sortierte Abfolge von GTMDs, die als Aufgabe1 bis AufgabeN etikettiert sind. Der Scheduler 610 ist so eingerichtet, dass er jeden GTMD empfängt und über den Verteiler einen entsprechenden Grid für die Ausführung in SM 630 terminiert. Der Verteiler dient allokierten Threads als CTAs innerhalb der SM 630. Der Fortsetzungsstatuspuffer, der in 5 erörtert wurde, kann im Fortsetzungsstatuspuffer 642 gespeichert werden, der sich in Speicher 640 befindet. In einem Ausführungsbeispiel umfasst der Scheduler 610 die Aufgabenverwaltungseinheit 300 aus 3A, der Verteiler 620 enthält die Arbeitsverteilungseinheit 340, SM 630 enthält SM 310 aus 3B, und der Speicher 640 umfasst PP-Speicher 204, Systemspeicher 104 oder eine Kombination aus beiden.
Wenn ein in SM 630 ausgeführter Thread einen Kind-CTA startet, wird für den Kind-CTA ein neuer GTMD erzeugt und in die GTMD-Warteschlange 652 zwecks späterer Ausführung gestellt. Der Scheduler 610 ist in der Lage, zwischen neuer Anwendungsarbeit 612, die von der GTMD-Warteschlange 650 kommt, und verschachtelter Prozessarbeit, die in GTMD-Warteschlange 652 eintrifft, zu unterscheiden, weil jede Menge von Arbeitslasten in jeweils separaten Warteschlangen gespeichert wird. Der Scheduler 610 kann den Arbeitslasten in jeder GTMD-Warteschlange zu unterschiedlichen Zeiten verschiedene Ausführungsprioritäten zuweisen und dabei jede technisch machbare Technik nutzen, die den vorwärts ablaufenden Ausführungsvorgang sicherstellt.
Wenn ein Grid, der in SM 630 ausgeführt wird, in Reaktion auf den Aufruf von cudaThreadSynchronize() angehalten wird, wird der Ausführungszustand in den Fortsetzungszustandspuffer 642 gesichert, und ein Scheduler-Kernel wird in die Warteschlange zur Ausführung gestellt. In einem Ausführungsbeispiel wird der Scheduler-Kernel zur Ausführung in eine Warteschlange QTMD 654 (QTMD: queue of task metadata descriptors) gestellt, in der sich Metadatendeskriptoren für Aufgaben befinden. Beispielsweise kann task7 einen Deskriptor für einen Scheduler-Kernel enthalten, der von einer CTA, die zuvor SM 630 verlassen hat, in die Warteschlange QTMD 654 gestellt worden ist, und taskP kann den aktuellsten Scheduler-Kernel enthalten, der von einer CTA, die erst jüngst SM 630 verlassen hat, zwecks Ausführung in die Warteschlange gestellt worden ist.
In einem Ausführungsbeispiel wird ein Wiederherstellungs-Kernel ausgeführt, um Ressourcen in SM 630 auf die Wiederaufnahme der Ausführung einer CTA vorzubereiten, die zuvor durch eine Ausführung von cudaThreadSynchronizer() den SM 630 verlassen hat. In bestimmten Implementierungen wird der Wiederherstellungs-Kernel zur Ausführung in QTMD 654 in eine Warteschlange gestellt. In alternativen Implementierungen wird der Wiederherstellungs-Kernel, beispielsweise taskQ, zur Ausführung in eine separate QTMD-Warteschlange 656 gestellt, um größere Flexibilität bei der Terminierung zu erhalten. Während die Ausführungswiederherstellung einer CTA oben in Begriffen eines Wiederherstellungs-Kernels beschrieben worden ist, lässt sich jede andere technisch machbare Technik, die Ausführung der CTA wiederherzustellen, ebenfalls implementieren, ohne von Umfang und Sinn der vorliegenden Erfindung abzuweichen.
Das Subsystem 600 für die Ausführung verschachtelter Kernels stellt ein System für das Aussetzen und Fortsetzen der Ausführung beliebiger Thread-Gruppen innerhalb eines Parallelverarbeitungssubsystems bereit, während es sich um Speicherkonsistenz und korrekte Ausführungssemantik für jeden Thread und jede Eltern-Tochter-Beziehung in einem Graphen für die hierarchische Ausführung kümmert.
Gemäß einem Ausführungsbeispiel in der vorliegenden Erfindung illustriert 7 einen solchen exemplarischen hierarchischen Ausführungsgraphen einschließlich der verknüpften Aufgaben-Metadaten-Warteschlangen und Aufgaben. Wie 7 zeigt, weist der hierarchische Ausführungsgraph die Thread-Gruppe 710 mit der Verschachtelungstiefe 0, Aufgaben-Metadaten-Deskriptor-Warteschlangen (TMDQs) 712, die Aufgaben 720 730 740, einen Ausführungsgraphen 780 auf der Verschachtelungstiefe 1 sowie einen Ausführungsgraphen 790 auf der Verschachtelungstiefe 2 auf.
Die Thread-Gruppe 710 mit der Verschachtelungstiefe 0 weist Threads auf, die von der CPU 102 erzeugt und verwaltet werden. Eine Thread-Gruppe weist jede Menge von Threads auf, darunter eine CTA, in der alle Threads auf der gleichen Verschachtelungstiefe vorliegen. Die Verschachtelungstiefe eines Threads ist gleich der Anzahl der Eltern-Grids über der Ebene des Threads. Ein CPU-Thread beispielsweise weist die Verschachtelungstiefe 0 auf, weil es über einem CPU-Thread keine Eltern-Grids gibt. Sollte dieser CPU-Thread einen Grid starten, dann wird diesem Grid die Verschachtelungstiefe 1 zugewiesen. Wenn ein Thread im Grid auf der Verschachtelungstiefe 1 einen neuen Grid startet, dann hat der neue Grid die Verschachtelungstiefe 2 und so weiter. Weil die Threads in Thread-Gruppe 710 CPU-Threads sind, weist jeder dieser Threads die Verschachtelungstiefe 0 auf.
Die TMDQs 712 enthalten Zeiger auf Datenstrukturen für anstehende Aufgaben, wie gleich beschrieben wird. Jede TMDQ 712 zeigt auf Aufgaben, die zu einem oder mehreren Streams gehören. Die TMDQ(0) 712(0) zeigt auf die Aufgabe 720(0), die mit einem ersten Stream verknüpft ist. Die TMDQ(1) 712(1) zeigt auf die Aufgaben 730(0) und 730(1), die mit einem zweiten Stream verknüpft sind. Die TMDQ(2) 712(2) zeigt auf die Aufgaben 734(0), 740(1) und 740(2), die mit einem dritten Stream verknüpft sind. Jede beliebige Anzahl von TMDQs 712 lässt sich definieren, solange jede TMDQ 712 eine beliebige Anzahl von Aufgaben enthält.
Die Aufgaben 720, 730, 740 enthalten Datenstrukturen, zu denen ein oder mehr Befehl(e) gehören, der/die auf der GPU auszuführen ist/sind. Die Aufgaben, die auf einer bestimmten TMDQ 712 gestartet werden, werden in sequentieller Reihenfolge nacheinander ausgeführt. Die Aufgabe 730(0) wird abgeschlossen, bevor Aufgabe 730(1) mit der Ausführung beginnt. In gleicher Weise wird Aufgabe 740(0) erst abgeschlossen, bevor Aufgabe 740(1) mit ihrer Ausführung beginnt, welche ihrerseits erst endet, bevor Aufgabe 740(2) starten kann. Eine Aufgabe am Kopf der Warteschlange TMDQ 712 beginnt mit ihrer Ausführung, sobald sie gestartet worden ist. Die Aufgaben 720(0), 730(0) und 740(0) werden also ausgeführt, sobald diese Aufgaben gestartet worden sind. Die Aufgaben in verschiedenen TMDQs 712 weisen keine sequentiellen Abhängigkeiten auf. Aufgabe 730(1) könnte beispielsweise entweder vor, nach oder während Aufgabe 740(1) ausgeführt werden.
Der Ausführungsgraph 780 auf der Verschachtelungstiefe 1 ist eine Thread-Gruppe, inklusive der zugehörigen TMDQs und Aufgaben. Sie wurde von einer der Aufgaben auf der Verschachtelungstiefe 0 gestartet. Jede Aufgabe darf einen oder mehr Grids starten, wobei solche Grids eine Verschachtelungstiefe aufweisen müssen, die um eins größer ist als bei der Aufgabe, die den Grid startete. Wie gezeigt wurde, hat Aufgabe 740(1), die auf der Verschachtelungstiefe 0 vorliegt, während der Ausführung der Aufgabe 740(1) einen Ausführungsgraphen 780 gestartet. Jede Aufgabe und TMDQ in diesem Ausführungsgraphen 780 fungieren im Wesentlichen ebenso wie Aufgaben und TMDQs auf der Verschachtelungstiefe 0. Sobald jede Aufgabe im Ausführungsgraphen 780 abgeschlossen ist und alle anderen Befehle in Aufgabe 740(1) ausgeführt worden sind, darf Aufgabe 740(2) mit ihrer Ausführung beginnen.
Der Ausführungsgraph 790 auf der Verschachtelungstiefe 2 ist eine Thread-Gruppe inklusive zugehöriger TMDQs und Aufgaben, welche von einer der Aufgaben auf der Verschachtelungstiefe 1 gestartet wurde. Jede Aufgabe und TMDQ im Ausführungsgraph 790 arbeitet im Wesentlichen ebenso wie die Aufgaben und TMDQs auf niedrigeren Verschachtelungstiefen. Sobald jede Aufgabe im Ausführungsgraph 790 abgeschlossen ist, darf die Startaufgabe abgeschlossen werden, sobald alle anderen Befehle in der Startaufgabe abgeschlossen worden sind. Auf diese Weise wird die sequentielle Ausführung in jedem Grid aufrechterhalten, und Grids dürfen zu jeder beliebigen Verschachtelungstiefe verschachtelt werden, solange sie die sequentielle Ausführung von Aufgaben in einem Stream beibehalten.
Threads innerhalb einer Thread-Gruppe werden innerhalb eines Kontextes definiert, wobei der Kontext aus dem Satz von Threads besteht, der Zugang zu den gleichen Ressourcen in Stream und TMDQ hat. Threads im gleichen Kontext dürfen TMDQs erzeugen und gemeinsam nutzen, solange die Threads die gleiche Verschachtelungstiefe aufweisen und sich auf dem gleichen Gerät (GPU oder CPU 102) befinden. Für CPU-Threads wird der Kontext als jener Satz von Threads festgelegt, der mit dem CUDA-Kontext verknüpft ist. Für GPU-Threads kann der Kontext ein Cooperative Thread Array (CTA) oder eine beliebige Menge von Threads darstellen, die auf der gleichen Verschachtelungstiefe vorliegen.
Erzeugt ein CPU-Thread einen neuen Stream, allokiert die CPU 102 automatisch Speicher, um die Verwaltung dieses Streams zu unterstützen. Sobald die Aufgaben des Streams abgeschlossen sind, wird der Stream gelöscht und die CPU gibt den zuvor dem Stream allokierten Speicher wieder frei. Die GPU ist typischerweise nicht in der Lage, Speicher dynamisch zu allokieren. Die GPU präallokiert daher Kontextdaten für jeden Kontext, der simultan ausgeführt werden könnte. Daraus ergibt sich, dass eine Thread-Gruppe, die mit einem GPU-Grid verknüpft ist, eine festgelegte Anzahl von TMDQs aufweist, die sich während der Ausführung des Grids nicht ändern darf. Ein neuer Stream wird in einem GPU-Grid mit Hilfe des Funktionsaufrufs cudaStreamCreate() erzeugt. Der Funktionsaufruf liefert einen Integerindex, der auf einen der zuvor allokierten TMDQs im Grid zeigt. Um den Stream zu erzeugen, ist also keine dynamische Zuweisung von Speicher vonnöten. Sobald alle Aufgaben in einem GPU-Stream abgeschlossen sind, wird der Stream mit dem Funktionsaufruf cudaStreamDestroy() gelöscht. Da dem GPU-Stream ja kein Speicher dynamisch zugewiesen wurde, braucht cudaStreamDestroy() keinen Speicher in einen allokationsfreien Speicherpool zurückzugeben und kehrt daher einfach zu seinem aufrufenden Programm zurück.
Sobald ein Stream erzeugt worden ist, lassen sich von jedem Thread im verknüpften Kontext neue Aufgaben in den Stream starten. Startet ein Thread eine Aufgabe in einen Stream, könnten sich keine Aufgaben im verknüpften TMDQ befinden oder alle vorhergehenden Aufgaben im TMDQ könnten ihre Ausführung vollendet haben. In einem solchen Fall darf die Aufgabe sofort mit ihrer Ausführung beginnen, sobald sie in den TMDQ geladen wurde. Alternativ könnte der TMDQ eine oder mehr anstehende Aufgaben aufweisen, die ihre Ausführung noch nicht abgeschlossen haben. In einem solchen Fall wird eine neue Aufgabe in den TMDQ geladen, doch die Aufgabe beginnt ihre Ausführung nicht, solange die anstehenden vorhergehenden Aufgaben nicht ihre Ausführung abgeschlossen haben. In beiden Fällen wird die neue Aufgabe mit Hilfe nichtsperrender Operationen, die keine Einmischung seitens der CPU 102 erfordern, in den TMDQ geladen.
8 zeigt gemäß einem weiteren Ausführungsbeispiel der vorliegenden Erfindung einen mit den obigen Ausführungen verbundenen hierarchischen Ausführungsgraphen einschließlich damit verknüpfter TMDQs und Aufgaben. Der hierarchische Ausführungsgraph enthält die Thread-Gruppe 810 auf der Verschachtelungstiefe 1, die TMDQs 812, die Aufgaben 820 830 840 850 860, einen Ausführungsgraphen 880 auf der Verschachtelungstiefe 2 und schließlich einen Ausführungsgraphen 890 auf der Verschachtelungstiefe 3. Die Komponenten des hierarchischen Ausführungsgraphen funktionieren im Wesentlichen wie es oben im Zusammenhang mit 7 beschrieben wurde, mit Ausnahme der folgenden Einzelheiten.
Wie in 8 gezeigt, weist jeder TMDQ 812 der Thread-Gruppe 810 eine oder mehr Aufgabe(n) auf. In einem Beispiel hätte Aufgabe 820(0), die mit Stream 870 verknüpft ist, in den TMDQ 812(0) gestartet werden können, doch Aufgabe 860(0), die mit dem Stream 875 verknüpft ist, wäre zu diesem Zeitpunkt noch nicht gestartet worden. Die Aufgaben 830, die mit einem Stream verknüpft sind, hätten in den TMDQ(1) 812(1) gestartet werden können. Gleichermaßen hätten die Aufgaben 840, die mit einem zweiten Stream verknüpft sind, in den TMDQ(2) 812(2) gestartet werden können, die Aufgaben 850, verknüpft mit einem dritten Stream, hätten in den TMDQ(N) 812(N) gestartet werden können, und alle dazwischen liegenden TMDQ 812 könnten ebenfalls eine oder mehr verknüpfte Aufgaben aufweisen. Zu einem solchen Zeitpunkt könnte ein Thread innerhalb der Thread-Gruppe 810 versuchen, einen neuen Stream 875 zu erzeugen. Die Thread-Gruppe 810 weist jedoch die Verschachtelungstiefe 1 auf und ist mit der GPU verbunden. Weil die GPU nicht in der Lage ist, Speicher dynamisch zu allokieren, könnte ein neuer TMDQ nicht erzeugt werden, um dem neuen Stream 875 Unterstützung zu geben. In einem solchen Fall könnten die Aufgaben 860, die mit dem neuen Stream 875 verknüpft sind, in den TMDQ(0), der gegenwärtig vom Stream 870 genutzt wird, gestartet werden. Der Stream 875 könnte die Aufgaben 860(0) und 860(1) in den TMDQ(0) 812(0) starten. Der Stream 870 könnte dann die Aufgabe 820(1) in den TMDQ(0) 812(0) starten. Der Stream 875 könnte dann die Aufgabe 860(2) in den TDMQ(0) 812(0) starten. Man beachte, dass dieses Verfahren zu unnötigen Abhängigkeiten führt. Obwohl die Streams 870 und 875 voneinander unabhängig sind, führt die sequentielle Natur bzw. Arbeitsweise von TMDQs dazu, dass die Aufgabe 860(0) vom Abschluss der Aufgabe 820(0) abhängt, die Aufgabe 820(1) vom Abschluss der Aufgabe 860(1) und so weiter. Obwohl als Ergebnis die Systemleistung sinken könnte, wird doch die sequentielle Abfolge der Aufgaben 820 in Stream 870 und der Aufgaben 860 in Stream 875 korrekt eingehalten.
9A zeigt gemäß einem Ausführungsbeispiel der vorliegenden Erfindung einen exemplarischen hierarischen Ausführungsgraphen 900 von Aufgaben, die Stream-übergreifende Abhängigkeiten aufweisen. Die Aufgaben 910 werden in einem ersten Stream in die Warteschlange gestellt, die Aufgaben 920 in einem zweiten Stream, die Aufgaben 930 in einem dritten und die Aufgaben 940 in einem vierten. Wie dargestellt hängt die Aufgabe B 910(1) sowohl von der Fertigstellung der Aufgabe A 910(0) ab, als auch von der der Aufgabe D 920(0), bevor sie ausgeführt werden kann. Die Abhängigkeit der Aufgabe B 910(1) von Aufgabe D 920(0) stellt eine Stream-übergreifende Abhängigkeit dar. Wie gezeigt weist auch die Aufgabe G 930(1) Stream-übergreifende Abhängigkeiten gegenüber den Aufgaben E 920(1) und Aufgabe 910(1) auf. In ähnlicher Weise weist auch die Aufgabe J 940(1) eine Stream-übergreifende Abhängigkeit von Aufgabe B 910(1) auf.
In einem Ausführungsbeispiel werden Aufgabenabhängigkeiten einschließlich Stream-übergreifender Abhängigkeiten unter Verwendung der CUDA-Funktionsaufrufe cudaEventCreate(), cudaEventRecord() und cudaStreamWaitEvent() spezifiziert. Die Funktion cudaEventCreate() erzeugt ein Ereignis-Objekt und weist diesem einen Identifikator zu. Auf das Event-Objekt lässt sich mit cudaStreamWaitEvent() als ein Event referenzieren, auf das gewartet werden soll. Ein Event lässt sich in einem entsprechenden Stream mit Hilfe von cudaEventRecord() aufzeichnen. Diese drei Funktionsaufrufe werden bislang in CUDA-Ausführungsmodellen implementiert, die von einem GPU-Treiber verwaltet werden, der in der CPU 102 läuft. In solchen Ausführungsmodellen vorhergehender Technik terminiert die CPU 102 explizit die Aufgabenausführung, um Blockadebedingungen zu verhindern, die als Folge von Stream-übergreifenden Abhängigkeiten auftreten können. Die Ausführung von Aufgaben mit Stream-übergreifenden Abhängigkeiten ohne die Einbindung der Verwaltung durch die CPU 102 ermöglicht hingegen eine höhere Gesamteffizienz des Systems, erfordert aber auch, dass die drei Funktionsaufrufe cudaEventCreate(), cudaEventRecord() und cudaStreamWaitEvent() analoge Implementierungen aufweisen, die im Parallelverarbeitungssubsystem 112 ausgeführt werden können. In einem Ausführungsbeispiel wird der hierarchische Ausführungsgraph in einen gleichwertigen hierarchischen Ausführungsgraphen transformiert, um die Implementierung der drei genannten Aufrufe zu erleichtern. Das wird im Folgenden eingehender beschrieben.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung zeigt 9B einen gleichwertigen hierarchischen Ausführungsgraphen 902 von Aufgaben und Ereignissen für das Durchsetzen der Ausführungsreihenfolge unter abhängigen Aufgaben in verschiedenen Streams. Stream-übergreifende Abhängigkeit wird mit Hilfe zweier neuer Konstrukte gehandhabt, einem Warten-Event (WE) und einem (feuernden) Signalisierungs-Event (SE). Jeder WE blockiert, bis alle Eingabebedingungen erfüllt sind. Jeder Feuern-Event erzeugt einen oder mehr Events, sobald er ausgelöst wird. Die Stream-übergreifende Abhängigkeit der Aufgabe B 910(1) von den Aufgaben A 910(0) und D 920(0) wird, wie gezeigt, mit Hilfe des WE 0 950(0) und des SE 0 952(0) dargestellt. In gleicher Weise wird die Stream-übergreifende Abhängigkeit der Aufgabe G 930(1) von den Aufgaben B 910(1), E 920(1) und F 930(0) mit Hilfe von WE 2(1) 950(2) und WE 1 950(1) dargestellt. Ein Fachmann wird erkennen, dass der hierarchische Ausführungsgraph 900 aus 9A und der gleichwertige hierarchische Ausführungsgraph 902 identische Aufgabenabhängigkeiten implementieren und demzufolge identische Ausführungsreihenfolgen erzwingen. Der hierarchische Ausführungsgraph 902 erfordert keine Sperren und darf daher effizient innerhalb des Parallelverarbeitungssubsystem 112 ohne Einmischung seitens der CPU 102 implementiert werden.
Um die Ausführungsreihenfolge zu erzwingen, kann ein WE in einem bestimmten Stream in eine Warteschlange gestellt werden. WE 0 950(0) wartet beispielsweise auf den Abschluss der Aufgabe 910(0) und darauf, dass SE 0 952(0) zu feuern aufhört. Weil SE 0 952(0) von der Fertigstellung der Aufgabe D 920(0) abhängt, wird die Stream-übergreifende Abhängigkeit der Aufgabe B 910(0) von sowohl Aufgabe A 910(0) als auch Aufgabe D 920(0) korrekt erzwungen.
Auf der CPU 102 darf Speicher mit minimalen ”Kosten” allokiert und freigegeben werden. Auf einer GPU jedoch darf Speicher nur allokiert oder freigegeben werden, wenn die GPU untätig ist. Diese Einschränkung hat Folgerungen für die Implementierung von cudaEventCreate() für die Ausführung auf der GPU. Um diese Einschränkungen bei der Speicherzuweisung zu überwinden, wird ein prä-allokierter Pool von Event-Strukturen eingerichtet, der sich für Signalisierungs- und Warte-Funktionen nutzen lässt, und die nach Gebrauch wieder in den frei verfügbaren Pool retourniert werden. Wenn die GPU versucht, einen Event zu allokieren, versucht ein Unter-Allokator, sich eine der prä-allokierten Event-Strukturen anzueignen. Ist die Allokation erfolgreich, wird ein Verweis zu dieser Struktur an den aufrufenden Thread geschickt. Schlägt die Allokation jedoch fehl, muss ein Fall von Speichermangel gehandhabt werden. In einem Speichermangel-Szenario wird die gesamte neue Arbeitslast in einem einzigen Stream serialisiert. Dieses Verfahren ist für verschachtelte Streams einschließlich Events sicher, denn der serialisierte Stream erfüllt die Semantik für korrekte Event-Reihenfolge und ist semantisch gleichwertig mit mehreren Streams, die Stream-übergreifende Abhängigkeiten aufweisen. Das Serialisieren aller Aufgaben in nur einem Stream wird korrekte Ergebnisse bei potentiell geringerer Leistung produzieren. Um ein Event-Objekt zu zerstören, wird es einfach wieder in den frei verfügbaren Pool retourniert.
Ein Thread kann cudaEventRecord() aufrufen, um einen SE in eine Stream-seitige Warteschlange zu stellen. Ist der Event bereits aufgezeichnet, muss ein neuer Event allokiert und als der jüngste Event-Record weiterverfolgt werden. Eine Datenstruktur wie etwa für einen Aufgabenstatus, die unten in 11 beschrieben wird, beschreibt den Status des Signalisierungs-Events. Die Datenstruktur lässt sich in eine Stream-seitige Warteschlange stellen. Jeder Signalisierungs-Event pflegt eine Liste von Events, die darauf warten, dass der Signalisierungs-Event feuert. Anfangs ist diese Liste leer, denn es wurde noch kein Warte-Event erzeugt, um auf den Signalisierungs-Event zu warten. Sobald der Signalisierungs-Event gefeuert hat (und vor ihm keine Stream-Arbeit liegt), wird eine verknüpfte Liste durchgegangen, um zu markieren, dass der Signalisierungs-Event gefeuert hat. Warte-Events werden zudem benachrichtigt, dass der Signalisierungs-Event gefeuert hat. Das Benachrichtigen von Warte-Events erfüllt die Forderung, dass ein oder möglicherweise mehr Events vorliegen, die der Signalisierungs-Event zur Fertigstellung benötigt. Das erlaubt es einem korrespondierenden Stream fortzufahren.
Ein Thread kann eine Stream-übergreifende Abhängigkeit erzeugen, indem er cudaStreamWaitEvent() aufruft, um einen WE in die Warteschlange eines beliebigen Streams zu stellen. Der Warte-Event muss zunächst aus dem frei verfügbaren Pool der Event-Strukturen allokiert werden. In einem Ausführungsbeispiel schlägt der WE den jüngsten Aufruf von cudaEventRecord() hinsichtlich des referenzierten Signalisierungs-Events nach. Wenn es nichts gibt, worauf zu warten ist, kann der Warte-Event abgeschlossen werden. Wird jedoch der jüngste Signalisierungs-Event gefunden, wird der WE atomisch zu einer Warteliste hinzugefügt, die mit einem entsprechenden Signalisierungs-Event-Objekt verknüpft ist. Ein WE, der dem selben Stream hinzugefügt worden ist wie der Stream des abzuwartenden Signalisierungs-Events, wird in semantischer Hinsicht als eine Null-Operation behandelt.
Wird ein Scheduler-Kernel nach der Fertigstellung eines Grid ausgeführt, sollte der Scheduler die nächste Aufgabe in einem zugeordneten Stream starten. Ist die nächste Aufgabe ein weiterer Grid, dann startet der Scheduler-Kernel einfach den Grid. Handelt es sich jedoch bei der nächsten Aufgabe um einen Event (WE oder SE), dann ist es nötig, dass der Scheduler-Kernel alle Events, die zum Auslösen bereitstehen (so etwa zum Feuern eines Signalisierungs-Events, der die Blockade eines oder mehrerer Warte-Events aufhebt, was wiederum die Blockade eines oder mehrerer Signalisierungs-Events aufheben mag), handhabt.
Zu jedem WE gehört eine Datenstruktur mit einem Abhängigkeitszähler, der angibt, auf wie viele Dinge, deren Fertigstellung ansteht, der WE wartet. In einem Ausführungsbeispiel kann der Zähler gleich Null, eins oder zwei sein. Ein Zähler von Null zeigt an, dass der Event ausgelöst worden ist. Ein Zähler von eins zeigt an, dass der WE auf eine Aufgabe oder einen Event wartet. Ein Zähler von zwei zeigt an, dass der Event sowohl auf einen Signalisierungs-Event als auch auf eine Aufgabe oder einen Event im gleichen Stream wartet. Dieser Zähler wird atomisch dekrementiert, sobald ein Scheduler-Kernel bestimmt, dass exakt eine Abhängigkeit erfüllt worden ist. Sobald der Zähler bis auf Null dekrementiert worden ist, kann in der Datenstruktur ein ”Stream-next”-Pointer (StreamNext) durchgegangen werden und damit verbundene Abhängigkeitszähler werden ebenfalls dekrementiert. Jeder Zugriff auf den Abhängigkeitszähler sollte atomisch erfolgen, um sicherzustellen, dass nur ein Scheduler-Kernel versucht, eine nächste Aufgabe in einen Stream zu terminieren oder eine abhängige Liste von WEs zu terminieren.
Wird ein Signalisierungs-Event ausgelöst, muss eine komplette verknüpfte eventWaitingList durchlaufen werden, wobei der mit dem Event verknüpfte Zähler, der die Abhängigkeiten jedes Warte-Events zählt, dekrementiert wird. Für Events, die keine Abhängigkeiten mehr aufweisen, muss jeder Stream-next-Pointer durchgegangen werden. Ein Stream-next-Pointer, der mit dem Signalisierungs-Event verknüpft ist, muss ebenfalls durchgegangen werden, denn der Signalisierungs-Event wird gerade fertiggestellt. Um die Notwendigkeit zu umgehen, einen Traversal-Status-Stapel zu errichten, wird der Tree-Walk-Algorithmus abgeflacht. Jede technisch realisierbare und non-rekursive Tree-walk-Technik lässt sich implementieren. Auf Maschinen, bei denen begrenzt vorhandener Thread-lokaler Stapelraum ein kleineres Problem als aktuelle PPU-Implementierungen darstellt, lässt sich jede technisch machbare, rekursive Tree-walk-Technik implementieren.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung zeigt 10 die Datenstruktur für den Kontext 1020 einer Thread-Gruppe, einschließlich Parametern und Kontextinformationen, die mit einer Thread-Gruppe verknüpft sind. Der Thread-Gruppenkontext 1020 weist, wie gezeigt, einen Pointer auf die letzte Aufgabe 1040 für jede TMDQ in der Thread-Gruppe sowie einen Arbeits-Pointer 1050 auf.
Der Pointer 1040 auf die letzte Aufgabe zeigt auf die letzte Aufgabe im zugeordneten TMDQ. Sobald eine neue Aufgabe in einen TMDQ gestartet wird, wird der Pointer auf die letzte Aufgabe 1040 durch eine atomische Operation aktualisiert, um anzuzeigen, dass die neue Aufgabe jetzt die letzte Aufgabe im TMDQ sei. Die untenstehende Tabelle 8 illustriert den Start einer neuen Aufgabe in einen TMDQ innerhalb eines exemplarischen CUDA-Programms.
Im Beispiel der Tabelle 8 wird der Pointer auf die letzte Aufgabe 1040 an der Adresse StreamEnd mit einem Pointer auf NewTask überschrieben, und der vorherige Wert im Pointer auf die letzte Aufgabe 1040 wird als FormerStreamEnd retourniert. Sollte FormerStreamEnd ungleich Null sein (also FormerStreamEnd ein Pointer auf eine Aufgabe sein), wird der StreamNext-Wert, der mit der Aufgabe verknüpft ist, aktualisiert, um auf die neu gestarteten Aufgaben zu zeigen. Ist FormerStreamEnd hingegen gleich Null, dann stehen im TMDQ keine Aufgaben zum Abschluss an und die neue Aufgabe darf sofort mit ihrer Ausführung beginnen.
Das Beispiel in Tabelle 8 wird in einem kritischen Operationsbereich ausgeführt, um eine Blockade an einer Stelle zu vermeiden, an der ein Thread eine Aufgabe in einen Stream gestartet hat, doch dann wurde die Aufgabe hinausgestellt, bevor die neue Aufgabe gestartet wurde. In solch einem Fall kann eine Blockade auftreten, falls einer hinausgestellten Aufgabe nicht erlaubt wird, wieder hereingestellt zu werden, bis die neue Aufgabe abgeschlossen worden ist. Die neue Aufgabe darf mit ihrer Ausführung jedoch nicht beginnen, weil die neue Aufgabe noch nicht aufgerufen worden ist.
Wenn eine Aufgabe abgeschlossen wird, liest ein auf dem Parallelverarbeitungssubsystem 112 ausgeführter Scheduler den Pointer auf den letzten Stream, welcher (der Pointer) mit dem TMDQ korrespondiert, der mit der abgeschlossenen Aufgabe verknüpft ist. Wenn der Pointer 1040 auf die letzte Aufgabe des verknüpften TMDQ nicht auf die abgeschlossene Aufgabe weisen sollte, dann ist die abgeschlossene Aufgabe nicht die letzte Aufgabe im TMDQ. In einem solchen Fall veranlasst der Scheduler die nächste Aufgabe im TMDQ dazu, mit ihrer Ausführung zu beginnen, wie unten in Verbindung mit 11 gezeigt. Wenn der Pointer auf die letzte Aufgabe 1040 des verknüpften TMDQ auf die abgeschlossene Aufgabe zeigt, dann ist diese abgeschlossene Aufgabe die letzte Aufgabe im TMDQ. In einem solchen Fall führt der Scheduler einen atomischen Vergleich und Tausch durch, um den Pointer 1040 auf die letzte Aufgabe auf einen Null-Pointer zu setzen und den Wert auszulesen, der aktuell im Pointer 1040 auf die letzte Aufgabe gespeichert ist. Der Scheduler führt einen Funktionsaufruf in der Form ”currentEnd = atomicCAS(&StreamEnd, finishedTask, NULL)” durch, bei dem ”StreamEnd” der Pointer 1040 auf die letzte Aufgabe des verknüpften TMDQ ist; ”finishedTask” ist ein Pointer auf die abgeschlossene Aufgabe, und ”NULL” ist der Null-Pointer. Die Funktion liefert atomisch den Wert, der jeweils im Pointer 1040 auf die letzte Aufgabe gespeichert ist, wie durch ”currentEnd” im Funktionsaufruf dargestellt.
Wenn der Wert von ”currentEnd” ein Pointer auf die abgeschlossene Aufgabe ist, dann sind alle Aufgaben im TMDQ abgeschlossen und keine neue Aufgabe wurde gestartet. Der Scheduler weiß, dass alle Aufgaben im Stream abgeschlossen wurden. Wenn der Wert ”currentEnd” hingegen kein Pointer auf die abgeschlossene Aufgabe ist, dann ist eine neue Aufgabe gestartet worden und der Thread-Gruppenkontext 1020 ist aktualisiert worden, um die Existenz der neuen Aufgabe widerzuspiegeln. In einem solchen Fall liest der Scheduler den Pointer StreamNext (der unten beschrieben wird), der mit der fertiggestellten Aufgabe verknüpft ist. Wenn der mit der fertiggestellten Aufgabe verknüpfte StreamNext-Pointer hingegen ungleich Null ist, dann veranlasst der Scheduler die Aufgabe an der Adresse StreamNext, mit ihrer Ausführung zu beginnen. Ist der Wert von StreamNext der Null-Pointer, dann ist eine neue Aufgabe zwar gestartet, der Aufgabenstatus aber noch nicht aktualisiert worden, um die Existenz der neuen Aufgabe zu reflektieren. In einem solchen Fall überwacht der Scheduler StreamNext, bis der Wert vom Null-Pointer zu Pointer auf die neue Aufgabe wechselt. Der Scheduler veranlasst anschließend die neue Aufgabe, auf die StreamNext zeigt, mit der Ausführung zu beginnen.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung stellt 11A eine Datenstruktur für den Aufgabenstatus 1120 einschließlich der Parameter, die mit einer Berechnungsaufgabe verknüpft sind, dar. Der Aufgabenstatus 1120 weist einen Aufgabenidentifikator (Aufgaben-Kennzeichen) 1140 auf, einen Pointer 1142 auf einen nächsten Stream, einen Identifikator 1144 für Thread-Gruppenkontext (Thread-Gruppenkontext-Kennzeichen) und andere Parameter, die mit der Aufgabe verknüpft sind (nicht dargestellt).
Das Aufgaben-Kennzeichen 1140 ist ein eindeutiger Identifikator, der auf die Aufgabe zeigt, die mit dem Aufgabenstatus 1120 verknüpft ist. Ein Aufgabenstatus 1120 wird für jede neue Aufgabe erzeugt, sobald Aufgaben erzeugt und auf einen TMDQ gestartet werden. Das Aufgaben-Kennzeichen ermöglicht es dem Scheduler, die Aufgabe, die jeweils mit dem Aufgabenstatus 1120 verknüpft ist, zu finden.
Der Pointer 1142 auf den nächsten Stream ist ein Pointer auf die nächste Aufgabe im TMDQ. Sobald eine Aufgabe abgeschlossen ist, liest der Scheduler den Pointer auf den nächsten Stream, um zu bestimmen, wo im TMDQ die nächste Aufgabe zu finden ist, die mit der Ausführung beginnen könnte. Der Scheduler veranlasst dann diejenige Aufgabe, die an der Adresse, auf die der Pointer 1142 auf den nächsten Stream zeigt, zu finden ist, mit ihrer Ausführung zu beginnen. Sollte die abgeschlossene Aufgabe hingegen die letzte Aufgabe im TMDQ sein, dann wird der Pointer 1142 auf den nächsten Stream auf einen Null-Pointer gesetzt.
Das Thread-Gruppenkontext-Kennzeichen 1120 ist ein eindeutiger Identifikator, der auf den Thread-Gruppenkontext 1020 zeigt, der mit dem Aufgabenstatus 1120 verknüpft ist. Wenn eine Aufgabe abgeschlossen wird, liest der Scheduler das Thread-Gruppenkontext-Kennzeichen 1120, um den Thread-Gruppenkontext 1020 zu finden. Daraufhin kann der Scheduler Schritte zur Fertigstellung der verknüpften Aufgabe ausführen, darunter die Aktualisierung des Arbeitszähler, der den TMDQ abschließt, sowie das Abschließen eines Kontexts, wie oben im Zusammenhang mit 10 beschrieben.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung stellt 11B eine Datenstruktur für den Status 1122 eines Signalisierungs-Events dar. Die Datenstruktur für den Status 1122 eines Signalisierungs-Events weist einen Event-Identifikator 1150, einen Pointer 1152 auf den nächsten Stream und einen Pointer 1154 auf den nächsten Event auf. Das Event-Kennzeichen 1150 identifiziert einen bestimmten Event eindeutig, von dem verschiedene Aufgaben abhängen können. StreamNext 1152 hat wesentliche identifizierende Bedeutung und Funktion im Hinblick auf StreamNext 1142 in 11A. EventNext 1154 ist ein Pointer auf einen abhängigen Event.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung stellt 11C eine Datenstruktur für einen Warte-Event-Status 1124 einschließlich Parametern, die mit einer Berechnungsaufgabe verknüpft sind, dar. Die Datenstruktur für einen Warte-Event-Status 1124 weist ein Event-Kennzeichen 1160, einen Pointer auf den nächsten Stream 1162 und einen Abhängigkeitszähler 1164 auf. Das Event-Kennzeichen 1160 wird definiert und operiert im Wesentlichen identisch zum Event-Kennzeichen 1150 in 11B. StreamNext wird definiert und operiert im Wesentlichen identisch zu StreamNext 1142 in 11A.
Die Datenstruktur zum Aufgabenstatus 1120 (11A), die Datenstruktur für den Signalisierungs-Event-Status 1122 (11B) und die Datenstruktur 1124 für den Warte-Event-Status weisen Artefakte auf, die sich zwecks Verarbeitung in eine Warteschlange innerhalb eines TMDQ stellen lassen. Aufgaben werden in eine Warteschlange gestellt, um eine korrespondierende Berechnung zu initiieren, wohingegen Events für die Berechnung die Aneinanderreihung unter einer Mehrzahl von Aufgaben in einer Warteschlange erzwingen. Bestimmte Anwendungen weisen Ausführungsabhängigkeiten unter verschiedenen Aufgaben auf, die erst einmal berücksichtigt sein müssen, bevor eine bestimmte Aufgabe ausgeführt werden darf.
Gemäß einem Ausführungsbeispiel der vorliegenden Erfindung ist 12 ein Flussdiagramm des Verfahrens 1200, das bestimmt, ob Ausführungsabhängigkeiten erfüllt worden sind oder nicht. Obwohl die Schritte des Verfahrens bereits in Verbindung mit den Systemen in den 1 bis sowie 6 beschrieben worden sind, wird ein Fachmann verstehen, dass jedes System, das zur Ausführung der Verfahrensschritte in beliebiger Reihenfolge eingerichtet ist, sich innerhalb des Umfangs der vorliegenden Erfindung befindet.
In einem Ausführungsbeispiel wird Verfahren 1200 von einem Scheduler-Kernel ausgeführt, der oben in 5 erörtert worden ist. Der Scheduler-Kernel wird ausgeführt, nachdem eine Aufgabe abgeschlossen worden, um zu bestimmen, welche Arbeit als nächste terminiert werden soll. Das Verfahren 1200 besitzt ein Artefakt am Kopf der Stream-Aufgaben-Warteschlange, wie etwa TMDQ, was oben beschrieben wurde.
Das Verfahren beginnt mit Schritt 1210, bei dem der Scheduler-Kernel ein nächstes Artefakt aus einem korrespondierenden TMDQ holt. In einem Ausführungsbeispiel kann es sich bei dem nächsten Artefakt um einen Aufgaben-Metadaten-Deskriptor handeln, einen Warte-Event-Deskriptor oder einen Signalisierungs-Event-Deskriptor. Wenn in Schritt 1220 das nächste Artefakt kein Deskriptor für eine Aufgabe ist, dann fährt das Verfahren bei Schritt 1230 fort. Wenn in Schritt 1230 der nächste Event kein Deskriptor für einen Warte-Event ist, fährt das Verfahren bei Schritt 1240 fort.
Wenn in Schritt 1240 das nächste Artefakt ein Signalisierungs-Event ist, fährt das Verfahren bei Schritt 1242 fort, wo der Scheduler-Kernel Abhängigkeitszählerwerte, die mit einem Event-Kennzeichen für den Signalisierungs-Event verknüpft sind, dekrementiert. Die Abhängigkeitszählerwerte weisen Einträge in einer Liste verbundener Events auf, die von sequentiellen Aufrufen mit cudaStreamWaitEvent() erzeugt worden sind. In einem Ausführungsbeispiel weist das nächste Artefakt eine Datenstruktur für einen Signalisierungs-Event-Status auf, so etwa den Signalisierungs-Event-Status 1122 in 11B, und das Event-Kennzeichen wird von der Datenstruktur für den Signalisierungs-Event-Status bereitgestellt. Die Liste der verbundenen Events weist eine verknüpfte Liste von Elementen des Warte-Event-Status 1124 auf, die durch den Event-Next-Pointer 1154 damit verknüpft sind. Um die Abhängigkeitszähler, die mit dem Event-Kennzeichen verknüpft sind, zu dekrementieren, durchläuft der Scheduler-Kernel die Liste verbundener Events und dekrementiert die Abhängigkeitszähler, die mit den Einträgen für Elementabhängigkeitszähler 1164 für den Warte-Event-Status 1124 verknüpft sind. Nachdem Schritt 1242 abgeschlossen worden ist, speichern die Abhängigkeitszählerwerte für Warte-Events, die direkt von dem Event-Kennzeichen abhängig sind, dekrementierte Werte. Der Vorgang, einen Abhängigkeitszähler zu dekrementieren, der mit einem Warte-Event verknüpft ist, kann bestimmte Aufgaben in die Lage versetzen, ausgeführt zu werden, oder eine Kaskade zusätzlicher Aktualisierungen auslösen. In Schritt 1244 durchläuft der Scheduler-Thread rekursiv StreamNext-Pointer, wie sie oben beschrieben wurden, die abhängig verbundenen Knoten zugeordnet sind, um sodann den Status bei jedem der abhängig verbundenen Knoten zu aktualisieren. Ein Fachmann wird erkennen, dass Verfahren 1200 rekursiv ausgeführt werden kann, um den Schritt 1244 fertigzustellen. Alternativ lässt sich eine ”abgeflachte”, nicht rekursive Traversaltechnik implementieren, während das Verfahren 1200 an jedem der abhängigen Knoten ausgeführt wird. In Schritt 1246 entfernt der Scheduler-Kernel die SE-Einträge aus dem TMDQ. Das Verfahren endet mit Schritt 1290.
Zurückkehrend zu Schritt 1230, wenn der nächste Event ein Deskriptor für einen Warte-Event ist, dann fährt das Verfahren bei Schritt 1232 fort, wo der Scheduler-Kernel einen Abhängigkeitszähler für den Warte-Event dekrementiert. In einem Ausführungsbeispiel wird der Abhängigkeitszähler mit einem Wert von zwei initialisiert, wenn der Warte-Event sowohl von einer Aufgabe als auch von einem Signalisierungs-Event abhängt. Falls in Schritt 1234 der Abhängigkeitszähler gleich Null ist, dann fährt das Verfahren bei Schritt 1236 fort, wo der Eintrag für den Warte-Event aus dem TMDQ entfernt wird, bevor das Verfahren wieder zu Schritt 1210 zurückkehrt.
Zurückkehrend zu Schritt 1220, wenn das nächste Artefakt ein Deskriptor für eine Aufgabe ist, fährt das Verfahren bei Schritt 1222 fort, wo der Scheduler-Kernel die dem Deskriptor zugeordnete Aufgabe zur Ausführung veranlasst. In einem Ausführungsbeispiel wird die Aufgabe zur Ausführung in einem SM 310 in 3B veranlasst.
Zurückkehrend zu Schritt 1240, wenn das nächste Artefakt kein Signalisierungs-Event ist, dann endet das Verfahren mit Schritt 1290.
Zurückkehrend zu Schritt 1234 – wenn der Abhängigkeitszähler ungleich Null ist, dann endet das Verfahren mit Schritt 1290.
Zusammenfassend lässt sich sagen, dass hier eine Technik für das Erzwingen Stream-übergreifender Abhängigkeiten offengelegt wird. Die Technik weist auf, dass Warte-Events in Warteschlangen gestellt werden, damit Stream-übergreifende Abhängigkeiten geschaffen werden, sowie Signalisierungs-Events, um den Warte-Events Abschlüsse anzuzeigen. Die Technik involviert keinen Sperrmechanismus und lässt sich auf einem Parallelverarbeitungssubsystem wie etwa einer GPU effizient ausführen. Ein Scheduler-Kernel untersucht eine Aufgabenstatus-Datenstruktur aus einem korrespondierenden Stream und aktualisiert Abhängigkeitszähler für Aufgaben und Events innerhalb des Streams. Sobald ein Abhängigkeitszähler den Wert Null erreicht, darf eine korrespondierende Aufgabe oder Event ausgeführt werden. Ein gegebener Thread kann Abhängigkeiten einrichten, indem er cudaEventCreate() aufruft, um ein Event-Objekt zu erzeugen, und cudaStreamWaitEvent(), um einen Warte-Event zu erzeugen, sowie cudaEventRecord(), um einen Signalisierungs-Event zu erzeugen.
Einer der Vorteile der hier offengelegten Technik besteht darin, dass eine GPU korrekt und effizient die Ausführungsreihenfolge in Aufgaben, die Stream-übergreifende Abhängigkeiten aufweisen, durchsetzen kann, ohne dass sich die CPU einmischt. Ein weiterer Vorteil liegt darin, dass sich ein gegebener hierarchischer Aufgabenausführungsgraph zusammensetzbar strukturieren lässt, wodurch er die korrekte Ausführungssemantik beibehält, die für viele allgemeine Berechnungsmodelle notwendig ist. Die offengelegte Technik ermöglicht somit vorteilhafterweise die Implementierung eines allgemeinen, zusammensetzbaren Ausführungsmodells auf Parallelverarbeitungssubsystemen wie etwa GPUs.
Während das oben Gesagte sich auf Ausführungsbeispiele der vorliegenden Erfindung bezieht, können andere und weitere Ausführungsbeispiele der Erfindung entworfen werden, ohne vom grundlegenden Umfang der Erfindung abzuweichen. So können beispielsweise Aspekte der vorliegenden Erfindung in Hardware oder Software oder Kombinationen davon implementiert werden. Ein Ausführungsbeispiel der vorliegenden Erfindung lässt sich als ein Programmprodukt zur Verwendung mit einem Computersystem implementieren. Das bzw. die Programm(e) des Programmprodukts definieren Funktionen der Ausführungsbeispiele (einschließlich der hier beschriebenen) und können auf einer Vielzahl von Computer-lesbaren Speichermedien enthalten sein. Beispielhafte Computer-lesbare Speichermedien können sein, ohne darauf beschränkt zu sein: (i) nicht-beschreibbare Speichermedien (wie etwa Nur-lese-Speichergeräte in einem Computer wie etwa CD-ROMs, die von einem CD-ROM-Laufwerk gelesen werden können; Flash-Speicher; ROM-Chips oder jede Art von Solid-State- und nicht-flüchtigem Halbleiterspeicher), auf denen Informationen dauerhaft gespeichert werden; (ii) beschreibbare Speichermedien (beispielsweise Floppy-Disks in einem Diskettenlaufwerk oder eine Festplatte oder jeder Typ von Solid-State-Halbleiterspeicher mit beliebigem Zugriff), auf dem veränderbare Informationen gespeichert werden. Wenn sie Computer-lesbare Anweisungen, die Funktionen der vorliegenden Erfindung steuern, tragen, dann sind solche Computer-lesbaren Speichermedien Ausführungsbeispiele der vorliegenden Erfindung.
Daher wird der Umfang der vorliegenden Erfindung durch die Ansprüche, die im Folgenden genannt werden, bestimmt.

Claims

Ein Computer-implementiertes Verfahren zur Verarbeitung einer Mehrzahl von Aufgaben verteilt über eine Gruppe von Threads, wobei das Verfahren aufweist Aufrufen eines ersten Elementes aus einer Warteschlange; Feststellen, dass das erste Element keine Aufgabe aufweist; Feststellen, ob das das erste Element einen Warte-Event oder einen Signalisierungs-Event aufweist; in Reaktion darauf, dekrementieren eines Zählers; und Entfernen des ersten Elementes aus der Warteschlange.
Das Verfahren gemäß Anspruch 1, wobei das erste Element einen Signalisierungs-Event aufweist und wobei das Dekrementieren eines Zählers ein Dekrementieren jedes Abhängigkeitszählers aufweist, der einem Event-Kennzeichen zugeordnet ist.
Das Verfahren gemäß Anspruch 1, wobei das Event-Kennzeichen die Fertigstellung einer gegeben Aufgabe repräsentiert und wobei jeder Abhängigkeitszähler eine unterschiedliche andere Aufgabe repräsentiert, die auf eine Fertigstellung gegebene Aufgabe wartet.
Das Verfahren gemäß Anspruch 3, weiterhin aufweisend ein rekursives Durchlaufen einer Mehrzahl von Pointern, die auf eine Mehrzahl von Knoten verweisen, wobei jeder Pointer auf einen anderen Knoten verweist, und wobei jeder Knoten einer der unterschiedlichen anderen Aufgaben, die auf eine Fertigstellung der gegeben Aufgaben warten, zugeordnet ist.
Das Verfahren gemäß Anspruch 1, wobei das erste Element einen Warte-Event aufweist und das Dekrementieren des Zählers ein Dekrementieren eines Abhängigkeitszählers aufweist, der dem Warte-Event zugeordnet ist.
Das Verfahren gemäß Anspruch 5, wobei der Warte-Event dahingehend wirkt, die Ausführung jeder Aufgabe, die auf eine Fertigstellung der gegeben Aufgabe wartet, zu blockieren.
Das Verfahren gemäß Anspruch 6, wobei alle Aufgabe, die der Warteschlange zugeordnet sind und die auf die Fertigstellung der gegebenen Aufgabe warten, in der Warteschlange hinter dem Warte-Event angeordnet sind.
Das Verfahren gemäß Anspruch 5, weiterhin aufweisend ein Feststellen, dass der Abhängigkeitszähler gleich Null ist.
Das Verfahren gemäß Anspruch 8, weiterhin aufweisend ein Aufrufen eines zweiten Elementes aus der Warteschlange.
Das Verfahren gemäß Anspruch 9, weiterhin aufweisend ein Feststellen, dass das zweite Element eine Aufgabe aufweist, und Veranlassen, dass die Aufgabe ausgeführt wird.
Eine Computervorrichtung aufweisend eine Zentralverarbeitungseinheit; und eine Parallelverarbeitungsuntereinheit, die mit der Zentralverarbeitungseinheit verbunden ist, aufweisend: eine Untersystem, welches einen Streaming-Multiprozessor aufweist, der konfiguriert ist zum: Aufrufen eines ersten Elementes aus einer Warteschlange; Feststellen, dass das erste Element keine Aufgabe aufweist; Feststellen, ob das das erste Element einen Warte-Event oder einen Signalisierungs-Event aufweist; in Reaktion darauf, dekrementieren eines Zählers; und Entfernen des ersten Elementes aus der Warteschlange.