DE202017103725U1

DE202017103725U1 - Blockoperationen für einen Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einem zweidimensionalen Schieberegister

Info

Publication number: DE202017103725U1
Application number: DE202017103725.8U
Authority: DE
Original assignee: Google LLC
Current assignee: Google LLC
Priority date: 2016-07-01
Filing date: 2017-06-22
Publication date: 2017-10-04
Anticipated expiration: 2027-06-23
Also published as: TW201921315A; JP2019521441A; US20200154072A1; TWI767190B; GB2560208A; KR102190318B1; US10334194B2; US11196953B2; TW202025081A; TW201826219A; JP6821715B2; GB2560208B; GB2577139B; GB201709786D0; EP3479340A1; TWI656508B; US9986187B2; US20180234653A1; CN107563953B; US20180007303A1

Abstract

Maschinenlesbares Speichermedium, das Programmcode enthält, der bei Verarbeitung durch einen Bildprozessor, der eine zweidimensionale Ausführungsbahnmatrix und eine zweidimensionale Schieberegistermatrix umfasst, den Bildprozessor veranlasst, ein Verfahren auszuführen mit wiederholtem Schieben von einem erstem Inhalt von mehreren Zeilen oder Spalten der zweidimensionalen Schieberegistermatrix und wiederholtes Ausführen von mindestens einem Befehl zwischen Schiebevorgängen, der an dem geschobenen ersten Inhalt und/oder zweiten Inhalt operiert, der an entsprechenden Orten der zweidimensionalen Schieberegistermatrix speicherresident ist, in die der geschobene erste Inhalt geschoben wurde.

Description

Gebiet der Erfindung
Das Gebiet der Erfindung betrifft generell Bildverarbeitung und insbesondere Blockoperationen für einen Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einem zweidimensionalen Schieberegister.
Allgemeiner Stand der Technik
Die Bildverarbeitung beinhaltet in der Regel die Verarbeitung von Bildpunktwerten, die in einer Matrix angeordnet sind. Hierbei erfasst eine räumlich gegliederte zweidimensionale Matrix die zweidimensionale Beschaffenheit der Bilder (zusätzliche Dimensionen können unter anderem Zeit (z. B. eine Sequenz von zweidimensionalen Bildern) und Datentyp (z. B. Farben) einschließen). In einem typischen Szenario werden die geordneten Bildpunktwerte von einer Kamera bereitgestellt, die ein Standbild oder eine Folge von Frames generiert hat, um Bewegungsbilder zu erfassen. Herkömmliche Bildprozessoren fallen in der Regel unter eines von zwei Extremen.
Ein erstes Extrem führt Bildverarbeitungstasks als Softwareprogramme aus, die auf einem Universalprozessor oder einem universell verwendbaren Prozessor (z. B. einem Universalprozessor mit Vektorbefehlserweiterungen) ausgeführt werden. Obwohl das erste Extrem in der Regel eine vielseitig einsetzbare Anwendungssoftware-Entwicklungsplattform bereitstellt, resultiert dessen Verwendung feinerer Datenstrukturen kombiniert mit den zugehörigen Verwaltungsdaten (z. B. Befehlsabruf und -decodierung, Handhabung von chipinternen und chipexternen Daten, spekulative Ausführung) letztendlich in einem Verbrauch größerer Energiemengen pro Dateneinheit während der Ausführung des Programmcodes.
Ein zweites, entgegengesetztes Extrem wendet stationäre fest verdrahtete Schaltungen auf viel größere Datenblöcke an. Die Verwendung von größeren (im Gegensatz zu feineren) Datenblöcken, die direkt auf benutzerdefinierte Schaltungen angewendet werden, verringert den Energieverbrauch pro Dateneinheit erheblich. Jedoch führt die Verwendung von benutzerdefinierten Schaltungen mit fester Funktion im Allgemeinen zu einer begrenzten Menge von Arbeitsschritten, die der Prozessor ausführen kann. Dementsprechend fehlt im zweiten Extrem die vielseitige Programmierumgebung (die mit dem ersten Extrem verbunden ist).
Eine Technologieplattform, die sowohl vielseitige Anwendungssoftware-Entwicklungsmöglichkeiten als auch eine verbesserte Energieeffizienz pro Dateneinheit bietet, bleibt eine wünschenswerte und dennoch fehlende Lösung.
Kurzdarstellung
Es wird ein Verfahren beschrieben, welches, auf einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einer zweidimensionalen Schieberegistermatrix, wiederholtes Schieben von erstem Inhalt von mehreren Reihen oder Spalten der zweidimensionalen Schieberegistermatrix und wiederholtes Ausführen von mindestens einem Befehl zwischen Schiebevorgängen umfasst, das an dem geschobenen ersten Inhalt und/oder zweiten Inhalt operiert, der an entsprechenden Orten der zweidimensionalen Schieberegistermatrix speicherresident ist, in die der geschobene erste Inhalt geschoben wurde.
Es wird eine Vorrichtung beschrieben, die Mittel umfasst, um auf einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einer zweidimensionalen Schieberegistermatrix wiederholt ersten Inhalt von mehreren Reihen oder Spalten der zweidimensionalen Schieberegistermatrix zu schieben und wiederholt mindestens einen Befehl zwischen Schiebevorgängen auszuführen, der an dem geschobenen ersten Inhalt und/oder zweiten Inhalt operiert, der an entsprechenden Orten der zweidimensionalen Schieberegistermatrix speicherresident ist, in die der geschobene erste Inhalt geschoben wurde.
Figurenverzeichnis
Die folgende Beschreibung und die begleitenden Zeichnungen dienen dazu, verschiedene Ausführungsformen der Erfindung zu veranschaulichen. In den Zeichnungen:
1 zeigt verschiedene Komponenten einer Technologieplattform;
2a zeigt eine Ausführungsform einer Anwendungssoftware, die mit Kerneln aufgebaut ist;
2b zeigt eine Ausführungsform der Struktur eines Kernels;
3 zeigt eine Ausführungsform des Betriebs eines Kernels;
Die 4a, 4b und 4c stellen verschiedene Aspekte eines Speichermodells eines virtuellen Prozessors zum Entwickeln von Kernel-Threads in einer höheren Anwendungssoftwareentwicklungsumgebung dar;
5a zeigt eine Ausführungsform eines Threads, der mit Ladebefehlen mit einem positionsrelativen Format geschrieben ist;
5b zeigt Bilder mit unterschiedlichen Bildpunktdichten;
6 zeigt eine Ausführungsform einer Anwendungssoftwareentwicklungs- und -simulationsumgebung;
7 zeigt eine Ausführungsform einer Bildprozessor-Hardwarearchitektur;
Die 8a, 8b, 8c, 8d und 8e zeigen das Parsen von Bilddaten in eine Zeilengruppe, das Parsen einer Zeilengruppe in ein Blatt und die an einem Blatt mit überlappenden Schablonen durchgeführte Operation;
9a zeigt eine Ausführungsform eines Schablonenprozessors;
9b zeigt eine Ausführungsform eines Befehlsworts des Schablonenprozessors;
10 zeigt eine Ausführungsform einer Datenberechnungseinheit innerhalb eines Schablonenprozessors;
Die 11a, 11b, 11c, 11d, 11e, 11f, 11g, 11h, 11i, 11j und 11k zeigen ein Beispiel für die Verwendung einer zweidimensionalen Schiebematrix und einer Ausführungsbahnmatrix, um ein Paar angrenzender Ausgabebildpunktwerte mit überlappenden Schablonen zu bestimmen;
12 zeigt eine Ausführungsform einer Einheitszelle für eine integrierte Ausführungsbahnmatrix und eine zweidimensionale Schiebematrix;
13 zeigt eine zweidimensionale Reihen-/Spaltensummieroperation;
Die 14a, 14b, 14c und 14d zeigen niedrigere Operationen zum Implementieren einer zweidimensionalen Zeilensummieroperation;
15 betrifft eine zweidimensionale Präfixsummieroperation;
Die 16a, 16b, 16c und 16d zeigen niedrigere Operationen zum Implementieren einer zweidimensionalen Präfixsummieroperation;
17 betrifft eine zweidimensionale Minimumfinden-Operation;
Die 18a, 18b, 18c und 18d zeigen niedrigere Operationen zum Implementieren einer zweidimensionalen Minimumfinden-Operation;
Die 19a und 19b zeigen eine Matrixmultiplizieroperation;
Die 20a, 20b, 20c, 20d und 20e zeigen niedrigere Operationen zum Implementieren einer Matrixmultiplizieroperation mit einem zweidimensionalen Schieberegister;
21 zeigt eine DFT-Operation;
Die 22a, 22b, 22c, 22d, 22e und 22f zeigen niedrigere Operationen zum Implementieren einer DFT-Operation mit einem zweidimensionalen Schieberegister;
23 zeigt eine Butterfly-Operation;
Die 24a, 24b und 24c zeigen mit einem zweidimensionalen Schieberegister implementierte Butterfly-Operationen;
25 zeigt ein Basisbild und ein alternatives Bild mit einem Blockbild;
Die 26a, 26b, 26c und 26d zeigen niedrigere Operationen zum Ausführen eines Blockabgleichalgorithmus;
27 zeigt eine Umgebung zum Erstellen von Programmcode, der an eine Hardwareplattform mit einer zweidimensionalen Ausführungsbahnmatrix und einer zweidimensionalen Schieberegistermatrix adressiert ist;
28 zeigt eine Ausführungsform eines Computersystems.
Ausführliche Beschreibung
i. Einführung
Die nachfolgende Beschreibung beschreibt zahlreiche Ausführungsformen, die eine neue Bildverarbeitungstechnologieplattform betreffen, welche eine vielseitig einsetzbare Anwendungssoftware-Entwicklungsumgebung bereitstellt, die größere Datenblöcke (z. B. Zeilengruppen und Blätter, wie weiter unten beschrieben) verwendet, um für eine verbesserte Energieeffizienz zu sorgen.
1.0 Anwendungssoftware-Entwicklungsumgebung
a. Anwendung und Struktur von Kerneln
1 zeigt eine umfassende Ansicht einer Bildprozessortechnologieplattform, die eine virtuelle Bildverarbeitungsumgebung 101, die eigentliche Bildverarbeitungshardware 103 und einen Compiler 102 beinhaltet, um einen für die virtuelle Verarbeitungsumgebung 101 geschriebenen höheren Code in Objektcode zu übersetzen, den die eigentliche Hardware 103 physisch ausgeführt. Wie nachfolgend näher beschrieben wird, ist die virtuelle Verarbeitungsumgebung 101 in Bezug auf die Anwendungen, die entwickelt werden können, vielseitig und auf eine einfache Visualisierung der einzelnen Prozesse einer Anwendung zugeschnitten. Nach Beendigung des Programmcode-Entwicklungsaufwands durch den Entwickler 104 übersetzt der Compiler 102 den Code, der in die virtuelle Verarbeitungsumgebung 101 geschrieben wurde, in einen Objektcode, der auf die eigentliche Hardware 103 zugeschnitten ist.
2a zeigt eine beispielhafte Ausführungsform der Struktur und Form, welche eine Anwendungssoftware, die in die virtuelle Umgebung geschrieben wird, annehmen kann. Wie in 2a dargestellt, kann erwartet werden, dass der Programmcode ein oder mehrere Frames der Eingabebilddaten 201 verarbeitet, um eine Gesamttransformation an den Eingabebilddaten 201 zu bewirken. Die Transformation wird mit dem Betrieb eines oder mehrerer Kernels des Programmcodes 202 realisiert, die an den Eingabebilddaten in einer abgestimmten Sequenz arbeiten, welche von dem Entwickler artikuliert wird.
Wie in 2a erkennbar, wird die Gesamttransformation durch eine erste Verarbeitung von jedem Eingabebild mit einem ersten Kernel K1 bewirkt. An den vom Kernel K1 erzeugten Ausgabebildern arbeitet dann der Kernel K2. An jedem von den vom Kernel K2 erzeugten Ausgabebildern arbeitet dann der Kernel K3_1 oder K3_2. An den von dem bzw. den Kerneln K3_1/K3_2 erzeugten Ausgabebildern arbeitet dann der Kernel K4. Die Kernel K3_1 und K3_2 können identische Kernel sein, die derart ausgelegt sind, dass sie die Gesamtverarbeitung beschleunigen, indem sie eine Parallelverarbeitung in der K3-Stufe auferlegen, oder sie können unterschiedliche Kernel sein (z. B. arbeitet der Kernel K3_1 an Eingabebildern eines ersten spezifischen Typs und der Kernel K3_2 arbeitet an Eingabebildern eines zweiten unterschiedlichen Typs).
Daher kann die größere Gesamtbildverarbeitungssequenz die Form einer Bildverarbeitungspipeline oder eines gerichteten azyklischen Graphen (DAG) annehmen und die Entwicklungsumgebung kann dafür ausgestattet sein, dem Entwickler tatsächlich eine Darstellung des Programmcodes als solches, der entwickelt wird, anzuzeigen. Kernels können von einem Entwickler einzeln entwickelt und/oder von einer Entität bereitgestellt werden, die jegliche darunterliegende Technologie liefert (wie z. B. die eigentliche Signalprozessorhardware und/oder ein Design derselben), und/oder durch einen Dritten (z. B. einen Anbieter von für die Entwicklungsumgebung geschriebene Kernelsoftware). Als solches wird davon ausgegangen, dass eine nominale Entwicklungsumgebung eine „Bibliothek“ von Kerneln einschließt, die von Entwicklern in verschiedenen Weisen „verbunden“ werden können, um den Gesamtablauf ihres größeren Entwicklungsaufwands herbeizuführen. Einige grundlegende Kernel, von denen erwartet wird, dass sie Teil einer solchen Bibliothek sind, können Kernel beinhalten, um eine oder mehrere der folgenden grundlegenden Bildverarbeitungstasks bereitzustellen: Faltungen, Entrauschen, Farbraumumwandlungen, Rand- und Eckenerfassung, Schärfen, Weißabgleich, Gammakorrektur, Tonemapping, Matrixmultiplikation, Bildregistrierung, Pyramidenbau, Wavelet-Transformation, blockweise diskrete Kosinus- und Fourier-Transformationen.
2b zeigt eine beispielhafte Darstellung der Struktur eines Kernels 203, der von einem Entwickler in Betracht gezogen werden kann. Wie in 2b dargestellt, kann der Kernel 203 als eine Anzahl von parallelen Threads von Programmcode („Threads“) 204 betrachtet werden, die jeweils auf einem entsprechenden zugrunde liegenden Prozessor 205 arbeiten, wobei jeder Prozessor 205 auf einen bestimmten Ort in einer Ausgabematrix 206 (z. B. einen bestimmten Bildpunktort in dem von dem Kernel generierten Ausgabebild) ausgerichtet ist. Der Einfachheit halber sind in 2b nur drei Prozessoren und entsprechende Threads dargestellt. Bei verschiedenen Ausführungsformen würde jeder dargestellte Ausgabematrixort einen eigenen fest zugeordneten Prozessor und einen entsprechenden Thread aufweisen. Das heißt, dass für jeden Bildpunkt in der Ausgabematrix ein separater Prozessor und Thread zugewiesen werden kann. Bei alternativen Herangehensweisen kann ein gleicher Thread Daten für mehr als Ausgabebildpunkte generieren und/oder zwei unterschiedliche Threads (z. B. in bestimmten begrenzten Fällen) können bei der Erzeugung der Daten für einen gleichen Ausgabebildpunkt zusammenwirken.
Wie nachstehend näher beschrieben, arbeiten eine Matrix von Ausführungsbahnen und entsprechenden Threads in der tatsächlichen zugrunde liegenden Hardware in verschiedenen Ausführungsformen im Gleichklang (z. B. in einer Single Instruction Multiple Data-(SIMD)-Weise), um für einen Teil einer „Zeilengruppe“ des gerade verarbeiteten Frames Ausgabebilddaten zu generieren. Eine Zeilengruppe ist ein zusammenhängender, beträchtlicher Abschnitt eines Bildframes. Bei verschiedenen Ausführungsformen kann sich der Entwickler im Klaren darüber sein, dass die Hardware auf Zeilengruppen arbeitet, oder die Entwicklungsumgebung kann eine Abstraktion darstellen, in der es einen separaten Prozessor und einen Thread für beispielsweise jeden Bildpunkt in dem Ausgabeframe gibt (z. B. jeder Bildpunkt in einem Ausgabeframe generiert durch einen eigenen fest zugeordneten Prozessor und Thread). Ungeachtet dessen versteht der Entwickler bei verschiedenen Ausführungsformen, dass der Kernel einen individuellen Thread für jeden Ausgabebildpunkt umfasst (egal, ob die Ausgabematrix als ein gesamter Ausgabeframe oder ein Abschnitt davon visualisiert wird).
Wie nachfolgend näher beschrieben wird, weisen die Prozessoren 205 bei einer Ausführungsform, die dem Entwickler in der virtuellen Umgebung präsentiert werden, eine Befehlssatzarchitektur (ISA) auf, die nicht nur Standard-(z. B. RISC)-Maschinenbefehle unterstützt, sondern auch speziell formatierte Datenzugriffsbefehle beinhaltet, die es dem Entwickler ermöglichen, die durchgeführte Bildpunkt-für-Bildpunkt-Verarbeitung leicht zu visualisieren. Die Fähigkeit, einen beliebigen Eingabematrixort in Kombination mit einer vollständigen ISA herkömmlicher mathematischer und Programmsteuerungsmaschinenbefehle leicht zu definieren/zu visualisieren, ermöglicht eine extrem vielseitige Programmierumgebung, die es einem Anwendungsprogrammentwickler im Wesentlichen ermöglicht, im Idealfall jede gewünschte Funktion, die ausgeführt werden soll, auf einer beliebig großen Bildfläche zu definieren. Zum Beispiel kann im Idealfall jede mathematische Operation leicht programmiert und auf jede Schablonengröße angewendet werden.
Hinsichtlich der Datenzugriffsbefehle beinhaltet die ISA der virtuellen Prozessoren („virtuelle ISA“) bei einer Ausführungsform einen speziellen Datenladebefehl und einen speziellen Datenspeicherbefehl. Der Datenladebefehl ist in der Lage, von jedem Ort innerhalb einer Eingabematrix von Bilddaten zu lesen. Der Datenspeicherbefehl ist in der Lage an einen beliebigen Ort innerhalb der Ausgabematrix von Bilddaten zu schreiben. Der letztgenannte Befehl ermöglicht es, mehrere Instanzen des gleichen Prozessors problemlos verschiedenen Ausgabebildpunktorten zuzuordnen (jeder Prozessor schreibt in einen anderen Bildpunkt in der Ausgabematrix). Dementsprechend kann beispielsweise die Schablonengröße selbst (z. B. als Bildpunktbreite und Bildpunkthöhe ausgedrückt) zu einem leicht programmierbaren Merkmal gemacht werden. Die Visualisierung der Verarbeitungsvorgänge wird weiter vereinfacht, wobei jede der speziellen Lade- und Speicherbefehle ein spezielles Befehlsformat aufweist, wodurch Zielmatrixorte vereinfacht als X- und Y-Koordinaten angegeben werden.
Unabhängig davon können durch Instanziieren eines getrennten Prozessors für jede von mehreren Orten in der Ausgabematrix die Prozessoren ihre jeweiligen Threads parallel ausführen, sodass z. B. die jeweiligen Werte für alle Orte in der Ausgabematrix gleichzeitig erzeugt werden. Es ist zu beachten, dass viele Bildverarbeitungsroutinen in der Regel die gleichen Vorgänge an verschiedenen Bildpunkten des gleichen Ausgabebildes ausführen. Dementsprechend wird bei einer Ausführungsform der Entwicklungsumgebung davon ausgegangen, dass jeder Prozessor identisch ist und den gleichen Thread-Programmcode ausführt. Daher kann die virtualisierte Umgebung als eine Art zweidimensionaler (2D) SIMD-Prozessor betrachtet werden, der aus einer 2D-Matrix von z. B. identischen Prozessoren besteht, die jeweils identischen Code im Sperrschritt ausführen.
3 zeigt eine ausführlichere beispielhafte Ausführungsform der Verarbeitungsumgebung für zwei virtuelle Prozessoren, die identischen Code für zwei unterschiedliche Bildpunktpositionen in einer Ausgabematrix verarbeiten. 3 zeigt eine Ausgabematrix 304, die einem Ausgabebild entspricht, das generiert wird. Hier verarbeitet ein erster virtueller Prozessor den Code des Threads 301, um einen Ausgabewert am Ort X1 der Ausgabematrix 304 zu generieren, und ein zweiter virtueller Prozessor verarbeitet den Code des Threads 302, um einen Ausgabewert am Ort X2 der Ausgabematrix 304 zu generieren. Bei verschiedenen Ausführungsformen würde der Entwickler wiederum verstehen, dass es einen separaten Prozessor und Thread für jede Bildpunktposition in der Ausgabematrix 304 gibt (der Einfachheit halber zeigt 3 lediglich zwei davon). Jedoch muss der Entwickler (aufgrund der SIMD-ähnlichen Beschaffenheit der Maschine) bei verschiedenen Ausführungsformen nur Code für einen Prozessor und Thread entwickeln.
Wie auf dem Fachgebiet bekannt, wird ein Ausgabebildpunktwert oft durch die Verarbeitung der Bildpunkte einer Eingabematrix bestimmt, welche die entsprechende Ausgabebildpunktposition beinhaltet und umgibt. Wie aus 3 ersichtlich, entspricht der Ort X1 der Ausgabematrix 304 beispielsweise dem Ort E der Eingabematrix 303. Die Schablone der Bildpunktwerte der Eingabematrix 303, die verarbeitet werden würden, um den Ausgabewert X1 zu bestimmen, würde daher den Eingabewerten ABCDEFGHI entsprechen. Ähnlich würde die Schablone der Eingabematrixbildpunkte, die verarbeitet werden würden, um den Ausgabewert X2 zu bestimmen, den Eingabewerten DEFGHIJKL entsprechen.
3 zeigt eine beispielhafte Ausführungsform eines entsprechenden Programmcodes einer virtuellen Umgebung für ein Thread-Paar 301, 302, das verwendet werden könnte, um die Ausgabewerte X1 bzw. X2 zu berechnen. Im Beispiel von 3 sind beide Code-Paare identisch und entsprechen durchschnittlich einer Schablone von neun Eingabematrixwerten, um einen entsprechenden Ausgabewert zu bestimmen. Der einzige Unterschied zwischen den beiden Threads sind die Variablen, die aus der Eingabematrix aufgerufen werden, und der Ort der Ausgabematrix, in die geschrieben wird. Speziell arbeitet der Thread, der in die Ausgabeposition X1 schreibt, an der Schablone ABCDEFGHI und der Thread, der in die Ausgabeposition X2 schreibt, an der Schablone DEFGHIJKL.
Wie aus dem jeweiligen Programmcode von dem Thread-Paar 301, 302 ersichtlich, beinhaltet jeder virtuelle Prozessor mindestens die internen Register R1 und R2 und unterstützt mindestens die folgenden Befehle: 1) einen LADE-Befehl aus der Eingabematrix in R1; 2) einen LADE-Befehl aus der Eingabematrix in R2; 3) einen ADDIER-Befehl, der den Inhalt von R1 und R2 addiert und die Resultante in R2 einfügt; 4) einen DIVIDIER-Befehl, der den Wert innerhalb R2 durch den unmittelbaren Operanden 9 teilt; und, 5) einen SPEICHER-Befehl, der den Inhalt von R2 in die Ausgabematrixposition speichert, zu der der Thread zugeordnet ist. Obwohl in 3 nur zwei Ausgabematrixpositionen und nur zwei Threads und entsprechende Prozessoren dargestellt sind, könnte jedem Ort in der Ausgabematrix durchaus ein virtueller Prozessor und ein entsprechender Thread zugewiesen werden, der diese Funktionen ausführt. Bei verschiedenen Ausführungsformen werden die mehreren Threads im Einklang mit der SIMD-artigen Beschaffenheit der Verarbeitungsumgebung isoliert voneinander ausgeführt. Das heißt, es gibt keine Thread-Thread-Kommunikation zwischen virtuellen Prozessoren (ein SIMD-Kanal verhindert den Übergang in einen anderen SIMD-Kanal).
b. Virtueller Prozessor-Speichermodell
Bei verschiedenen Ausführungsformen ist ein zugehöriges Merkmal der virtuellen Prozessoren ihr Speichermodell. Wie auf dem Fachgebiet bekannt, liest ein Prozessor Daten aus einem Speicher, arbeitet an diesen Daten und schreibt neue Daten in den Speicher zurück. Ein Speichermodell ist die Perspektive oder Sicht, die ein Prozessor von der Art und Weise hat, in der Daten im Speicher organisiert sind. Die 4a bis 4c betreffen eine Ausführungsform des Speichermodells für die virtuellen Prozessoren der Entwicklungsumgebung. Es wird eine simplifizierende Umgebung, die nur drei virtuelle Prozessoren und entsprechende Threads 401 einbezieht, für Beispielzwecke verwendet. Wie nachfolgend ausführlicher beschrieben wird, achtet das Speichermodell der virtuellen Prozessoren darauf, eine SIMD-Semantik zu erhalten, während es gleichzeitig Skalaroperationen und eigenen Zwischenwertspeicherplatz für jeden virtuellen Prozessor bereitstellt.
Wie in 4a ersichtlich, ist bei einer Ausführungsform die Speicherregion 420, aus der heraus jeder virtuelle Prozessor operiert, basierend auf der gespeicherten Informationsart in sechs unterschiedliche Partitionen organisiert. Speziell existiert dort: 1) eine eigene Notizblockspeicherregion 402; 2) eine globale Eingabedatenmatrixregion 403; 3) eine globale Ausgabedatenmatrixregion 404; 4) eine globale Nachschlagetabelleninformationsregion 405; 5) eine globale atomare Statistikregion 406; und 6) eine globale Konstantentabelleninformationsregion 407.
Die Partitionen wie in 4a dargestellt versuchen diejenigen Speicherregionen, die unter virtuellen Prozessoren geteilt werden oder „global“ sind, in Einklang mit der SIMD-artigen Beschaffenheit der gesamten Verarbeitungsumgebung zu visualisieren. Desgleichen versucht 4a auch andere Speicherregionen zu visualisieren, die unter virtuellen Prozessoren nicht geteilt werden oder einem bestimmten virtuellen Prozessor „eigen“ sind. Speziell sind wie in 4a ersichtlich alle Speicherzuordnungsbereiche global mit Ausnahme einer Notizblockspeicherregion 402, die jedem virtuellen Prozessor eigen ist. Mehrere unterschiedliche Speicherregionen weisen auch unterschiedliche Speicheradressierungsschemen auf, wie es nachfolgend näher beschrieben wird.
In Bezug auf die Zwischenregisterregion 402 ist es nicht ungewöhnlich, Zwischeninformationen im Laufe der Ausführung eines aufwendigen Bildverarbeitungsalgorithmus zu speichern (z. B., und dann die Informationen zurückzulesen und sie zu einer späteren Zeit zu verwenden). Zusätzlichist es bei solchen Informationen nicht ungewöhnlich, dass sie sich über Threads unterscheiden (unterschiedliche Eingabewerte können unterschiedliche Zwischenwerte bewirken). Das Speichermodell umfasst daher pro Prozessor eigene Notizblockspeicherregionen 402 zum Speichern solcher Zwischeninformationen durch den entsprechenden Thread jedes virtuellen Prozessors. Bei einer Ausführungsform wird auf die Zwischenregisterregion für einen bestimmten Prozessor von diesem Prozessor durch einen typischen (z. B. linearen) Direktzugriffsspeicher zugegriffen 409 und sie ist eine Lese-/Schreibspeicherregion (d. h., ein virtueller Prozessor kann sowohl Informationen vom eigenen Speicher lesen als auch Informationen in den eigenen Speicher schreiben). Ausführungsformen des ISA-Befehlsformats des virtuellen Prozessors zum Zugreifen auf die Notizblockspeicherregion werden nachfolgend ausführlicher beschrieben.
Der Eingabematrixabschnitt 403 enthält die Reihe von Eingabedaten, die in die Reihe von Threads gerufen werden 408, um Ausgabedaten zu erzeugen. Bei einer typischen Situation entspricht die Eingabematrix einem Bild (z. B. einem Frame) oder Abschnitt eines Bildes, an dem oder innerhalb dessen jeder Thread operiert. Das Eingabebild kann ein wahrer Eingang sein, wie beispielsweise die durch eine Kamera bereitgestellte Bildpunktinformation, oder eine Form von Zwischenbild, wie die durch einen vorhergehenden Kernel in einer größeren gesamten Bildverarbeitungsreihenfolge bereitgestellten Informationen sein. Virtuelle Prozessoren konkurrieren in der Regel nicht bezüglich den gleichen Eingabedatenelementen, da sie während desselben Zyklus an verschiedenen Bildpunktpositionen der Eingabebilddaten arbeiten.
Bei einer Ausführungsform wird ein neuartiges Speicheradressierschema verwendet, um zu definieren, welche bestimmten Eingabewerte von der Eingabematrix 403 aufgerufen werden sollen. Genauer gesagt wird ein „positionsbezogenes“ Adressierungsschema verwendet, das die gewünschten Eingabedaten anstatt einer herkömmlichen linearen Speicheradresse mit X, Y-Koordinaten definiert. Dementsprechend beinhaltet der Ladebefehl der Befehlssatzarchitektur der virtuellen Prozessoren ein Befehlsformat, das eine spezifische Speicherposition innerhalb der Eingabematrix mit einer X-Komponente und einer Y-Komponente identifiziert. Als solches wird ein zweidimensionales Koordinatensystem verwendet, um Speicher für Eingabewerte zu adressieren, die von der Eingabematrix 403 gelesen werden.
Die Verwendung eines positionsbezogenen Speicheradressierungsansatzes ermöglicht es, dass der Bereich eines Bildes, an dem ein virtueller Prozessor arbeitet, für einen Entwickler leichter identifizierbar ist. Wie vorstehend erwähnt, ermöglicht die, eine beliebige Eingangsmatrixposition in Kombination mit einer vollständigen ISA herkömmlicher mathematischer und Programmsteuerungsmaschinenbefehle leicht zu definieren/zu visualisieren, ermöglicht eine extrem vielseitige Programmierumgebung, die es einem Anwendungsprogrammentwickler im Wesentlichen ermöglicht, im Idealfall jede gewünschte Funktion, die ausgeführt werden soll, auf einer beliebig großen Bildfläche leicht zu definieren. Verschiedene Befehlsformat-Ausführungsformen für Befehle, die ein positionsbezogenes Adressierungsschema übernehmen, sowie Ausführungsformen anderer Merkmale der unterstützten ISA werden nachfolgend näher beschrieben.
Die Ausgabematrix 404 enthält die Ausgabebilddaten, für deren Generierung die Threads verantwortlich sind. Die Ausgabebilddaten können endgültige Bilddaten, wie z. B. die tatsächlichen Bilddaten, sein, die auf einem Display dargestellt werden, das der Gesamtbildverarbeitungssequenz folgt, oder es können Zwischenbilddaten sein, die ein nachfolgender Kernel der Gesamtbildverarbeitungssequenz als Eingabebilddateninformationen verwendet. Virtuelle Prozessoren konkurrieren wiederum in der Regel jedoch nicht um Ausgabedatenelemente, da sie während eines gleichen Zyklus in verschiedene Bildpunktpositionen der Ausgabebilddaten schreiben.
Bei einer Ausführungsform wird das positionsbezogene Adressierungsschema auch für Schreibvorgänge in die Ausgabematrix verwendet. Dementsprechend beinhaltet die ISA für jeden virtuellen Prozessor einen Speicherbefehl, dessen Befehlsformat anstelle einer herkömmlichen Speicherung mit wahlfreiem Zugriff einen adressierten Schreibort im Speicher als eine zweidimensionale X, Y-Koordinate definiert. Mehr Details bezüglich Ausführungsformen der positionsrelativen Befehle der virtuellen ISA, werden nachfolgend bereitgestellt.
4a zeigt auch jeden virtuellen Prozessor, der ein Nachschlagen 410 in einer Nachschlagetabelle 411 ausführt, die innerhalb der Nachschlagetabellenspeicherregion 405 gehalten wird. Nachschlagetabellen werden häufig von Bildverarbeitungsaufgaben verwendet um z. B. Filter zu erlangen oder Koeffizienten für unterschiedliche Matrixorte umzuwandeln, komplexe Funktionen (z. B. Gammakurven, Sinus, Cosinus) zu implementieren, wobei die Nachschlagetabelle die Funktionsausgabe für einen Eingabeindexwert bereitstellt, usw. Hier wird erwartet, dass SIMD-Bildverarbeitungssequenzen häufig ein Nachschlagen in einer gleichen Nachschlagetabelle während eines gleichen Taktzyklus ausführen. Als solches kann auf die Nachschlagetabellenregion 405, wie auf die Eingabe- und Ausgabematrixspeicherregionen 403, 404, global durch jeden virtuellen Prozessor zugegriffen werden. 4a zeigt desgleichen jeden der drei virtuellen Prozessoren, die Informationen von einer gleichen Nachschlagetabelle 411 effektiv nachschlagen, die in der Speicherregion der Nachschlagetabelle 405 gehalten werden.
Da Indexwerte typischerweise verwendet werden, um einen gewünschten Nachschlagetabelleneintrag zu definieren, wird bei einer Ausführungsform auf die Nachschlagetabelleninformationsregion unter Verwendung eines normalen linear zugreifenden Schemas zugegriffen. Bei einer Ausführungsform ist die Nachschlagspeicherregion nur Lesen (d. h., der Prozessor kann Informationen in einer Nachschlagetabelle nicht ändern und ihm ist nur erlaubt, Informationen daraus zu lesen). Der Einfachheit halber zeigt 4a, dass nur eine Nachschlagetabelle innerhalb der Nachschlagetabellenregion 405 speicherresident ist, aber die virtuelle Umgebung erlaubt es, dass während der simulierten Laufzeit mehrere unterschiedliche Nachschlagetabellen speicherresident sind. Ausführungsformen des virtuellen ISA-Befehlsformats für Befehle, die ein Nachschlagen in der Nachschlagetabelle ausführen, werden nachfolgend bereitgestellt.
4b zeigt jeden der drei virtuellen Prozessoren, die in die atomare Statistikregion 406 schreiben 413. Es ist für Bildprozesse nicht ungewöhnlich, zu „aktualisieren“ oder eine einfache Änderung vorzunehmen, um Informationen auszugeben. Die aktualisierten Informationen können dann für andere nachgeschaltete Prozesse verwendet werden, die von den aktualisierten Informationen Gebrauch machen. Beispiele von solchen Aktualisierungen oder einfachen Änderungen umfassen einfache Hinzufügungen eines festen Offsets zu Ausgabedaten, einfache Multiplikation eines Multiplikanden mit Ausgabedaten oder Minimum- oder Maximumvergleiche von Ausgabedaten mit einem Schwellenwert.
Wie in 4b ersichtlich, kann in diesen Sequenzen an Ausgabedaten, die gerade durch die individuellen Threads 401 berechnet wurden, operiert werden und die Resultate können in die atomare Statistikregion 406 geschrieben 413 werden. Abhängig von Implementierungssemantiken können die Ausgabedaten, an denen durch eine atomare Operation gearbeitet wird, durch den Prozessor intern gehalten oder von der Ausgabematrix aufgerufen 412 werden. 4b zeigt Letzteres. Bei verschiedenen Ausführungsformen umfassen die atomaren Operationen, die an den Ausgabedaten ausgeführt werden können, Addieren, Multiplizieren, min und max. Bei einer Ausführungsform wird auf die atomare Statistikregion 406 unter Verwendung eines positionsrelativen Adressierungschemas zugegriffen (wie mit Eingabe- und Ausgabematrixzugriffen), vorausgesetzt, dass Aktualisierungen von Ausgabedaten in einer gleichen zweidimensionalen Matrix wie die Ausgabedaten selbst logisch organisiert werden würden. Ausführungsformen des virtuellen ISA-Befehlsformats zum Ausführen einer atomaren Operation an Ausgabedaten und Schreiben der Resultanten in die Statistikregion 406 werden nachfolgend ausführlicher beschrieben.
4c zeigt jeden der virtuellen Prozessoren, die einen Konstantenwert von einer Konstantennachschlagetabelle 415 innerhalb der Konstantenspeicherregion 407 lesen 414. Hier wird z. B. erwartet, dass unterschiedliche Threads 401 eine gleiche Konstante oder einen anderen Wert im gleichen Taktzyklus benötigen können (z. B. einen bestimmten Multiplikator, der auf ein gesamtes Bild anzuwenden ist). Daher geben Zugriffe auf die Konstantennachschlagetabelle 415 einen gleichen Skalarwert an jeden der virtuellen Prozessoren zurück, wie es in 4c dargestellt ist. Da auf Nachschlagetabellen typischerweise mit einem Indexwert zugegriffen wird, wird bei einer Ausführungsform auf die Konstantennachschlagetabellen-Speicherregion mit einer linearen Random Access Memory-Adresse zugegriffen. Bei einer Ausführungsform ist die Konstantenspeicherregion nur Lesen (d. h., dass der Prozessor keine Informationen in einer Konstantentabelle ändern kann und ihm nur erlaubt ist, Informationen daraus zu lesen). Der Einfachheit halber zeigt 4c nur eine einzelne Konstantenachschlagetabelle 415 in der Konstantenspeicherregion 407. Da Threads von mehr als einer solchen Tabelle Gebrauch machen können, ist die Speicherregion 407 derart konfiguriert, dass sie ausreichend groß ist, um so viele Konstantentabellen wie benötigt/verwendet zu halten.
c. Virtueller ISA-Prozessor
Wie in mehreren vorstehenden Fällen darauf hingewiesen kann der virtuelle ISA-Prozessor eine Anzahl entsprechender Merkmale umfassen. Einige davon werden unmittelbar nachstehend ausführlich beschrieben.
Bei verschiedenen Ausführungsformen verwendet das Befehlsformat von dem ISA eines jeden virtuellen Prozessors eine relative Positionierungsherangehensweise zum Definieren einer X, Y-Koordinate für jedes von Folgendem: 1) einen LOAD-Befehl, der Eingabebilddaten von der Eingabematrixspeicherregion liest; 2) einen STORE-Befehl, der Ausgabebilddaten in die Ausgabematrix schreibt; und 3) eine atomare Aktualisierung an der Statistikregion eines Speichers.
Die Fähigkeit, eine beliebige Eingabematrixposition in Kombination mit einer vollständigen ISA von herkömmlichem Datenzugriff, herkömmlicher mathematischer und Programmsteuerungsmaschinenbefehle leicht zu definieren, ermöglicht eine extrem vielseitige Programmierumgebung, die es einem Anwendungsprogrammentwickler im Wesentlichen ermöglicht, im Idealfall jede gewünschte Funktion, die ausgeführt werden soll, auf einer beliebig großen Bildfläche zu definieren. Zum Beispiel kann im Idealfall jede mathematische Operation leicht programmiert und auf jede Schablonengröße angewendet werden.
Bei einer Ausführungsform weisen Befehle zum Laden von/Speichern in den Ein-/Ausgangmatrizen das folgende Format auf [OPCODE] LINEGROUP_(name)[(((X·XS + X0)/XD); ((Y·YS + Y0)/YD); Z] wobei [OPCODE] die spezifische Art der Operation ist (LADEN (LOAD) von der Eingabematrix, SPEICHERN (STORE) in die Ausgabematrix) und LINEGROUP_(name) der Name ist, der einem bestimmten Abschnitt eines bestimmten Bildes (z. B. eine Zeilengruppe für einen Frame von Bilddaten) innerhalb der Eingabe- oder Ausgabematrixspeicherregion zugewiesen ist. Da an unterschiedlichen Zeilengruppen separat operiert wird, sind den unterschiedlichen Zeilengruppen hier unterschiedliche Namen gegeben, sodass diese eindeutig identifiziert werden können bzw. eindeutig darauf zugegriffen werden kann (z. B. LINEGROUP_1, LINEGROUP_2 usw.). Zeilengruppen des gleichen Namens können sowohl in der Eingabematrixspeicherregion als auch in der Ausgabematrixspeicherregion existieren. Der Ursprung irgendeiner Zeilengruppe kann z. B. deren untere linke Ecke innerhalb ihrer entsprechenden Speicherregion sein.
Im Fall von Befehlen, die Aktualisierungen an der atomaren Statistiktabelle ausführen, nimmt bei einer Ausführungsform das Befehlsformat die folgende ähnliche Struktur an [OPCODE] STATS_(name)[(((X·XS + X0)/XD); ((Y·YS + Y0)/YD); Z] wobei der erwähnenswerte Unterschied ist, dass die Eingabeoperandinformationen eine Position innerhalb einer bestimmten Statistiktabelle (STATS_(name)) anstatt einer bestimmten Zeilengruppe innerhalb der Eingabe- oder Ausgabematrix definiert. Wie auch bei Zeilengruppen werden unterschiedlichen Statistiktabellen unterschiedliche Namen gegeben, sodass ein Thread eindeutig an unterschiedlichen Statistiktabellen im Laufe seiner Operation operieren kann. Der [OPCODE] spezifiziert die bestimmte auszuführende atomare Operation (z. B. STAT_ADD; STAT_MUL; STAT_MIN; STAT_MAX).
Entweder für Ein-/Ausgabematrixzugriffe oder für atomare Statistiktabellenzugriffe definiert der Z-Operand des Befehls, welcher Kanal einer benannten Zeilengruppe oder Statistiktabelle durch den Befehl adressiert wird. Typischerweise wird hier ein Einzelbild mehrere Kanäle aufweisen. Videobilder weisen beispielsweise typischerweise einen roten Kanal (R), einen grünen Kanal (G) und einen blauen Kanal (B) für einen gleichen Frame des Videostreams auf. Ein komplettes Bild kann gewissermaßen als separate R-, G- und B-Kanalbilder angesehen werden, die aufeinandergestapelt sind. Der Z-Operand definiert, welcher von diesen durch den Befehl adressiert wird (z. B. Z = 0 entspricht dem roten Kanal, Z = 1 entspricht dem blauen Kanal und Z = 2 entspricht dem grünen Kanal). Jede Zeilengruppe und Statistiktabelle ist daher strukturiert, den Inhalt jedes Kanals für das bestimmte Bild zu beinhalten, das verarbeitet wird.
Der Operand (X·XS + X0)/XD definiert den X-Ort innerhalb einer benannten Zeilengruppe oder Statistiktabelle, die durch den Befehl adressiert wird, und der Operand (Y·YS + Y0)/YD definiert den Y-Ort innerhalb einer benannten Zeilengruppe oder Statistiktabelle, die durch den Befehl adressiert wird. Die XS- und XD-Termini für den X-Ort und die YS- und YD-Termini für den Y-Ort werden für das Skalieren zwischen Eingabe- und Ausgabebildern mit unterschiedlichen Bildpunktdichten verwendet. Skalieren wird nachfolgend ausführlicher beschrieben.
Bei einem einfachsten Fall gibt es kein Skalieren zwischen Eingabe- und Ausgabebildern und die X- und Y-Komponenten des Befehlsformats nehmen einfach die Form X + X0 und Y + Y0 an, wobei X0 und Y0 Positionsoffsets relativ zur Position des Threads sind. Ein Thread wird als der Position innerhalb der Zeilengruppe der Ausgabematrix zugewiesen angesehen, in die sein Ausgabewert geschrieben wird. Eine entsprechende gleiche Position ist ohne Weiteres in der Zeilengruppe der Eingabematrix und irgendeiner Statistiktabelle identifizierbar.
Wenn dem Thread als Beispiel ein spezifischer X, Y-Ort in einer Ausgabematrix LINEGROUP_1 zugewiesen wird, würde der Befehl LOAD LINEGROUP_1[(x – 1); (Y – 1); Z] von LINEGROUP_1 einen Wert laden, der sich ein Bildpunktort nach links und ein Bildpunktort nach unten von dem gleichen X,Y-Ort innerhalb der Eingabematrix befindet.
Ein einfacher Unschärfekernel, der Durchschnitte für die Bildpunktwerte des X,Y-Orts zusammen mit ihren linken und rechten Nachbarn bildet, kann daher wie in 5a dargestellt in Pseudocode geschrieben werden. Wie in 5a ersichtlich, entspricht der Ort ((X); (Y)) dem Ort des virtuellen Prozessors, der in die Ausgabematrix schreibt. Im vorstehenden Pseudocode entspricht LOAD dem Operationscode für ein Laden von der Eingabematrix und STORE dem Operationscode für das Speichern in die Ausgabematrix. Es ist zu beachten, dass eine LINEGROUP_1 in der Eingabematrix und eine LINEGROUP_1 in der Ausgabematrix existiert.
5b stellt skalierte Bilder zum Zwecke des Erklärens der Skaliermerkmale des relativen Positionierungs-Laden- und Speichern-Befehlsformats dar. Downsampling bezeichnet die Transformation eines Bildes mit einer höheren Auflösung in ein Bild mit einer niedrigeren Auflösung durch Bereitstellen von weniger als allen der Bildpunkte, die im Eingabebild existieren, im Ausgabebild. Upsampling bezeichnet die Transformation eines Bildes mit einer niedrigeren Auflösung in ein Bild mit einer höheren Auflösung durch Erzeugen von mehr Bildpunkten im Ausgabebild, als im Eingabebild existieren.
Wenn beispielsweise unter Bezugnahme auf 5b das Bild 501 das Eingabebild darstellt und das Bild 502 das Ausgabebild darstellt, wird Downsampling ausgeführt, da es im Ausgabebild weniger Bildpunkte gibt als im Eingabebild. Für jeden Bildpunkt im Ausgabebild schreiten hier die entsprechenden Bildpunkte im Eingabebild, die den Ausgabewert für einen Ausgabebildpunkt bestimmen, „weiter fort“ vom Ausgabebildpunktort, indem sie sich entlang einer Achse im Ausgabebild bewegen. Für ein 3:1-Downsampling-Verhältnis entspricht der erste Bildpunkt im Ausgabebild entlang einer Achse den ersten, zweiten und dritten Bildpunkten entlang der gleichen Achse im Ausgabebild, der zweite Bildpunkt im Ausgabebild den vierten, fünften und sechsten Bildpunkten im Eingabebild usw. Daher weist der erste Ausgabebildpunkt einen entsprechenden Bildpunkt am dritten Ort auf, während der zweite Ausgabebildpunkt einen entsprechenden Bildpunkt am sechsten Ort aufweist.
Als solches werden die XS- und YS-Multiplikandentermini im relativen Positionierungs-Befehlsformat verwendet, um Downsampling zu implementieren. Wenn der Unschärfepseudocode der 5a für ein 3:1-Downsampling entlang beider Achsen neu geschrieben werden müsste, würde der Code neu geschrieben werden als:
R1 <= LOAD LINEGROUP_1[((3X) – 1); 3(Y); 0]
R2 <= LOAD LINEGROUP_1[3(X); 3(Y); 0]
R3 <= LOAD LINEGROUP_1[((3X) + 1); 3(Y); 0]
R2 <= ADD R1, R2
R2 <= ADD R2, R3
R2 <= DIV R2, 3
STORE LINEGROUP_1[(X); (Y); (0)]; R2.
Im Vergleich dazu würden im Fall eines 1:3-Upsamplings (z. B. ist Bild 502 das Eingabebild und Bild 501 ist das Ausgabebild) die XD- und YD-Divisoren dazu verwendet werden, um drei Ausgabebildpunkte für jeden Eingabebildpunkt entlang einer Achse zu erzeugen. Als solches würde der Unschärfecode neu geschrieben werden als:
R1 <= LOAD LINEGROUP_1[(X – 1)/3; (Y)/3; 0]
R2 <= LOAD LINEGROUP_1[(X)/3; (Y)/3; 0]
R3 <= LOAD LINEGROUP_1[(X + 1)/3; (Y)/3; 0]
R2 <= ADD R1, R2
R2 <= ADD R2, R3
R2 <= DIV R2, 3
STORE LINEGROUP_1[(X); (Y); (0)]; R2
Bei verschiedenen Ausführungsformen umfasst das Befehlsformat für Befehle, die auf die eigenen, Konstanten- und Nachschlagabschnitte des Speichers zugreifen, einen Operanden, der auch die Form a·b + c annimmt, wobei a eine Grundposition ist, b ein Skalierbegriff ist und c ein Offset ist. Hier wird jedoch eine Herangehensweise linearer Adressierung verwendet, bei der der Term a·b + c – im Wesentlichen einem linearen Index entspricht, der auf die adressierte Tabelle angewandt wird. Jeder dieser Befehle umfasst auch im Operationscode und einem Bezeichner der Speicherregion, auf die zugegriffen wird. Ein Befehl, der ein Nachschlagen in der Nachschlagetabellenspeicherregion ausführt, kann beispielsweise ausgedrückt werden als LOAD LKUP_(name)[(A·B + C)]. wobei LOAD der Operationscode ist, der eine Ladeoperation identifiziert und LKUP_(name) den Namen der Nachschlagetabelle in der Nachschlagetabellenspeicherregion spezifiziert, auf die zugegriffen wird. Es können erneut mehrere Nachschlagetabellen durch einen Thread verwendet werden, und daher wird ein Bennenungsschema verwendet, um die geeignete eine von den mehr als einen zu identifizieren, die in der Nachschlagetabellenspeicherregion existieren.
Ein ähnliches Format mit einem in ähnlicher Weise gesonnenen Operationscode kann für Befehle verwendet werden, welche die Konstanten- und die eigenen Speicherregionen adressieren (z. B. LOAD CNST_(name)[(A·B + C)]; LOAD PRVT_(name)[(A·B + C)]. Bei einer Ausführungsform sind die Nachschlagetabellen- und Konstantentabellenzugriffe nur lesen (ein Prozessor kann die Daten nicht verändern, die dort abgelegt wurden). Als solches existieren keine STORE-Befehle für diese Speicherregionen. Bei einer Ausführungsform ist die eigene Speicherregion lesen/schreiben. Als solches existiert ein Speicherbefehl für diese Speicherregion (z. B. STORE PRVT[(A·B + C)].
Bei verschiedenen Ausführungsformen umfasst jeder virtuelle Prozessor Allzweckregister, die Ganzzahl-, Fließkomma- oder Festpunktwerte enthalten können. Zusätzlich können die Allzweckregister Dateninhalte mit konfigurierbarer Bitbreite wie 8-, 16- oder 32-Bit-Werte enthalten. Daher können die Bilddaten an jeder Bildpunktposition in einer Eingabematrix oder Ausgabematrix eine Datengröße von 8, 16 oder 32 Bits aufweisen. Hier kann ein virtueller Prozessor für einen Ausführungsmodus konfiguriert sein, der die Bitgröße und das Zahlenformat der Werte innerhalb der Allzweckregister etabliert. Befehle können auch unmittelbare Operanden spezifizieren (die Eingabeoperanden sind, deren Eingabewerte direkt im Befehl selbst ausgedrückt werden, als dass sie in einem spezifizierten Register gefunden werden). Unmittelbare Operanden können auch konfigurierbare 8-, 16- oder 32-Bit-Breiten aufweisen.
Bei einer erweiterten Ausführungsform ist jeder virtuelle Prozessor auch fähig, in einem Skalarmodus oder einem SIMD-Modus intern von sich selbst zu operieren. D. h., dass die Daten innerhalb eines spezifischen Matrixorts als ein Skalarwert oder als ein Vektor mit mehreren Elementen betrachtet werden können. Eine erste Konfiguration kann beispielsweise eine Skalaroperation von 8 Bits etablieren, wobei jede Bildmatrixposition einen Skalar mit einem 8-Bit-Wert hält. Im Vergleich dazu kann eine andere Konfiguration ein parallele/SIMD-Operation mit 32 Bits etablieren, wobei von jedem Bildanordungsort angenommen wird, dass er vier 8-Bit-Werte für eine gesamte Datengröße von 32 Bit pro Matrixort hält.
Bei verschiedenen Ausführungsformen umfasst jeder virtuelle Prozessor auch Register zum Halten von Prädikatwerten. Ein einzelner Prädikatwert ist häufig nur ein Bit lang und drückt eine Resultante von einem Operationscode aus, der einen Richtig/Falsch- oder Größer/Kleiner-Test an vorhandenen Daten ausführt. Prädikatwerte werden verwendet, um z. B. während der Ausführung Zweigrichtungen durch den Code zu bestimmen (und werden daher als Operanden bei bedingten Zweigbefehlen verwendet). Prädikatwerte können auch als ein Direktoperand in einem Befehl ausgedrückt werden.
Bei verschiedenen Ausführungsformen umfasst jeder virtuelle Prozessor Register zum Halten von Skalarwerten. Skalarwerte werden hier in den Partitionsraum des Speichermodells, der für Konstanten reserviert ist, gespeichert und daraus gelesen (wie vorstehend in Bezug auf 4c beschrieben). Jeder virtuelle Prozessor einer Gruppe von virtuellen Prozessoren, die ein gleiches Bild verarbeiten, verwendet hier den gleichen Skalarwert des Konstantenspeicherraums. Bei erweiterten Ausführungsformen existieren auch Skalarprädikate. Diese sind in Registerraum gehaltene Werte, die der Definition von sowohl eines Prädikats als auch eines Skalars entsprechen.
Bei verschiedenen Ausführungsformen ist jeder virtuelle Prozessor als ein RISC-artiger Befehlssatz ausgelegt, dessen unterstützte arithmetische Operationscodes irgendeine brauchbare Kombination von Folgendem umfassen: 1) ADD (Addition der Operanden A und B); 2) SUB (Subtraktion der Operanden A und B); 3) MOV (Operand von einem Register zu einem anderen Register bewegen); 4) MUL (Operanden A und B multiplizieren); 5) MAD (die Operanden A und B multiplizieren und C zu der Resultante addieren); 6) ABS (den Absolutwert des Operanden A zurückgeben); 7) DIV (Operand A durch Operand B teilen); 8) SHL (Operand A nach links schieben); 9) SHR (Operand A nach rechts schieben); 10) MIN/MAX (zurückgeben, welcher der Operanden A und B größer ist); 11) SEL (spezifizierte Bytes des Operanden A auswählen); 12) AND (das logische UND der Operanden A und B zurückgeben); 13) OR (das logische ODER der Operanden A und B zurückgeben); 14) XOR (das logische exklusive ODER der Operanden A und B zurückgeben); 15) NOT (die logische Umkehrung des Operanden A zurückgeben).
Der Befehlssatz umfasst auch Standardprädikatoperationen wie: 1) SEQ (gibt 1 zurück, wenn A gleich B ist); 2) SNE (gibt 1 zurück, wenn A ungleich B ist); 3) SLT (gibt 1 zurück, wenn A kleiner B ist); 4) SLE (gibt 1 zurück, wenn A kleiner oder gleich B ist). Steuerablaufbefehle wie JMP (springe) und BRANCH sind auch beinhaltet, wobei jeder von beiden nominale Variablen oder Prädikate als Operanden umfassen kann.
d. Anwendungssoftwareentwicklung und Simulationsumgebung
6 stellt eine Anwendungssoftwareentwicklung und Simulationsumgebung 601 dar. Wie vorstehend in Bezug auf 2 beschrieben, kann ein Entwickler eine umfangreiche Bildverarbeitungsfunktion (wie z. B. eine Bildverarbeitungspipeline, wobei jede Stufe der Pipeline eine zugehörige Bildverarbeitungsaufgabe ausführt, einige andere DAG-vorgeschriebene Routinensätze usw.) durch Anordnen von Kerneln in einer strategischen Sequenz, die mit der gesamten beabsichtigten Bildveränderung konsistent ist, entwickeln. Kernel können von einer Bibliothek 602 abgerufen werden und/oder der Entwickler kann einen oder mehrere benutzerdefinierte Kernel entwickeln.
Kernel innerhalb der Bibliothek 602 können durch einen Drittanbieter von Kerneln und/oder einen Provider von irgendeiner zugrundeliegenden Technologie bereitgestellt werden (z. B. ein Lieferant einer Hardwareplattform, die den adressierten Hardwarebildprozessor umfasst, oder ein Lieferant des adressierten Hardwarebildprozessors (z. B. als ein dazu bereitgestelltes Design oder als tatsächliche Hardware)).
Im Fall von benutzerdefiniert entwickelten Kerneln muss der Entwickler in vielen Situationen nur den Programmcode für einen einzelnen Thread 603 schreiben. D. h., dass der Entwickler nur Programmcode schreiben muss, der einen einzelnen Ausgabebildpunktwert durch Referenzieren von Eingabebildpunktwerten relativ zum Ausgabebildpunktort (z. B. mit dem vorstehend beschriebenen positionsrelativen Speicherzugriffsbefehlsformat) bestimmt. Nach Erfüllen der Operation des einzelnen Threads 603 kann die Entwicklungsumgebung dann automatisch mehrere Instanzen des Threadcodes auf einem entsprechenden virtuellen Prozessor instanziieren, um einen Kernel auf einer Matrix von Prozessoren zu bewirken, die an einem Bildflächenbereich operieren. Der Bildflächenbereich kann ein Abschnitt eines Bildframes (wie eine Zeilengruppe) sein.
Bei verschiedenen Ausführungsformen wird der benutzerdefinierte Threadprogrammcode in den Objektcode des ISA des virtuellen Prozessors (oder eine höhere Sprache, die in den ISA-Objektcode des virtuellen Prozessors abwärtskompiliert wird) geschrieben. Die Simulation der Ausführung des benutzerdefinierten Programmcodes des Kernels kann in einer simulierten Laufzeitumgebung ausgeführt werden, die einen virtuellen Prozessor umfasst, der auf einen gemäß dem Speichermodell organisierten Speicher zugreift. Hier werden Softwaremodelle (objektorientiert oder anderweitig) eines virtuellen Prozessors 604 und ein Speicher 605, der das Modell einbezieht, instanziiert.
Das virtuelle Prozessormodell 604 simuliert dann die Ausführung des Threadcodes 603. Nach Erfüllen der Ausführung eines Threads, wird ein größerer Kernel und irgendeine größere Funktion, zu welcher der Kernel gehört, in den tatsächlichen Objektcode der darunterliegenden Hardware kompiliert. Die Gesamtheit der Simulationsumgebung 601 kann als Software implementiert sein, die auf einem Computersystem (z. B. einer Arbeitsstation) 606 läuft.
2.0 Hardwarearchitektur-Ausführungsformen
a. Bildprozessor-Hardwarearchitektur und -Betrieb
7 zeigt eine Ausführungsform einer Architektur 700 für einen in Hardware implementierten Bildprozessor. Der Bildprozessor kann z. B. von einem Compiler angesteuert werden, der den Programmcode, der für einen virtuellen Prozessor geschrieben wurde, in einer simulierten Umgebung in Programmcode umwandelt, der von dem Hardwareprozessor tatsächlich ausgeführt wird. Wie in 7 ersichtlich umfasst die Architektur 700 mehrere Zeilenpuffereinheiten 701_1 bis 701_M (im Folgenden „Zeilenpuffer“: „Zeilenpuffereinheiten“ oder dergleichen), die miteinander zu mehreren Schablonen-Prozessoreinheiten 702_1 bis 702_N (im Folgenden: „Schablonenprozessoren“: „Schablonen-Prozessoreinheiten“ oder dergleichen) und entsprechende Blattgeneratoreinheiten 703_1 bis 703_N (im Folgenden „Blattgenerator“: „Blattgeneratoreinheiten“ oder dergleichen) durch ein Netzwerk 704 (z. B. ein Network-On-Chip (NOC) einschließlich eines chipinternen Wählnetzwerks, eines chipinternen Ringnetzwerks oder einer anderen Art von Netzwerk) verbunden sind. Bei einer Ausführungsform kann jede Zeilenpuffereinheit mit jedem Blattgenerator und einem entsprechenden Schablonenprozessor über das Netzwerk 704 verbunden sein.
Bei einer Ausführungsform wird der Programmcode kompiliert und auf einen entsprechenden Schablonenprozessor 702 geladen, um die zuvor von einem Softwareentwickler definierten Bildverarbeitungsoperationen auszuführen (der Programmcode kann je nach Auslegung und Implementierung auch auf den zugehörigen Blattgenerator des Schablonenprozessors 703 geladen werden). In zumindest einigen Fällen kann eine Bildverarbeitungspipeline realisiert werden, indem ein erstes Kernelprogramm für eine erste Pipelinestufe in einen ersten Schablonenprozessor 702_1 geladen, ein zweites Kernelprogramm für eine zweite Pipelinestufe in einen zweiten Schablonenprozessor 702_2 geladen wird usw., wobei der erste Kernel die Funktionen der ersten Pipelinestufe durchführt, der zweite Kernel die Funktionen der zweiten Pipelinestufe durchführt usw., und zusätzliche Steuerablaufverfahren installiert werden, um Ausgabebilddaten von einer Pipelinestufe zur nächsten Pipelinestufe weiterzugeben.
Bei anderen Konfigurationen kann der Bildprozessor als eine parallele Maschine mit zwei oder mehr Schablonenprozessoren 702_1, 702_2 realisiert sein, die den gleichen Kernelprogrammcode ausführen. Zum Beispiel kann ein hochgradig dichter und hoher Datenratenstrom von Bilddaten verarbeitet werden, indem Frames über mehrere Schablonenprozessoren verteilt werden, von denen jeder dieselbe Funktion ausführt.
Bei noch anderen Konfigurationen kann im Wesentlichen jeder DAG von Kerneln auf den Hardwareprozessor geladen werden, indem jeweilige Schablonenprozessoren mit deren eigenen jeweiligen Kernel von Programmcode konfiguriert und geeignete Steuerablauf-Hooks in die Hardware konfiguriert werden, um Ausgabebilder von einem Kernel an den Eingang eines nächsten Kernels im DAG-Design zu leiten.
Als allgemeiner Ablauf werden Frames der Bilddaten von einer Makro-E/A-Einheit 705 empfangen und zu einer oder mehreren der Zeilenpuffereinheiten 701 auf einer Frame-per-Frame-Basis weitergegeben. Eine bestimmte Zeilenpuffereinheit parst ihren Frame aus Bilddaten in einen kleineren Bereich von Bilddaten, der als „Zeilengruppe“ bezeichnet wird, und gibt dann die Zeilengruppe durch das Netzwerk 704 an einen bestimmten Blattgenerator weiter. Eine vollständige oder „volle“ singuläre Zeilengruppe kann sich beispielsweise aus den Daten mehrerer zusammenhängender vollständiger Zeilen oder Spalten eines Frames zusammensetzen (der Einfachheit halber bezieht sich die vorliegende Beschreibung hauptsächlich auf zusammenhängende Zeilen). Der Blattgenerator parst die Zeilengruppe von Bilddaten weiter in einen kleineren Bereich von Bilddaten, der als „Blatt“ bezeichnet wird, und präsentiert das Blatt seinem entsprechenden Schablonenprozessor.
Im Falle einer Bildverarbeitungspipeline oder eines DAG-Ablaufs mit einem einzigen Eingang werden im Allgemeinen Eingabeframes an die gleiche Zeilenpuffereinheit 701_1 geleitet, welche die Bilddaten in Zeilengruppen parst und die Zeilengruppen zu dem Blattgenerator 703_1 leitet, dessen entsprechender Schablonenprozessor 702_1 den Code des ersten Kernels in der Pipeline/dem DAG ausführt. Nach der Fertigstellung der Operationen durch den Schablonenprozessor 702_1 an den Zeilengruppen, die er verarbeitet, sendet der Blattgenerator 703_1 Ausgabezeilengruppen zu einer „nachgeschalteten“ Zeilenpuffereinheit 701_2 (in einigen Anwendungsfällen kann die Ausgabezeilengruppe zur gleichen Zeilenpuffereinheit 701_1 zurückgesendet werden, die zuvor die Eingabezeilengruppen gesendet hat).
Ein oder mehrere „Abnehmer“-Kernel, welche die nächste Phase/Operation in der Pipeline/dem DAG darstellen, die auf deren eigenen jeweiligen anderen Blattgenerator und Schablonenprozessor (z. B. Blattgenerator 703_2 und Schablonenprozessor 702_2) ausgeführt werden, empfangen dann die von dem ersten Schablonenprozessor 702_1 generierten Bilddaten von der nachgeschalteten Zeilenpuffereinheit 701_2. Auf diese Weise werden die Ausgabedaten eines „Erzeuger“-Kernels, der auf einem ersten Schablonenprozessor betrieben wird, an einen „Abnehmer“-Kernel weitergeleitet, der auf einem zweiten Schablonenprozessor betrieben wird, wobei der Abnehmerkernel nach dem Erzeugerkernel den nächsten Satz von Arbeitsschritten gemäß des Designs der gesamten Pipeline oder des DAG ausführt.
Ein Schablonenprozessor 702 ist dafür ausgelegt, gleichzeitig an mehreren überlappenden Schablonen von Bilddaten zu arbeiten. Die mehreren überlappenden Schablonen und die interne Hardwareverarbeitungskapazität des Schablonenprozessors bestimmen effektiv die Größe eines Blattes. Hier arbeiten innerhalb eines Schablonenprozessors 702 Matrizen von Ausführungsbahnen zusammen, um den Bilddatenoberflächenbereich gleichzeitig zu bearbeiten, der von den mehreren überlappenden Schablonen abgedeckt ist.
Wie nachstehend näher beschrieben, werden in verschiedenen Ausführungsformen, Blätter von Bilddaten in eine zweidimensionale Registermatrixstruktur innerhalb des Schablonenprozessors 702 geladen. Es wird davon ausgegangen, dass die Verwendung von Blättern und die zweidimensionale Registermatrixstruktur für effektive Energieverbrauchsverbesserungen sorgen, indem eine große Datenmenge in einen großen Registerraum bewegt wird, wie beispielsweise eine einzelne Ladeoperation mit direkt an den Daten ausgeführten Verarbeitungsschritten, die unmittelbar danach durch eine Ausführungsbahnmatrix erfolgen. Zudem stellt die Verwendung einer Ausführungsbahnmatrix und einer entsprechenden Registermatrix verschiedene Schablonengrößen bereit, die leicht programmierbar/konfigurierbar sind.
Die 8a bis 8e veranschaulichen umfassend Ausführungsformen sowohl der Parsing-Aktivität einer Zeilenpuffereinheit 701, der feineren Parsing-Aktivität einer Blattgeneratoreinheit 703, als auch der Schablonenverarbeitungsaktivität des Schablonenprozessors 702, der mit der Blattgeneratoreinheit 703 gekoppelt ist.
8a zeigt eine Ausführungsform eines Eingabeframes der Bilddaten 801. 8a zeigt zudem einen Umriss von drei überlappenden Schablonen 802 (wobei jede Schablone eine Dimension von 3 Bildpunkten × 3 Bildpunkten aufweist), für die der Schablonenprozessor ausgelegt ist, daran zu arbeiten. Der Ausgabebildpunkt, für den jede Schablone jeweils die Ausgabebilddaten generiert, ist in reinem Schwarz hervorgehoben. Der Einfachheit halber sind die drei überlappenden Schablonen 802 als nur in vertikaler Richtung überlappend dargestellt. Es ist relevant, zu erkennen, dass ein Schablonenprozessor in Wirklichkeit derart ausgelegt sein kann, dass dieser sowohl in vertikaler als auch in horizontaler Richtung überlappende Schablonen aufweist.
Aufgrund der vertikal überlappenden Schablonen 802 innerhalb des Schablonenprozessors, wie in 8a dargestellt, gibt es ein breites Band von Bilddaten innerhalb des Frames, an denen ein einzelner Schablonenprozessor betreiben kann. Wie nachfolgend näher beschrieben, verarbeiten die Schablonenprozessoren bei einer Ausführungsform innerhalb ihrer überlappenden Schablonen Daten von links nach rechts über die Bilddaten (und wiederholen den Vorgang dann für die nächste Reihe von Zeilen in der Reihenfolge von oben nach unten). Daher nimmt, während die Schablonenprozessoren mit ihrer Operation fortfahren, die Anzahl der schwarzen Ausgabebildpunktblöcke horizontal nach rechts zu. Wie vorstehend beschrieben, ist eine Zeilenpuffereinheit 701 für das Parsen einer Zeilengruppe von Eingabebilddaten aus einem eingehenden Frame verantwortlich, der für die Schablonenprozessoren ausreichend ist, um daran eine ausgedehnte Anzahl von kommenden Zyklen zu arbeiten. Eine beispielhafte Darstellung einer Zeilengruppe ist als schattierter Bereich 803 veranschaulicht. Bei einer weiteren Ausführungsform kann wie nachstehend weiter beschrieben die Zeilenpuffereinheit 701 unterschiedliche Dynamiken zum Senden/Empfangen einer Zeilengruppe an einen/von einem Blattgenerator umfassen. Beispielsweise werden gemäß einem Modus, der als „vollständige Gruppe“ bezeichnet wird, die Gesamtbreitenzeilen von Bilddaten zwischen einer Zeilenpuffereinheit und einem Blattgenerator weitergegeben. Gemäß einem zweiten Modus, der als „virtuell groß“ bezeichnet wird, wird eine Zeilengruppe zunächst mit einer Teilmenge von Zeilen mit voller Breite weitergegeben. Die verbleibenden Zeilen werden dann nacheinander in kleineren Stücken (mit weniger als voller Breite) weitergegeben.
Wenn die Zeilengruppe 803 der Eingabebilddaten durch die Zeilenpuffereinheit definiert und an die Blattgeneratoreinheit weitergegeben wurde, parst die Blattgeneratoreinheit die Zeilengruppe weiter in feinere Blätter, die an die Hardwarebeschränkungen des Schablonenprozessors präziser angepasst sind. Insbesondere besteht bei einer Ausführungsform, wie nachfolgend näher beschrieben, jeder Schablonenprozessor aus einer zweidimensionalen Schieberegistermatrix. Die zweidimensionale Schieberegistermatrix schiebt im Wesentlichen Bilddaten „unterhalb“ einer Matrix von Ausführungsbahnen, wobei das Muster des Schiebens bewirkt, dass jede Ausführungsbahn innerhalb ihrer eigenen jeweiligen Schablone an Daten arbeitet (d. h. jede Ausführungsbahn ihre eigene Schablone von Informationen verarbeitet, um eine Ausgabe für diese Schablone zu generieren). Bei einer Ausführungsform sind Blätter Flächenbereiche von Eingabebilddaten, die die zweidimensionale Schieberegistermatrix „ausfüllen“ oder anderweitig in diese geladen werden.
Wie in 8b dargestellt, parst der Blattgenerator ein Anfangsblatt 804 von der Zeilengruppe 803 und stellt es dem Schablonenprozessor bereit (hier entspricht das beispielhafte Datenblatt dem schattierten Bereich von fünf mal fünf, der generell mit der Bezugsnummer 804 gekennzeichnet ist). Wie in den 8c und 8d dargestellt, arbeitet der Schablonenprozessor an dem Blatt der eingegebenen Bilddaten durch effektives Bewegen der überlappenden Schablonen 802 in einer Weise von links nach rechts über das Blatt. Bezüglich 8d ist die Bildpunktanzahl, für die ein Ausgabewert berechnet werden könnte (neun in einer abgedunkelten 3 mal 3-Matrix) von den Daten innerhalb des Blattes erschöpft (keine anderen Bildpunktpositionen können einen Ausgabewert aufweisen, der von den Informationen innerhalb des Blattes bestimmt wird). Zur Vereinfachung wurden die Randbereiche des Bildes ignoriert.
Wie in 8e ersichtlich, stellt der Blattgenerator dann ein nächstes Blatt 805 für den Schablonenprozessor bereit, sodass er die Operationen daran fortsetzt. Zu beachten ist, dass die Anfangspositionen der Schablonen der nächsten Progression vom Erschöpfungspunkt nach rechts auf dem ersten Blatt entsprechen (wie zuvor in 8d dargestellt), wenn sie mit der Operation am nächsten Blatt beginnen. Mit dem neuen Blatt 805 bewegen sich die Schablonen einfach weiter nach rechts, während der Schablonenprozessor an dem neuen Blatt auf die gleiche Weise arbeitet wie bei der Verarbeitung des ersten Blattes.
Zu beachten ist, dass zwischen den Daten des ersten Blattes 804 und den Daten des zweiten Blattes 805 aufgrund der Randbereiche der Schablonen, die einen Ausgabebildpunktort umgeben, eine gewisse Überlappung vorhanden ist. Die Überlappung könnte einfach gehandhabt werden, indem der Blattgenerator die überlappenden Daten zweimal überträgt. Bei alternativen Implementierungen kann, um dem Schablonenprozessor ein nächstes Blatt zuzuführen, der Blattgenerator damit fortfahren, ausschließlich neue Daten an den Schablonenprozessor zu senden, und der Schablonenprozessor verwendet die überlappenden Daten aus dem vorhergehenden Blatt erneut.
b. Schablonenprozessordesign und -betrieb
9a zeigt eine Ausführungsform einer Schablonenprozessoreinheitarchitektur 900. Wie in 9a dargestellt, beinhaltet der Schablonenprozessor eine Datenberechnungseinheit 901, einen Skalarprozessor 902 und einen zugehörigen Speicher 903, sowie eine E-/A-Einheit 904. Die Datenberechnungseinheit 901 beinhaltet eine Matrix von Ausführungsbahnen 905, eine zweidimensionale Schiebematrixstruktur 906 und getrennte jeweilige Direktzugriffsspeicher 907, die mit bestimmten Zeilen oder Spalten der Matrix verbunden sind.
Die E-/A-Einheit 904 ist verantwortlich für das Laden von „eingegebenen“ Datenblättern, die von dem Blattgenerator empfangen wurden, in die Datenberechnungseinheit 901, sowie das Speichern der von dem Schablonenprozessor „ausgegebenen“ Datenblätter in den Blattgenerator. Bei einer Ausführungsform erfordert das Laden von Blattdaten in die Datenberechnungseinheit 901 das Parsen eines empfangenen Blattes in Zeilen/Spalten von Bilddaten sowie das Laden der Zeilen/Spalten von Bilddaten in die zweidimensionale Schieberegisterstruktur 906 oder in die jeweiligen Direktzugriffsspeicher 907 der Zeilen/Spalten der Ausführungsbahnmatrix (wie nachfolgend näher beschrieben). Wird das Blatt anfänglich in die Speicher 907 geladen, können die einzelnen Ausführungsbahnen innerhalb der Ausführungsbahnmatrix 905 dann die Blattdaten, sofern geeignet (z. B. als Ladebefehl kurz vor der Operation der Blattdaten), in die zweidimensionale Schieberegisterstruktur 906 der Direktzugriffsspeicher 907 laden. Nach Beendigung des Ladens eines Blattes in die Registerstruktur 906 (ob direkt aus einem Blattgenerator oder aus den Speichern 907) arbeiten die Ausführungsbahnen der Ausführungsbahnmatrix 905 an den Daten und „schreiben“ letztendlich die fertigen Daten als ein Blatt direkt „zurück“ in den Blattgenerator oder in die Direktzugriffsspeicher 907. Wenn die Ausführungsbahnen in die Random Access Memories 907 zurückschreiben, ruft die E/A-Einheit 904 die Daten von den Random Access Memories 907 ab, um ein Ausgabeblatt zu bilden, das dann an den Blattgenerator weitergeleitet wird.
Der Skalarprozessor 902 beinhaltet einen Programmcontroller 909, der die Befehle des Programmcodes des Schablonenprozessors aus dem Skalarspeicher 903 liest und die Befehle an die Ausführungsbahnen in der Ausführungsbahnmatrix 905 ausgibt. Bei einer Ausführungsform wird ein einzelner gleicher Befehl auf alle Ausführungsbahnen innerhalb der Matrix 905 übertragen, um ein SIMD-(Single Instruction Multiple Data)-ähnliches Verhalten der Datenberechnungseinheit 901 zu bewirken. Bei einer Ausführungsform beinhaltet das Befehlsformat der Befehle, die aus dem Skalarspeicher 903 gelesen und an die Ausführungsbahnen der Ausführungsbahnmatrix 905 ausgegeben werden, ein sehr langes Befehlswortformat (VLIW), welches mehr als einen Operationscode pro Befehl beinhaltet. Bei einer weiteren Ausführungsform beinhaltet das VLIW-Format sowohl einen ALU-Operationscode, der eine mathematische Funktion anweist, die von der ALU einer Ausführungsbahn ausgeführt wird (wobei, wie nachstehend beschrieben, in einer Ausführungsform mehr als eine herkömmliche ALU-Operation spezifizieren kann), als auch einen Speicheroperationscode (der eine Speicheroperation für eine spezifische Ausführungsbahn oder eine Gruppe von Ausführungsbahnen anweist).
Der Begriff „Ausführungsbahn“ bezieht sich auf eine Gruppe von einer oder mehreren Ausführungseinheiten, die einen Befehl ausführen können (z. B. Logikschaltungen, die einen Befehl ausführen können). Eine Ausführungsbahn kann bei verschiedenen Ausführungsformen jedoch prozessorähnlichere Funktionalität als nur Ausführungseinheiten beinhalten. Beispielsweise kann eine Ausführungsbahn neben einer oder mehreren Ausführungseinheiten auch Logikschaltungen beinhalten, die einen empfangenen Befehl decodieren, oder im Fall von mehr MIMD-ähnlichen Designs Logikschaltungen beinhalten, die einen Befehl abrufen und decodieren. In Bezug auf MIMD-ähnliche Ansätze kann, obwohl ein zentraler Programmsteuerungsansatz hier weitgehend beschrieben wurde, auch ein verteilterer Ansatz in verschiedenen alternativen Ausführungsformen (z. B. unter anderem auch einschließlich Programmcode und eines Programmcontrollers innerhalb jeder Ausführungsbahn der Matrix 905) implementiert werden.
Die Kombination einer Ausführungsbahnmatrix 905, eines Programmcontrollers 909 und einer zweidimensionalen Schieberegisterstruktur 906 stellt eine weitgehend anpassbare/konfigurierbare Hardware-Plattform für ein breites Spektrum programmierbarer Funktionen bereit. Beispielsweise können Anwendungssoftwareentwickler in der Lage sein, Kernels mit einem breiten Spektrum unterschiedlicher Funktionsfähigkeiten sowie Dimensionen (z. B. Schablonengrößen) zu programmieren, da die einzelnen Ausführungsbahnen in der Lage sind, eine breite Palette von Funktionen auszuführen und ohne Weiteres auf Eingabebilddaten in der Nähe eines beliebigen Ausgabematrixorts zuzugreifen.
Abgesehen davon, dass die Random Access Memories 907 als ein Datenspeicher für Bilddaten agieren, die durch die Ausführungsbahnmatrix 905 bearbeitet werden, können diese auch eine oder mehrere Nachschlagetabellen halt wie irgendwelche Nachschlagetabellen, die in der Nachschlagetabellenkomponente des im Abschnitt 1.0 vorstehend beschriebenen virtuellen Verarbeitungsspeichers gehalten werden. Bei verschiedenen Ausführungsformen können eine oder mehrere skalare Nachschlagetabellen auch innerhalb des skalaren Speichers 903 instanziiert werden. Die eine oder die mehreren Skalarnachschlagetabellen können irgendwelche Skalarnachschlagetabellen sein, die in der Nachschlagetabellenskalarkomponente des in Abschnitt 1.0 vorstehend beschriebenen Speichermodells beinhaltet sind.
Ein skalarer Nachschlagevorgang beinhaltet das Weitergeben des gleichen Datenwertes aus der gleichen Nachschlagetabelle von dem gleichen Index an sämtliche Ausführungsbahnen innerhalb der Ausführungsbahnmatrix 905. Bei verschiedenen Ausführungsformen wird das oben beschriebene VLIW-Befehlsformat erweitert, um auch einen skalaren Operationscode einzuschließen, der eine vom Skalarprozessor ausgeführte Nachschlageoperation in eine skalare Nachschlagetabelle leitet. Der für die Verwendung mit dem Operationscode angegebene Index kann ein unmittelbarer Operand sein oder von einem anderen Datenspeicherort abgerufen werden. Unabhängig davon umfasst in einer Ausführungsform ein Nachschlagevorgang in einer skalaren Nachschlagetabelle innerhalb des skalaren Speichers im Wesentlichen das Senden des gleichen Datenwertes an alle Ausführungsbahnen innerhalb der Ausführungsbahnmatrix 905 während des gleichen Taktzyklus. Weitere Details zur Verwendung und Operation von Nachschlagetabellen werden nachfolgend bereitgestellt.
9b fasst die oben beschriebene bzw. beschriebenen VLIW-Befehlswort-Ausführungsformen zusammen. Wie in 9b dargestellt, beinhaltet das VLIW-Befehlswortformat Felder für drei separate Befehle: 1) einen Skalarbefehl 951, der durch den Skalarprozessor ausgeführt wird; 2) einen ALU-Befehl 952, der von den jeweiligen ALUs innerhalb der Ausführungsbahnmatrix in SIMD-Weise gesendet und ausgeführt wird; und 3) einen Speicherbefehl 953, der auf partielle SIMD-Weise gesendet und ausgeführt wird (z. B. wenn Ausführungsbahnen entlang derselben Zeile in der Ausführungsbahnmatrix einen gleichen Direktzugriffsspeicher teilen, anschließend wird der Befehl von einer Ausführungsbahn aus jeder der verschiedenen Zeilen tatsächlich ausgeführt (das Format des Speicherbefehls 953 kann einen Operanden beinhalten, der identifiziert, welche Ausführungsbahn von jeder Zeile den Befehl ausführt).
Ein Feld 954 für einen oder mehrere unmittelbare Operanden ist ebenfalls enthalten. Welche der Befehle 951, 952, 953 welche unmittelbaren Operandeninformationen verwenden, kann im Befehlsformat identifiziert sein. Jeder der Befehle 951, 952, 953 beinhaltet zudem seine eigenen Eingabeoperanden sowie resultierende Informationen (z. B. lokale Register für ALU-Operationen und ein lokales Register sowie eine Speicheradresse für Speicherzugriffsbefehle). Bei einer Ausführungsform wird der Skalarbefehl 951 durch den Skalarprozessor ausgeführt, bevor die Ausführungsbahnen innerhalb der Ausführungsbahnmatrix einen der anderen Befehle 952, 953 ausführen. Das heißt, die Ausführung des VLIW-Wortes beinhaltet einen ersten Zyklus, bei dem der Skalarbefehl 951 ausgeführt wird, gefolgt von einem zweiten Zyklus, bei dem die anderen Befehle 952, 953 ausgeführt werden können (es ist zu beachten, dass bei verschiedenen Ausführungsformen die Befehle 952 und 953 parallel ausgeführt werden können).
Bei einer Ausführungsform beinhalten die Skalarbefehle, die von dem Skalarprozessor 902 ausgeführt werden, Befehle, die an den Blattgenerator 703 ausgegeben werden, um Blätter von/in die Speicher oder 2D-Schieberegister 906 der Datenberechnungseinheit 901 zu laden/zu speichern. Hier kann der Betrieb des Blattgenerators von der Operation der Zeilenpuffereinheit 701 oder von anderen Variablen abhängig sein, die ein Vorlaufzeitverständnis der Anzahl an Zyklen verhindern, die der Blattgenerator 703 benötigen wird, um einen Befehl zu beenden, der von dem Skalarprozessor 902 ausgegeben wurde. Dementsprechend beinhaltet bei einer Ausführungsform jedes beliebige VLIW-Wort, dessen Skalarbefehl 951 dem Blattgenerator 703 entspricht oder anderweitig bewirkt, dass ein Befehl an diesen ausgegeben wird, auch Nicht-Operations-(NOOP)-Befehle in den anderen zwei Befehlsfeldern 952, 953 beinhaltet. Der Programmcode tritt dann in eine Schleife von NOOP-Befehlen für die Befehlsfelder 952, 953 ein, bis der Blattgenerator sein Laden von/Speichern unter der Datenberechnungseinheit beendet. Hier kann bei der Ausgabe eines Befehls an den Blattgenerator der Skalarprozessor ein Bit eines Verriegelungsregisters setzen, das der Blattgenerator nach Beendigung des Befehls zurücksetzt. Während der NOOP-Schleife überwacht der Skalarprozessor das Bit des Verriegelungsbits. Sobald der Skalarprozessor erkennt, dass der Blattgenerator seinen Befehl beendet hat, beginnt die normale Ausführung erneut.
10 zeigt eine Ausführungsform einer Datenberechnungseinheit 1001. Wie in 10 dargestellt, beinhaltet die Datenberechnungskomponente 1001 eine Matrix von Ausführungsbahnen 1005, die logisch „oberhalb“ einer zweidimensionalen Schieberegistermatrixstruktur 1006 positioniert sind. Wie vorstehend beschrieben, wird bei verschiedenen Ausführungsformen ein von einem Blattgenerator bereitgestelltes Bilddatenblatt in das zweidimensionale Schieberegister 1006 geladen. Die Ausführungsbahnen arbeiten dann an den Blattdaten aus der Registerstruktur 1006.
Die Ausführungsbahnmatrix 1005 und die Schieberegisterstruktur 1006 sind in Bezug zueinander in der Position fixiert. Die Daten innerhalb der Schieberegistermatrix 1006 schieben sich jedoch in einer strategischen und koordinierten Weise, um zu bewirken, dass jede Ausführungsbahn in der Ausführungsbahnmatrix eine andere Schablone innerhalb der Daten verarbeitet. Demgemäß bestimmt jede Ausführungsbahn den Ausgabebildwert für einen anderen Bildpunkt in dem generierten Ausgabeblatt. Aus der Architektur von 10 sollte hervorgehen, dass überlappende Schablonen nicht nur vertikal, sondern auch horizontal angeordnet sind, da die Ausführungsbahnmatrix 1005 vertikal angrenzende Ausführungsbahnen als auch horizontal angrenzende Ausführungsbahnen aufweist.
Einige zu beachtende architektonische Merkmale der Datenberechnungseinheit 1001 beinhalten die Schieberegisterstruktur 1006, die breitere Dimensionen als die Ausführungsbahnmatrix 1005 aufweist. Das heißt, es gibt einen „Ring“ von Registern 1009 außerhalb der Ausführungsbahnmatrix 1005. Obwohl der Ring 1009 auf zwei Seiten der Ausführungsbahnmatrix dargestellt ist, kann der Ring je nach Implementierung auf weniger (einer) oder mehr (drei oder vier) Seiten der Ausführungsbahnmatrix 1005 existieren. Der Ring 1005 dient dazu, einen „Überlaufraum“ für Daten bereitzustellen, die außerhalb der Grenzen der Ausführungsbahnmatrix 1005 überlaufen, wenn die Daten „unterhalb“ der Ausführungsbahnen 1005 geschoben werden. In einem einfachen Fall benötigt eine 5 × 5-Schablone, die am rechten Rand der Ausführungsbahnmatrix 1005 zentriert ist, vier Ringregisterorte weiter nach rechts, wenn die linksseitigen Bildpunkte der Schablone verarbeitet werden. Zur Vereinfachung der Zeichnung zeigt 10 die Register der rechten Seite des Rings mit nur horizontalen Schiebeverbindungen und Register der Unterseite des Rings mit nur vertikalen Schiebeverbindungen, wenn in einer nominalen Ausführungsform die Register auf beiden Seiten (rechts, unten) sowohl horizontale als auch vertikale Verbindungen aufweisen würden.
Zusätzlicher Überlaufraum wird durch Direktzugriffsspeicher 1007 bereitgestellt, die mit jeder Zeile und/oder jeder Spalte in der Matrix oder Teilen davon gekoppelt sind (z. B. kann ein Direktzugriffsspeicher einer „Region“ der Ausführungsbahnmatrix zugewiesen werden, die reihenweise 4 Ausführungsbahnen und spaltenweise 2 Ausführungsbahnen überspannt. Zur Vereinfachung bezieht sich der Rest der Anwendung hauptsächlich auf zeilen- und/oder spaltenbasierte Zuordnungsschemen). Wenn die Kerneloperationen einer Ausführungsbahn es erfordern, Bildpunktwerte außerhalb der zweidimensionalen Schieberegistermatrix 1006 zu verarbeiten (was einige Bildverarbeitungsroutinen ggf. erfordern), kann die Ebene der Bilddaten z. B. vom Ringbereich 1009 in den Direktzugriffsspeicher 1007 weiter überlaufen. Betrachten wir zum Beispiel eine 6X6-Schablone, bei der die Hardware einen Ringbereich von nur vier Speicherelementen rechts von einer Ausführungsbahn am rechten Rand der Ausführungsbahnmatrix beinhaltet. In diesem Fall müssten die Daten vom rechten Rand des Rings 1009 weiter nach rechts geschoben werden, um die Schablone vollständig zu verarbeiten. Daten, die außerhalb des Ringbereichs 1009 geschoben werden, würden dann in den Direktzugriffsspeicher 1007 überlaufen. Andere Anwendungen der Direktzugriffsspeicher 1007 und des Schablonenprozessors aus 3 werden nachfolgend bereitgestellt.
Die 11a bis 11k zeigen ein Ausführungsbeispiel für die Art und Weise, in der Bilddaten wie vorstehend angedeutet innerhalb der zweidimensionalen Schieberegistermatrix „unterhalb“ der Ausführungsbahnmatrix geschoben werden. Wie in 11a dargestellt, sind die Dateninhalte der zweidimensionalen Schiebematrix in einer ersten Matrix 1107 dargestellt, während die Ausführungsbahnmatrix durch einen Frame 1105 dargestellt ist. Außerdem sind innerhalb der Ausführungsbahnmatrix vereinfacht zwei angrenzende Ausführungsbahnen 1110 dargestellt. In dieser vereinfachten Darstellung 1110 beinhaltet jede Ausführungsbahn ein Register R1, das Daten aus dem Schieberegister akzeptieren kann, Daten von einer ALU-Ausgabe akzeptieren kann (z. B. sich über Zyklen als Akkumulator verhalten) oder Ausgabedaten in ein Ausgabeziel schreiben kann.
Jede Ausführungsbahn verfügt zudem in einem lokalen Register R2 über den Inhalt „darunter“ in der zweidimensionalen Schiebematrix. Somit ist R1 ein physisches Register der Ausführungsbahn, während R2 ein physisches Register der zweidimensionalen Schieberegistermatrix ist. Die Ausführungsbahn beinhaltet eine ALU, die mit Operanden arbeiten kann, die von R1 und/oder R2 bereitgestellt werden. Wie weiter nachfolgend näher beschrieben wird, wird bei einer Ausführungsform das Schieberegister tatsächlich mit mehreren (einer „Tiefe“ von) Speicher-/Registerelementen pro Matrixort implementiert, die Schiebeaktivität ist jedoch auf eine Ebene von Speicherelementen begrenzt (z. B. kann nur eine Ebene von Speicherelementen pro Zyklus geschoben werden). Die 11a bis 11k stellen einen dieser tieferen Registerorte dar, wie sie verwendet werden, um das resultierende X aus den jeweiligen Ausführungsbahnen zu speichern. Zur Veranschaulichung ist das tiefere resultierende Register neben anstatt unter dessen Gegenstückregister R2 gezeichnet.
Die 11a bis 11k konzentrieren sich auf die Berechnung von zwei Schablonen, deren zentrale Position mit dem Paar von Ausführungsbahnpositionen 1111 ausgerichtet ist, die in der Ausführungsbahnmatrix 1105 dargestellt sind. Zur Vereinfachung der Darstellung ist das Paar von Ausführungsbahnen 1110 als horizontal aneinander angrenzend gezeichnet, obgleich sie nach dem folgenden Beispiel tatsächlich vertikal aneinander angrenzend sind.
Wie anfangs in 11a ersichtlich, sind die Ausführungsbahnen 1111 auf ihren zentralen Schablonenorten zentriert. 11b zeigt den Objektcode, der von den beiden Ausführungsbahnen 1111 ausgeführt wird. Wie in 11b dargestellt, bewirkt der Programmcode beider Ausführungsbahnen 1111, dass die Daten innerhalb der Schieberegistermatrix 1107 eine Position nach unten und eine Position nach rechts geschoben werden. Damit werden beide Ausführungsbahnen 1111 an der oberen linken Ecke ihrer jeweiligen Schablonen ausgerichtet. Der Programmcode bewirkt dann, dass die Daten (in R2) an ihren jeweiligen Orten in R1 geladen werden.
Wie in 11c ersichtlich, bewirkt der Programmcode als Nächstes, dass das Paar von Ausführungsbahnen 1111 die Daten innerhalb der Schieberegistermatrix 1107 um eine Einheit nach links schiebt, wodurch bewirkt wird, dass der Wert rechts von der jeweiligen Position jeder Ausführungsbahn in die Position jeder Ausführungsbahn geschoben wird. Der Wert in R1 (vorheriger Wert) wird dann mit dem neuen Wert addiert, der in die Position der Ausführungsbahn (in R2) geschoben wurde. Das Ergebnis wird in R1 geschrieben. Wie in 11d ersichtlich, wird der gleiche Prozess wie vorstehend für 11c beschrieben wiederholt, wodurch bewirkt wird, dass die Resultante R1 jetzt den Wert A + B + C in der oberen Ausführungsbahn und F + G + H in der unteren Ausführungsbahn beinhaltet. An diesem Punkt haben beide Ausführungsbahnen 1111 die obere Reihe ihrer jeweiligen Schablonen verarbeitet. Zu beachten ist der Überlauf in einen Ringbereich auf der linken Seite der Ausführungsbahnmatrix 1105 (falls einer auf der linken Seite existiert) oder in einen Direktzugriffsspeicher, wenn auf der linken Seite der Ausführungsbahnmatrix 1105 kein Ringbereich existiert.
Wie in 11e dargestellt, bewirkt der Programmcode als Nächstes, dass sich Daten innerhalb der Schieberegistermatrix eine Einheit nach oben geschoben werden, wodurch bewirkt wird, dass beide Ausführungsbahnen 1111 am rechten Rand der mittleren Reihe ihrer jeweiligen Schablonen ausgerichtet werden. Register R1 der beiden Ausführungsbahnen 1111 beinhaltet derzeit die Summe der oberen Zeile der Schablone und den Wert ganz rechts der mittleren Zeile. Die 11f und 11g zeigen weitere Fortschritte beim Bewegen nach links über die mittlere Reihe der Schablonen der beiden Ausführungsbahnen. Die kumulative Addition setzt sich derart fort, dass am Ende der Verarbeitung von 11g beide Ausführungsbahnen 1111 die Summe der Werte der obersten Reihe und der mittleren Reihe ihrer jeweiligen Schablonen beinhalten.
11h zeigt einen weiteren Schiebevorgang, um jede Ausführungsbahn an der untersten Zeile ihrer entsprechenden Schablonen auszurichten. Die 11i und 11j zeigen einen fortlaufenden Schiebvorgang zum Abschließen der Verarbeitung über den Verlauf der Schablonen der beiden Ausführungsbahnen. 11k zeigt einen zusätzlichen Schiebevorgang, um jede Ausführungsbahn an ihrer korrekten Position in der Datenmatrix auszurichten und die Resultante dort hinein zu schreiben.
In dem Beispiel der 11a bis 11k ist zu beachten, dass der Objektcode für die Schiebeoperationen ein Befehlsformat beinhalten kann, das die Richtung und die Größe des in (X, Y) Koordinaten ausgedrückten Schiebevorgangs identifiziert. Der Objektcode für einen Schiebevorgang um einen Ort nach oben kann in Objektcode beispielsweise als SHIFT 0, +1 ausgedrückt werden. Als weiteres Beispiel kann einen Schiebevorgang um einen Ort nach rechts im Objektcode als SHIFT +1, 0 ausgedrückt werden. Darüber hinaus können in verschiedenen Ausführungsformen Schiebevorgänge größerer Größe in Objektcode (z. B. SHIFT 0, +2) spezifiziert werden. Wenn die 2D-Schieberegister-Hardware hier nur Verschiebungen um einen Ort pro Zyklus unterstützt, kann der Befehl durch die Maschine so interpretiert werden, dass eine Mehrfachzyklusausführung erforderlich ist, oder die 2D-Schieberegister-Hardware kann so ausgelegt sein, dass sie Schiebevorgänge um mehr als einen Ort pro Zyklus unterstützt. Ausführungsformen der letzteren Variante werden nachfolgend näher beschrieben.
12 zeigt eine weitere, detailliertere Darstellung der Einheitszelle für die Ausführungsbahnmatrix und Schieberegisterstruktur (Register im Ringbereich beinhalten keine entsprechende Ausführungsbahn). Die Ausführungsbahn und der Registerraum, der mit jedem Ort in der Ausführungsbahnmatrix verknüpft ist, sind bei einer Ausführungsform implementiert, indem die in 12 ersichtlichen Schaltungen an jedem Knoten der Ausführungsbahnmatrix instanziiert wird. Wie in 12 ersichtlich, beinhaltet die Einheitszelle eine Ausführungsbahn 1201, die mit einer Registerdatei 1202 gekoppelt ist, die aus vier Registern R2 bis R5 besteht. Während eines Zyklus kann die Ausführungsbahn 1201 von einem Register R1 bis R5 lesen oder darin schreiben. Für Befehle, die zwei Eingabeoperanden erfordern, kann die Ausführungsbahn beide Operanden von irgendeinem von R1 bis R5 abrufen.
Bei einer Ausführungsform ist die zweidimensionale Schieberegisterstruktur implementiert, indem sie während eines einzigen Zyklus erlaubt, dass der Inhalt eines beliebigen von (nur) einem der Register R2 bis R4 zu einer seiner angrenzenden Registerdateien durch den Ausgabemultiplexer 1203 „heraus“-geschoben wird, und der Inhalt eines beliebigen von (nur) einem der Register R2 bis R4 durch den Inhalt ersetzt wird, der von einem entsprechenden Nachbarn durch den Eingabemultiplexer 1204 „hinein“-geschoben wird, sodass Verschiebungen zwischen den Nachbarn in gleicher Richtung (z. B. alle Ausführungsbahnen nach links, alle Ausführungsbahnen nach rechts usw.) erfolgen. Obwohl es für ein gleiches Register üblich sein kann, dass sein Inhalt herausgeschoben und durch den Inhalt ersetzt wird, der in dem gleichen Zyklus hereingeschoben wird, erlaubt die Multiplexeranordnung 1203, 1204 unterschiedliche Schiebequellen- und Schiebezielregister innerhalb einer gleichen Registerdatei während eines gleichen Zyklus.
Wie in 12 dargestellt, ist zu beachten, dass während einer Schiebesequenz eine Ausführungsbahn den Inhalt aus ihrer Registerdatei 1202 zu deren jeweiligen linken, rechten, oberen und unteren Nachbarn schiebt. In Verbindung mit der gleichen Schiebesequenz schiebt die Ausführungsbahn zudem den Inhalt in deren Registerdatei von einem bestimmten seiner linken, rechten, oberen und unteren Nachbarn. Das Herausschiebeziel und die Hereinschiebequelle sollten wiederum mit einer gleichen Schieberichtung für alle Ausführungsbahnen übereinstimmen (wenn z. B. das Herausschieben zum rechten Nachbarn erfolgt, sollte das Hereinschieben vom linken Nachbarn erfolgen).
Obwohl in einer Ausführungsform der Inhalt von nur einem Register pro Ausführungsbahn pro Zyklus geschoben werden darf, können andere Ausführungsformen zulassen, dass der Inhalt von mehr als einem Register hinein-/herausgeschoben wird. Beispielsweise kann der Inhalt von zwei Registern während eines gleichen Zyklus hinaus/herein geschoben werden, wenn eine zweite Instanz der in 12 ersichtlichen Multiplexerschaltung 1203, 1204 in das Design von 12 eingefügt wird. Selbstverständlich können bei Ausführungsformen, bei denen der Inhalt von nur einem Register pro Zyklus geschoben werden kann, Schieben von mehreren Registern zwischen mathematischen Operationen stattfinden, indem mehr Taktzyklen für Schieben zwischen mathematischen Operationen verbraucht werden (z. B. kann der Inhalt von zwei Registern zwischen mathematischen Operationen geschoben werden, indem zwei Schiebeoperationen zwischen den mathematischen Operationen verbraucht werden).
Falls weniger als der gesamte Inhalt der Registerdateien einer Ausführungsbahn während einer Schiebesequenz geschoben wird, ist zu beachten, dass der Inhalt der nicht geschobenen Register jeder Ausführungsbahn an Ort und Stelle bleibt (nicht geschoben wird). Dementsprechend bleibt jeder nicht geschobene Inhalt, der nicht durch hereingeschobenen Inhalt ersetzt wird, über den Schiebezyklus hinweg lokal auf der Ausführungsbahn. Die in jeder Ausführungsbahn beobachtete Speichereinheit („M“) wird verwendet, um Daten von dem/in den Direktzugriffsspeicherraum zu laden/zu speichern, der mit der Zeile und/oder Spalte der Ausführungsbahn innerhalb der Ausführungsbahnmatrix verbunden ist. Hier fungiert die M-Einheit als Standard-M-Einheit, indem sie häufig zum Laden/Speichern von Daten verwendet wird, die nicht von/in den eigenen Registerraum der Ausführungsbahn geladen/gespeichert werden können. Bei verschiedenen Ausführungsformen besteht die primäre Operation der M-Einheit darin, Daten von einem lokalen Register in den Speicher zu schreiben und Daten aus dem Speicher zu lesen und in ein lokales Register zu schreiben.
In Bezug auf die von der ALU-Einheit der Hardware-Ausführungsbahn 1201 unterstützten ISA-Operationscodes sind in verschiedenen Ausführungsformen die von der Hardware-ALU unterstützten mathematischen Operationscodes ganz (d. h. im Wesentlichen gleich) mit den von den unterstützten mathematischen Operationscodes verbunden, die eine virtuelle Ausführungsbahn unterstützen (z. B. ADD, SUB, MOV, MUL, MAD, ABS, DIV, SHL, SHR, MIN/MAX, SEL, ODER, XOR, NOT). Wie vorstehend beschrieben, können Speicherzugriffsbefehle von der Ausführungsbahn 1201 ausgeführt werden, um Daten von/in ihrem zugehörigen Direktzugriffsspeicher abzurufen/zu speichern. Zudem unterstützt die Hardware-Ausführungsbahn 1201 Schiebeoperationsbefehle (nach rechts, links, oben, unten), um Daten innerhalb der zweidimensionalen Schieberegisterstruktur zu schieben. Wie vorstehend beschrieben, werden Programmsteuerbefehle weitgehend durch den Skalarprozessor des Schablonenprozessors ausgeführt.
3.0 In einem Bildprozessor implementierte Blockoperationen
a. Zeilen-/Spaltensummieroperation
Die 13 und 14a, b betreffen eine Zeilen-/Spaltensummieroperation. Zeilen-/Spaltensummieroperationen sind besonders für Statistikrechnung, Rauschunterdrückung und Grossformat-Downsampling nützlich. Zeilensummen- und/oder Spaltensummieroperationen können in einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und entsprechender zweidimensionaler Schieberegistermatrix implementiert werden, wie beispielsweise Ausführungsformen des in den vorhergehenden Abschnitten beschriebenen Bildprozessors.
Wie in 13 ersichtlich, fügt eine Zeilensummieroperation alle Werte in einer gleichen Zeile einer Matrix hinzu und eine Spaltensummieroperation fügt alle Werte in einer gleichen Spalte einer Matrix hinzu. Zusätzlich können mit einer Zeilensummieroperation die Werte von allen Zeilen wie dargestellt innerhalb ihrer entsprechenden Zeilen gleichzeitig hinzugefügt werden. Desgleichen können mit einer Spaltensummieroperation die Werte von allen Spalten innerhalb ihrer entsprechenden Spalten gleichzeitig hinzugefügt werden. Aufgrund der Vielseitigkeit der Schieberegistermatrix ist das Summieren über alle Zeilen oder Spalten jedoch keine Voraussetzung. D. h., dass weniger als alle Zeilen in einer Matrix gleichzeitig summiert werden können, oder dass weniger als alle Spalten in einer Matrix gleichzeitig summiert werden können.
Die 14a bis 14d zeigen eine Ausführungsform einer Maschinenebenenoperation zum Implementieren einer Zeilensummieroperation. Der Einfachheit halber ist nur eine einzelne Zeile gezeigt. Der Leser wird verstehen, dass die in den 14a bis 14d dargestellten Operationen auch auf eine Spalte angewendet werden können. Zusätzlich kann für beide Zeilen- oder Spaltenoperationen die Sequenz der 14a bis 14d auch gleichzeitig für mehrere Zeilen oder Spalten der Matrix ausgeführt werden. Die Dimension der Zeile ist nur als 8 Stellen breit gezeigt (wohingegen bei tatsächlichen Implementierungen die Ausführungsbahnmatrix und die Schieberegistermatrix 16 × 16 oder sogar größer sein können).
Bei einer Implementierung ist wie in 14a ersichtlich die zweidimensionale Schieberegistermatrix derart ausgelegt, dass sie direkte logische Schiebevorgänge zwischen gegenüberliegenden Enden der Matrix 1401 unterstützt. D. h., dass das Schieberegister seinen Inhalt zwischen einer Spaltenmatrix ganz rechts und einer Spaltenmatrix ganz links „rollen“ oder „schleifen“ oder „umbrechen“ kann, wenn Schiebevorgänge entlang einer Zeilenachse ausgeführt werden, und/oder seinen Inhalt zwischen einer obersten Zeilenmatrix und einer untersten Zeilenmatrix rollt oder schleift, wenn Schiebevorgänge entlang einer Spaltenachse ausgeführt werden. Bei verschiedenen Ausführungsformen kann das Schieberegister zusätzlich mehrere Registerplatzsprünge in einem einzelnen Befehl unterstützen (z. B. spezifiziert Operationscode und/oder eine Variable, die mit einem Schiebebefehl verbunden ist, ob der Schiebebetrag +/–1, +/–2, +/–3 oder +/–4 Registerorte entlang der horizontalen und/oder vertikalen Achse ist). Schiebeabstände, die von der Hardware nicht unterstützt werden, können durch den Compiler emuliert werden.
Wie in 14a ersichtlich wird die Zeile anfänglich mit Datenwerten A0 bis A7 an den entsprechenden Registerplätzen R0 und R1 von jeder Ausführungsbahn geladen. Wie in 14b ersichtlich, werden dann in einer ersten Iteration von Maschinenoperationen die R1-Registerraumorte um einen Ort nach links geschoben und die R1- und R0-Inhalte jeder Ausführungsbahn werden mit der in R0 und R1 zurückgeschriebenen Resultante summiert. Dies erzeugt eine erste akkumulierte Partialsumme in R1, die wie in der folgenden Erörterung eindeutiger ausgeführt als ein Akkumulator für die gesamte Additionsoperation agiert.
Bei einer nächsten Iteration von Maschinenoperationen, die in 14c dargestellt sind, werden die R1-Registerraumorte um zwei Orte nach links geschoben und die R1- und R0-Inhalte jeder Ausführungsbahn werden summiert. Die Resultante wird erneut in R0 und R1 gehalten. Letztendlich werden in einer dritten Iteration von Maschinenoperationen, die in 14d dargestellt sind, die R1-Registerraumorte um vier Orte nach links geschoben, und die Resultate werden in einen oder in beide Registerplätze R0 und R1 geschrieben. In Anbetracht der Iterationen der Summieroperation über die 14b, 14c und 14d ist daher zu beachten, dass sich der Schiebebetrag mit jeder Iteration verdoppelt, und dass das komplette Resultat nach nur drei Iterationen (d. h. nach 14c) in allen Zeilen vorhanden ist. Im Fall einer 16 breiten Zeile wäre der Schiebebetrag acht Orte für die vierte Iteration und die Operation wäre nach der vierten Iteration abgeschlossen.
b. Zeilen/Spalten-Präfixsummieroperation
Die 15 und 16a bis 16d betreffen eine Zeilen/Spalten-Präfixsummieroperation, die auch in einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einem entsprechenden zweidimensionalen Schieberegister ausgeführt werden kann. Zeilen/Spalten-Präfixsummieroperationen sind für ganzheitliche Bilder, Beschleunigen von Mittelwertfiltern und Berechnen von Adressen für Verdichtungsoperationen besonders nützlich. Unter Bezugnahme auf 15 ist der Einfachheit halber nur eine einzelne Zeile gezeigt. Die Operation kann jedoch auch für Spalten ausgeführt werden. Obwohl nur eine Zeile bei verschiedenen Implementierungen dargestellt ist, wie bei der unmittelbar zuvor beschriebenen Zeilen-/Spaltensummieroperation, kann außerdem jegliche Anzahl von Zeilen (oder Spalten) in der Registermatrix (einschließlich aller Zeilen/Spalten) gleichzeitig summiert werden. Zusätzlich ist die Dimension der Zeile in 15 nur als 8 Orte breit gezeigt, wohingegen bei tatsächlichen Implementierungen die Ausführungsbahn und die Schieberegistermatrizen 16 × 16 oder noch größer sein können.
Wie in 15 ersichtlich, ist die Resultante einer Zeilenpräfixoperation für jeden Ort innerhalb einer Zeile die Summe der Werte, die bis zu diesem Ort führen. Desgleichen ist der Wert einer Spaltenpräfixoperation für irgendeinen Ort innerhalb einer Spalte die Summe der Werte, die bis zu diesem Ort führen.
Die 16a bis 16d zeigen eine Ausführungsform einer Maschinenebenenoperation für eine Zeilenpräfixoperation. Wie bei der Zeilensummieroperation im vorhergehenden Abschnitt ausführlich beschrieben, verwendet die Zeilenpräfixoperation ein zweidimensionales Schieberegister, das Schiebevorgänge zwischen Matrixrandstellen 1601 rollen kann.
Wie in 16a ersichtlich, ist eine Zeile anfänglich mit Datenwerten A0 bis A7 im entsprechenden R0-Registerplatz jeder Ausführungsbahn geladen. Außerdem ist eine Null („0“) in den R2-Registerraum jeder Ausführungsbahn geladen.
Bei einer ersten Iteration einer Maschinenebenenoperation, die in 16b dargestellt ist, werden die R0-Registerraumorte um einen Ort nach rechts in den R1-Registerraum des Zielorts geschoben. Eine anschließende ADD-Operation fügt den R0-Inhalt abhängig von dem Ort der Bahn relativ zu der Iterationszahl entweder dem R1-Inhalt oder dem R2-Inhalt hinzu. Der erste Zeilenort wählt speziell die Null in R2 aus (und nicht den geschobenen Inhalt in R1), da sein Ort (0) gleich oder kleiner als 2^N – 1 ist, wobei N die Iterationszahl (2⁰ – 1 = 0) ist, addiert die Null zum Inhalt in R0 und speichert die Resultante zurück in R0. Der erste Zeilenort erhält daher einen Wert von A0 in R0 aufrecht.
Im Vergleich dazu wird jeder der anderen Zeilenorte, da ihr Ort größer als 2^N – 1 ist, den geschobenen Inhalt in R1 auswählen (anstatt den Nullwert in R2), ihn dem Inhalt in R0 hinzufügen und die Resultante in R0 speichern. Daher wird jeder andere Zeilenort wie der erste Zeilenort die Summe seines ursprünglichen Inhalts und seinen ganz linken Nachbarn in R0 halten, wohingegen der erste Zeilenort einfach nur seinen ursprünglichen Inhalt in R0 halten wird.
Bei einer zweiten Iteration von Maschinenebenenoperationen, die in 16c dargestellt sind, wird der Inhalt der R0-Registerraumorte um zwei Orte nach rechts in den R1-Registerraum des Ziels geschoben. Daher verdoppelt sich wie bei den unmittelbar vorstehend beschriebenen Zeilen/Säulenoperationen mit jeder nächsten Iteration von Maschinenoperationen der Schiebebetrag. Wie bei der ersten Iteration fügt eine anschließende ADD-Operation abhängig von dem Ort der Bahn relativ zur gegenwärtigen Iterationszahl den R0-Inhalt entweder dem R1-Inhalt oder dem R2-Inhalt hinzu. In diesem Fall wählen die ersten und zweiten Zeilenorte die Null in R2 aus (und nicht den geschobenen Inhalt in R1), da deren Ort (0, 1) kleiner oder gleich 2^N – 1 (2¹ – 1 = 1) ist. Diese Bahnen fügen daher den Nullwert dem Inhalt in R0 hinzu und speichern die Resultante zurück in R0. Der erste Zeilenort wird daher seinen Wert von A0 in R0 aufrechterhalten und der zweite Zeilenort wird einen Wert von A0 + A1 in R0 aufrechterhalten.
Im Vergleich dazu wird jeder der anderen Zeilenorte, da sein Ort größer als 2^N – 1 ist, den geschobenen Inhalt in R1 (anstatt den Nullwert in R2) auswählen, den Inhalt von R1 dem Inhalt in R0 hinzufügen und die Resultante in R0 speichern. Daher wird jeder andere Zeilenort als die ersten und zweiten Zeilenorte basierend auf seinem ursprünglichen Inhalt und seinem gesamten geschobenen Inhalt in R0 eine akkumulierte Summe halten.
Bei einer dritten Iteration von Maschinenebenenoperationen, die in 16d dargestellt sind, werden die R0-Registerraumorte um vier Orte nach rechts in den R1-Registerraum des Ziels geschoben. Daher verdoppelt sich erneut mit jeder nächsten Iteration von Maschinenoperationen der Schiebebetrag. Wie auch bei den früheren Iterationen fügt eine anschließende ADD-Operation abhängig von dem Ort der Bahn relativ zur gegenwärtigen Durchlaufzahl den R0-Inhalt entweder dem R1-Inhalt oder dem R2-Inhalt hinzu. In diesem Fall werden die ersten bis vierten Zeilenorte (0 bis 3) die Null in R2 auswählen (und nicht den geschobenen Inhalt in R1), da deren Ort kleiner oder gleich 2^N – 1 (2² – 1 = 3) ist. Jede dieser Bahnen fügt daher den Nullwert dem Inhalt in R0 hinzu und speichert die Resultante zurück in R0. Der erste Zeilenort wird daher seinen Anfangswert von A0 in R0 aufrechterhalten, der zweite Zeilenort wird einen Wert von A0 + A1 in R0 aufrechterhalten, der dritte Zeilenort wird einen Wert von A0 + A1 + A2 in R0 aufrechterhalten und der vierte Zeilenort wird einen Wert von A0 + A1 + A2 + A3 in R0 aufrechterhalten.
Im Vergleich dazu wird jeder der anderen Zeilenorte, da sein Ort größer als 2^N – 1 bleibt, den geschobenen Inhalt in R1 (anstatt den Nullwert in R2) auswählen, den Inhalt von R1 dem Inhalt in R0 hinzufügen und die Resultante in R0 speichern. Daher wird jeder andere Zeilenort als die ersten, zweiten und dritten Zeilenorte basierend auf seinem ursprünglichen Inhalt und seinem gesamten geschobenen Inhalt eine akkumulierte Summe halten.
Nach der dritten Iteration ist die Präfixsummieroperation abgeschlossen. Wenn die Reihe die Dimension Sechzehn hätte, wäre nur ein weiterer Operationssatz erforderlich, der den R1-Registerinhalt um acht Orte schiebt, um die Addition zu akkumulieren, die für alle 16 unterschiedlichen Orte in allen sechzehn Ausführungsbahnen einzigartig ist.
c. Zeilen-/Spaltenminimumfinden-Operation
Die 17 und 18a bis 18d betreffen eine Zeilen-/Spaltenminimumfinden-Operation, die auch auf einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einem entsprechenden zweidimensionalen Schieberegister ausgeführt werden kann. Zeilen-/Spaltenminimumfinden-Operationen sind für Statistikrechnung und Blockabgleichnachverarbeitung besonders nützlich.
Unter Bezugnahme auf 17 ist der Einfachheit halber nur eine einzelne Zeile gezeigt. Bei der praktischen Ausführung kann jedoch eine Spaltenminimumfinden-Operation auf einem gleichen Prozessor implementiert sein, der eine Zeilen Minimumfinden-Operation implementiert. Zusätzlich kann das Minimum für eine Zeile (oder Spalte) gleichzeitig für jegliche Anzahl von Zeilen (oder Spalten) in der Registermatrix gefunden werden (einschließlich bis zu allen Zeilen/Spalten). Zusätzlich ist die Dimension der Zeile/Spalte nur als 8 Orte breit gezeigt, wohingegen bei tatsächlichen Implementierungen die Ausführungsbahn und die Schieberegistermatrizen 16 × 16 oder noch größer sein können.
Wie in 17 ersichtlich entspricht die Resultante einer Zeilenminimumfinden-Operation dem kleinsten Wert unter allen Werten innerhalb einer gleichen Zeile und dessen Ort/Position (der auch als ihr Index bezeichnet wird) in der Zeile. Desgleichen entspricht die Resultante einer Spaltenminimumfinden-Operation dem kleinsten Wert unter allen Werten innerhalb einer gleichen Spalte und dessen Ort/Position innerhalb der Spalte. Wie bei der Zeilensummen- und Präfixsummieroperation in den vorhergehenden Abschnitten ausführlich beschrieben, verwendet die Zeilen/Spalten-Minimumfinden-Operation ein zweidimensionales Schieberegister, das Schiebevorgänge zwischen den Matrixrandstellen 1701 rollen kann.
Die 18a bis 18d zeigen eine Ausführungsform von Maschinenebenenoperationen für eine Zeilenpräfixoperation. Wie in 18a ersichtlich, ist eine Zeile anfänglich mit Dateninhalten A0 bis A7 am entsprechenden R0-Registerplatz jeder Ausführungsbahn geladen. Außerdem ist der Index jedes Zeilenorts in den R1-Registerraum jeder Ausführungsbahn geladen.
Bei einer ersten Iteration von Maschinenebenenoperationen, die in 18b dargestellt sind, wird der Inhalt der R0- und R1-Registerplätze um eine Einheit in den R2- und R3-Registerraum einer angrenzenden Ausführungsbahn geschoben. Die entsprechenden Werte der R0- und R2-Register werden dann innerhalb jeder Ausführungsbahn verglichen. Der Mindestwert des Vergleichs und sein entsprechender Index werden in den R0- und R1-Registerraum gespeichert. D. h., dass die R0- und R1-Register ihren ursprünglichen Inhalt aufrechterhalten, wenn der R0-Wert kleiner ist als der R2-Wert. Wohingegen, wenn der R2-Wert kleiner ist als der R0-Wert, der R2-Wert in R0 geschrieben wird und der R3-Wert in R1 geschrieben wird. Dies hat den Effekt des Haltens des Mindestwerts des Vergleichs in R0 und dessen Index in R1.
Bei einer zweiten Iteration von Maschinenebenenoperationen, die in 18c dargestellt sind, wird der Inhalt der R0- und R1-Registerplätze um zwei Einheiten in den R2- und R3-Registerraum einer nachgeschalteten Ausführungsbahn geschoben. Daher verdoppelt sich der Schiebebetrag erneut bei einer nächsten Iteration. Die entsprechenden Werte der R0- und R2-Register werden dann innerhalb jeder Ausführungsbahn verglichen. Der Mindestwert des Vergleichs und sein entsprechender Index werden in den R0- und R1-Registerraum gespeichert. Dies hat den Effekt des Haltens des beobachteten Mindestwerts von beiden Vergleichen, die in jeder Ausführungsbahn in R0 und ihrem Index in R1 ausgeführt wurden.
Bei einer dritten Iteration von Maschinenebenenoperationen, die in 18d ersichtlich sind, wird der Inhalt der R0- und R1-Registerplätze um vier Einheiten in den R2- und R3-Registerraum einer nachgeschalteten Ausführungsbahn geschoben. Die entsprechenden Werte der R0- und R2-Register werden dann innerhalb jeder Ausführungsbahn verglichen. Der Mindestwert des Vergleichs und sein entsprechender Index werden in den R0- und R1-Registerraum gespeichert. Dies hat den Effekt des Haltens des beobachteten Mindestwerts von allen drei Vergleichen, die in jeder Ausführungsbahn in R0 und ihrem Index in R1 ausgeführt wurden.
Nach der dritten Iteration ist die Minimumfinden-Operation abgeschlossen, da jede Ausführungsbahn den Mindestwert der gesamten Zeile in deren R0-Registerraum und dessen entsprechenden Index in deren R1-Raum aufweisen wird (der Zeilenort, der ursprünglich mit dem Mindestwert versehen wurde, wird seinen eigenen Zeilenort in seinem R1-Registerraum identifiziert finden). Wenn die Zeile die Dimension sechzehn hätte, dann wäre nur ein weiterer Satz an Operationen basierend auf einem Schieben des R0- und R1-Registerinhalts nachgeschaltet in das Schieberegister um acht Orte erforderlich, um das Minimum von allen sechzehn unterschiedlichen Orten in allen sechzehn Ausführungsbahnen bereitzustellen.
Zu beachten ist, dass eine Maximumfindenoperation auch unter Verwendung der gleichen vorstehend beschriebenen Prinzipien implementiert werden könnte, außer, dass die Kernrechenoperation das Maximumfinden anstatt des Minimumfindens umfasst.
d. Matrixmultplizieren
Die 19a, b und 20 betreffen Matrixmultiplizieren. Matrixmultiplizieren ist für diskrete Fourier- oder Kosinustransformationen (die gemeinsame Grundelemente beim Komprimieren und Filtern sind) und zum Ausdrücken größerer Matrix-/Multiplizieroperationen (die bei der Bildanalyse gewöhnlich verwendet werden) besonders nützlich. Ein Matrixmultiplizieren von zwei Matrizen A und B wird durch Summieren von Produkten der Elemente in der Zeile des Koordinatenorts und ihrer entsprechenden Elemente in der Spalte des Koordinatenorts für jeden Matrixkoordinatenort in der Resultante ausgeführt. 19b zeigt die resultierende Matrix X für das Matrixmultiplizieren der Matrizen A und B in 19a. Hier sind speziell zugehörige Teilprodukttermini der Matrizen A und B für zwei Koordinatenorte 1901, 1902 in der resultierenden Matrix C gezeigt.
Insbesondere ist zum Beispiel das resultierende Produkt für den Koordinatenort c12 in der resultierenden Matrix C von 19b (a11·b11) + (a12·b21) + (a13·b31) + (a14·b41) das den Elementen der Zeile 1901 in Matrix A entspricht, das mit den entsprechenden Elementen der Spalte 1902 der Matrix B in 19a multipliziert ist.
Desgleichen ist das resultierende Produkt für den Koordinatenort c22 in der resultierenden Matrix C von 19b (a21·b12) + (a22·b22) + (a23·b32) + (a24·b42) das den Elementen der Zeile 1903 in Matrix A entspricht, das mit den entsprechenden Elementen der Spalte 1902 der Matrix B in 19a multipliziert ist. Wie aus diesen zwei Beispielen zu ersehen ist, kann die Resultante für irgendeinen Koordinatenort x,y in der resultierenden Matrix C ausgedrückt werden als: ∑_{k=1 bis 4}(a_x,k)·(b_k,y).
Die 20a bis 20e zeigen eine Ausführungsform einer Matrixmultiplizieroperation, die eine zweidimensionale Ausführungsbahnmatrix und die einzigartigen Schiebevorteile einer entsprechenden zweidimensionalen Schieberegistermatrix verwendet.
Wie in 20a ersichtlich werden anfänglich zwei Matrizen A und B in die zweidimensionale Schieberegistermatrix geladen. Die Werte der Matrix A können beispielsweise in den R0-Registerraum des zweidimensionalen Schieberegisters geladen werden und die Werte der Matrix B können in den R1-Registerraum der zweidimensionalen Schieberegistermatrix geladen werden, sodass jede Matrixkoordinate dem zugehörigen Registerraum einer unterschiedlichen Ausführungsbahn entspricht. Erneut sind zur leichteren Veranschaulichung 4 × 4-Matrizen dargestellt, obwohl bei tatsächlichen Implementierungen größere Matrizen mit einer entsprechend größer dimensionierten Schieberegistermatrix multipliziert werden können.
Wie in 20b ersichtlich, wird ein Rotationsscheralgorithmusschiebevorgang sowohl auf Matrizen mit einer zeilenweisen Rotationsscheralgorithmusschiebesequenz auf Matrix A als auch auf eine spaltenweise Rotationsscheralgorithmusschiebesequenz auf Matrix B angewandt. Wie auf dem Fachgebiet bekannt erhöht ein Rotationsscheralgorithmus den Schiebebetrag um N – 1, wobei N die Position in der Matrix ist.
Wie in 20b ersichtlich, wird daher die erste Zeile der Matrix A überhaupt nicht geschoben, die zweite Zeile der Matrix A wird um eine Einheit geschoben, die dritte Zeile der Matrix A wird um zwei Einheiten geschoben und die vierte Zeile der Matrix A wird um drei Einheiten geschoben. Ähnlich wird die erste Spalte der Matrix B überhaupt nicht geschoben, die zweite Spalte der Matrix B wird um eine Einheit geschoben, die dritte Spalte der Matrix B wird um zwei Einheiten geschoben und die vierte Spalte der Matrix B wird um drei Einheiten geschoben. Hier wird die zweidimensionale Schieberegistermatrix als fähig verstanden, Elemente an Matrixrändern sowohl für zeilenorientierte Schiebevorgänge als auch für spaltenorientierte Schiebevorgänge zu rollen.
Zu beachten ist, dass infolge der vielseitigen Art der zweidimensionalen Schieberegistermatrix und der entsprechenden Ausführungsbahnmatrix, bei der der durch jede Ausführungsbahn ausgeführte entsprechende Schiebebefehl unterschiedliche Eingabeoperandendaten spezifizieren kann, die zweidimensionale Schieberegistermatrix fähig ist, unterschiedliche Zeilen durch unterschiedliche Horizontalschiebebeträge zu schieben und unterschiedliche Spalten durch unterschiedliche Vertikalschiebebeträge zu schieben, während die Verschiebebefehle über alle Ausführungsbahnen für eine gleiche Matrix gleichzeitig ausgeführt werden (bei verschiedenen Ausführungsformen muss ein gleicher Horizontalschiebebetrag für Bahnen in einer gleichen Zeile spezifiziert werden und ein gleicher Vertikalschiebebetrag für Bahnen in einer gleichen Spalte spezifiziert werden). Daher ist es denkbar, dass das Scheren sowohl der A- als auch der B-Matrizen in nur zwei Zyklen abgeschlossen werden kann (d. h., dass alle Schiebevorgänge für eine Matrix unter der Annahme, dass das Schieberegister mehrere Sprungschiebevorgänge in einem Einzelzyklus implementieren kann, in einem Zyklus ausgeführt werden).
Mit den in den Matrizen A und B von den Scheralgorithmen neu ausgerichteten Daten wird wie in 20b ersichtlich eine Multiplizieroperation ausgeführt, bei der jede Ausführungsbahn die A- und B-Werte in ihrem entsprechenden zweidimensionalen Schieberegisterraum multipliziert. Es sei hier daran erinnert, dass A-Werte z. B. im R0-Raum und B-Werte im R1-Raum gehalten werden. Die Resultante der Multiplikation wird in lokalem R2-Raum gespeichert. Nullwerte können als Anfangszustand in den R3-Raum geladen werden und die Resultante der Multiplikation in R2 wird zu den Inhalten von R3 hinzugefügt. Die Resultante der Addition wird in R3 zurückgespeichert. Aus den folgenden Lehren wird offensichtlich, dass R3 die Rolle eines Akkumulators annimmt, der die Addition von Teilprodukttermini im Verlauf der Matrixmultiplizieroperation akkumuliert.
Zum besseren Verständnis zeigt 20b explizit die Inhalte im resultierenden R3-Raum nach der ersten Iteration für die Koordinatenorte c11 und c22 in der resultierenden Matrix C, die ursprünglich in 19b dargestellt sind. Hier enthält nach der ersten Iteration von 20b der R3-Registerraum an Ort C11 den Teilproduktterm a11·b11 und der R3-Registerraum an Ort C22 den Teilproduktterm (a21·b12).
Wie in 20c ersichtlich, werden nach einer nächsten Iteration die Matrix-A-Daten enthaltenden R0-Register dann horizontal um eine Einheit geschoben und die Matrix-B-Daten enthaltenden R1-Register werden vertikal um eine Einheit geschoben. Die gerade zuvor beschriebenen mathematischen Operationen werden dann in Bezug auf 20b wiederholt. Hier enthält jedoch der R3-Registerraum jeder Bahn anfänglich den Teilproduktterm der ersten Iteration (z. B. a11·b11 in R3 von Ort c11 und a21·b12 in R3 von Ort c22). Als solches wird am Ende der Iteration das Register R3 die akkumulierte Summe von beiden Teilprodukten enthalten, die bis jetzt berechnet wurden. D. h., dass R3 am Ort c11 die Summe (a11·b11) + (a12·b21) aufweisen wird und R3 am Ort c22 die Summe (a21·b12) + (a22·b22) aufweisen wird.
Wie in den 20d und 20e ersichtlich, wird nach zwei weiteren Iterationen einer einzelnen Schiebeeinheit für beide Matrizen gefolgt von einem Multiplizieren-Addieren das Matrixmultiplizieren abgeschlossen sein, wobei jeder Registermatrixort die korrekte Addition von Teilprodukten in seinem lokalen R3-Registerraum aufweist.
e. Zweidimensionale diskrete Fouriertransformation (2D DFT)
21 stellt eine zweidimensionale DFT (2D-DFT) dar. Eine 2D-DFT ist besonders für eine Rauschminderung und beschleunigende Faltungen nützlich. Wie in 21 ersichtlich, kann eine 2D-DFT als die Addition über den zweidimensionalen Raum des Produkts zweier komplexer Termini 2101, 2102 ausgedrückt werden. Ein erster von den komplexen Termini 2101 entspricht einem Phasor, dessen Größenordnung und Phase eine Funktion von Zeit und Frequenz sind. Bei der Implementierung wird der erste komplexe Term 2101 ausdrücklich als eine erste Koeffizientenmatrix berechnet. Ein zweiter von den komplexen Termini 2102 entspricht dem Signal, das von der räumlichen Domäne in den Frequenzbereich transformiert wird.
Der Einfachheit halber stellt 21 den ersten komplexen Term 2101 als Re1 + jIm1 dar und den zweiten komplexen Term 2102 als Re2 + jIm2 dar. Wie auf dem Fachgebiet bekannt kann der Realteil von (Re1 + jIm1)·(Re2 + jIm2) ausgedrückt werden als (Re1·Re2) – (Im1·Im2), während der Imaginärteil als j((Re1·Im2) + (Re2·Im1)) ausgedrückt werden kann. Die Additionen der 2D-DFT über 2D-Raum, fügen genau wie das unmittelbar zuvor ausführlich beschriebene Matrixmultiplizieren die Produkte von Elementen in einer Zeile eines Koordinatenorts durch entsprechende Elemente in der Spalte des Koordinatenorts hinzu.
Daher kann der Realteil der DFT-Resultante durch Ausführen eines Matrixmultiplizierens an einer Matrix von Re1-Werten und einer Matrix von Re2-Werten und Subtrahieren des Resultats eines Matrixmultiplizierens an einer Matrix von Im1-Werten und Im2-Werten von der resultierenden Matrix berechnet werden. Ähnlich kann der Imaginärteil der DFT-Resultante durch Ausführen eines Matrixmultiplizierens an einer Matrix von Re1-Werten und einer Matrix von Im2-Werten und Hinzufügen der resultierenden Matrix zum Resultat eines Matrixmultiplizierens an einer Matrix von Re2-Werten und Im1-Werten berechnet werden.
22a zeigt Maschinenoperationen zum Berechnen des Realteils der DFT. Wie in 22a ersichtlich wird eine Matrix von Re1-Werten in den R0-Registerraum der zweidimensionalen Schieberegistermatrix geladen, eine Matrix von Re2-Werten wird in den R1-Registerraum der zweidimensionalen Schieberegistermatrix geladen, eine Matrix von Im1-Werten wird in den R2-Registerraum der zweidimensionalen Schieberegistermatrix geladen und eine Matrix von Im2-Werten wird in den R3-Registerraum der zweidimensionalen Schieberegistermatrix geladen. Wie in 22b ersichtlich wird dann ein Rotationsscheralgorithmus auf jeden der Matrixwerte angewandt, wobei die Re1- und Im1-Werte horizontal geschert werden und die Re2- und Im2-Werte vertikal geschert werden.
Wie in 22c ersichtlich wird dann ein Re1·Re2-Matrixmultiplizieren und ein Im1·Im2-Matrixmultiplizieren ausgeführt, wobei die Resultanten entsprechend im R0- und R2-Registerraum gehalten werden. Der Inhalt des R2-Registerraums wird dann vom R0-Registerraum subtrahiert, wobei die Resultante der Subtraktion den Realteil der DFT-Transformation wie in 22d ersichtlich im R0-Registerraum hinterlässt.
Bei einer Ausführungsform wird die Realteilresultante im R0-Registerraum zum R4-Registerraum bewegt (wenn dieser existiert) oder in den lokalen Speicher, der mit der Schieberegistermatrix gekoppelt ist, ausgeschrieben. Dann werden die ursprünglichen Re1- und Im1-Werte in den R0- und R1-Registerraum (z. B. vom gleichen lokalen Speicher) zurückgeschrieben und horizontal geschert, sodass der Registerinhalt von 22a in der zweidimensionalen Schieberegistermatrix neu erzeugt wird.
Ein Re1·Im2-Matrixmultiplizieren und ein Re2·Im1-Matrixmultiplizieren wird dann entsprechend mit den im R0- und R2-Registerraum gehaltenen Resultanten ausgeführt. D. h., dass ein Matrixmultiplizieren auf den Inhalten von R0 und R3 ausgeführt wird, wobei die Resultante zu R0 zurückgeschrieben wird, und ein Matrixmultiplizieren an den Inhalten von R2 und R1 ausgeführt wird, wobei die Inhalte in R2 geschrieben werden. Die resultierenden Matrizen in R0 und R2 sind in 22e gezeigt. Der Inhalt des R0-Registerraums wird dann zum Inhalt des R2-Registerraums hinzugefügt und zu R0 zurückgeschrieben. Dies hinterlässt den Imaginärteil der DFT-Transformation wie in 22f dargestellt im R0-Registerraum.
f. FFT-Butterfly
Wie auf dem Fachgebiet bekannt ist eine schnelle Fourier-Transformation (FFT) eine schnellere weniger rechenintensive Herangehensweise an eine DFT. FFTs hängen von speziellen wirkungsvollen Algorithmen, um Zeit- oder Raumbereichsdaten schnell in Frequenzbereichsdaten umzuwandeln. Eine kritische Komponente solcher Algorithmen ist ein Butterfly-Algorithmus. Ein beispielhafter Butterfly-Algorithmus ist in 23 dargestellt. Hier ist eine Butterfly-Operation durch eine Stufe definiert, wobei die spezifische Stufe einen Betrag an Tauschen bestimmt, der sich zwischen Elementen einer gleichen Zeile oder Spalte ereignet. Wie auf dem Fachgebiet bekannt umfasst eine komplette FFT das Ausführen mathematischer Operationen an dem getauschten Inhalt von mehreren Butterfly-Operationen unterschiedlicher Stufen zwischen den Butterfly-Operationen.
23 zeigt jede der 1-, 2- und 4-Stufen-Butterfly-Operationen. Im Fall des 1-Stufen-Butterfly werden angrenzende Elemente getauscht. Im Fall des 2-Stufen-Butterfly werden angrenzende Paare von Elementen getauscht. Im Fall des 4-Stufen-Butterfly werden Gruppen von 4 angrenzenden Elementen getauscht. Im Fall einer zweidimensionalen FFT werden Elemente in einer ersten Matrix von Signaldaten gemäß jeder von mehreren Butterflystufen mit mathematischen Operationen getauscht, die an den getauschten Signaldatenelementen ausgeführt werden.
24a zeigt Maschinenebenenschiebeoperationen, die verwendet werden können, um einen 1-Stufe-1-Schmetterling über eine Matrix von Signaldaten zu bewirken, die im zweidimensionalen R0-Registerraum gespeichert sind. Der Einfachheit halber ist nur eine Zeile gezeigt. Es ist entsprechend zu erkennen, dass mehrere (z. B. alle) Zeilen oder Spalten einer Matrix in einem zweidimensionalen Schieberegister gemäß der bestimmten Zeile, die beobachtet wird, gleichzeitig verarbeitet werden können.
Hier ist das Schieberegister anfänglich mit A0 bis A7 über jede seiner Zeilen in seinem R0-Registerraum geladen. Das Schieberegister schiebt dann den R0-Registerinhalt um eine Einheit nach rechts und jede Ausführungsbahn speichert die Daten, die gerade geschoben wurden, in ihren Ort in ihrem R1-Registerraum. Dann schiebt das Schieberegister den R0-Registerinhalt um zwei Einheiten nach links (wobei die Resultante als R0' bezeichnet wird). Jede Ausführungsbahn führt dann eine Auswahloperation aus, bei der entweder deren lokaler R0-Inhalt oder deren lokaler R1-Inhalt abhängig vom Ort der Bahn in R0 gespeichert wird (d. h., wenn R0 ausgewählt ist, bleibt der Inhalt von R0 unverändert). Bei der bestimmten Ausführungsform von 24a wählen gerade Bahnen R0 aus, während ungerade Bahnen R1 wählen. Nach der Auswahl befinden sich die korrekt getauschten Werte im R0-Registerraum von jeder Ausführungsbahn (bezeichnet als R0'').
Die 24b und 24c stellen 2-Stufen- und 4-Stufen-Butterfly-Operationen dar. Die Verarbeitung ist der gerade zuvor beschriebenen 1-Stufen-Butterfly-Operation ähnlich, außer dass im Fall des 2-Stufenschmetterlings die Schieberegistermatrix den R0-Registerinhalt um zwei Einheiten nach rechts und dann um vier Einheiten nach links schiebt. Jede andere gerade Bahn und deren Nachbar ganz rechts wählen von einem von R1 und R0'' aus, wohingegen die anderen geraden Bahnen und deren Nachbar ganz rechts von dem anderen von R1 und R0'' auswählen. Im Fall des 4-Stufenschmetterlings, der in 25c dargestellt ist, schiebt die Schieberegistermatrix die Inhalte von R0 um vier Einheiten nach rechts und wählt dann alle Werte, wie sie sich in R0 befinden (als R0' bezeichnet).
g. Blockabgleich
25 zeigt eine Darstellung eines Blockabgleichs. Ein Blockabgleich ist besonders für Bewegungseinschätzung (z. B. für Videokompression), Bildstabilisierung und für Bildfusionsalgorithmen (z. B. Belichtungsverschmelzung und zeitliche Geräuschverschmelzung) nützlich. Hier ein spezifisches Merkmal in einem ersten Basis-(z. B. vorhergehenden)-Bild, nach dem in einem alternativen (z. B. späteren) Bild gesucht wird. Im Beispiel von 26 muss der Block ABCD, der im Basisbild vorhanden ist, im alternativen Bild gefunden werden.
Die 26a bis 26d zeigen ein Verfahren zum Ausführen eines Blockabgleichs auf einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einer zweidimensionalen Schieberegistermatrix. Hier werden die Bildpunkte des alternativen Bilds im R0-Registerraum des zweidimensionalen Schieberegisters gehalten. Ein erster Bildpunkt des Merkmals im Basisbild, nach dem gesucht wird (z. B.: „A”) wird zu allen Ausführungsbahnen übertragen und eine absolute Differenz wird an dem Inhalt im R0-Registerraum und dem Wert „A” genommen und die Resultante in R1 gespeichert (der Einfachheit halber wird angenommen, dass keine anderen Bildpunkte im alternativen Bild irgendwelche der gesuchten Werte des Merkmals (A, B, C, D) aufweisen. Wenn ein Bildpunktwert im alternativen Bild mit dem übertragenen Bildpunktwert A übereinstimmt (oder annähernd übereinstimmt), dann sollte die Resultante in R1 bei (oder nahe bei) Null sein. Alle anderen Resultanten in R1 sollten im Wesentlichen nicht Null sein.
Unter Bezugnahme auf 26b wird das alternative Bild dann um eine Einheit im zweidimensionalen Schieberegister geschoben, ein nächster Bildpunktwert „B“, nach dem in dem Merkmal gesucht wird, wird zu allen Ausführungsbahnen übertragen und die absolute Differenz erneut genommen und in R1 gespeichert. Hier sollte der bestimmte Ausführungsbahnmatrixort 2601, der zwei aufeinanderfolgende Bildpunktübereinstimmungen aufwies, den niedrigsten resultierenden Wert in R1 aufweisen. Dieser bestimmte Ausführungsbahnort ist der Ort, an dem die bestimmte Reihenfolge der Bildpunktwerte, die übertragen werden, sich mit der bestimmten Scanning-/Schiebebewegung des alternativen Bildes in R0 der Schieberegistermatrix ausrichtet.
Wie in den 26c und 26d ersichtlich, fährt der Prozess mit jeder nächsten Iteration fort, die einen nächsten Bildpunktwert überträgt und die alternativen Bilddaten in einer bestimmten Scanreihenfolge schiebt. Nach einer Anzahl von Iterationen, die zum Übertragen aller gesuchten Bild-Bildpunkte ausreichend sind, und um das alternative Bild in R0 entsprechend über einen Bereich zu schieben, welcher der Größe des gesuchten Bildes entspricht, wird idealerweise nur ein Ort 2601 in der Ausführungsbahnmatrix eine Übereinstimmung bei jedem Iterationszyklus erfahren haben. Als solches sollte dieser bestimmte Matrixort 2601 einen Nullwert oder einen dem Nullwert nahen Wert in seinem R1-Registerraum aufrechterhalten können (oder mindestens einen kleineren Wert als die anderen Ausführungsbahnen), nachdem alle Iterationen abgeschlossen wurden.
Es wird eine Minimumfinden-Operation wie vorstehend beschrieben am R1-Registerraum über alle Zeilen in der Matrix ausgeführt. Der niedrigste entsprechende Wert in jeder Zeile wird dann an jedem Ort seiner Zeile gehalten. Es wird dann eine Minimumfinden-Operation am R1-Registerraum über alle Spalten in der Matrix ausgeführt. Die Resultante sollte den Bildpunkt, der mit dem Basisbild übereinstimmte, bei jeder Iteration identifizieren, die wiederum verwendet werden kann, um den exakten Ort des gesuchten Bildes in der alternativen Matrix zu identifizieren.
4.0 Aufbau von niederem Programmcode
27 zeigt eine Vorlaufzeitentwicklungsumgebung, in der ein Programmierer eine höhere Bildverarbeitungsfunktion ausgelegt, und die Anwendungsentwicklungsumgebung (welche die vorstehend ausführlich unter Bezug auf die 1 bis 6 beschriebene virtuelle ISA-Umgebung umfassen kann) stellt eine/alle der vorstehend beschriebenen Sonderoperationen bereit, sodass der Entwickler diese nicht vollständig neu schreiben muss.
Der Entwickler kann hier speziell irgendeine der oben ausführlich beschriebenen Operationen aufrufen und/oder die Entwicklungsumgebung stellt diese als Reaktion automatisch von einer Bibliothek 2701 bereit. Alternativ oder in Kombination kann der Bedarf des Entwicklers an solchen Operationen impliziert oder abgeleitet werden (wie z. B. ein Matrixmultiplizieren im Fall von einer 2D-DFT) und die Entwicklungsumgebung fügt automatisch Programmcode von der Bibliothek 2701 ein, der diese Funktionen ausführt (z. B. als Teil eines Kompilierungsprozesses).
Der Programmcode, der die vorstehend beschriebenen Operationen oder alternativen Ausführungsformen davon ausführt, kann daher in höherem (z. B. virtuelle ISA) Programmcode oder in höherem Objektcode ausgedrückt werden. Bei verschiedenen Ausführungsformen kann der höhere virtuelle ISA-Code zu bearbeitende Dateninhalte als Speicherlesevorgänge mit x,y-Adresskoordinaten spezifizieren, während der Objektcode stattdessen diese Datenzugriffe als zweidimensionale Schieberegisteroperationen umfassen kann (wie z. B. einige der vorstehend beschriebenen Schiebeoperationen oder ähnliche Ausführungsformen). Ein Compiler kann die x,y-Lesevorgänge in der Entwicklungsumgebung in entsprechende Schiebevorgänge des zweidimensionalen Schieberegisters umwandeln, die spezifizierter Objektcode sind (z. B. kann ein Lesen in der Entwicklungsumgebung mit x,y-Koordinaten (+2, +2) in Objektcode als ein Schieben nach links um zwei Räume und ein Verschieben nach unten um zwei Räume realisiert werden). Abhängig von der Umgebung kann der Entwickler in beide dieser Ebenen (oder z. B. nur die höhere VISA-Ebene) Einsicht haben. Bei noch anderen Ausführungsformen können solche vorher geschriebenen Routinen während der Laufzeit (z. B. durch einen Just-In-Time-Compiler) anstatt der Vorlaufzeit aufgerufen werden.
5.0 Abschließen von Aussagen
Aus den vorhergehenden Abschnitten ist entsprechend zu erkennen, dass die im Abschnitt 1.0 beschriebene virtuelle Umgebung auf einem Computersystem instanziiert werden kann. Desgleichen kann ein Bildprozessor wie vorstehend im Abschnitt 2.0 beschrieben in Hardware auf einem Computersystem verkörpert sein (z. B. als Teil eines System-On-Chips (SOC) einer handgehaltenen Vorrichtung, der Daten von der Kamera der handgehaltenen Vorrichtung verarbeitet).
Es ist wichtig, darauf hinzuweisen, dass die vorstehend beschriebenen verschiedenen Merkmale der Bildprozessorarchitektur nicht zwangsläufig auf die Bildverarbeitung im herkömmlichen Sinne beschränkt sind und daher auf andere Anwendungen angewendet werden können, die ggf. veranlassen können, dass der Bildprozessor neu charakterisiert wird oder auch nicht. Wenn beispielsweise eines der vorstehend beschriebenen verschiedenen Merkmale der Bildprozessorarchitektur bei der Erstellung und/oder Erzeugung und/oder Wiedergabe von Animationen anstatt bei der Verarbeitung von tatsächlichen Kamerabildern verwendet werden soll, kann der Bildprozessor als Grafikprozessor charakterisiert sein. Zudem können die oben beschriebenen Architekturmerkmale des Bildprozessors in anderen technischen Anwendungen, wie in der Videoverarbeitung, Bildverarbeitung, Bilderkennung und/oder dem maschinellen Lernen, angewendet werden. Auf diese Weise kann der Bildprozessor (z. B. als Koprozessor) in einen allgemeineren Universalprozessor (z. B. als Teil einer CPU des Computersystems) mit integriert werden oder ein eigenständiger Prozessor innerhalb eines Computersystems sein.
Die vorstehend beschriebenen Hardware-Ausführungsformen können in einem Halbleiterchip und/oder als Beschreibung eines Schaltungsdesigns zur letztendlichen Ausrichtung auf einen Halbleiterherstellungsprozess verkörpert sein. Im letzteren Fall können diese Schaltkreisbeschreibungen die Form einer (z. B. VHDL oder Verilog) Beschreibung einer Registerüberleitungsschaltung (RTL), einer Torschaltung, einer Transistorschaltung oder einer Maske oder verschiedener Kombinationen derselben annehmen. Schaltungsbeschreibungen sind in der Regel auf einem computerlesbaren Speichermedium (wie z. B. einer CD-ROM oder einer anderen Art von Speichertechnologie) enthalten.
Aus den vorangehenden Abschnitten ist zu erkennen, dass ein Bildprozessor wie vorstehend beschrieben in der Hardware auf einem Computersystem (z. B. als Teil eines Handgerätsystems on Chip (SOC), das Daten von der Kamera des Handgerätes verarbeitet) enthalten sein kann. In Fällen, in denen der Bildprozessor als Hardware-Schaltung verkörpert ist, ist zu beachten, dass die Bilddaten, die von dem Bildprozessor verarbeitet werden, direkt von einer Kamera empfangen werden können. Hier kann der Bildprozessor Teil einer diskreten Kamera oder Teil eines Computersystems mit einer integrierten Kamera sein. Im letzteren Fall können die Bilddaten direkt von der Kamera oder aus dem Systemspeicher des Computersystems empfangen werden (z. B. sendet die Kamera ihre Bilddaten anstatt an den Bildprozessor an den Systemspeicher). Zu beachten ist auch, dass viele der in den vorangehenden Abschnitten beschriebenen Merkmale auf eine Bildprozessoreinheit (zur Darstellung von Animationen) anwendbar sind.
28 zeigt eine beispielhafte Darstellung eines Computersystems. Viele der Komponenten des nachstehend beschriebenen Computersystems sind auf ein Computersystem mit einer integrierten Kamera und einem zugehörigen Bildprozessor (z. B. einem Handgerät, wie etwa einem Smartphone oder Tablet-Computer) anwendbar. Ein Durchschnittsfachmann wird leicht zwischen beiden unterscheiden können.
Wie in 28 dargestellt, kann das grundlegende Computersystem eine zentrale Verarbeitungseinheit 2801 (die beispielsweise eine Vielzahl von Universal-Verarbeitungskernen 2815_1 bis 2815_N und einen auf einem Multikernprozessor oder einem Anwendungsprozessor angeordneten Hauptspeichercontroller 2817 beinhalten kann), Systemspeicher 2802, ein Display 2803 (z. B. Touchscreen, Flachbildschirm), eine lokal verdrahtete Punkt-zu-Punkt-Verbindung (z. B. eine USB-Schnittstelle) 2804, verschiedene Netzwerk-E-/A-Funktionen 2805 (wie z. B. eine Ethernet-Schnittstelle und/oder ein Mobilfunkmodem-Teilsystem), ein drahtloses lokales Netzwerk (z. B. WLAN) 2806, eine drahtlose Punkt-zu-Punkt-Verbindung (z. B. Bluetooth-Schnittstelle) 2807 und eine globale Positionierungssystemschnittstelle 2808, verschiedene Sensoren 2809_1 bis 2809_N, eine oder mehrere Kameras 2810, eine Batterie 2811, eine Energieverwaltungssteuereinheit 2824, einen Lautsprecher und ein Mikrofon 2813 sowie einen Audio-Codierer/Decodierer 2814 beinhalten.
Ein Anwendungsprozessor oder Multikernprozessor 2850 kann einen oder mehrere Universalprozessorkerne 2815 innerhalb seiner CPU 2801, eine oder mehrere grafische Verarbeitungseinheiten 2816, eine Speicherverwaltungsfunktion 2817 (z. B. einen Speichercontroller), eine E-/A-Steuerfunktion 2818 und eine Bildverarbeitungseinheit 2819 beinhalten. Die Universalprozessorkerne 2815 führen typischerweise das Betriebssystem und die Anwendungssoftware des Computersystems aus. Die Grafikprozessoren 2816 führen typischerweise grafikintensive Funktionen aus, um z. B. Grafikinformationen, die auf dem Display 2803 dargestellt werden, zu generieren. Die Speichersteuerungsfunktion 2817 verbindet mit dem Systemspeicher 2802, um Daten in den Systemspeicher 2802 zu schreiben bzw. davon lesen. Die Leistungsmanagementsteuereinheit 2824 steuert generell den Energieverbrauch des Systems 2800.
Die Bildverarbeitungseinheit 2819 kann gemäß einer der in den vorangehenden Abschnitten beschriebenen Ausführungsformen der Bildverarbeitungseinheit implementiert sein. Alternativ dazu oder in Kombination kann die IPU 2819 mit einer oder sowohl der GPU 2816 als auch der CPU 2801 als Koprozessor derselben gekoppelt sein. Darüber hinaus kann in verschiedenen Ausführungsformen die GPU 2816 mit einem der vorstehend beschriebenen Prozessormerkmale implementiert sein.
Das Touchscreen-Display 2803, die Kommunikationsschnittstellen 2804–2807, die GPS-Schnittstelle 2808, die Sensoren 2809, die Kamera 2810 und der Lautsprecher/Mikrofon-Codec 2813, 2814 können alle als unterschiedliche Formen von E/A (Eingabe und/oder Ausgabe) in Bezug auf das gesamte Computersystem betrachtet werden, darunter auch gegebenenfalls ein integriertes Peripheriegerät (z. B. die eine oder die mehreren Kameras 2810). Abhängig von der Implementierung können verschiedene dieser I/O-Komponenten im Anwendungsprozessor/Mehrkernprozessor 2850 integriert sein oder sie können sich außerhalb des Chips oder Pakets des Anwendungsprozessors/Mehrkernprozessors 2850 befinden.
Bei einer Ausführungsform beinhalten eine oder mehrere Kameras 2810 eine Tiefenkamera, die in der Lage ist, die Tiefe zwischen der Kamera und einem Objekt in dessen Sichtfeld zu messen. Anwendungssoftware, Betriebssystemsoftware, Gerätetreibersoftware und/oder Firmware, die auf einem universellen CPU-Kern (oder einem anderen Funktionsblock mit einer Befehlsausführungspipeline zum Ausführen eines Programmcodes) eines Anwendungsprozessors oder eines anderen Prozessors ausgeführt werden, können irgendwelche der vorstehend beschriebenen Funktionen ausführen.
Ausführungsformen der Erfindung können, wie vorstehend dargelegt, verschiedene Prozesse beinhalten. Die Prozesse können in maschinenausführbaren Befehlen verkörpert sein. Die Befehle können verwendet werden, um einen Allzweck- oder Spezialprozessor zu veranlassen, bestimmte Prozesse auszuführen. Alternativ können diese Prozesse durch spezifische Hardwarekomponenten ausgeführt werden, die festverdrahtete Logik enthalten, um die Prozesse auszuführen, oder durch jede Kombination von programmierten Computerkomponenten und benutzerdefinierten Hardwarekomponenten.
Elemente der vorliegenden Erfindung können ebenfalls als ein maschinenlesbares Medium zum Speichern der maschinenausführbaren Befehle bereitgestellt werden. Das maschinenlesbare Medium kann umfassen, ist aber nicht beschränkt auf, Disketten, optische Disks, CD-ROMs und magnetooptische Disks, FLASH-Speicher, ROMs, RAMs, EPROMs, EEPROMs, magnetische oder optische Karten, Ausbreitungsmedien oder andere Arten von Medien/maschinenlesbaren Medien, die zum Speichern elektronischer Befehle geeignet sind. Die Elemente können beispielsweise als ein Computerprogramm heruntergeladen werden, das von einem dezentralen Computer (z. B. einem Server) mittels eines in einer Trägerwelle oder in einem anderen Ausbreitungsmedium enthaltenen Datensignals an einen anfordernden Computer (z. B. einen Client) über eine Kommunikationsverbindung (z. B. ein Modem oder eine Netzwerkverbindung) übertragen werden kann.
In der vorhergehenden Beschreibung wurden spezifische beispielhafte Ausführungsformen beschrieben. Es ist jedoch offensichtlich, dass verschiedene Modifikationen und Änderungen daran vorgenommen werden können, ohne von dem in den beigefügten Ansprüchen dargelegten Sinn und Umfang der Erfindung abzuweichen. Die Beschreibung und die Zeichnungen sind daher in einem veranschaulichenden und nicht in einem einschränkenden Sinne zu betrachten.

Claims

Maschinenlesbares Speichermedium, das Programmcode enthält, der bei Verarbeitung durch einen Bildprozessor, der eine zweidimensionale Ausführungsbahnmatrix und eine zweidimensionale Schieberegistermatrix umfasst, den Bildprozessor veranlasst, ein Verfahren auszuführen mit wiederholtem Schieben von einem erstem Inhalt von mehreren Zeilen oder Spalten der zweidimensionalen Schieberegistermatrix und wiederholtes Ausführen von mindestens einem Befehl zwischen Schiebevorgängen, der an dem geschobenen ersten Inhalt und/oder zweiten Inhalt operiert, der an entsprechenden Orten der zweidimensionalen Schieberegistermatrix speicherresident ist, in die der geschobene erste Inhalt geschoben wurde.
Maschinenlesbares Speichermedium nach Anspruch 1, wobei bei Ausführung des Verfahrens, eine Zeilenadditions- oder Spaltenadditions- oder Präfixadditionsoperation implementiert ist, die an mehreren entsprechenden Zeilen oder Spalten einer Matrix gleichzeitig ausgeführt wird.
Maschinenlesbares Speichermedium nach Anspruch 2, wobei die Operation eine Zeilen- oder Spaltenadditionsoperation ist, um irgendeines auszuführen von: a) Berechnen von Statistiken; b) Entstören; c) Downsampling.
Maschinenlesbares Speichermedium nach Anspruch 2, wobei die Operation eine Zeilen- oder Spaltenpräfixadditionsoperation ist, um irgendeines auszuführen von: a) Bestimmen eines ganzheitlichen Bildes; b) Beschleunigen eines Mittelwertfilters; c) Berechnen einer Adresse für eine Verdichtungsoperation.
Maschinenlesbares Speichermedium nach einem der Ansprüche 2 bis 4, wobei der Inhalt der zweidimensionalen Schieberegistermatrix in gegenüberliegende Ränder der zweidimensionalen Schieberegistermatrix umgebrochen ist.
Maschinenlesbares Speichermedium nach einem der Ansprüche 2 bis 5, wobei die Operation über alle Zeilen oder Spalten der Schieberegistermatrix summiert.
Maschinenlesbares Speichermedium nach Anspruch 1, das beim Ausführen des Verfahrens, eine Minimum- oder Maximumfindenoperation implementiert ist, die an mehreren entsprechenden Zeilen oder Spalten einer Datenmatrix gleichzeitig ausgeführt wird, die einen Indexwert eines Minimum- oder Maximumwertes über eine entsprechende Reihe oder Spalte wiederholt.
Maschinenlesbares Speichermedium nach Anspruch 7, wobei die Operation irgendeines ausführt von: a) Berechnen von Statistiken; b) Nachbearbeiten eines Blockabgleichalgorithmus.
Maschinenlesbares Speichermedium nach Anspruch 7 oder 8, wobei der Inhalt der zweidimensionalen Schieberegistermatrix um gegenüberliegende Ränder der zweidimensionalen Schieberegistermatrix umgebrochen ist.
Maschinenlesbares Speichermedium nach einem der Ansprüche 7 bis 9, wobei die Operation ein Minimum oder ein Maximum über alle Zeilen oder Spalten der Schieberegistermatrix findet.
Maschinenlesbares Speichermedium nach Anspruch 1, wobei bei Ausführung des Verfahrens eine Matrixmultiplizieroperation implementiert ist.
Maschinenlesbares Medium nach Anspruch 11, wobei die Matrixmultiplizieroperation irgendeines ausführt von: a) Berechnen einer diskreten Fouriertransformation; b) Berechnen einer diskreten Kosinustransformation.
Maschinenlesbares Medium nach Anspruch 11, wobei die Matrixmultiplizieroperation irgendeines ausführt von: a) Komprimieren von Daten; b) Filtern von Daten.
Maschinenlesbares Medium nach einem der Ansprüche 11 bis 13, wobei vor dem wiederholten Schieben und dem wiederholten Ausführen des mindestens einen Befehls ein Rotationsscheralgorithmus über Zeilen oder Spalten von Daten in der zweidimensionalen Schieberegistermatrix angewandt wird.
Maschinenlesbares Medium nach einem der Ansprüche 11 bis 14, wobei der mindestens eine Befehl eine Multiplizier-Addier-Operation implementiert.
Maschinenlesbares Speichermedium, wobei bei Ausführung des Verfahrens eine zweidimensionale diskrete Fouriertransformation implementiert ist.
Maschinenlesbares Medium nach Anspruch 16, wobei die zweidimensionale diskrete Fouriertransformation irgendeines ausführt von: a) Reduzieren von Rauschen; b) Beschleunigen einer Faltung.
Maschinenlesbares Medium nach Anspruch 16 oder 17, wobei vor dem wiederholten Schieben und dem wiederholten Ausführen des mindestens einen Befehls ein Rotationsscheralgorithmus über Zeilen oder Spalten von Daten in der zweidimensionalen Schieberegistermatrix angewandt wird.
Maschinenlesbares Medium nach einem der Ansprüche 16 bis 18, wobei der eine oder die mehreren Befehle gemeinsam auf eine erste Schicht von realen Werten innerhalb der zweidimensionalen Schieberegistermatrix und eine zweite Schicht von imaginären Werten innerhalb der zweidimensionalen Schieberegistermatrix verweisen.
Maschinenlesbares Speichermedium, wobei bei Ausführung des Verfahrens eine Butterfly-Operation für eine schnelle Fouriertransformation implementiert ist.
Maschinenlesbares Speichermedium, wobei bei Ausführung des Verfahrens Blockabgleichoperation zu implementiert ist.
Maschinenlesbares Speichermedium nach Anspruch 21, wobei die Blockabgleichoperation irgendeines ausführt von: a) Abschätzen von Bewegung; b) Stabilisieren eines Bildes; c) Implementieren eines Bildfusionsalgorithmus.
Maschinenlesbares Medium nach Anspruch 21 oder 22, wobei das Verfahren weiter das wiederholte Übertragen eines nächsten gesuchten Bildpunktwerts und des einen oder der mehreren Befehle umfasst, die einen Unterschied zwischen dem nächsten gesuchten Bildpunktwert und dem geschobenen ersten Inhalt bestimmen.
Maschinenlesbares Speichermedium, das Programmcode enthält, der bei Verarbeitung durch eine Verarbeitungseinheit bewirkt, dass ein Verfahren ausgeführt wird, umfassend: Artikulieren eines Verfahrens in einem höheren Programmcode, der auf Daten mit Speicherlese- und -schreibbefehlen mit einem Befehlsformat zugreift, das eine Speicherstelle in x,y-Koordinaten spezifiziert, wobei der höhere Programmcode in einen Objektcode zu kompilieren ist, der auf einem Bildprozessor mit einer zweidimensionalen Ausführungsbahnmatrix und einer zweidimensionalen Schieberegistermatrix ausführbar ist, und der Objektcode ausgeführt ist, den Bildprozessor zu veranlassen, ein Verfahren auszuführen, mit wiederholtem Schieben von einem erstem Inhalt von mehreren Zeilen oder Spalten der zweidimensionalen Schieberegistermatrix und wiederholtes Ausführen von mindestens einem Befehl zwischen Schiebevorgängen, der an dem geschobenen ersten Inhalt und/oder zweiten Inhalt operiert, der an entsprechenden Orten der zweidimensionalen Schieberegistermatrix speicherresident ist, in die der geschobene erste Inhalt geschoben wurde.
Maschinenlesbares Speichermedium nach Anspruch 24, wobei der Objektcode ausgeführt ist, irgendeines auszuführen von: a) Berechnen von Statistiken; b) Entstören; c) Downsampling; d) Bestimmen eines ganzheitlichen Bildes; e) Beschleunigen eines Mittelwertfilters; f) Nachbearbeiten eines Blockabgleichalgorithmus; g) Berechnen einer diskreten Fouriertransformation; h) Berechnen einer diskreten Kosinustransformation. i) Komprimieren von Daten; j) Filtern von Daten; k) Reduzieren von Rauschen; l) Beschleunigen einer Faltung; m) Abschätzen von Bewegung; n) Stabilisieren eines Bildes; o) Implementieren eines Bildfusionsalgorithmus.