DE102015224026A1

DE102015224026A1 - Indirektes Erfassen von Sampledaten zur Durchführung mehrfacher Faltungsoperationen in einem Parallelverarbeitungssystem

Info

Publication number: DE102015224026A1
Application number: DE102015224026.6A
Authority: DE
Inventors: John Clifton Woolley jun.; John Tran
Original assignee: Nvidia Corp
Current assignee: Nvidia Corp
Priority date: 2014-12-04
Filing date: 2015-12-02
Publication date: 2016-06-09
Also published as: US20160162402A1; CN105678378A; CN105678378B; US20190220731A1; US10255547B2

Abstract

Bei einer erfindungsgemäßen Ausführungsform konfiguriert eine Faltungsmaschine eine Parallelverarbeitungs-Pipeline, um mehrfache Faltungsoperationen durchzuführen. Genauer gesagt konfiguriert die Faltungsmaschine die Parallelverarbeitungs-Pipeline, um unabhängig individuelle Bildabschnitte zu erzeugen und zu verarbeiten. Im Betrieb berechnet die Pipeline für jeden Bildabschnitt Ursprungsstellen, welche in einem Eingabebildstapel enthalten sind, abhängig von einer oder von mehreren Startadressen und von einem oder von mehreren Offsets. Anschließend kopiert die Pipeline Daten von den Ursprungsstellen in den Bildabschnitt. Die Pipeline führt dann Matrixmultiplikationsoperationen zwischen dem Bildabschnitt und einem Filterabschnitt durch, um einen Anteil des Bildabschnitts für eine Ausgabematrix zu erzeugen. Um den Umfang eines eingesetzten Speichers zu optimieren, erzeugt die Pipeline jeden Bildabschnitt bei Bedarf in einem Gemeinschaftsspeicher. Um den Durchsatz der Matrixmultiplikationsoperationen zu optimieren, werden darüber hinaus die Werte der Offsets durch einen Faltungspräprozessor vorberechnet.

Description

QUERVERWEIS ZU VERWANDTEN ANMELDUNGEN
Diese Anmeldung bezieht sich auf die vorläufige US-Patentanmeldung mit der Nummer 62/087,681 (Anwaltsaktenzeichen 14-SC-0391-USL), die am 4. Dezember 2014 eingereicht wurde. Der Gegenstand dieser Anmeldung wird hier per Referenz aufgenommen.
HINTERGRUND DER ERFINDUNG
Bereich der Erfindung
Erfindungsgemäße Ausführungsformen betreffen im Allgemeinen eine Verarbeitung mittels Computer und insbesondere ein indirektes Erfassen von Sampledaten, um mehrfache Faltungsoperationen in einem Parallelverarbeitungssystem durchzuführen.
Beschreibung des Stands der Technik
Faltende neuronale Netzwerke („Convolution Neural Networks” (CNNs)) werden oft eingesetzt, um effizient und zuverlässig einen weiten Bereich von Interferenzproblemen zu lösen. Zum Beispiel sind CNNs in vielen Bilderkennungs-, Handschrifterkennungs- und Sprachübersetzungs-Algorithmen vorhanden. Im Betrieb können CNNs Fehlerraten im Vergleich zu vielen einfacheren maschinellen Lerntechniken wesentlich verringern. Jedoch übersteigt die Zeit, welche zur Ausführung von CNNs erforderlich ist, gewöhnlicherweise die Zeit, welche zur Ausführung von einfacheren maschinellen Lerntechniken erforderlich ist. Daher können zeitkritische Anwendungen derart aufgebaut sein, dass sie einfachere maschinelle Lerntechniken auf Kosten von schlechteren Ergebnissen einsetzen.
Im Allgemeinen wird die Zeit, welche zur Ausführung eines CNNs erforderlich ist, durch die Zeit dominiert, welche für das CNN erforderlich ist, um mehrfache Faltungsoperationen bzw. Mehrfachfaltungsoperationen („Multi-Convolution” Operations) durchzuführen. Eine mehrfache Faltungsoperation ist eine verallgemeinerte Form einer mehrdimensionalen Faltungsoperation zwischen Sampledaten, wie beispielsweise einem Bild und einem Filter. Die mehrfache Faltungsoperation wird oft unter Verwendung einer Schablonen basierten Technik oder unter Verwendung einer schnellen Fouriertransformation (FFT) implementiert. Während Schablonen basierte Techniken und FFT-basierte Techniken ermöglichen, dass einige mehrfache Faltungsoperationen effizienter implementiert werden, sind solche Techniken normalerweise nicht in der Lage, mehrfache Faltungsoperationen effizient über den gesamten Bereich von Dimensionen bzw. Ausmaßen und zusätzlichen Parametern, welche typischerweise standardisierten CNNs zugeordnet sind, auszuführen.
Diesbezüglich weist ein CNN typischerweise mehrere „Faltungsebenen” bzw. „Faltungsschichten” („Convolution Layers”) auf, wobei jede Faltungsebene Faltungsoperationen über mehrere Dimensionen eines Sampledaten-Stapels und mehrere Dimensionen eines Filterstapels ausführt. Zum Beispiel ist der Sampledaten-Stapel für ein vierdimensionales CNN, welches Bildsample umfasst, ein Stapel von Bildern, und die vier Dimensionen des Bildstapels umfassen die Bildbreite, die Bildhöhe, die Anzahl von Farbebenen pro Bild und die Anzahl der Bilder in dem Bildstapel. Die vier Dimensionen des Filterstapels umfassen die Filterbreite, die Filterhöhe, die Anzahl von Merkmalsebenen pro Filter und die Anzahl von Filtern in dem Filterstapel. Zusätzliche Parameter können die mehrfachen Faltungsoperationen weiter individualisieren. Zum Beispiel können eine horizontale Filterschrittweite und eine vertikale Filterschrittweite die gesamte Rechenbelastung verringern, indem die Größe der Teilmenge der Pixel, welche in die Faltungsoperation involviert ist, verringert wird. Es sei angemerkt, dass die Dimensionen des Bildstapels und des Filterstapels wie auch die zusätzlichen Parameter oft zwischen den Faltungsebenen variieren.
Schablonen basierte Techniken werden typischerweise eingestellt, um mehrfache Faltungsoperationen über eine relativ kleine Teilmenge von Dimensionen und Parametern zu optimieren. Die Leistungsfähigkeit von Schablonen basierten Techniken über andere Dimensionen und Parameter übersteigt jedoch gewöhnlicherweise die Zeit, welche erforderlich ist, um einfachere maschinelle Lerntechniken auszuführen. Daher ist, wie es vorab erwähnt wird, die Zeit, welche erforderlich ist, um viele CNNs unter Verwendung von Schablonen basierten Techniken auszuführen, typischerweise unakzeptabel lang. Wie ebenfalls vorab erwähnt wird, variiert die Zeit, welche erforderlich ist, um viele CNNs unter Verwendung von FFT-basierten Ansätzen auszuführen, auch dramatisch abhängig von den Werten der Parameter.
Ein Ansatz, um die Zeit zu verringern, welche erforderlich ist, um CNNs über einen großen Bereich von Parameterwerten auszuführen, umfasst die Beobachtung, dass eine Faltung ein linearer Operator ist und daher auf eine Matrixmultiplikation reduziert werden kann. Ein solcher Ansatz erfordert, dass die Sampledaten in die erforderliche Matrixform expandiert werden. Genauer gesagt wandelt die Faltungsmaschine bei solchen Implementierungen den Bildstapel in eine spaltengewichtete Bildmatrix um und formuliert den Filterstapel als eine Filtermatrix. Daher führt die Faltungsmaschine Matrixmultiplikationsoperationen zwischen der Bildmatrix und dem Filterstapel aus. Es sei angemerkt, dass die Dimensionen der Bildmatrix und der Filtermatrix Produkten von Teilmengen der unabhängigen Parameter des CNNs entsprechen anstatt den individuellen Parametern. Daher weisen Matrix basierte Techniken relativ gleichförmige Leistungseigenschaften über verschiedene vorgegebene Dimensionen und Parameter auf. Da Codebibliotheken, welche für jede von vielen Arten von Verarbeitungseinheiten geschrieben sind, optimierte Matrixmultiplikationsroutinen aufweisen, kann die Zeit, welche erforderlich ist, um ein CNN mittels des vorab beschriebenen Ansatzes auszuführen, wesentlich geringer sein, als die Zeit, welche erforderlich ist, um das CNN unter Verwendung einer Schablonen basierten oder FFT-basierten Technik auszuführen.
Ein Nachteil einer Implementierung von solchen Matrix basierten Operationen bei einer Faltungsmaschine ist, dass die Faltungsmaschine als Teil eines Expandierens des Bildstapels zum geeigneten Bilden der Matrixmultiplikationsoperationen die Bilddaten an viele Stellen in der Bildmatrix kopieren muss. Daher kann die Größe der Bildmatrix bis zu einem Punkt ansteigen, ab welchem der verfügbare Speicher vollständig aufgebraucht wird. Es sei angenommen, dass die Bildbreite zum Beispiel W, die Bildhöhe H, die Anzahl der Farbebenen pro Bild C und die Anzahl der Bilder in dem Bildstapel N sei. Darüber hinaus sei angenommen, dass die Ausmaße von jedem der Ausgangsbilder (P × Q) sei. In einem solchen Szenario wären die Ausmaße der Bildmatrix (N × P × Q) × (C × R × S). Bei vielen Systemen kann der Platz, welcher erforderlich ist, um Bildmatrizen dieser Größe zu speichern, den verfügbaren Speicherplatz übersteigen.
Im Bemühen, einen Speicherverbrauch zu verringern, während eine mehrfache Faltung mittels einer optimierten Matrixmultiplikationsroutine ausgeführt wird, kann eine Abschnitt basierte Faltungsmaschine implementiert werden, welche eine Parallelverarbeitungs-Pipeline konfiguriert, um unabhängig bestimmte Abschnitte der Bildmatrix zu expandieren und zu verarbeiten. Bei einem solchen Ansatz führt die Parallelverarbeitungs-Pipeline Adressberechnungen durch, um jeden Abschnitt der Bildmatrix in einen Gemeinschaftsspeicher auf einer „bei Bedarf”-Basis zu expandieren. Die Parallelverarbeitungs-Pipeline führt dann Matrixmultiplikationsoperationen zwischen dem Bildabschnitt und dem Filterstapel durch. Da die expandierte Bildmatrix direkt als ein Abschnitt zu jedem Zeitpunkt in den Gemeinschaftsspeicher expandiert wird, wird die Matrix niemals vollständig gespeichert, und der Umfang des zur Parallelverarbeitung benötigten Speichers kann dramatisch im Vergleich zu typischen Matrix basierten Faltungsmaschinen verringert werden.
Ein Nachteil der Abschnitt basierten Faltungsmaschinen ist jedoch, dass das Berechnen der Adressfolge, welche benötigt wird, um die Bilddaten in der richtigen Reihenfolge zu laden, um einen Abschnitt der expandierten Bildmatrix zu expandieren, die Durchführung einer Reihe von unabhängigen Ganzzahl-Operationen umfasst. Diese Reihe von Ganzzahl-Operationen erfordert typischerweise eine relativ große Anzahl von Taktzyklen zur Durchführung. Bisweilen übersteigt die Anzahl von Taktzyklen, welche erforderlich sind, um die Ganzzahl-Operationen durchzuführen, die Anzahl von Taktzyklen, welche erforderlich sind, um die Matrixmultiplikationsoperationen durchzuführen. Daher können die Vorteile der optimierten Matrixmultiplikationsroutine nicht vollständig realisiert werden, und die Gesamtzeit, um die CNN auszuführen, kann unakzeptabel lang sein.
Genauer gesagt ist typischerweise jede Schleifeniteration bei einer Matrixmultiplikation für eine bestimmte Anzahl von Fließkomma-Rechenoperationen ausgelegt, um die Speicherwartezeit der Ladevorgänge abzudecken. Zum Beispiel kann eine Implementierung 100 Rechenoperationen für 10 Speicherladevorgänge aufweisen. Typischerweise werden diese 10 Speicherladevorgänge relativ rasch abgearbeitet und beantwortet, wenn die 100 Rechenoperationen abgeschlossen sind. Wenn jedoch jede solche Speicheroperation 10 zusätzliche Ganzzahloperationen aufweist, welche jeweils abhängig von der vorherigen Operation eine Wartezeit von 10 Zyklen aufweisen, dann betragen die Kosten, um die 10 Adressen zu erzeugen, 100 Zyklen – das stimmt mit der Anzahl von Rechenoperationen überein, bevor die Speicherwartezeit zur Unterstützung solcher Speicherladevorgänge berücksichtigt wird. Wenn solche Speicherladevorgänge selbst im Mittel 10 Zyklen benötigen, dann benötigen wir nun 200 Zyklen, um den Speicher zu laden, gegenüber 100 Zyklen, um die Fließkomma-Rechenoperationen zu berechnen, was zu 100 Zyklen führt, bei welchen keine nützliche Rechenoperation ausführbar ist, um die Speicherwartezeit abzudecken, was für die Gesamteffizienz von Nachteil ist.
Aus dem vorab beschriebenen geht hervor, dass nach dem Stand der Technik ein effektiverer Ansatz zur Durchführung von mehrfachen Faltungsoperationen erforderlich ist.
ZUSAMMENFASSUNG DER ERFINDUNG
Eine erfindungsgemäße Ausführungsform stellt ein in einem Computer implementiertes Verfahren zur Durchführung einer mehrfachen Faltungsoperation dar. Das Verfahren weist auf ein Auswählen einer ersten Startadresse abhängig von einer ersten Zieladresse, welche in einem ersten Bildabschnitt vorhanden ist, welcher in einem ersten Speicher gespeichert ist; ein Identifizieren eines ersten Offsets abhängig von der ersten Zieladresse; ein Berechnen einer ersten Ursprungsadresse, welche in einem Bildstapel vorhanden ist, welcher in einem zweiten Speicher gespeichert ist, abhängig von der ersten Startadresse und dem ersten Offset; ein Kopieren von Daten von der ersten Ursprungsadresse zu der ersten Zieladresse; und nach dem Kopieren der Daten ein Durchführen von einer oder von mehreren Matrixmultiplikationsoperationen zwischen dem ersten Bildabschnitt und einem ersten Filterabschnitt.
Weitere Ausführungsformen stellen neben anderen Dingen ein nicht flüchtiges von einem Computer lesbares Medium und ein System bereit, welches ausgestaltet ist, um das vorab beschriebene Verfahren zu implementieren.
Ein Vorteil der offenbarten Techniken ist, dass Anwendungen mehrfache Faltungsoperationen mittels einer optimierten Matrixmultiplikationsroutine durchführen können, während die Verwendung eines Speichers bei der Parallelverarbeitung optimiert wird. Insbesondere verringert eine vorherige Berechnung von Offsets die Wartezeit, welche einer Berechnung von Adressen zugeordnet ist, während jeder Bildabschnitt einer virtuellen Bildmatrix unmittelbar expandiert wird.
KURZE BESCHREIBUNG DER ZEICHNUNGEN
Um die vorab beschriebenen Merkmale der vorliegenden Erfindung im Detail zu verstehen, kann eine genauere Beschreibung der Erfindung, welche vorab kurz umrissen wurde, mit Bezug zu Ausführungsformen, von welchen einige in den beigefügten Zeichnungen dargestellt sind, folgen. Es sei jedoch angemerkt, dass die beigefügten Zeichnungen nur typische erfindungsgemäße Ausführungsformen darstellen und daher nicht als den Umfang der Erfindung einschränkend anzusehen sind, da die Erfindung auch durch andere gleich effektive Ausführungsformen realisiert werden kann.
1 ist ein Blockdiagramm, welches ein Computersystem darstellt, das ausgestaltet ist, um einen oder mehrere Aspekte der vorliegenden Erfindung zu implementieren.
2 ist gemäß verschiedenen erfindungsgemäßen Ausführungsformen ein Blockdiagramm einer Parallelverarbeitungseinheit, welche in dem Parallelverarbeitungs-Teilsystem der 1 vorhanden ist.
3 ist gemäß verschiedenen erfindungsgemäßen Ausführungsformen ein Blockdiagramm eines allgemeinen Verarbeitungsclusters, welches in der Parallelverarbeitungseinheit der 2 vorhanden ist.
4 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen einen Bildstapel, einen Filterstapel und einen Ausgabestapel, welche einer mehrfachen Faltungsoperation zugeordnet sind, dar.
5 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen die Beziehung zwischen dem Bildstapel der 4 und einer virtuellen Bildmatrix dar.
6 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen die Beziehungen zwischen dem Bildstapel der 4, einer Offsetfolge und der virtuellen Bildmatrix der 5 dar.
7 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen den Faltungspräprozessor der 1 dar, welcher ausgestaltet ist, um die Offsetfolge der 6 zu erzeugen.
8 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen die Faltungsmaschine der 1 dar, welche ausgestaltet ist, um eine mehrfache Faltungsoperation durchzuführen.
9 ist gemäß verschiedenen erfindungsgemäßen Ausführungsformen ein Flussplan von Verfahrensschritten, um eine mehrfache Faltungsoperation in einem Parallelverarbeitungssystem durchzuführen.
DETAILLIERTE BESCHREIBUNG
In der folgenden Beschreibung werden zahlreiche bestimmte Details zum besseren Verständnis der vorliegenden Erfindung dargelegt. Der Fachmann erkennt jedoch, dass die vorliegende Erfindung auch ohne eines oder mehrerer dieser bestimmten Details ausgeführt werden kann.
Systemüberblick
1 ist ein Blockdiagramm, welches ein Computersystem 100 darstellt, welches ausgestaltet ist, um einen oder mehrere Aspekte der vorliegenden Erfindung auszuführen. Wie dargestellt ist, weist das System 100, ohne Einschränkung, eine Zentral-Verarbeitungseinheit (CPU) 102 und einen Systemspeicher 104 auf, welche über eine Speicher-Brücke 105 und einen Kommunikationspfad 113 mit einem Parallelverarbeitungs-Teilsystem 112 gekoppelt sind. Die Speicher-Brücke 105 ist darüber hinaus über einen Kommunikationspfad 106 mit einer I/O-(Eingabe-/Ausgabe-)-Brücke 107 gekoppelt, und die I/O-Brücke 107 ist wiederum mit einem Switch 116 gekoppelt.
Im Betrieb ist die I/O-Brücke 107 ausgestaltet, um eine Benutzer-Eingabe-Information von Eingabe-Geräten 108, wie beispielsweise einer Tastatur oder einer Maus, zu erfassen, und um diese Eingabe-Information über den Kommunikationspfad 106 und die Speicher-Brücke 105 an die CPU 102 zur Verarbeitung weiterzuleiten. Der Switch 116 ist ausgestaltet, um Verbindungen zwischen der I/O-Brücke 107 und anderen Komponenten des Computersystems 100, wie beispielsweise einem Netzwerk-Adapter 118 und verschiedenen Zusatzkarten, bereitzustellen.
Wie auch dargestellt ist, ist die I/O-Brücke 107 mit einer Systemplatte 114 gekoppelt, welche ausgestaltet sein kann, um Inhalte und Anwendungen und Daten zur Verarbeitung durch die CPU 102 und das Parallelverarbeitungs-Teilsystem 112 zu speichern. Im Allgemeinen stellt eine Systemplatte 114 einen nicht flüchtigen Speicher für Anwendungen und Daten bereit und kann fest eingebaute oder entfernbare Festplatten, Flash-Speicher und CD-ROM-(„Compact Disk Read-Only Memory”), DVD-ROM-(„Digital Versstile Disc-ROM”), Blu-ray-, HD-DVD-(„High Definition DVD”) oder andere magnetische, optische oder Festspeicher-Geräte aufweisen. Schließlich können, obwohl es nicht ausdrücklich dargestellt ist, andere Komponenten, wie beispielsweise ein „Universal Serial Bus” oder andere Anschlussverbindungen, Compact-Disk-Laufwerke, Digital-Versatile-Disc-Laufwerke, Filmaufnahme-Geräte und Ähnliches auch mit der I/O-Brücke 107 verbunden sein.
Bei verschiedenen Ausführungsformen kann die Speicher-Brücke 105 ein Northbridge-Chip und die I/O-Brücke 107 ein Southbridge-Chip sein. Darüber hinaus können die Kommunikationspfade 106 und 113 genauso wie die anderen Kommunikationspfade innerhalb des Computersystems 100 unter Verwendung jedes technisch geeigneten Protokolls ausgeführt werden, was ohne Einschränkung AGP-(„Accelerated Graphics Port”), HyperTransport- oder irgendein anderes Bus- oder Punkt-zu-Punkt-Kommunikations-Protokoll, welche nach dem Stand der Technik bekannt sind, einschließt.
Bei einigen Ausführungsformen umfasst das Parallelverarbeitungs-Teilsystem 112 ein Grafik-Teilsystem, welches einem Anzeigegerät 110 Pixel zuliefert, welches irgendeine herkömmliche Kathodenstrahlröhre, Flüssigkristallanzeige, Leuchtdioden-Anzeige oder Ähnliches sein kann. Bei solchen Ausführungsformen umfasst das Parallelverarbeitungs-Teilsystem 112 eine Schaltung, welche für Bild- und Video-Verarbeitung optimiert ist, was zum Beispiel eine Videoausgabe-Schaltung einschließt. Wie im Folgenden in 2 im Detail beschrieben wird, kann eine solche Schaltung in einem oder in mehreren Parallelverarbeitungs-Einheiten (PPU („Parallel Processing Unit”)) enthalten sein, welche in dem Parallelverarbeitungs-Teilsystem 112 enthalten sind. Bei anderen Ausführungsformen umfasst das Parallelverarbeitungs-Teilsystem 112 eine Schaltung, welche für eine allgemeine Verarbeitung und/oder Berechnung („general purpose and/or compute processing”) optimiert ist. Wiederum kann eine solche Schaltung in einem oder in mehreren PPUs, welche in dem Parallelverarbeitungs-Teilsystem 112 enthalten sind, enthalten sein, welche ausgestaltet sind, um solche Operationen zur allgemeinen Verarbeitung und/oder Berechnung durchzuführen. Bei noch anderen Ausführungsformen kann die eine oder können die mehreren PPUs, welche in dem Parallelverarbeitungs-Teilsystem 112 enthalten sind, ausgestaltet sein, um Operationen zur Bildverarbeitung, zur allgemeinen Verarbeitung und zur allgemeinen Berechnung durchzuführen.
Wie dargestellt ist, weist der Systemspeicher 104 mindestens einen Gerätetreiber 175 und ein Faltungsteilsystem 180 auf. Der Gerätetreiber 175 ist ausgestaltet, um die Verarbeitungsoperationen der einen oder der mehreren PPUs in dem Parallelverarbeitungs-Teilsystem 112 zu verwalten. Das Faltungsteilsystem 180 weist ohne Beschränkung einen Faltungspräprozessor 182 und eine Faltungsmaschine 184 auf. Der Faltungspräprozessor 182 führt Berechnungen durch, welche entworfen sind, um die Effizienz der Faltungsmaschine 184 zu erhöhen, und die Faltungsmaschine 184 ist ausgestaltet, um mehrfache Faltungsoperationen durchzuführen.
Der Faltungspräprozessor 182 kann auf der CPU 120, dem Parallelverarbeitungs-Teilsystem 112 oder irgendeiner Kombination davon ausgeführt werden. Die Faltungsmaschine 184 wird auf dem Parallelverarbeitungs-Teilsystem 112 ausgeführt, und das Parallelverarbeitungs-Teilsystem 112 führt eine optimierte Matrixmultiplikationsroutine, welche in einer Bibliothek vorhanden ist, aus. Es sei angemerkt, dass solche mehrfachen Faltungsoperationen die Zeit dominieren, welche erforderlich ist, um faltende neuronale Netze (CNNs) auszuführen. Obwohl es nicht dargestellt ist, weist der Systemspeicher 104 irgendeine Anzahl von Softwareapplikationen auf, welche auf der CPU 102 ausgeführt werden, kann Kommandos absetzen, welche den Betrieb der PPUs steuern, und kann das Faltungsteilsystem 180 beeinflussen, um die CNN effizient auszuführen.
Bei verschiedenen Ausführungsformen kann das Parallelverarbeitungs-Teilsystem 112 mit einem oder mit mehreren der anderen Elemente der 1 integriert sein, um ein Einzelsystem auszubilden. Zum Beispiel kann das Parallelverarbeitungs-Teilsystem 112 mit der CPU 102 und einer anderen Verbindungsschaltung auf einem einzigen Chip integriert sein, um ein Ein-Chip-System („system an chip” (SOC)) auszubilden.
Es sei angemerkt, dass das System hier veranschaulichend dargestellt ist und dass Variationen und Veränderungen möglich sind. Die Verbindungstopologie einschließlich der Anzahl und der Anordnung von Brücken, der Anzahl von CPUs 102 und der Anzahl von Parallelverarbeitungs-Teilsystemen 112 kann wie gewünscht modifiziert werden. Zum Beispiel kann bei einigen Ausführungsformen der Systemspeicher 104 direkt mit der CPU 102, anstatt über die Speicher-Brücke 105, verbunden sein, und andere Bauelemente können über die Speicher-Brücke 105 und die CPU 102 mit dem Systemspeicher 104 kommunizieren. Bei anderen alternativen Topologien kann das Parallelverarbeitungs-Teilsystem 112 mit der I/O-Brücke 107 oder direkt mit der CPU 102, anstelle der Speicher-Brücke 105, verbunden sein. Bei noch anderen Ausführungsformen können die I/O-Brücke 107 und die Speicher-Brücke 105 in einem Einzelchip integriert sein, anstatt dass sie auf einem oder mehreren diskreten Bauelementen vorhanden sind. Schließlich können bei bestimmten Ausführungsformen eine oder mehrere in 1 dargestellte Komponenten nicht vorhanden sein. Zum Beispiel kann der Switch 116 weggelassen werden, und der Netzwerk-Adapter 118 und die Zusatzkarten 120, 121 können direkt mit der I/O-Brücke 107 verbunden sein.
2 ist gemäß verschiedenen erfindungsgemäßen Ausführungsformen ein Blockdiagramm einer Parallelverarbeitungs-Einheit (PPU) 202, welche in dem Parallelverarbeitungs-Teilsystem 112 der 1 vorhanden ist. Obwohl in 2 eine PPU 202 dargestellt ist, wie es vorab beschrieben ist, kann das Parallelverarbeitungs-Teilsystem 112 irgendeine Anzahl von PPUs aufweisen. Wie dargestellt ist, ist die PPU 202 mit einem lokalen Parallelverarbeitungs-(PP-)Speicher 204 gekoppelt. Die PPU 202 und der PP-Speicher 204 können mittels einer integrierten Schaltung oder mittels mehrerer integrierter Schaltungen ausgeführt sein, wie beispielsweise mittels programmierbarer Prozessoren, Anwendungs-spezifischen integrierten Schaltungen (ASICs) oder Speicherbauelementen oder in irgendeiner anderen technisch machbaren Weise.
Bei einigen Ausführungsformen umfasst die PPU 202 eine Grafik-Verarbeitungseinheit („Graphics Processing Unit” (GPU)), welche ausgestaltet sein kann, um eine Grafik-Render-Pipeline zu implementieren, welche verschiedene Operationen bezüglich einer Pixel-Erzeugung abhängig von Grafikdaten, welche durch die CPU 102 und/oder den Systemspeicher 104 zugeführt werden, ausführt. Wenn Grafikdaten bearbeitet werden, kann der PP-Speicher 204 als Bildspeicher eingesetzt werden, welcher einen oder mehrere herkömmliche Bildpufferspeicher speichert und, wenn es erforderlich ist, auch ein oder mehrere Render-Targets. Neben anderen Dingen kann der PP-Speicher 204 eingesetzt werden, um Pixeldaten zu speichern und zu aktualisieren und um einem Anzeigegerät 110 zur Anzeige endgültige Pixeldaten zuzuführen oder auf diesem Einzelbilder darzustellen. Bei einigen Ausführungsformen kann die PPU 202 auch für Operationen zur allgemeinen Verarbeitung und Berechnung ausgestaltet sein.
Im Betrieb ist die CPU 102 der Haupt-Prozessor des Computersystems 100 und steuert und koordiniert Operationen der anderen Systemkomponenten. Insbesondere gibt die CPU 102 Befehle aus, welche den Betrieb der PPU 202 steuern. Bei einigen Ausführungsformen schreibt die CPU 102 eine Folge von Befehlen für die PPU 202 in eine Datenstruktur (weder in 1 noch in 2 explizit dargestellt), welche sich in dem Systemspeicher 104, in dem PP-Speicher 204 oder in irgendeiner anderen Speicherstelle, welche sowohl von der CPU 102 als auch von der PPU 202 zugreifbar ist, befindet. Ein Zeiger dieser Datenstruktur wird in einen Schiebe-Pufferspeicher („pushbuffer”) geschrieben, um eine Verarbeitung der Folge der Befehle in der Datenstruktur zu initiieren. Die PPU 202 liest Befehlsfolgen von dem Schiebe-Pufferspeicher und führt dann die Befehle asynchron relativ zu dem Betrieb der CPU 102 aus. Bei Ausführungsformen, bei denen mehrere Schiebe-Pufferspeicher erzeugt werden, können Ausführungsprioritäten für jeden Schiebe-Pufferspeicher durch ein Applikationsprogramm über eine Geräte-Steuerung 103 spezifiziert werden, um eine Ablaufsteuerung der verschiedenen Schiebe-Pufferspeicher zu steuern.
Wie dargestellt ist, weist die PPU 202 eine I/O(Eingangs-/Ausgabe-)Einheit 205 auf, welche mit dem Rest des Computersystems 100 über den Kommunikationspfad 113 und die Speicher-Brücke 105 kommuniziert. Die I/O-Einheit 205 erzeugt Pakete (oder andere Signale) zur Übertragung auf dem Kommunikationspfad 113 und empfängt auch alle eingehenden Pakete (oder anderen Signale) von dem Kommunikationspfad 113 und leitet die eingehenden Pakete an die entsprechenden Komponenten der PPU 202 weiter. Zum Beispiel können Befehle, welche Verarbeitungsaufgaben betreffen, direkt an eine Host-Schnittstelle 206 geleitet werden, während Befehle, welche Speicheroperationen (z. B. Lesen von dem oder Schreiben auf den PP-Speicher 204) betreffen, an eine Koppelfeld-Einheit 210 geleitet werden. Die Host-Schnittstelle 206 liest von jedem Schiebe-Pufferspeicher und überträgt die Befehlsfolge, welche in dem Schiebe-Pufferspeicher gespeichert ist, an ein Frontend 212.
Wie bereits vorab in Verbindung mit 1 erwähnt wurde, kann die Verbindung der PPU 202 mit dem Rest des Computersystems 100 variieren. Bei einigen Ausführungsformen ist das Parallelverarbeitungs-Teilsystem 112, welches mindestens eine PPU 202 aufweist, als eine Zusatzkarte implementiert, welche in einen Erweiterungsschlitz des Computersystems 100 eingeführt sein kann. Bei anderen Ausführungsformen kann die PPU 202 mit einer Bus-Brücke, wie beispielsweise einer Speicher-Brücke 105 oder einer I/O-Brücke 107, auf einem Einzel-Chip integriert sein. Wiederum können bei noch anderen Ausführungsformen einige oder alle Elemente der PPU 202 zusammen mit der CPU 102 in einer einzigen integrierten Schaltung oder auf einem Ein-Chip-System (SoC) vorhanden sein.
Im Betrieb überträgt das Frontend 212 Verarbeitungsaufgaben, welche es von der Host-Schnittstelle 206 empfängt, an eine Arbeitsverteilungseinheit (nicht dargestellt) in der Aufgaben-/Arbeits-Einheit („task/work unit”) 207. Die Arbeitsverteilungseinheit empfängt Zeiger auf Verarbeitungsaufgaben, welche als Aufgaben-Metadaten („task metadata” (TMD)) codiert und in einem Speicher gespeichert sind. Die Zeiger auf die TMDs sind in einer Befehlsfolge enthalten, welche als ein Schiebe-Pufferspeicher gespeichert ist und von der Host-Schnittstelle 206 durch die Frontend-Einheit 212 empfangen wird. Verarbeitungsaufgaben, welche als TMDs codiert sind, weisen Indizes, welche den zu verarbeitenden Daten zugeordnet sind, wie auch Zustandsparameter und Befehle, welche definieren, wie die Daten zu verarbeiten sind, auf. Zum Beispiel können die Zustandsparameter und Befehle das Programm definieren, welches mit den Daten ausgeführt wird. Die Aufgaben-/Arbeits-Einheit 207 empfängt Aufgaben von dem Frontend 212 und stellt sicher, dass GPCs 208 mit einem gültigen Zustand konfiguriert sind, bevor die Verarbeitungsaufgabe, welche durch die jeweiligen TMDs spezifiziert ist, begonnen wird. Eine Priorität kann für die jeweiligen TMD spezifiziert werden, welche eingesetzt wird, um die Ausführung der Verarbeitungsaufgabe festzulegen. Verarbeitungsaufgaben werden auch von der Verarbeitungs-Clusteranordnung 230 empfangen. Optional kann die TMD einen Parameter aufweisen, welcher steuert, ob die TMD dem Anfang oder dem Ende einer Liste von Verarbeitungsaufgaben (oder einer Liste von Zeigern auf die Verarbeitungsaufgaben) zugeordnet ist, wodurch eine andere Stufe einer Steuerung über die Ausführungspriorität bereitgestellt wird.
Die PPU 202 implementiert vorteilhafterweise eine hochparallele Verarbeitungsarchitektur, welche auf einer Verarbeitungs-Clusteranordnung 230 basiert, die eine Gruppe von C allgemeinen Verarbeitungs-Clustern („general processing cluster” (GPC)”) beinhaltet, wobei C ≥ 1 gilt. Jedes GPC 208 ist in der Lage eine große Anzahl (z. B. hunderte oder tausende) von Threads gleichzeitig zu verarbeiten, wobei jeder Thread eine Instanz eines Programms ist. Bei verschiedenen Anwendungen können verschiedene GPCs 208 zur Bearbeitung von verschiedenen Typen von Programmen oder zur Durchführung von verschiedenen Typen von Berechnungen zugewiesen werden. Die Zuweisung von GPCs 208 kann sich abhängig von der Arbeitslast, welche für jeden Typ eines Programms oder Berechnung auftritt, verändern.
Die Speicher-Schnittstelle 214 weist eine Gruppe von D Partitions-Einheiten 215 auf, wobei D ≥ 1 gilt. Jede Partitions-Einheit 215 ist mit einem oder mit mehreren dynamischen Direktzugriffsspeichern (DRAMs) 220, welche sich innerhalb des PPM-Speichers 204 befinden, gekoppelt. Bei einer Ausführungsform entspricht die Anzahl der Partitions-Einheiten 215 der Anzahl der DRAMs 220, und jede Partitions-Einheit 215 ist mit einem anderen DRAM 220 gekoppelt. Bei anderen Ausführungsformen kann sich die Anzahl der Partitions-Einheiten 215 von der Anzahl der DRAMs 220 unterscheiden. Der Fachmann erkennt, dass ein DRAM 220 durch jedes andere technisch geeignete Speicherbauelement ersetzt werden kann. Im Betrieb können verschiedene Render-Targets, wie beispielsweise Texturabbildungen und Einzelbild-Speicherbereiche, über die DRAMs 220 gespeichert werden, was ermöglicht, dass die Partitions-Einheiten 215 Teile von jedem Render-Target parallel schreiben, um effizient die verfügbare Bandbreite des PP-Speichers 204 zu nutzen.
Ein bestimmtes GPC 208 kann Daten verarbeiten, welche zu jedem der DRAMs 220 in dem PP-Speicher 204 zu schreiben sind. Die Koppelfeld-Einheit 210 ist ausgestaltet, um die Ausgabe von jedem GPC 208 zu dem Eingang von jeder Partitions-Einheit 215 oder zu jedem anderen GPC 208 zur weiteren Verarbeitung zu routen. Die GPCs kommunizieren über die Koppelfeld-Einheit 210 mit der Speicher-Schnittstelle 214, um von verschiedenen DRAMs 220 zu lesen oder auf diese zu schreiben. Bei einer Ausführungsform weist die Koppelfeld-Einheit 210 eine Verbindung zu der I/O-Einheit 205, zusätzlich eine Verbindung mittels der Speicher-Schnittstelle 214 zu dem PP-Speicher 204 auf, wodurch ermöglicht wird, dass die verarbeitenden Kerne innerhalb der verschiedenen GPCs 208 mit dem Systemspeicher 104 oder mit einem anderen Speicher, welcher sich nicht lokal auf der PPU 202 befindet, kommunizieren. Bei der in 2 dargestellten Ausführungsform ist die Koppelfeld-Einheit 210 direkt mit der I/O-Einheit 205 verbunden. Bei verschiedenen Ausführungsformen kann die Koppelfeld-Einheit 210 virtuelle Kanäle verwenden, um Verkehrsströme zwischen den GPCs 208 und den Partitions-Einheiten 215 zu trennen.
Wiederum können die GPCs 208 programmiert werden, um Verarbeitungsaufgaben, welche eine große Vielzahl von Anwendungen betreffen, was ohne Einschränkung lineare und nicht-lineare Daten-Transformationen, ein Filtern von Video- und/oder Audio-Daten, Modellierungs-Operationen (z. B. Anwenden von physikalischen Gesetzen, um eine Position, Geschwindigkeit und andere Attribute von Objekten zu bestimmen), Bild-Render-Operationen (z. B. Tessellations-Shader-, Vertex-Shader-, Geometrie-Shader- und/oder Pixel-/Fragment-Shader-Programme), allgemeine Berechnungs-Operationen usw. einschließt. Im Betrieb ist die PPU 202 ausgestaltet, um Daten von dem Systemspeicher 104 und/oder dem PP-Speicher 204 zu einer oder zu mehreren On-Chip-Speichereinheiten zu übertragen, die Daten zu verarbeiten und die sich ergebenden Daten zurück zu dem Systemspeicher 104 und/oder dem PP-Speicher 204 zu schreiben. Auf die sich ergebenden Daten kann dann durch andere Systemkomponenten, was die CPU 102, eine andere PPU 202 innerhalb des Parallelverarbeitungs-Teilsystems 112 oder ein anderes Parallelverarbeitungs-Teilsystem 112 innerhalb des Computersystems 100 einschließt, zugegriffen werden.
Wie bereits vorab angemerkt ist, kann irgendeine Anzahl von PPUs 202 in einem Parallelverarbeitungs-Teilsystem 112 enthalten sein. Zum Beispiel können mehrere PPUs 202 auf einer einzigen Zusatzkarte vorhanden sein, oder mehrere Zusatzkarten können mit einem Kommunikationspfad 113 verbunden sein, oder eine oder mehrere PPUs 202 können in einem Brücken-Chip integriert sein. Die PPUs 202 in einem Multi-PPU-System können identisch oder verschieden voneinander sein. Zum Beispiel können unterschiedliche PPUs 202 eine unterschiedliche Anzahl von verarbeitenden Kernen und/oder einen unterschiedlich großen PP-Speicher 204 aufweisen. Bei Ausführungsformen, bei welchen mehrere PPUs 202 vorhanden sind, können solche PPUs parallel arbeiten, um Daten mit einem höheren Durchsatz zu verarbeiten, als es mit einer einzigen PPU 202 möglich wäre. Systeme, welche eine oder mehrere PPUs 202 umfassen, können in einer Vielzahl von Konfigurationen und Formfaktoren ausgeführt sein, was ohne Einschränkung Desktop-Computer, Laptops, tragbare Computer oder andere tragbare Geräte, Server, Workstations, Spielekonsolen, eingebettete Systeme und Ähnliches einschließt.
3 ist gemäß verschiedenen erfindungsgemäßen Ausführungsformen ein Blockdiagramm eines GPC 208, welches in der PPU 202 der 2 enthalten ist. Im Betrieb kann das GPC 208 ausgestaltet sein, um eine große Anzahl von Threads parallel auszuführen, um eine Bildverarbeitung, eine allgemeine Verarbeitung und/oder Rechenoperationen auszuführen. Ein „Thread” bezeichnet dabei eine Instanz eines bestimmten Programms, welche auf einen bestimmten Satz von Eingangsdaten ausgeführt wird. Bei einigen Ausführungsformen werden SIMD-Anweisungserstellungstechniken („Single-Instruction, Multiple-Data”) eingesetzt, um eine parallele Ausführung einer großen Anzahl von Threads zu unterstützen, ohne mehrere unabhängige Anweisungs-Einheiten bereitzustellen. Bei anderen Ausführungsformen werden SIMT-Techniken („Single-Instruction, Multiple-Tread”) eingesetzt, um eine parallele Ausführung einer großen Anzahl von im Allgemeinen synchronisierten Threads zu unterstützen, wobei eine gemeinsame Anweisungs-Einheit eingesetzt wird, welche ausgestaltet ist, um Anweisungen an eine Gruppe von Verarbeitungs-Maschinen innerhalb des GPC 208 auszugeben. Im Gegensatz zu einem SIMD-Ausführungssystem, bei welchem alle Verarbeitungs-Maschinen typischerweise identische Anweisungen ausführen, ermöglicht eine SIMT-Ausführung unterschiedlichen Threads einfacher divergierenden Ausführungs-Pfaden durch ein bestimmtes Programm zu folgen. Dem Fachmann ist bekannt, dass ein SIMD-Verarbeitungssystem eine funktionale Untermenge eines SIMT-Verarbeitungssystems repräsentiert.
Der Betrieb eines GPC 208 wird mittels eines Pipeline-Managers 305 gesteuert, welcher Verarbeitungs-Aufgaben, welche von einer Arbeitsverteilungs-Einheit (nicht dargestellt) innerhalb der Aufgaben-/Arbeits-Einheit 207 empfangen werden, an einen oder an mehrere Streaming-Multiprozessoren (SM) 310 verteilt. Der Pipeline-Manager 305 kann auch ausgestaltet sein, um ein Arbeitsverteilungs-Koppelfeld 330 zu steuern, indem Ziele für die verarbeitete Datenausgabe von den SMs 310 spezifiziert werden.
Bei einer Ausführungsform weist das GPC 208 einen Satz von M von SMs 310 auf, wobei M ≥ 1 gilt. Auch jeder SM 310 weist einen Satz von funktionalen Ausführungs-Einheiten (in 3 nicht dargestellt), wie beispielsweise Ausführungs-Einheiten und Lade-Speicher-Einheiten, auf. Verarbeitungsoperationen, welche für irgendeine der funktionalen Ausführungs-Einheiten bestimmt sind, können hintereinander ausgeführt werden („pipelined”), wodurch eine neue Anweisung zur Ausführung ausgegeben werden kann, bevor bei einer vorherigen Anweisung die Ausführung abgeschlossen worden ist. Jede Kombination von funktionalen Ausführungs-Einheiten innerhalb eines bestimmten SM 310 kann bereitgestellt werden. Bei verschiedenen Ausführungsformen können die funktionalen Ausführungs-Einheiten ausgestaltet sein, um eine Vielzahl von unterschiedlichen Operationen zu unterstützen, was Ganzzahl- und Fließkomma-Arithmetik (z. B. Addition und Multiplikation), Vergleichs-Operationen, Boolsche Operationen (AND, OR, XOR), Bit-Schiebe-Operationen und die Berechnung von verschiedenen algebraischen Funktionen (z. B. planare Interpolation und trigonometrische, exponentielle und logarithmische Funktionen, usw.) einschließt. Vorteilhafterweise kann dieselbe funktionale Ausführungseinheit ausgestaltet sein, um verschiedene Operationen auszuführen.
Im Betrieb ist jeder SM 310 ausgestaltet, um eine oder mehrere Thread-Gruppen zu verarbeiten. Eine „Thread-Gruppe” oder „Warp” bezeichnet dabei eine Gruppe von Threads, welche gleichzeitig dasselbe Programm mit verschiedenen Eingabedaten ausführen, wobei ein Thread der Gruppe einer anderen Ausführungs-Einheit innerhalb eines SM 310 zugewiesen ist. Eine Thread-Gruppe kann weniger Threads als die Anzahl von Ausführungs-Einheiten innerhalb des SM 310 aufweisen, wobei in diesem Fall die Ausführung während Zyklen inaktiv sein kann, wenn diese Thread-Gruppe verarbeitet wird. Eine Thread-Gruppe kann auch mehr Threads als die Anzahl der Ausführungs-Einheiten innerhalb des SM 310 aufweisen, wobei in diesem Fall eine Verarbeitung über aufeinanderfolgende Taktzyklen erfolgen kann. Da jeder SM 310 bis zu G Thread-Gruppen gleichzeitig unterstützen kann, folgt, dass bis zu G·M Thread-Gruppen zu einer bestimmten Zeit in einem GPC 208 ausgeführt werden können.
Darüber hinaus kann eine Mehrzahl von in Beziehung stehenden Thread-Gruppen zur selben Zeit innerhalb eines SM 310 aktiv (in verschiedenen Phasen der Ausführung) sein. Diese Sammlung von Thread-Gruppen wird hier als eine kooperative Thread-Anordnung („Cooperative Thread Array” (CTA)) oder eine „Thread-Anordnung” bezeichnet. Die Größe einer bestimmten CTA ist gleich m·k, wobei k die Anzahl von gleichzeitig laufenden Ausführungs-Threads in einer Thread-Gruppe ist, was typischerweise ein ganzzahliges Mehrfaches der Anzahl der Ausführungs-Einheiten innerhalb des SM 310 ist, und wobei m die Anzahl der Thread-Gruppen ist, welche gleichzeitig aktiv innerhalb des SM 310 sind.
Wie dargestellt ist, enthält, ohne Einschränkung, jeder SM 310 einen Gemeinschaftsspeicher 382 und einen Level-1-Cache bzw. L1-Cache 384. der Gemeinschaftsspeicher 382 ist typischerweise eine relativ kleine Sektion eines statischen Direktzugriffsspeichers („Static Random-Access Memory” (SRAM)), welcher sich lokal bei dem SM 310 befindet. Ein oder mehrere Abschnitte des Gemeinschaftsspeichers 382 werden gemeinsam von den Threads in einer CTA geteilt. Der L1-Cache 384 unterstützt, neben anderen Dingen, Lade- und Speicher-Operationen, welche durch die Ausführungseinheiten ausgeführt werden.
Jeder SM 310 weist auch einen Zugang zu Level-2-Caches bzw. L2-Caches (nicht dargestellt) auf, welche sich alle GPCs 208 in einer PPU 202 teilen. Die L2-Caches können eingesetzt werden, um Daten zwischen Threads zu übertragen. Schließlich weisen die SMs 310 auch einen Zugang zu einem externen (off-chip) Speicher auf, welcher den PP-Speicher 204 (auch als „globaler” Speicher bekannt) und/oder Systemspeicher 104 umfassen kann. Darüber hinaus kann, wie es in 3 dargestellt ist, ein Level-1.5-Cache bzw. L1.5-Cache 335 in dem GPC 208 enthalten und ausgestaltet sein, um Daten, welche durch den SM 310 über die Speicher-Schnittstelle 214 von einem Speicher angefordert werden, zu empfangen und zu halten. Solche Daten können ohne Einschränkung Anweisungen, einheitliche Daten und konstante Daten beinhalten. Bei Ausführungsformen, welche mehrere SMs 310 innerhalb des GPC 208 aufweisen, teilen sich die SMs 310 vorteilhafterweise gemeinsame Anweisungen und Daten, welche in dem L1.5-Cache 335 gecacht sind.
Jedes GPC 208 kann eine zugeordnete Speicher-Verwaltungseinheit („Memory Management Unit” (MMU)) 320 aufweisen, welche ausgestaltet ist, um virtuelle Adressen in physikalische Adressen umzusetzen. Bei verschiedenen Ausführungsformen kann die MMU 320 entweder in dem GPC 208 oder in der Speicher-Schnittstelle 214 liegen. Die MMU 320 umfasst einen Satz von Seiten-Tabellen-Einträgen („Page Table Entries” (PTE)), welche eingesetzt werden, um eine virtuelle Adresse in eine physikalische Adresse eines Tiles (z. B. Cache-Abschnitts) oder einer Speicherseite und optional einen Cache-Speicher-Zeilenindex umzusetzen. Die MMU 320 kann Adress-Übersetzungs-Pufferspeicher („Translation Lookaside Buffers” (TLB)) oder Cache-Speicher enthalten, welche in den SMs 310, in einem oder in mehreren L1.5-Cache-Speichern oder innerhalb des GPC 208 angeordnet sein können.
Bei Grafik- und Berechnungs-Anwendungen kann das GPC 208 ausgestaltet sein, so dass jeder SM 310 mit einer Textur-Einheit 315 gekoppelt ist, um Textur-Abbildungs-Operationen, wie beispielsweise die Bestimmung der Textur-Sample-Positionen, das Lesen von Textur-Daten und das Filtern von Textur-Daten, durchzuführen.
Im Betrieb überträgt jeder SM 310 eine verarbeitete Aufgabe zu einem Arbeitsverteilungs-Koppelfeld 330, um die verarbeitete Aufgabe einem anderen GPC 208 zur weiteren Verarbeitung bereitzustellen und um die verarbeitete Aufgabe in einem L2-Cachespeicher (nicht dargestellt), einem Parallelverarbeitungs-Speicher 204 oder einem Systemspeicher 104 über die Koppelfeld-Einheit 210 zu speichern. Darüber hinaus ist eine Einheit 325 für Vorrasterungs-Operationen („pre-raster operations” (preROP)) ausgestaltet, um Daten von dem SM 310 zu empfangen, Daten zu einer oder zu mehreren Einheiten für Rasterungs-Operationen innerhalb der Partitions-Einheiten 215 weiterzuleiten, um Optimierungen für Farbverläufe durchzuführen, um Pixel-Farbdaten zu organisieren und um Adress-Übersetzungen durchzuführen.
Es sei angemerkt, dass die Kernarchitektur hier illustrativ dargestellt ist, und dass viele Variationen und Änderungen möglich sind. Neben anderen Dingen kann irgendeine Anzahl von Verarbeitungs-Einheiten, wie beispielsweise SMs 310, Textur-Einheiten 315 oder preROP-Einheiten 325, in dem GPC 208 enthalten sein. Darüber hinaus kann die PPU 202, wie es im Zusammenhang mit 2 beschrieben ist, irgendeine Anzahl von GPCs 208 aufweisen, welche ausgestaltet sind, um funktional ähnlich zueinander zu sein, so dass das Ausführungsverhalten nicht davon abhängt, von welchem GPC 208 eine bestimmte Verarbeitungs-Aufgabe erhalten wird. Darüber hinaus arbeitet jedes GPC 208 unabhängig von den anderen GPCs 208 in der PPU 202, um Aufgaben für ein oder mehrere Anwendungsprogramme auszuführen. Unter Berücksichtigung des vorab Stehenden erkennt der Fachmann, dass die Architektur, welche in 1–3 beschrieben ist, auf keine Weise den Umfang der vorliegenden Erfindung beschränkt.
Erzeugen von Bildabschnitten
Im Allgemeinen kann der SM 310 ausgestaltet sein, um eine große Anzahl von Threads parallel auszuführen, um Grafik-, allgemeine Verarbeitungs- und/oder Rechen-Operationen auszuführen. Es sei angemerkt, dass die Gleichzeitigkeit und die bestimmten Speichermittel, welche von dem SM 310 bereitgestellt werden, typischerweise dem SM 310 ermöglichen, die Ausführung von rechenintensiven Operationen zu optimieren. Eine rechenintensive Operation, welche zur Ausführung durch den SM 310 besonders geeignet ist, ist die mehrfache Faltungsoperation. Typischerweise führen die SMs 310 bei herkömmlichen Techniken, welche Parallelverarbeitungs-Teilsysteme wirksam einsetzen, um mehrfache Faltungsoperationen auszuführen, optimierte Matrixmultiplikationsroutinen, welche in Bibliotheken vorhanden sind, aus.
Eine Einschränkung von solchen Matrix basierten Ansätzen zur Ausführung von mehrfachen Faltungsoperationen ist, dass der Speicher, welcher erforderlich ist, um effiziente Matrixmultiplikationsoperationen aufzubauen, den verfügbaren PP-Speicher 204 beansprucht. Genauer gesagt ist die Bildmatrix, welche die Eingabe für die Matrixmultiplikation darstellt, eine expandierte Version – die in erheblicher Weise redundante Daten enthält – des Bildstapels, welcher die Eingabe für das Mehrfachfaltungsbild darstellt. Im Betrieb führt der SM 310 die Matrixmultiplikationsoperationen auf Teilmatrizen, welche hier als Abschnitte (Tiles) bezeichnet werden, des Bildstapels aus. Um die optimierte Matrixmultiplikationsroutine auszunutzen, ohne den PP-Speicher 204 zu beanspruchen, erzeugt dementsprechend das Faltungsteilsystem 180 für jeden Bildabschnitt („image tile”) den Bildabschnitt, wenn er benötigt wird, verarbeitet den Bildabschnitt und verwirft dann den Bildabschnitt. Vorteilhafterweise wird nur ein Teil der Bildmatrix in dem Gemeinschaftsspeicher 382 zu einem bestimmten Zeitpunkt gespeichert. Bei alternativen Ausführungsformen kann das Faltungsteilsystem 180 jeden Typ von Eingangsdaten, was hier auch als „Samples” bezeichnet wird, anstatt von Bilddaten bearbeiten.
4 stellt einen Bildstapel 410, einen Filterstapel 440 und einen Ausgabestapel 470, welche einer mehrfachen Faltungsoperation zugeordnet sind, gemäß verschiedenen erfindungsgemäßen Ausführungsformen dar. Mit Bezug auf 4 ist der Streaming-Multiprozessor (SM) 310 ausgestaltet, um eine mehrfache Faltungsoperation zwischen dem Bildstapel 410 und dem Filterstapel 440 auszuführen, um den Ausgabestapel 470 zu erzeugen. Die mehrfache Faltungsoperation entspricht der vorherrschenden Rechnung, welche in eine Ausführung einer bestimmten Faltungsebene, die in einem CNN enthalten ist, einbezogen ist.
Wie dargestellt ist, weist der Bildstapel 410, ohne Einschränkung, irgendeine Anzahl von Eingabebildern 420(0: N – 1) auf. Zur Erläuterung sind mehrere Instanzen von ähnlichen Objekten mit Bezugszeichen bezeichnet, welche das Objekt identifizieren, und Zahlen in Klammern identifizieren die Instanz, wo dies erforderlich ist. Darüber hinaus ist ein Umfang von „X” ähnlichen Objekten mit einem Bereich in Klammern (d. h. (0: X – 1)) bezeichnet. Jedes der Eingabebilder 420 weist, ohne Einschränkung, irgendeine Anzahl von Farbebenen 430(0: C – 1) auf. Zum Beispiel kann jedes der Eingabebilder 420 drei Farbebenen 430: die Farbebene 430(0) „rot”, die Farbebene 430(1) „grün”, und die Farbebene 430(2) „blau” aufweisen. Jedes der Eingabebilder 420 ist einer Bildhöhe, welche als „H” dargestellt ist, und einer Bildbreite, welche als „W” dargestellt ist, zugeordnet. Es sei angemerkt, dass die Bildhöhe und die Bildbreite die Ausmaße von jeder der Farbebenen 430 definieren. Dementsprechend weist der Bildstapel 410 (N × C × H × W) Einzelwerte auf.
In ähnlicher Weise weist der Filterstapel 440, ohne Einschränkung, irgendeine Anzahl von Filtern 450(0: K – 1) auf. Bei einigen Ausführungsformen kann jedes der Filter 450 ein auslösendes Suchelement repräsentieren, welches der Ebene des CNN zugeordnet ist. Zum Beispiel kann das CNN in einem Gesichtserkennungsalgorithmus vorhanden sein, und das Filter 450(0) kann ein Ohr repräsentieren. Jedes der Filter 450 weist, ohne Einschränkung, Merkmalebenen 460(0: C – 1) auf, wobei die Anzahl der Merkmalebenen 460 gleich der Anzahl der Farbebenen 430 ist. Jedes der Filter 450 ist einer Filterhöhe, welche als „R” dargestellt ist, und einer Filterbreite, welche als „S” dargestellt ist, zugeordnet. Die Filterhöhe und die Filterbreite definieren die Ausmaße von jeder Merkmalebene 460. Daher weist der Filterstapel 440 (K × C × R × S) Einzelwerte auf.
Wie auch dargestellt ist, sind zahlreiche Parameter 465 der mehrfachen Faltungsoperation zugeordnet. Die Ausmaße des Bildstapels 410 und des Filterstapels 440 repräsentieren fünf unabhängige Parameter der mehrfachen Faltungsoperation: N (die Anzahl der Eingabebilder 420 in dem Bildstapel 410), C (die Anzahl der Farbebenen 430 in jedem der Eingabebilder 420 und die Anzahl der Merkmalebenen 460 in jedem der Filter 450), H (die Bildhöhe), W (die Bildbreite), K (die Anzahl der Filter 450 in dem Filterstapel 440), R (die Filterhöhe) und S (die Filterbreite). Die Parameter 465 umfassen auch, ohne Einschränkung, V (eine horizontale Filterschrittweite („filter stride”)) und (eine vertikale Filterschrittweite). Die horizontale Filterschrittweite und die vertikale Filterschrittweite verringern die Rechenlast, indem die Größe der Teilmenge von Pixeln, welche in die mehrfache Faltungsoperation einbezogen wird, verringert wird. Es sei angemerkt, dass die horizontale Filterschrittweite und die vertikale Filterschrittweite nicht nur die Zeit, welche erforderlich ist, um die mehrfache Faltungsoperation durchzuführen, reduzieren, sondern auch die Größe des Ausgabestapels 470 reduzieren, welcher durch die mehrfache Faltungsoperation erzeugt wird.
Bei alternativen Ausführungsformen können zusätzliche Parameter 465 zusätzlichen Ausmaßen, Schrittweiten, Optimierungs-, Formatierungs- und/oder anderen Konfigurations-Optionen entsprechen. Zum Beispiel können bei einigen Ausführungsformen die Parameter 465 eine Füllhöhe und eine Füllbreite umfassen. Die Füllhöhe und die Füllbreite fügen Reihen von Nullen und Spalten von Nullen den Ausgabebildern 480, welche in dem Ausgabestapel 470 vorhanden sind, aus irgendeinem technischen Grund, beispielsweise zur Formatierung für zukünftige Operationen, hinzu.
Der Ausgabestapel 470 weist, ohne Einschränkung, die Ausgabebilder 480(0: N – 1) auf, wobei die Anzahl der Ausgabebilder 480 der Anzahl der Eingabebilder 420 entspricht. Jedes der Ausgabebilder 480 weist, ohne Einschränkung, Merkmalkarten 490(0: K – 1) auf, wobei die Anzahl der Merkmalkarten 490 der Anzahl der Filter 450 entspricht. Jedes der Ausgabebilder 480 ist einer Ausgabehöhe, welche als „P” dargestellt ist, und einer Ausgabebreite, welche als „Q” dargestellt ist, zugeordnet. Die Ausgabehöhe und die Ausgabebreite definieren die Ausmaße der Merkmalkarten 490. Dementsprechend weist der Ausgabestapel 470 (N × K × P × Q) Einzelwerte auf.
Wie hier vorab beschrieben ist, setzt das Faltungsteilsystem 180 die optimierten Matrixmultiplikationseigenschaften des SM 310 wirksam ein, um die mehrfache Faltungsoperation effizient auszuführen. Der Fachmann erkennt, dass die mehrfache Faltungsoperation zwischen dem Eingabestapel 410 und dem Filterstapel 440 in Matrixmultiplikationsoperationen zwischen einer Bildmatrix und einer Filtermatrix umgesetzt werden kann. Die Umsetzoperationen sind nach dem Stand der Technik gut bekannt und führen zu deterministischen Beziehungen zwischen den Werten, welche in dem Eingabestapel 410 vorhanden sind, und den Werten, welche in der Bildmatrix vorhanden sind. In einer komplementären Weise ergeben die Umsetzoperationen deterministische Beziehungen zwischen den Werten, welche in dem Filterstapel 440 vorhanden sind, und den Werten, welche in der Filtermatrix vorhanden sind. Um den Einsatz des PP-Speichers 204 zu optimieren, speichert das Faltungsteilsystem 180 die Bildmatrix nicht in dem PP-Speicher 204. Stattdessen konfiguriert das Faltungsteilsystem 180 den SM 310 abhängig von diesen deterministischen Beziehungen – expandiert Bildabschnitte, welche in einer „virtuellen” Bildmatrix vorhanden sind, direkt.
5 stellt die Beziehung zwischen dem Bildstapel 410 der 4 und einer virtuellen Bildmatrix 510 gemäß verschiedenen erfindungsgemäßen Ausführungsformen dar. 5 stellt auch Beziehungen zwischen dem Filterstapel 440 der 4 und einer virtuellen Filtermatrix 540 dar. Für eine beispielhafte Erläuterung gilt für die Parameter 465 und damit für die Ausmaße des Bildstapels 410, der virtuellen Bildmatrix 510, des Filterstapels 440 und der virtuellen Filtermatrix 540: N = 1, C = 2, H = 3, W = 3, K = 2, R = 2, S = 2, U = 1 und V = 1. Darüber hinaus gilt entsprechend dem Speicher-Fußabdruck des Bildstapels 410 für die Bilddaten-Schrittweiten 475: Schrittweite C = 9, Schrittweite H = 3 und Schrittweite W = 1.
Als Teil der Erweiterung des Bildstapels 410 in die virtuelle Bildmatrix 510 ist jede der vier Spalten in der virtuellen Bildmatrix 510 den Werten zugeordnet, welche in dem Bildstapel 410 vorhanden sind und welche erforderlich sind, um eine Spalte von Werten in einer Ausgabematrix (nicht dargestellt) zu berechnen. Eine solche Erweiterung umfasst eine Vervielfältigung von einigen der Werte, welche in dem Bildstapel 410 vorhanden sind. Zum Beispiel wird, wie es für den Wert „D4” dargestellt ist, das Zentrum von jeder der 3×3-Farbebenen 410 viermal verwendet, um jede der vier Spalten in der Ausgabematrix zu berechnen, und folglich ist jeder der Zentrumswerte (z. B. der „D4”-Werte) vier getrennten Spalten der virtuellen Bildmatrix 510 zugeordnet. Für eine beispielhafte Erläuterung werden Werte, welche in dem Bildstapel 410 vorhanden sind, als die Daten bei den „Ursprungs”-Adressen bezeichnet. In ähnlicher Weise werden die Werte, welche in der virtuellen Bildmatrix 510 vorhanden sind, hier als die Daten bei den entsprechenden „virtuellen” Adressen, bezeichnet. Dementsprechend sind mehrere virtuelle Adressen in der virtuellen Bildmatrix 510 einer einzigen Ursprungsadresse, welche in dem Bildstapel 410 vorhanden ist, zugeordnet. In einer komplementären Weise enthält jede der Zeilen der virtuellen Filtermatrix 540 die Werte, welche in dem Filterstapel 440 vorhanden sind, und welche erforderlich sind, um einen oder mehrere der Abschnitte in einer Ausgabematrix zu berechnen.
Wenn die Ausmaße des Eingabestapels 410 (N × C × H × W) entsprechen, gelten für die Ausmaße des Filterstapels 440 im Allgemeinen (K × C × R × S) und für die Ausmaße des Ausgabestapels 470 (N × K × P × Q), wobei dann die Ausmaße der virtuellen Bildmatrix 510 (C × R × S) × (N × P × Q), die Ausmaße der virtuellen Filtermatrix 540 k × (C × R × S) und die Ausmaße der Ausgabe Matrix K × (N × P × Q) entsprechen. Für das dargestellte Beispiel entsprechen die Ausmaße des Eingabestapels 410 (1 × 3 × 3 × 3), die Ausmaße des Filterstapels 440 (2 × 3 × 2 × 2) und die Ausmaße des Ausgabestapels 470 (1 × 2 × 2 × 2). Dementsprechend entsprechen die Ausmaße der virtuellen Bildmatrix 510 (12 × 4), die Ausmaße der virtuellen Filtermatrix 540 (2 × 12) und die Ausmaße der Ausgabematrix (2 × 4).
Es sei angemerkt, dass, da die Ausmaße der virtuellen Bildmatrix 510 Produkten der unabhängigen Parameter, welche der mehrfachen Faltungsoperation zugeordnet sind, entsprechen, die Matrix basierte mehrfache Faltungsoperation ein relativ gleichartiges Verhalten über sich verändernde Parameter aufweist. Obwohl die Parameter C, R und S sich individuell dramatisch über die mehrfachen Faltungsoperationen, welche den verschiedenen Ebenen eines bestimmten CNN zugeordnet sind, verändern können, verändern sich die Produkte der Parameter C, R und S typischerweise nicht dramatisch über die mehrfachen Faltungsoperationen. Dementsprechend ist das optimierte Leistungsverhalten der Matrix basierten mehrfachen Faltungsoperation relativ konsistent gegenüber Veränderungen bei den Werten der individuellen Parameter.
Wie die Ausmaße (C × R × S) × (N × P × Q) der virtuellen Bildmatrix 510 zeigen, kann das gleichzeitige und redundante Speichern der Werte, welche den Adressen zugeordnet sind, die in der virtuellen Matrix 510 enthalten sind, den PP-Speicher 204 belasten. Daher ist das Faltungsteilsystem 180 des SM 310 ausgestaltet, um die virtuelle Bildmatrix 510 in einer hinterher speichernden Weise („lazy” manner) zu manifestieren und zu verarbeiten. Insbesondere unterteilt das Faltungsteilsystem 180 die virtuelle Bildmatrix 510 in getrennte Bildabschnitte 542 und konfiguriert den SM 310, um die Bildabschnitte 542 zu verarbeiten. Darüber hinaus ordnet das Faltungsteilsystem 180 jede der „Ziel”-Adressen in jedem der Bildabschnitte 542 einer virtuellen Adresse zu, welche in der virtuellen Bildmatrix 510 enthalten ist. Zum Beispiel ordnet die Faltungsmaschine 125, wie es in 5 dargestellt ist, die sechzehn Zieladressen, welche in dem Bildabschnitt 542(0) enthalten sind, den sechzehn virtuellen Adressen zu, welche in den ersten vier Zeilen der virtuellen Bildmatrix 510 enthalten sind.
Jede der virtuellen Adressen in der virtuellen Bildmatrix 510 bezieht sich deterministisch auf eine Ursprungsadresse, welche in dem Bildstapel 410 vorhanden ist. Daher bezieht sich jede der Zieladressen in den Bildabschnitten 542 deterministisch auf eine Ursprungsadresse, welche in dem Bildstapel 410 enthalten ist. Dementsprechend kann das Faltungsteilsystem 180 Adressberechnungen durchführen, welche es dem Faltungsteilsystem 180 ermöglichen, die geeigneten Daten von dem Bildstapel 410 direkt zu jeder Zieladresse zu kopieren, welche in jedem der Bildabschnitte 542 enthalten ist, ohne die virtuelle Eingabematrix 510 zu erzeugen.
Um die Einführung von ganzzahligen Wartezeiten zu vermeiden, welche der Durchführung von Adressberechnungen während der Erzeugung der Bildabschnitte 542 zugeordnet sind, ist der Faltungspräprozessor 182 ausgestaltet, um ein konsistentes Muster von Ursprungsadressen wirksam einzusetzen, welches sich inhärent in der virtuellen Bildmatrix 510 befindet. Genauer gesagt sind, während jede Spalte der virtuellen Bildmatrix 510 einer anderen Folge von Ursprungsadressen zugeordnet ist, welche einem gewundenen Pfad durch den Bildstapel 410 folgen, die Folgen (in einer multidimensionalen Weise) affin.
Zum Beispiel ist die erste Spalte der virtuellen Bildmatrix 510 der ersten Ursprungsadressfolge D0, D1, D3 und D4 für jede der drei Farbebenen 430 zugeordnet. Die Addition von 1 zu jedem Element in dieser ersten Folge ergibt D1, D2, D4 und D5 für jede der drei Farbebenen 430 – entsprechend der Ursprungsadressfolge, welche der zweiten Spalte der virtuellen Bildmatrix 510 zugeordnet ist. In ähnlicher Weise ergibt die Addition von 3 zu jedem Element in der ersten Folge die Ursprungsadressfolge, welche der dritten Spalte der virtuellen Bildmatrix 510 zugeordnet ist, usw. Ein Beispiel, wie der Faltungspräprozessor 182 dieses konsistente Muster wirksam einsetzt, wird mit mehr Details in 6 beschrieben.
Als Teil einer Verarbeitung jedes der Bildabschnitte 542 lädt der SM 310 Daten von dem Bildstapel 410, um den Bildabschnitt 542 auszubilden, und lädt Daten von dem Filterstapel 440, um den entsprechenden Filterabschnitt 544 auszubilden. Der SM 310 führt dann Matrixmultiplikationsoperationen zwischen dem Bildabschnitt 542 und dem Filterabschnitt 544 durch, speichert das Ergebnis als einen Ausgabeabschnitt in dem PP-Speicher 204 und verwirft die Daten dann, welche in dem Bildabschnitt 542 und dem Filterabschnitt 544 vorhanden sind. Daher weist zu jedem bestimmten Zeitpunkt der Gemeinschaftsspeicher 382 die Bildabschnitte 542 auf, welche der SM 310 gerade bearbeitet, aber weist nicht notwendigerweise die Bildabschnitte 542 auf, welche der SM 310 bereits verarbeitet hat oder noch nicht begonnen hat zu verarbeiten.
Das Faltungsteilsystem 180 kann die Größe des Bildabschnitts 542 in jeder technisch möglichen Weise einstellen, welche die Fähigkeiten des SM 310 optimiert. Zum Beispiel kann das Faltungsteilsystem 180 die Größe des Bildabschnitts 542 abhängig von irgendeiner Anzahl und Kombination aus der Größe des Gemeinschaftsspeichers 382, der Anzahl der Threads in jeder Threadgruppe und so weiter einstellen. Bei alternativen Ausführungsformen kann das Faltungsteilsystem 180 die Größe des Bildabschnitts 542 als eine Hilfseingabe für die mehrfache Faltungsoperation erfassen. Das Faltungsteilsystem 180 stellt die Größe des Filterabschnitts 544 abhängig von der Größe des Bildabschnitts 542 ein. Genauer gesagt stellt das Faltungsteilsystem 180 die Größe des Filterabschnitts 545 ein, so dass die Matrixmultiplikation zwischen jedem Bildabschnitt 542 und dem entsprechenden Filterabschnitt 544 die Daten erzeugt, welche geeignet einen Ausgabeabschnitt füllen.
Bei alternativen Ausführungsformen kann das Faltungsteilsystem 180 den SM 310 abhängig von irgendeiner technisch geeigneten Implementierung der virtuellen Bildmatrix 510 und der virtuellen Filtermatrix 540 konfigurieren, welche die Durchführung der mehrfachen Faltungsoperation über Matrixmultiplikationsoperationen erleichtert. Darüber hinaus kann das Faltungsteilsystem 180 die Daten, welche in der virtuellen Bildmatrix 510 und der virtuellen Filtermatrix 540 enthalten sind, in Bildabschnitte 542 und Filterabschnitte 544 in irgendeiner technisch geeigneten, konsistenten Weise unterteilen.
Erzeugen der Offsetfolge
6 stellt die Beziehungen zwischen dem Bildstapel 410 der 4, einer Offsetfolge 640 und der virtuellen Bildmatrix 510 der 5 gemäß verschiedenen erfindungsgemäßen Ausführungsformen dar. Zur beispielhaften Erläuterung ist der Bildstapel 410 als Bilddaten 610 und als Speicheroffsets 620 dargestellt. Wie dargestellt ist, hängt der Speicheroffset 620 für einen bestimmten Wert der Bilddaten 610 von der Farbebene 430, der vertikalen Lage innerhalb des Bildes, der horizontalen Lage innerhalb des Bildes und der Datengröße, welche dem Wert der Bilddaten 610 zugeordnet ist, ab. Zum Beispiel ist der Speicheroffset 620 für D4 (1·2·2)·4 Byte = 16.
Wie mit Bezug zu 5 erläutert ist, repräsentiert das gewundene Muster eine gleichförmige Folge von Offsets für jede Zeile der virtuellen Bildmatrix 510, während das gewundene Muster jeder Spalte gegenüber dem gewundenen Muster der anderen Spalten versetzt ist. Zum Beispiel ist für jede Spalte der virtuellen Bildmatrix 510 die Ursprungsadresse, welche der zweiten Zeile zugeordnet ist, um 4 größer als die Ursprungsadresse, welche der ersten Zeile zugeordnet ist. Im Betrieb erzeugt der Faltungspräprozessor 182 die Offsetfolge 640 abhängig von diesem gewundenen Muster. Die Offseffolge 640 umfasst einen Offset 642 für jede Zeile der virtuellen Bildmatrix 510. Der Faltungspräprozessor 182 kann die Offseffolge 640 auf jede technisch mögliche Weise erzeugen, welche jedes Muster inhärent in der virtuellen Bildmatrix 510 erfasst.
Darüber hinaus ist das gewundene Muster von jeder Spalte versetzt gegenüber dem gewundenen Muster der anderen Spalten. Zum Beispiel ist die erste Spalte der virtuellen Bildmatrix 510 der Ursprungsadressfolge 0, 4, 12, 16, 26, 40, 48, 52, 72, 76, 84 und 88 zugeordnet. Eine Addition von 4 zu jeder Ursprungsadresse, welche in dieser Folge enthalten ist, ergibt die Ursprungsadressfolge 4, 8, 16, 20, 40, 44, 52, 56, 76, 80, 88 und 92, welche der zweiten Spalte der virtuellen Bildmatrix 510 zugeordnet ist. Der Fachmann erkennt, dass ein Spaltenoffset 632, welcher eine Differenz zwischen der ersten Spalte der virtuellen Bildmatrix 510 und einer bestimmten Spalte der virtuellen Bildmatrix 510 spezifiziert, der Ursprungsadresse entspricht, welche der ersten Zeile der bestimmten Spalte zugeordnet ist.
7 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen den Faltungspräprozessor 182 der 1 dar, welcher ausgestaltet ist, um die Offseffolge 640 der 6 zu erzeugen. Wie dargestellt ist, stellt 7 einen Offsetfolgegenerator 720 als einen Pseudocode dar, welcher den Faltungspräprozessor 182 implementieren kann. Der Offsetfolgegenerator 720 erzeugt die Offseffolge 640, um die Faltungsmaschine 184 von zeitaufwändigen Adressberechnungen zu entlasten. Solch eine Entlastung ermöglicht es dem SM 310, eine optimierte Matrixmultiplikationsroutine effizient auszuführen, um Faltungsoperationen auszuführen.
Im Allgemeinen kann der Offsetfolgegenerator 720 auf irgendeine Art und irgendeiner Anzahl von Verarbeitungseinheiten ausgeführt werden und die Offseffolge 640 in irgendeinem zugreifbaren Speicher speichern. Zum Beispiel kann der Offsetfolgegenerator 720 auf der CPU 102 ausgeführt werden und die Offseffolge 640 in dem Systemspeicher 104 speichern. Daher kann der Gerätetreiber 175 die Offsetfolge 640 in den PP-Speicher 204 kopieren. Wenn die Faltungsmaschine 184 die Offsetfolge 640 erfasst, kann die Offsetfolge 640 in den L1-Cache 384 geladen werden. Bei anderen Ausführungsformen kann der Offsetfolgegenerator 720 auf dem Parallelverarbeitungs-Teilsystem 112 ausgeführt werden. Bei solchen Ausführungsformen kann der Offsetfolgegenerator 720 die Offsetfolge 640 direkt in den PP-Speicher 204 speichern.
Wie dargestellt ist, erzeugt der Offsetfolgegenerator 720 (C·R·S) Offsets 642 abhängig von der Farbebene (C) 430, der Filterhöhe (R), der Filterbreite (S), der Bildhöhen-Schrittweite (Schrittweite H) und der Bildbreiten-Schrittweite (Schrittweite W). Bei alternativen Ausführungsformen kann der in 7 dargestellte Pseudocode modifiziert werden, um irgendeine Anzahl von Dimensionen zu repräsentieren. Darüber hinaus kann, wie es in dem Pseudocode dargestellt ist, der Offsetfolgegenerator 720 ausgestaltet sein, um eine Offsetfolge 640 zu erzeugen, welche eine Kreuzkorrelationsoperation anstelle einer Faltungsoperation darstellt.
Wie vorab beschrieben ist, spezifiziert die Offsetfolge 640 eine Zuordnung von Ursprungsadressen in dem Bildstapel 410 zu virtuellen Adressen in der virtuellen Bildmatrix 510. Da die Faltungsmaschine 125 jeder Zieladresse in jedem Bildabschnitt 542 eine virtuelle Adresse in der virtuellen Bildmatrix 510 zuordnet, setzt die Faltungsmaschine 125 die Offsetfolge 640 wirksam ein, um die Bildabschnitte 542 geeignet zu füllen. Im Allgemeinen können der Faltungspräprozessor 182 und die Faltungsmaschine 125 ausgestaltet sein, um irgendein Muster, welches inhärent in der virtuellen Bildmatrix 510 vorhanden ist, abhängig von irgendeiner Anzahl von Offsets 642 und irgendeiner Anzahl von Offsetfolgen 640 auszunutzen.
Durchführen von Matrix basierten mehrfachen Faltungsoperationen
8 stellt gemäß verschiedenen erfindungsgemäßen Ausführungsformen die Faltungsmaschine 184 der 1 dar, welche ausgestaltet ist, um eine mehrfache Faltungsoperation durchzuführen. Mit Bezug zu 8 konfiguriert die Faltungsmaschine 184 funktionale Einheiten (z. B. Ausführungseinheiten, Lade-Speicher-Einheiten, usw.), welche in dem Streaming-Multiprozessor (SM) 310 enthalten sind, um Operationen auszuführen, welche mehrfache Faltungsoperationen implementieren. Zur beispielhaften Erläuterung werden Operationen, welche durch den SM 310 ausgeführt werden, was die funktionalen Ausführungseinheiten beinhaltet, die durch die Faltungsmaschine 184 konfiguriert werden, hier auch als Operationen bezeichnet, die durch die Faltungsmaschine 184 ausgeführt werden.
Im Betrieb weist, um Parallelverarbeitungseigenschaften des SM 310 auszunutzen, die Faltungsmaschine 184 die Verarbeitung von jedem der Bildabschnitte 542 einer Threadgruppe oder einer Threadanordnung zu. Der Fachmann erkennt, dass die Faltungsmaschine 184 irgendeine Anzahl von Bildabschnitten einer einzigen Threadgruppe zuordnen kann und/oder irgendeine Anzahl von Operationen oder Zieladressen in den Bildabschnitten 542 einem einzigen Thread zuweisen kann. Wenn einer Threadgruppe eine Bearbeitung von mehreren Bildabschnitten 542 zugewiesen wird, dann kann die Threadgruppe die zugewiesenen Bildabschnitte 542 nacheinander verarbeiten oder kann die Verarbeitung in irgendeiner technisch möglichen Weise zwischen den Threads, welche in der Threadgruppe enthalten sind, aufteilen. Wenn einem Thread zugewiesen wird, mehrere Zieladressen in dem Bildabschnitt 542 zu bearbeiten, dann kann der Thread die zugewiesenen Zieladressen sequenziell verarbeiten.
Vorteilhafterweise kann die Faltungsmaschine 184 den SM 310 konfigurieren, um die Bearbeitung der Bildabschnitte 542 hintereinander auszuführen, um die Wartezeit, welche einem Zugriff auf die Eingabedaten, die in dem PP-Speicher 210 enthalten sind, zu minimieren. Genauer gesagt, kann die Faltungsmaschine 184 den SM 310 konfigurieren, um Daten, welche in dem Bildstapel 410 bzw. dem Filterstapel 440 enthalten sind, in den Bildabschnitt 542(0) bzw. den Filterabschnitt 144(0) zu kopieren. Die Faltungsmaschine 184 kann den SM 310 konfigurieren, um dann Matrixmultiplikationsoperationen zwischen dem Bildabschnitt 542(0) und dem Filterabschnitt 544(0) auszuführen und um, im Wesentlichen parallel, Daten, welche in dem Bildstapel 410 bzw. dem Filterstapel 440 vorhanden sind, in den Bildabschnitt 542(1) bzw. den Filterabschnitt 544(1) zu kopieren. Bei alternativen Ausführungsformen kann die Faltungsmaschine 184 irgendeine Art von Pipeline-Verarbeitung in irgendeiner technisch möglichen Weise einstellen. Zum Beispiel kann, ohne Einschränkung, die Faltungsmaschine 184 die Verarbeitung der Bildabschnitte 542 strategisch Threadgruppen zuweisen, um eine zweistufige Pipeline-Verarbeitung (ein Laden von Daten und ein Durchführen von Matrixmultiplikationsoperationen) zu ermöglichen.
Wie dargestellt ist weist die Faltungsmaschine 184, ohne Einschränkung, Spaltenoffsets 632 und Startadressen 634 auf. Der PP-Speicher 204 weist, ohne Einschränkung, den Filterstapel 440, den Bildstapel 410, die Offsetfolge 640 und eine Ausgabematrix 860 auf. Wie es als gestrichelter Kasten dargestellt ist, erfasst die Faltungsmaschine 184 typischerweise die Offsetfolge 640 durch den L1-Cache 384. Der Gemeinschaftsspeicher 382 weist, ohne Einschränkung, den Bildabschnitt 542 und den Filterabschnitt 544 auf. Für jede Threadgruppe bestimmt die Faltungsmaschine 184 die Ursprungsadressen in dem Bildstapel 410 entsprechend der Zieladressen in dem Bildabschnitt 542, welcher der Threadgruppe zugewiesen ist.
Genauer gesagt berechnen die Threads in der zugewiesenen Threadgruppe für jeden Bildabschnitt 542 zusammenarbeitend die Spaltenoffsets 632, welche den Spalten zugewiesen sind, die in dem Bildabschnitt 542 enthalten sind. Zum Beispiel würden die Threads für einen 4·4-Bildabschnitt 542 vier Spaltenoffsets 632 berechnen, welche den vier Spalten entsprechen. Genauer gesagt stellt die zugewiesene Threadgruppe für jede der relevanten Spalten den Spaltenoffset 632 auf die Differenz zwischen der Ursprungsadresse, welche der ersten Zeile der relevanten Spalte zugeordnet ist, und der Ursprungsadresse, welche der ersten Zeile der ersten Spalte zugeordnet ist, ein.
Daher summiert die zugewiesene Threadgruppe für jeden der Spaltenoffsets 632 die Basisadresse des Bildstapels 410 und den Spaltenoffset 632, um die Startadresse 634 zu erzeugen. Der Fachmann erkennt, dass die Startadressen 634 den Ursprungsadressen in dem Bildstapel 410 entsprechen, welche der ersten Zeile der virtuellen Bildmatrix 510 zugeordnet sind. Nachdem die Startadressen 634 für den Bildabschnitt 542 berechnet sind, führen die Threads in der zugewiesenen Threadgruppe eine Bilddatenadressierung 850 durch, um den Bildabschnitt 542 zu erzeugen.
Zuerst erfasst die Threadgruppe den Teil der Offsetfolge 640, welcher den Zeilen der virtuellen Bildmatrix 510 zugeordnet ist, welche dem Bildabschnitt 542 entsprechen. Wie es vorab beschrieben ist, ist jede Zieladresse, welche in dem Bildabschnitt 542 enthalten ist, einer virtuellen Adresse zugeordnet, welche in der virtuellen Bildmatrix 510 enthalten ist. Dementsprechend füllt die Threadgruppe den Bildabschnitt 542 abhängig von der Zuordnung zwischen der Ursprungsadresse, welche in dem Bildstapel 410 enthalten ist, und der virtuellen Adresse, welche in der virtuellen Bildmatrix 510 enthalten ist. Insbesondere berechnet die Threadgruppe die Ursprungsadresse, welche einer Spalte der virtuellen Bildmatrix 510 und einer Zeile der virtuellen Bildmatrix 510 entspricht, als die Summe der Startadresse 634, welche der Spalte zugeordnet ist, und des Offsets 642, welcher der Zeile zugeordnet ist. Nachdem jede Ursprungsadresse berechnet worden ist, kopiert die Threadgruppe die Bilddaten 610, welche bei der Ursprungsadresse (welche in dem Bildstapel 410 enthalten ist) gespeichert ist, an die entsprechende Zieladresse, welche in dem Bildabschnitt 542 enthalten ist.
Da die Faltungsmaschine 184 die Bildabschnitte 542 auf Anforderung erzeugt, speichert die Faltungsmaschine 184 zu einem bestimmten Zeitpunkt nur einen Teil der virtuellen Bildmatrix 510 in dem Gemeinschaftsspeicher 382. Die Zeit, welche für die Threads erforderlich ist, um die Bilddatenadressierung 850 auszuführen, ist wesentlich kürzer als die Zeit, welche erforderlich ist, um die zahlreichen, abhängigen Ganzzahladressierungsoperationen auszuführen, welche in einigen herkömmlichen Abschnitt basierten Faltungsmaschinen implementiert sind.
Die Threadgruppe erzeugt auch den Filterabschnitt 544, welcher dem Bildabschnitt 542 zugeordnet ist. Die Threadgruppe kann Daten von dem Filterstapel 440, welcher in dem PP-Speicher 204 enthalten ist, in den Filterabschnitt 544, welcher in dem Gemeinschaftsspeicher 382 enthalten ist, auf irgendeine technisch mögliche Weise kopieren, welche konsistent zu den Daten ist, die in dem Bildabschnitt 542 enthalten sind. Zum Beispiel kann die Threadgruppe eine lineare Zuordnung zwischen dem Filterstapel 440 und dem Filterabschnitt 544 abhängig von den Ursprungsadressen, welche dem Bildabschnitt 542 zugeordnet sind, implementieren.
Nachdem jede Threadgruppe die Erzeugung des zugewiesenen Bildabschnitts 542 und des entsprechenden Filterabschnitts 544 abgeschlossen hat, arbeitet die Threadgruppe mit einer Fließkommaeinheit, welche in dem SM 310 enthalten ist, wobei die Funktionalität einer Matrixmultiplikation 855 pro Abschnitt implementiert wird. Genauer gesagt konfiguriert jede der Threadgruppen die Fließkommaeinheit, um Matrixmultiplikationsoperationen zwischen dem zugewiesenen Bildabschnitt 542 und dem entsprechenden Filterabschnitt 544 durchzuführen. Die Threadgruppe konfiguriert darüber hinaus die Fließkommaeinheit, um die Ausgabematrix 860 abhängig von den Ergebnissen der Matrixmultiplikation zu aktualisieren.
Nachdem die Threadgruppen eine Erzeugung eines Ausgabeabschnitts, welcher in der Ausgabematrix 860 vorhanden ist, abgeschlossen haben, setzen die Threadgruppen die Ausgabematrix 860 in den Ausgabestapel 470 (in 8 nicht dargestellt), welcher auch in dem PP-Speicher 204 enthalten ist, um. Die Threadgruppen können irgendeine Anzahl von Formatierungsoperationen implementieren, welche den Ausgabestapel 470 abhängig von irgendeiner Organisation oder irgendeiner Teilmenge oder Obermenge der Daten, welche in der Ausgabematrix 860 enthalten ist, erzeugen. Typischerweise wird der Ausgabestapel 470 in einem Format implementiert, welches mit dem Format des Bildstapels 410 konsistent ist, wodurch es möglich ist, dass der Ausgabestapel 470 als der Eingabestapel 410 für die mehrfache Faltungsoperation verwendet wird, welche für die nächste Faltungsebene, die in dem CNN enthalten ist, implementiert ist.
Im Allgemeinen können Komponenten, welche in dem Computersystem 100 enthalten sind, den Bildstapel 410, den Filterstapel 440, die Offsetfolge 640 und/oder die Ausgabematrix 860 in irgendeiner Art einer Speicherstruktur, welche in dem PP-Speicher 204 enthalten ist, speichern. Zum Beispiel kann irgendeine Anzahl, was Null einschließt, von dem Bildstapel 410, dem Filterstapel 440, der Offsetfolge 640 und/oder der Ausgabematrix 860 in einem Bildspeicher enthalten sein. Bei anderen Ausführungsformen können Komponenten, welche in dem Computersystem 100 enthalten sind, den Bildstapel 410, den Filterstapel 440, die Offsetfolge 640 und/oder die Ausgabematrix 860 in irgendeiner Art von Speicher, anstelle des PP-Speichers 204, speichern. In ähnlicher Weise kann bei alternativen Ausführungsformen die Faltungsmaschine 125 die Bildabschnitte 542 und die Filterabschnitte 544 in irgendeiner Art von Speicher, anstelle des Gemeinschaftsspeichers 382, speichern.
Bei alternativen Ausführungsformen kann das Faltungsteilsystem 180 die Offsetfolge 640 konfigurieren, um irgendeine Anzahl von Offsets 642 abhängig von jeder Anzahl der Parameter 465 und jeder Anzahl der Bilddaten-Schrittweiten 475 in irgendeiner Kombination zu umfassen. Zum Beispiel können sich die Anzahl der Ausmaße bzw. Dimensionen, welche dem Bildstapel 410 zugeordnet sind, die Anzahl von Ausmaßen bzw. Dimensionen, welche dem Filterstapel 440 zugeordnet sind, die Bilddaten-Schrittweiten 475 und die Anzahl von Zeilen, welche in der virtuellen Bildmatrix 510 enthalten sind, über verschiedene Ausführungsformen hinweg unterscheiden. Darüber hinaus kann bei einigen Ausführungsformen die Offsetfolge 640 nur eine Teilmenge von Dimensionen bzw. Ausmaßen umfassen. Zum Beispiel kann der Faltungspräprozessor 182 die Offsets 642 für jede der (R·S) Zeilen der virtuellen Bildmatrix 510 abhängig von der Bildhöhen-Schrittweite und der Bildbreiten-Schrittweite berechnen. In einer komplementären Weise kann die Faltungsmaschine 184 konfiguriert sein, um die Offsets 642 wiederholt – einmal für jede der C Farbebenen 430 – für jede der Spalten der virtuellen Bildmatrix 510 anzuwenden, um die Ursprungsadressen zu bestimmen, welche jeder der (C·R·S) Zeilen der virtuellen Bildmatrix 510 zugeordnet sind.
In 9 ist ein Flussplan von Verfahrensschritten zur Durchführung einer mehrfachen Faltungsoperation in einem Parallelverarbeitungssystem gemäß verschiedenen erfindungsgemäßen Ausführungsformen dargestellt. Obwohl die Verfahrensschritte mit Bezug zu den Systemen der 1–8 beschrieben sind, erkennt der Fachmann, dass jedes System, welches ausgestaltet ist, um die Verfahrensschritte in irgendeiner Reihenfolge auszuführen, unter den Umfang der vorliegenden Erfindung fällt.
Wie dargestellt ist, beginnt ein Verfahren 900 bei Schritt 902, bei welchem das Faltungsteilsystem 180 den Bildstapel 410 und den Filterstapel 440 empfängt. Bei Schritt 904 berechnet der Faltungspräprozessor 182 die Offsetfolge 640, welche der virtuellen Bildmatrix 510 zugeordnet ist, und speichert dann die Offsetfolge 640 in den PP-Speicher 204. Der Faltungspräprozessor 182 kann die Offsetfolge 640 in irgendeiner technisch möglichen Weise berechnen, welche ein Muster repräsentiert, welches eine Teilmenge von Ursprungsadressen in dem Bildstapel 410 entsprechenden virtuellen Adressen entlang der Bildmatrixspalten 544 der virtuellen Bildmatrix 510 abhängig von der Zeile zuordnet. Zum Beispiel sei angenommen, dass die Dimensionen bzw. Ausmaße des Eingabestapels 410 (N × C × H × W) und dass die Dimensionen bzw. Ausmaße des Filterstapels 440 (K × C × R × S) betragen. Bei einem solchen Beispiel kann der Faltungspräprozessor 182 die Offsets 642 für jede der (C·R·S) Zeilen der virtuellen Bildmatrix 510 abhängig von der Schrittweite C, Schrittweite H und der Schrittweite W berechnen.
Der Faltungspräprozessor 182 kann auf irgendeiner Verarbeitungseinheit (z. B. der CPU 182, dem SM 310, usw.) ausgeführt werden und die Offsetfolge 640 in dem PP-Speicher 204 auf irgendeine technisch mögliche Weise speichern. Zum Beispiel kann der Faltungspräprozessor 182 bei einigen Ausführungsformen auf der CPU 182 ausgeführt werden und die Offsetfolge 640 in dem Systemspeicher 104 speichern. Daher kann der Gerätetreiber 175 die Offsetfolge 640 von dem Systemspeicher 104 in den PP-Speicher 204 kopieren. Wenn der SM 310 anfänglich die Offsetfolge 640 erfasst, kann der SM 310 die Offsetfolge 640 in den L1-Cache 384 laden.
Bei Schritt 906 bestimmt das Faltungsteilsystem 180 die Größe des Bildabschnitts 542 und definiert dann die Bildabschnitte 542 – wobei Zieladressen in jedem der Bildabschnitte 542 virtuellen Adressen in der virtuellen Bildmatrix 510 zugeordnet werden. Bei Schritt 908 weist die Faltungsmaschine 184 die Verarbeitung von jedem der Bildabschnitte 542 einer Threadgruppe zu. Die Faltungsmaschine 184 konfiguriert dann den SM 310, um die Threadgruppen auszuführen.
Bei Schritt 910 berechnet für jeden der Bildabschnitte 542 die zugewiesene Threadgruppe die Startadressen 634, welche den Spalten, die in dem Bildabschnitt 542 enthalten sind, zugeordnet sind. Genauer gesagt stellt die zugewiesene Threadgruppe für jede Spalte die Startadressen 634 auf die Ursprungsadresse ein, welche in dem Bildstapel 410 enthalten ist und welche der Spalte der virtuellen Bildmatrix 510 und der ersten Zeile der virtuellen Bildmatrix 510 zugeordnet ist. Bei Schritt 912 berechnet dann die Threadgruppe für jeden Bildabschnitt 542 die Ursprungsadressen in dem Bildstapel 410 und kopiert die Daten, welche bei den Ursprungsadressen gespeichert sind, zu den entsprechenden Zieladressen in dem zugewiesenen Bildabschnitt 542, welcher in dem Gemeinschaftsspeicher 382 gespeichert ist. Es sei angemerkt dass die Threads in der zugewiesene Threadgruppe die Ursprungsadresse, welche einer bestimmten Spalte der virtuellen Bildmatrix 510 und einer bestimmten Zeile der virtuellen Bildmatrix 510 entspricht, als die Summe der Startadresse 634, welche der Spalte zugeordnet ist, und des Offsets 642, welcher der Zeile zugeordnet ist, berechnen.
Bei Schritt 916 führt die Threadgruppe für jeden Bildabschnitt 542 Matrixmultiplikationsoperationen zwischen dem Bildabschnitt 514 und dem entsprechenden Filterabschnitt 544 durch. Nachdem alle Threadgruppen die Verarbeitung aller zugewiesener Bildabschnitte 514 abgeschlossen haben, ist die Ausgabematrix 860 vollständig, und Threads, welche durch das Faltungsteilsystem 180 konfiguriert sind, kopieren die Daten, welche in der Ausgabematrix 860 enthalten sind, in den Ausgabestapel 470, welcher in dem PP-Speicher 204 enthalten ist.
Zusammengefasst ermöglichen die offenbarten Techniken einem Faltungsteilsystem, effizient mehrfache Faltungsoperationen in einem Parallelverarbeitungssystem durchzuführen. Im Allgemeinen implementiert das Faltungsteilsystem eine virtuelle Bildmatrix, welche mit einem Spaltenhauptformat übereinstimmt, das eine Matrix basierte Faltungsoperation ermöglicht. Das Faltungsteilsystem weist einen Faltungspräprozessor, welcher auf der CPU und/oder einem Streaming-Multiprozessor (SM), welche in dem Parallelverarbeitungs-Teilsystem enthalten sind, ausgeführt wird, und eine Faltungsmaschine, welche auf einem SM ausgeführt wird, welcher in dem Parallelverarbeitungs-Teilsystem enthalten ist, auf.
Im Betrieb berechnet ein Offsetfolgegenerator, welcher in dem Faltungspräprozessor enthalten ist, eine Offsetfolge vor, welche ein gewundenes Muster von Ursprungsspeicheradressen reflektiert, welches einer Spalte der virtuellen Bildmatrix zugeordnet ist. Es sei angemerkt, dass das Muster über Spalten relativ zu der ersten Ursprungsspeicheradresse, die jeder Spalte zugeordnet ist, konsistent ist. Ein Treiber führt dann Kopieroperationen aus, welche die Offsetfolge in den Parallelverarbeitungsspeicher speichern. Die Faltungsmaschine unterteilt die virtuelle Bildmatrix in getrennte Bildabschnitte und weist dann die Bearbeitung von jedem Bildabschnitt einer anderen Threadgruppe zu.
Für jede Threadgruppe berechnen die Threads, welche in der jeweiligen Threadgruppe enthalten sind, zusammenwirkend Startadressen, wobei jede Startadresse der ersten Ursprungsadresse entspricht, die einer Spalte zugeordnet ist, welche in dem zugewiesenen Bildabschnitt repräsentiert wird. Um den Bildabschnitt zu füllen, greifen daher die Threads abhängig von den Startadressen und der Offsetfolge indirekt auf die entsprechenden Bilddaten zu, welche in dem Bildstapel enthalten sind, welcher in dem Parallelverarbeitungsspeicher gespeichert ist. Die Threads führen dann Matrixmultiplikationsoperationen zwischen dem Bildabschnitt und einem entsprechenden Filterabschnitt durch, um Teilergebnisse für die Ausgabematrix zu erzeugen. Es sei angemerkt, dass der Thread, da jeder Thread zugewiesene Bildabschnitte bearbeitet, durch die Offsetfolge schreitet. Nachdem die Threadgruppen die Bearbeitung von allen Bildabschnitten abgeschlossen haben, konfiguriert die Faltungsmaschine die Threads, um die Daten, welche in der Ausgabematrix enthalten sind, in einen Ausgabestapel zu kopieren.
Zumindest ein Vorteil des offenbarten Ansatzes ist, dass das Faltungsteilsystem die Vorteile, welche inhärent in Parallelverarbeitungssystemen vorhanden sind, vollständig ausnutzt, um die hohe Genauigkeit, welche durch die CNNs bereitgestellt wird, zu erzielen, während eine Ausführungsgeschwindigkeit und der Umfang von eingesetztem Parallelverarbeitungsspeicher optimiert wird. Genauer gesagt entkoppelt die Faltungsmaschine rechenintensive Adressoperationen, welche mit dem Füllen der Abschnitte der virtuellen Bildmatrix verbunden sind, von den performancekritischen Matrixmultiplikationsoperationen, indem die CPU konfiguriert wird, die Folgesequenz vorher zu berechnen. Darüber hinaus befindet sich zu einem bestimmten Zeitpunkt nur ein Teil der virtuellen Bildmatrix in dem Gemeinschaftsspeicher und die gesamte Bildmatrix wird nicht in dem Parallelverarbeitungsspeicher gespeichert. Daher realisiert die Parallelverarbeitungspipeline die Vorteile einer optimierten Matrixmultiplikation, während der Umfang von eingesetztem Parallelverarbeitungsspeicher minimiert wird.
Die Beschreibung der verschiedenen Ausführungsformen ist nur zum Zweck der Darstellung präsentiert worden und soll nicht vollständig sein oder auf die offenbarten Ausführungen beschränkt sein. Viele Modifikationen und Variationen sind dem Fachmann bekannt, ohne dass dadurch der Umfang und der Geist der beschriebenen Ausführungsformen verlassen werden.
Aspekte der vorliegenden Ausführungsformen können als ein System, ein Verfahren oder ein Computerprogrammprodukt ausgestaltet werden. Dementsprechend nehmen Aspekte der vorliegenden Offenbarung die Form einer Ausführungsform vollständig in Hardware, einer Ausführungsform vollständig in Software (einschließlich Firmware, residenter Software und Mikroprogrammcode, usw.) oder eine Ausführungsform, welche Software- und Hardware-Aspekte kombiniert, an, was im Allgemeinen hier als „Schaltung”, „Modul”, oder „System” bezeichnet wird. Darüber hinaus können Aspekte der vorliegenden Offenbarung die Form eines Computerprogrammprodukts annehmen, welches in einem oder in mehreren von einem Computer lesbaren Medien, ausgestaltet ist, in welchen ein von einem Computer lesbarer Programmcode enthalten ist.
Jede Kombination von einem oder von mehreren von einem Computer lesbaren Medien kann eingesetzt werden. Das von einem Computer lesbare Medium kann ein von einem Computer lesbares Signal-Medium oder ein von einem Computer lesbares Speichermedium sein. Ein von einem Computer lesbares Speichermedium kann zum Beispiel sein ein(e) elektronische(s), ein(e) magnetische(s), ein(e) optische(s), ein(e) elektromagnetische(s), ein(e) Infrarot- oder ein(e) Halbleiter-System, Vorrichtung oder Gerät oder irgendeine geeignete Kombination von den vorab Stehenden sein, ist aber nicht darauf eingeschränkt. Speziellere Beispiele (eine nicht vollständige Liste) der von einem Computer lesbaren Speichermedien würde Folgendes umfassen: eine elektrische Verbindung mit einem oder mit mehreren Leitungen, eine tragbare Computerdiskette, eine Festplatte, einen Direktzugriffspeicher (RAM), einen Festwertspeicher (ROM), einen löschbaren programmierbaren Nur-Lese-Speicher (EPROM oder Flash-Speicher), eine optische Faser, eine tragbare CD-ROM, eine optische Speichervorrichtung, eine magnetische Speichervorrichtung oder jede geeignete Kombination des vorab Stehenden. Im Kontext dieses Dokuments kann ein von einem Computer lesbares Speichermedium jedes zugreifbare Medium sein, welches ein Programm zum Einsatz mit oder in Verbindung mit einem/r Instruktions-Ausführungs-System, Vorrichtung oder Gerät enthält oder speichert.
Aspekte der vorliegenden Offenbarung sind vorab mit Bezug zu Flussplan-Darstellungen und/oder Blockdiagrammen von Verfahren, Vorrichtungen (Systemen) und Computerprogrammprodukten gemäß Ausführungsformen der Offenbarung beschrieben worden. Es sei darauf hingewiesen, dass jeder Block der Flussplan-Darstellungen und/oder Blockdiagramme und Kombinationen von Blöcken in den Flussplan-Darstellungen und/oder Blockdiagrammen durch Computerprogramm-Anweisungen implementiert werden kann bzw. können. Diese Computerprogramm-Anweisungen können von einem Prozessor eines Computers für allgemeine Zwecke, eines Computers für spezielle Zwecke oder einer anderen programmierbaren Datenverarbeitungsvorrichtung bereitgestellt werden, um eine Maschine zu erzeugen, so dass die Anweisungen, welche mittels des Prozessors des Computers oder der anderen programmierbaren Datenverarbeitungsvorrichtung ausgeführt werden, die Implementierung der Funktionen/Vorgänge ermöglichen, welche in dem Flussplan und/oder dem Blockdiagrammblock oder Blockdiagrammblöcken spezifiziert sind. Solche Prozessoren können ohne Einschränkung Prozessoren für einen allgemeinen Zweck, Prozessoren für spezielle Zwecke, Anwendungs-spezifische Prozessoren oder feldprogrammierbare Prozessoren oder Gate Arrays sein.
Der Flussplan und die Blockdiagramme in den Figuren stellen die Architektur, Funktionalität und den Betrieb von möglichen Implementierungen von Systemen, Verfahren und Computerprogrammprodukten gemäß verschiedener Ausführungsformen der vorliegenden Offenbarung dar. Diesbezüglich kann jeder Block in dem Flussplan oder Blockdiagramm ein Modul, ein Segment oder einen Codeabschnitt repräsentieren, welcher eine oder mehrere ausführbare Anweisungen umfasst, um die spezifizierten logischen Funktion(en) zu implementieren. Es sei auch angemerkt, dass bei einigen alternativen Ausführungsformen die Funktionen, welche in dem Block beschrieben sind, in einer anderen Reihenfolge, als in den Figuren dargestellt, auftreten können. Zum Beispiel können zwei Blöcke, welche nacheinander dargestellt sind, tatsächlich im Wesentlichen gleichzeitig ausgeführt werden, oder die Blöcke können bisweilen in der umgekehrten Reihenfolge, abhängig von der entsprechenden Funktionalität, ausgeführt werden. Es sei auch angemerkt, dass jeder Block der Blockdiagramme und/oder Flussplan-Darstellung und Kombinationen der Blöcke in den Blockdiagrammen und/oder Flussplan-Darstellung durch Hardware-basierte Systeme für einen speziellen Zweck implementiert werden können, welche die spezifizierten Funktionen oder Vorgänge ausführen, oder durch Kombinationen von Hardware für einen speziellen Zweck und Computer-Anweisungen.
Während das vorab Stehende auf Ausführungsformen der vorliegenden Offenbarung gerichtet ist, können andere und weitere Ausführungsformen der Offenbarung formuliert werden, ohne den wesentlichen Umfang davon zu verlassen, und der Umfang davon wird durch die folgenden Ansprüche bestimmt.

Claims

Computer-implementiertes Verfahren zur Durchführung einer mehrfachen Faltungsoperation, wobei das Verfahren umfasst: Auswählen einer ersten Startadresse abhängig von einer ersten Zieladresse, welche in einem ersten Bildabschnitt enthalten ist, welcher in einem ersten Speicher gespeichert ist; Identifizieren eines ersten Offsets abhängig von der ersten Zieladresse; Berechnen einer ersten Ursprungsadresse, welche in einem Bildstapel enthalten ist, welcher in einem zweiten Speicher gespeichert ist, abhängig von der ersten Startadresse und dem ersten Offset; Kopieren von Daten von der ersten Ursprungsadresse zu der ersten Zieladresse; und nach dem Kopieren der Daten Durchführen einer oder mehrerer Matrixmultiplikationsoperationen zwischen dem ersten Bildabschnitt und einem ersten Filterabschnitt.
Computer-implementiertes Verfahren nach Anspruch 1, wobei der erste Speicher einen Gemeinschaftsspeicher umfasst, und wobei der zweite Speicher einen Parallelverarbeitungsspeicher umfasst.
Computer-implementiertes Verfahren nach Anspruch 1 oder 2, wobei der erste Filterabschnitt in dem ersten Speicher gespeichert ist, und darüber hinaus umfassend: Berechnen einer Filter-Ursprungsadresse abhängig von der ersten Zieladresse; und Kopieren von Daten, welche in einem Filterstapel bei der Filter-Ursprungsadresse gespeichert sind, zu einer Filter-Zieladresse, welche in dem ersten Filterabschnitt enthalten ist.
Computer-implementiertes Verfahren nach einem der vorhergehenden Ansprüche, wobei das Auswählen der ersten Startadresse umfasst: Zuordnen der ersten Zieladresse zu einer Spalte einer virtuellen Bildmatrix; und Durchführen einer oder mehrerer Operationen, welche die Spalte einer Adresse, welche in dem Bildstapel enthalten ist, zuordnen.
Computer-implementiertes Verfahren nach einem der vorhergehenden Ansprüche, wobei das Identifizieren des ersten Offsets umfasst: Zuordnen der ersten Zieladresse zu einer Zeile einer virtuellen Bildmatrix; und Abfragen eines Werts, welcher in einer Offsetfolge enthalten ist, abhängig von der Zeile.
Computer-implementiertes Verfahren nach Anspruch 5, darüber hinaus umfassend Erzeugen der Offsetfolge abhängig von einer deterministischen Beziehung zwischen dem Bildstapel und der virtuellen Bildmatrix.
Computer-implementiertes Verfahren nach einem der vorhergehenden Ansprüche, darüber hinaus umfassend ein Zuweisen des ersten Bildabschnitts zu einer ersten Threadgruppe, und Konfigurieren mindestens eines Threads in der Threadgruppe, um die erste Ursprungsadresse zu berechnen.
Computer-implementiertes Verfahren nach Anspruch 7, darüber hinaus umfassend ein Zuweisen eines zweiten Bildabschnitts zu einer zweiten Threadgruppe, und Konfigurieren mindestens eines Threads in der zweiten Threadgruppe, um eine zweite Ursprungsadresse, welche in dem Bildstapel enthalten ist, abhängig von einer zweiten Startadresse und dem ersten Offset zu berechnen.
Computer-implementiertes Verfahren nach Anspruch 8, wobei die erste Ursprungsadresse und die zweite Ursprungsadresse im Wesentlichen parallel berechnet werden.
Nicht flüchtiges, von einem Computer lesbares Speichermedium, welches Anweisungen aufweist, die, wenn sie durch einen Prozessor ausgeführt werden, bewirken, dass der Prozessor eine mehrfache Faltungsoperation durchführt, indem folgende Schritte durchgeführt werden: Auswählen einer ersten Startadresse abhängig von einer ersten Zieladresse, welche in einem ersten Bildabschnitt enthalten ist, welcher in einem ersten Speicher gespeichert ist; Identifizieren eines ersten Offsets abhängig von der ersten Zieladresse; Berechnen einer ersten Ursprungsadresse, welche in einem Bildstapel enthalten ist, welcher in einem zweiten Speicher gespeichert ist, abhängig von der ersten Startadresse und dem ersten Offset; Kopieren von Daten von der ersten Ursprungsadresse zu der ersten Zieladresse; und nach dem Kopieren der Daten Durchführen einer oder mehrerer Matrixmultiplikationsoperationen zwischen dem ersten Bildabschnitt und einem ersten Filterabschnitt.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach Anspruch 10, wobei der erste Speicher einen Gemeinschaftsspeicher umfasst, und wobei der zweite Speicher einen Parallelverarbeitungsspeicher umfasst.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach Anspruch 10 oder 11, wobei der erste Filterabschnitt in dem ersten Speicher gespeichert ist, und darüber hinaus umfassend: Berechnen einer Filter-Ursprungsadresse abhängig von der ersten Zieladresse; und Kopieren von Daten, welche in einem Filterstapel unter der Filter-Ursprungsadresse gespeichert sind, zu einer Filter-Zieladresse, welche in dem ersten Filterabschnitt enthalten ist.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach einem der Ansprüche 10 bis 12, wobei das Auswählen der ersten Startadresse umfasst: Zuordnen der ersten Zieladresse zu einer Spalte einer virtuellen Bildmatrix; und Durchführen einer oder mehrerer Operationen, welche die Spalte einer Adresse, welche in dem Bildstapel enthalten ist, zuordnen.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach einem der Ansprüche 11 bis 13, wobei das Identifizieren des ersten Offsets umfasst: Zuordnen der ersten Zieladresse zu einer Zeile einer virtuellen Bildmatrix; und Abfragen eines Werts, welcher in einer Offsetfolge enthalten ist, abhängig von der Zeile.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach Anspruch 14, darüber hinaus umfassend Erzeugen der Offseffolge abhängig von einer deterministischen Beziehung zwischen dem Bildstapel und der virtuellen Bildmatrix.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach einem der Ansprüche 10 bis 15, darüber hinaus umfassend ein Konfigurieren mindestens eines Threads in einer zweiten Threadgruppe, um eine zweite Ursprungsadresse, welche in dem Bildstapel enthalten ist, abhängig von der ersten Startadresse und einem zweiten Offset zu berechnen.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach einem der Ansprüche 10 bis 16, darüber hinaus umfassend Durchführen von einer oder von mehreren Ausgabeformatierungsoperationen abhängig von der Ausgabematrix, um einen Ausgabestapel zu erzeugen.
Nicht flüchtiges von einem Computer lesbares Speichermedium nach Anspruch 17, wobei eine erste Schicht, welche in einem faltenden neuronalen Netz enthalten ist, mindestens den Bildstapel aufweist, und wobei ein zweite Schicht, welche in dem faltenden neuronalen Netz enthalten ist, mindestens den Ausgabestapel aufweist.
System, welches ausgestaltet ist, um eine mehrfache Faltungsoperation auszuführen, wobei das System umfasst: einen ersten Speicher; einen zweiten Speicher; und eine Faltungsmaschine, welche sowohl mit dem ersten Speicher als auch mit dem zweiten Speiche gekoppelt ist, und welche ausgestaltet ist, um: einen ersten Offset zu identifizieren, welcher in einer Offsetfolge enthalten ist, abhängig von einer ersten Zieladresse, welche in einem ersten Bildabschnitt enthalten ist, welcher in dem ersten Speicher gespeichert ist; eine erste Ursprungsadresse zu berechnen, welche in einem Bildstapel enthalten ist, welcher in dem zweiten Speicher gespeichert ist, abhängig von dem ersten Offset; um Daten von der ersten Ursprungsadresse zu der ersten Zieladresse zu kopieren; und um nach dem Kopieren der Daten eine oder mehrere Matrixmultiplikationsoperationen zwischen dem ersten Bildabschnitt und einem ersten Filterabschnitt auszuführen.
System nach Anspruch 19, wobei der erste Speicher einen Gemeinschaftsspeicher umfasst, und wobei der zweite Speicher einen Parallelverarbeitungsspeicher umfasst.