DE102013018380A1 - System for compilation or run-time execution of data-parallel program in single-instruction multi-strand processor, has strand association unit connected to partition unit and determining strand from parent group to run as parent strand - Google Patents
System for compilation or run-time execution of data-parallel program in single-instruction multi-strand processor, has strand association unit connected to partition unit and determining strand from parent group to run as parent strand Download PDFInfo
- Publication number
- DE102013018380A1 DE102013018380A1 DE201310018380 DE102013018380A DE102013018380A1 DE 102013018380 A1 DE102013018380 A1 DE 102013018380A1 DE 201310018380 DE201310018380 DE 201310018380 DE 102013018380 A DE102013018380 A DE 102013018380A DE 102013018380 A1 DE102013018380 A1 DE 102013018380A1
- Authority
- DE
- Germany
- Prior art keywords
- strand
- parent
- parallel
- function
- worker
- Prior art date
- Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
- Pending
Links
Images
Classifications
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F8/00—Arrangements for software engineering
- G06F8/40—Transformation of program code
- G06F8/41—Compilation
- G06F8/45—Exploiting coarse grain parallelism in compilation, i.e. parallelism between groups of instructions
- G06F8/456—Parallelism detection
Landscapes
- Engineering & Computer Science (AREA)
- General Engineering & Computer Science (AREA)
- Theoretical Computer Science (AREA)
- Software Systems (AREA)
- Physics & Mathematics (AREA)
- General Physics & Mathematics (AREA)
- Stored Programmes (AREA)
Abstract
Description
QUERVERWEIS AUF VERWANDTE ANMELDUNGCROSS-REFERENCE TO RELATED APPLICATION
Diese Anmeldung beansprucht die Priorität der vorläufigen US-Anmeldung mit der Seriennummer 61/722661, die von Lin et al. am 5. November 2012 eingereicht wurde mit dem Titel „AUSFÜHRUNG EINES SEQUENZIELLEN CODES UNTER ANWENDUNG EINER GRUPPE AUS STRÄNGEN” und der US-Anmeldung mit der Seriennummer 13/724359, die von Lin, et al. am 21. Dezember 2012 eingereicht wurde mit dem Titel „SYSTEM UND VERFAHREN ZUR KOMPILIERUNG ODER LAUFZEITAUSFÜHRUNG EINES DATENPARALLELEN PROGRAMMS MIT AUFTEILUNG-VEREINIGUNG MIT FUNKTIONSAUFRUFEN IN EINEM EINZELBEFEHL-MULTI-STRANG-PROZESSOR”, die beide die gleiche Anmelderin wie diese Anmeldung haben und hierin durch Bezugnahme mit eingeschlossen sind.This application claims priority to US Provisional Application Ser. No. 61 / 722,661, which is assigned to Lin et al. filed November 5, 2012, entitled "IMPLEMENTATION OF A SEQUENCIAL CODE USING A GROUP OF STRANDS" and US Serial No. 13 / 724,359 filed by Lin, et al. on December 21, 2012, entitled "SYSTEM AND METHOD FOR COMPILING OR RUNNING A DATA PARALLEL PROGRAM WITH DISTRIBUTION ASSOCIATION WITH FUNCTION CALLS IN A SINGLE COMMAND MULTI-STRESS PROCESSOR", both of which are assigned to the same assignee as this application and herein Reference are included.
TECHNISCHES GEBIETTECHNICAL AREA
Diese Anmeldung betrifft generell parallele Prozessoren und insbesondere ein System und ein Verfahren zur Kompilierung oder Laufzeitausführung eines datenparallelen Programms mit Aufteilung-Vereinigung mit Funktionsaufrufen in einem Einzelbefehl-Multi-Strang-(SIMT-)Prozessor.This application relates generally to parallel processors and, more particularly, to a system and method for compiling or executing a data parallelled partition-sharing program with function calls in a single-instruction multi-strand (SIMT) processor.
HINTERGRUNDBACKGROUND
Wie der Fachmann auf diesem Gebiet weiß, können Anwendungen oder Programme parallel ausgeführt werden, um ihr Leistungsverhalten zu verbessern. Datenparallele Programme führen den gleichen Prozess gleichzeitig an unterschiedlichen Daten aus. Aufgabenparallele Programme führen unterschiedliche Prozesse gleichzeitig an den gleichen Daten aus. Statische parallele Programme sind Programme mit einem Grad an Parallelität, der vor der Ausführung bestimmt werden kann. Im Gegensatz dazu kann die Parallelität, die von dynamischen parallelen Programmen erreichbar ist, nur ermittelt werden, wenn sie ausgeführt werden. Unabhängig davon, ob das Programm datenparallel oder aufgabenparallel oder statisch oder dynamisch parallel ist, kann es in einer Pipeline bzw. Parallelverarbeitungslinie ausgeführt werden, was häufig der Fall ist für graphische Programme.As those skilled in the art know, applications or programs may be executed in parallel to improve their performance. Data-parallel programs execute the same process simultaneously on different data. Task-parallel programs execute different processes simultaneously on the same data. Static parallel programs are programs with a degree of parallelism that can be determined before execution. In contrast, the parallelism achievable by dynamic parallel programs can only be determined when they are executed. Regardless of whether the program is data-parallel or task-parallel or static or dynamically parallel, it can be executed in a pipeline or parallel-processing line, which is often the case for graphical programs.
Ein SIMT-Prozessor ist besonders geschickt bei der Ausführung datenparalleler Programme. Eine Steuereinheit in dem SIMT-Prozessor erzeugt Gruppen aus Strängen zur Ausführung und disponiert diese für die Ausführung, während welcher alle Stränge in der Gruppe den gleichen Befehl gleichzeitig ausführen. In einem speziellen Prozessor hat jede Gruppe oder „Wölbung bzw. Kette” 32 Stränge, die 32 Ausführungs-Pipelines oder Bahnen in dem SIMT-Prozessor entsprechen.A SIMT processor is particularly adept at executing data-parallel programs. A controller in the SIMT processor generates strings of strings for execution and schedules them for execution, during which all threads in the group execute the same instruction concurrently. In a particular processor, each group or "warp" has 32 strands corresponding to 32 execution pipelines or lanes in the SIMT processor.
Ein datenparalleles Programm mit Aufteilung-Vereinigung beginnt mit einem Hauptprogramm, das nur einen Strang aufweist. Das Programm ist in dieser Phase in einer sequenziellen Phase oder einem sequenziellen Bereich. Bei einem gewissen Punkt während der Ausführung des Hauptprogramms trifft der Haupt- oder „Master-”Strang auf eine Sequenz aus parallelen Phasen oder Bereichen. Jeder parallele Bereich hat einen unabhängigen Datensatz und kann von mehreren Strängen gleichzeitig ausgeführt werden. Die Anzahl an gleichzeitigen Aufgaben in jedem parallelen Bereich wird bestimmt, wenn der parallele Bereich beginnt, und sich während des parallelen Bereichs nicht ändert. Wenn ein paralleler Bereich angetroffen wird, führt der Haupt-Strang eine Aufteilung in eine Gruppe aus Strängen (die als Arbeiter-Stränge bezeichnet werden) durch, um die parallelen Bereiche parallel abzuarbeiten. Das Programm tritt dann in den parallelen Bereich ein. Wenn ein Arbeiter-Strang einen neuen parallelen Bereich antrifft, wird der neue parallele Bereich serialisiert, d. h. der parallele Bereich wird von dem eintreffenden Arbeiter-Strang selbst ausgeführt. Der Haupt-Strang wartet, bis der parallele Bereich beendet ist. Beim Austritt aus dem parallelen Bereich vereinigen sich die Arbeiter-Stränge mit dem Haupt-Strang, der dann die Ausführung des Hauptprogramms wieder fortgesetzt, wobei dann das Programm einen sequenziellen Bereich betritt.A data-parallel program with split-union begins with a main program, which has only one strand. The program is in a sequential or sequential area at this stage. At some point during execution of the main program, the master or "master" thread encounters a sequence of parallel phases or regions. Each parallel area has an independent data set and can be executed by several strands simultaneously. The number of concurrent tasks in each parallel area is determined when the parallel area starts and does not change during the parallel area. When a parallel region is encountered, the main strand splits into a group of strands (referred to as worker strands) to work the parallel regions in parallel. The program then enters the parallel area. When a worker strand encounters a new parallel region, the new parallel region is serialized, i. H. the parallel area is executed by the incoming worker strand itself. The main line waits until the parallel area is finished. Upon exiting the parallel area, the worker strands unite with the main strand, which then resumes execution of the main program, and then the program enters a sequential area.
Die nachfolgende Tabelle 1 gibt ein Beispiel für ein datenparalleles Programm mit Aufteilung-Vereinigung an.Table 1 below gives an example of a data-parallel program with split-union.
Tabelle 1 – ein Beispiel eines datenparallelen Programms mit Aufteilung-Vereinigung Table 1 - an example of a data-parallel program with split-union
Zum Zwecke des Verständnisses der Tabelle 1 und des Restes dieser Offenbarung sind die Begriffe „foo” und „bar” willkürliche Namen von Funktionen. Es kann daher eine beliebige Funktion anstelle von „foo” oder „bar” verwendet werden.For purposes of understanding Table 1 and the remainder of this disclosure, the terms "foo" and "bar" are arbitrary names of functions. It is therefore possible to use any function instead of "foo" or "bar".
Das Datenparallele Modell mit Aufteilung-Verzweigung wird häufig in der parallelen Programmierung eingesetzt. Beispielsweise verwendet der OpenMP-Standard dieses Modell als ein grundlegendes Strang-Ausführungsmodell. Der OpenACC-Standard verwendet dieses Modell für die Arbeiter-Stränge in einer Gruppe, die als eine „Arbeitsgruppe” bezeichnet wird.The data-parallel split-split model is often used in parallel programming. For example, the OpenMP standard uses this model as a basic thread execution model. The OpenACC standard uses this model for worker strands in a group called a "workgroup".
ÜBERBLICKOVERVIEW
Ein Aspekt stellt ein System zur Kompilierung oder Laufzeitausführung eines datenparallelen Programms mit Aufteilung-Vereinigung mit Funktionsaufrufen bereit. In einer Ausführungsform umfasst das System: (1) eine Partitionseinheit, die ausgebildet ist, Ketten in eine übergeordnete Kette und mindestens eine Arbeiter-Kette zu unterteilen, und (2) eine Strang-Zuweisungseinheit, die mit der Partitionseinheit verbunden und ausgebildet ist, nur einen einzelnen Strang aus der übergeordneten Kette zur Ausführung und alle Stränge in der mindestens einen Arbeiter-Kette zur Ausführung anzuweisen bzw. zu bestimmen.One aspect provides a compile-time or run-time execution system of a data parallelled split-function program with function calls. In one embodiment, the system comprises: (1) a partition unit configured to partition chains into a parent chain and at least one worker chain, and (2) a thread assignment unit connected and configured with the partition unit only to designate a single strand from the parent chain for execution and all strands in the at least one worker chain for execution.
Ein weiterer Aspekt stellt ein Verfahren zur Kompilierung oder Laufzeitausführung eines datenparallelen Programms mit Aufteilung-Vereinigung mit Funktionsaufrufen bereit. In einer Ausführungsform umfasst das Verfahren: (1) Unterteilen von Ketten in eine übergeordnete Kette und mindestens eine Arbeiter-Kette und (2) Anweisen bzw. Bestimmen nur eines einzelnen Strangs aus der übergeordneten Kette zur Ausführung und aller Stränge in der mindestens einen Arbeiter-Kette für die Ausführung.Another aspect provides a method for compiling or executing runtime a data parallels program with partitioning association with function calls. In one embodiment, the method comprises: (1) dividing chains into a parent chain and at least one worker chain, and (2) instructing only a single strand from the parent chain to execute and all strands in the at least one worker chain Chain for the execution.
KURZE BESCHREIBUNG SHORT DESCRIPTION
Es wird nun auf die folgenden Beschreibungen in Verbindung mit den begleitenden Zeichnungen verwiesen, in denen:Reference is now made to the following descriptions taken in conjunction with the accompanying drawings, in which:
DETAILLIERTE BESCHREIBUNGDETAILED DESCRIPTION
In einem SIMT-Prozessor werden mehrere Ausführungs-Stränge in Gruppen unterteilt. Alle Stränge in der Gruppe führen den gleichen Befehl zur gleichen Zeit aus. In grafischen Verarbeitungseinheiten (GPUs), die kommerziell von der Firma Nvidia, Santa Clara, Kalifornien erhältlich sind, die eine Art von SIMT-Prozessoren sind, werden Gruppen als „Wölbungen bzw. Ketten” bezeichnet, und sie werden in Blöcken ausgeführt.In a SIMT processor, multiple execution strands are divided into groups. All strands in the group execute the same command at the same time. In graphical processing units (GPUs), commercially available from Nvidia, Santa Clara, California, which are a type of SIMT processor, groups are referred to as "bumps" and are executed in blocks.
Eine Pipeline-Steuereinheit des SIMT-Prozessors erzeugt, verwaltet, disponiert, führt aus und stellt bereit einen Mechanismus, um Gruppen zu synchronisieren. Nvidia-GPUs stellen einen bar.sync-Befehl zum Synchronisieren von Gruppen bereit. Nvidia-GPUs unterstützen ferner die Ausführung einer „divergenten” bedingten Verzweigung durch eine Gruppe; einige Stränge der Gruppe müssen die Verzweigung nehmen (da die Verzweigungsbedingungs-Vorbestimmung ein „wahr” ermittelt), und die anderen Stränge müssen auf den nächsten Befehl springen (da die Verzweigungsbedingungs-Vorbestimmung ein „falsch” ermittelt). Die Pipeline-Steuereinheit überwacht aktive Stränge in der Gruppe. Sie führt zuerst einen der Pfade (Verzweigung genommen oder Verzweigung nicht genommen) aus und führt dann den anderen Pfad aus; die entsprechenden Stränge werden in jedem Pfad aktiviert.A SIMT processor pipeline controller generates, manages, schedules, executes, and provides a mechanism to synchronize groups. Nvidia GPUs provide a bar.sync command to synchronize groups. Nvidia GPUs also support the execution of a "divergent" conditional branch by a group; some strands of the group must take the branch (since the branch condition predetermination determines a "true") and the other strands must jump to the next command (since the branch condition predetermination determines a "false"). The pipeline control unit monitors active threads in the group. It first takes one of the paths (taken branch or branch not taken) and then executes the other path; the corresponding strands are activated in each path.
Es wird hierin erkannt, dass, während alle Stränge innerhalb eines GPU-Strangblocks an der gleichen Programmadresse starten, die Pipeline-Steuereinheit von einem Software-Mechanismus profitieren würde, der die Stränge in einen Haupt-Strang und Arbeiter-Stränge unterteilt und disponiert, so dass diese in dem Aufteilung-Vereinigungs-Modell ausführbar sind.It is recognized herein that while all threads within a GPU thread block start at the same program address, the pipeline control unit would benefit from a software mechanism that divides and schedules the threads into a main thread and worker threads that these are executable in the split-union model.
Es wird ferner hierin erkannt, dass gewisse Ausführungsformen des Software-Mechanismus die Stränge als Gruppen verwalten und synchronisieren sollten, da die Pipeline-Steuereinheit die Stränge als Gruppen verwaltet.It is further appreciated herein that certain embodiments of the software mechanism should manage and synchronize the strands as groups because the pipeline control unit manages the strands as groups.
Es wird ferner hierin erkannt, dass, da das Hauptprogramm ein Einzel-Strang-Programm in dem Aufteilung-Vereinigungs-Modell ist, gewisse Ausführungsformen des Software-Mechanismus eine Semantik eines sequenziellen Bereichs erreichen sollten, ohne dass Nebenwirkungen eingeführt werden. Beispiele von Nebenwirkungen hervorrufenden Befehlen sind jene, die gemeinsam benutzte Ressourcen verwenden, etwa einen Lesebefehl oder einen Schreibbefehl für einen gemeinsam benutzten Speicher oder eine beliebige Code-Operation, die eine gemeinsam genutzte Ausnahmebehandlung (beispielsweise Division) aufrufen kann.It is further recognized herein that since the main program is a single-strand program in the split-merge model, certain embodiments of the software mechanism should achieve semantics of a sequential range without introducing side-effects. Examples of side-effecting commands are those that use shared resources, such as a read command or a shared memory write command, or any code operation that can call a shared exception handler (for example, division).
Es wird ferner hierin erkannt, dass gewisse Ausführungsformen des Software-Mechanismus Funktionen unterstützen sollten, die innerhalb des sequenziellen Bereichs und innerhalb des parallelen Bereichs aufgerufen werden können. Und derartige Funktionen können selbst parallele Konstrukte enthalten. Es wird ferner hierin erkannt, dass gewisse Ausführungsformen des Software-Mechanismus Funktionsaufrufe unterstützen sollten, die parallele Bereiche aufteilen-vereinigen können.It is further recognized herein that certain embodiments of the software mechanism should support functions that may be invoked within the sequential area and within the parallel area. And such functions may even contain parallel constructs. It is further appreciated herein that certain embodiments of the software mechanism should support function calls that may divide-unite parallel portions.
Weiterhin wird hierin erkannt, dass gewisse Ausführungsformen des Software-Mechanismus außenstehende bzw. äußere Funktionen unterstützen sollten, das heißt Funktionen, die nicht von dem gleichen Compiler wie das Programm kompiliert werden. Beispielsweise mathematische Funktionen in bestehenden GPU-Mathematik-Bibliotheken, und Systemfunktionen wie malloc, free und print. In gewissen Ausführungsformen sollten sowohl der übergeordnete Strang in den sequenziellen Bereichen als auch die Arbeiter-Stränge in dem parallelen Bereich in der Lage sein, eine äußere Funktion aufzurufen.Further, it is recognized herein that certain embodiments of the software mechanism should support outside functions, that is, functions that are not compiled by the same compiler as the program. For example, mathematical functions in existing GPU math libraries, and system functions such as malloc, free, and print. In certain embodiments, both the parent strand in the sequential regions and the worker strands in the parallel region should be able to invoke an outside function.
Folglich sind hierin diverse Ausführungsformen eines Systems und eines Verfahrens zur Kompilierung und Ausführung von datenparallelen Programmen mit Aufteilung-Vereinigung und mit Funktionsaufrufen in einem SIMT-Prozessor, etwa einer GPU, beschrieben. Thus, various embodiments of a system and method for compiling and executing data parallels with partition sharing and function calls in a SIMT processor, such as a GPU, are described herein.
Vor der Beschreibung gewisser Ausführungsformen des Systems und des Verfahrens wird ein SIMT-Prozessor beschrieben, der ausgebildet ist, ein System oder ein Verfahren zur Kompilierung oder zur Laufzeitausführung von datenparallelen Programmen mit Aufteilung-Vereinigung mit Funktionsaufrufen zu enthalten oder auszuführen.Prior to describing certain embodiments of the system and method, a SIMT processor is described that is configured to include or execute a system or method for compiling or executing data-parallel split-union programs with function calls.
Der SIMT-Prozessor
Die Ausführungsform aus
Nach der Beschreibung eines SIMT-Prozessors, in welchem das System oder das Verfahren, wie sie hierin eingeführt sind, enthalten oder ausgeführt werden kann, werden nunmehr diverse Ausführungsformen des Systems und des Verfahrens beschrieben.Having described a SIMT processor in which the system or method as incorporated herein can be incorporated or embodied, various embodiments of the system and method will now be described.
Eine Ausführungsform des hierin eingeführten Systems umfasst einen Compiler und eine Geräte-Laufzeitbibliothek. Die Geräte-Laufzeitbibliothek realisiert die Funktion der Strang- und Gruppenverwaltung. Der Compiler übersetzt ein datenparalleles Programm mit Aufteilung-Vereinigung in ein Haupt-Strangprogramm und eine Gruppe aus äußeren Funktionen, wovon jede einem parallelen Konstrukt entspricht. Der übersetzte Code nimmt Aufrufe von Funktionen der Geräte-Laufzeitbibliothek vor, um die Strang- und Gruppenverwaltung auszuführen.One embodiment of the system introduced herein includes a compiler and a device runtime library. The device runtime library implements the function of thread and group management. The compiler translates a data-parallel splitting-union program into a main thread program and a set of outer functions, each corresponding to a parallel construct. The translated code makes calls to device runtime library functions to perform thread and group management.
Die nachfolgende Tabelle 2 zeigt ein Beispielprogramm, um die Compiler-Übersetzung und die Realisierung der Geräte-Laufzeit darzustellen.Table 2 below shows an example program to illustrate the compiler translation and device runtime implementation.
Tabelle 2 – Beispielprogramm für die Compiler-Übersetzung und die Realisierung der Geräte-Laufzeit Table 2 - Example program for the compiler translation and the realization of the device runtime
Der Ablauf des main()-Programms aus Tabelle 2 beginnt mit dem einzelnen übergeordneten Strang. Der übergeordnete Strang ruft die Funktion foo() auf, die einen Körper aufweist, der für diesen Compiler transparent ist und von diesem übersetzt wird. Der übergeordnete Strang ruft dann eine Funktion ext() auf, d. h. eine externe oder außen liegende Funktion mit einem Körper, der für diesen Compiler unsichtbar ist. Aufrufe von außen liegenden Funktionen werden übersetzt, wie sie sind, ohne dass eine spezielle Handhabung durch den Compiler erfolgt. Der übergeordnete Strang trifft dann auf den ersten parallelen Bereich. Arbeiter-Stränge führen den parallelen Bereich aus, während der übergeordnete Strang auf deren Beendigung wartet. Innerhalb des parallelen Bereichs ruft jeder Arbeiter-Strang die Funktion foo() und bar() auf. Die Funktion bar() enthält einen weiteren parallelen Bereich; jedoch liegt bar() bereits innerhalb eines parallelen Bereichs. Da bar() bereits innerhalb eines Parallelenbereichs liegt, wird der parallele Bereich innerhalb von bar() sequenziell von jedem Arbeiter-Strang ausgeführt.The flow of the main () program from Table 2 begins with the single parent thread. The parent thread calls foo (), which has a body that is transparent to and translated by this compiler. The parent thread then calls an ext () function, d. H. an external or external function with a body that is invisible to this compiler. External function calls are translated as they are without any special handling by the compiler. The parent strand then encounters the first parallel region. Worker strands execute the parallel area while the parent strand is waiting to complete. Within the parallel area, each worker strand calls foo () and bar (). The function bar () contains another parallel area; however, bar () is already within a parallel range. Since bar () is already within a parallel region, the parallel region within bar () is executed sequentially by each worker strand.
Nach dem ersten parallelen Bereich trifft der übergeordnete Strang auf einen zweiten parallelen Bereich. Innerhalb des zweiten parallelen Bereichs ruft jeder Arbeiter-Strang die externe außen liegende Funktion ext() auf. Nach dem zweiten parallelen Bereich ruft der übergeordnete Strang die Funktion bar() auf. Innerhalb von bar() trifft der übergeordnete Strang auf einen dritten parallelen Bereich, der wiederum von den Arbeiter-Strängen abgearbeitet wird.After the first parallel region, the parent strand strikes a second parallel region. Within the second parallel area, each worker thread calls the external external function ext (). After the second parallel area, the higher-level thread calls the bar () function. Within bar (), the parent strand encounters a third parallel region, which in turn is processed by the worker strands.
Die Funktion main() ist als eine Eintrittsfunktion bekannt, da dies der Punkt ist, an welchem das Programm startet. Funktionen, etwa foo() und bar() sind Nicht-Eintrittsfunktionen. The function main () is known as an entry function, since this is the point where the program starts. Functions such as foo () and bar () are non-entry functions.
Für eine Eintrittsfunktion stellt der Compiler zuerst eine geklonte Kopie, die als main_core() bezeichnet wird. Die geklonte Kopie wird dann als eine Nicht-Eintrittsfunktion verarbeitet, wie nachfolgend beschrieben ist. Für die main()-Funktion erzeugt der Compiler einen Code, der in Tabelle 3 nachfolgend gezeigt ist, wobei groupID() die ID der Stranggruppe zurückgibt, die einen Befehl ausführenden Strang enthält. threadID() gibt die ID des Strangs zurück. init(), signal_done() und scheduler() sind Funktionen in der Geräte-Laufzeitbibliothek.For an entry function, the compiler first makes a cloned copy called main_core (). The cloned copy is then processed as a non-entry function, as described below. For the main () function, the compiler generates a code shown in Table 3 below, where groupID () returns the ID of the thread group containing a thread executing command. threadID () returns the ID of the thread. init (), signal_done () and scheduler () are functions in the device runtime library.
Tabelle 3 – Compiler-erzeugter Beispiels-Code Table 3 - Compiler Generated Example Code
Wenn ein GPU-Strangblock startet, führen alle Stränge innerhalb des Blocks main() aus; jedoch nehmen sie unterschiedliche Pfade. Der Strang 0 ist der übergeordnete Strang und führt init(), main_core() und signal_done() aus. Andere Stränge in der Gruppe 0 gehen geradewegs zum Ende der main()-Funktion und warten dort. Die Stränge in den verbleibenden Gruppen führen scheduler() aus.When a GPU thread block starts, all strands within the block execute main (); however, they take different paths. String 0 is the parent thread and runs init (), main_core (), and signal_done (). Other strands in group 0 go straight to the end of the main () function and wait there. The strands in the remaining groups execute scheduler ().
Für eine Nicht-Eintrittsfunktion wie foo(), bar() und main core() übersetzt der Compiler den Code so, als ob kein paralleles Konstrukt existieren würde. Wenn eine Nicht-Eintrittsfunktion ein paralleles Konstrukt enthält, dann erzeugt der Compiler für jedes parallele Konstrukt eine Funktion, die den Körper des parallelen Konstrukts enthält (eine ausgelagerte Funktion), und erzeugt dann eine bedingte Verzweigung, die prüft, ob der gerade ausgeführte Strang der übergeordnete Strang ist. In der falschen Verzweigung fügt der Compiler Code ein, der die Schleife ausführt. In der wahren Verzweigung fügt der Compiler Aufrufe in die Geräte-Laufzeitbibliothek ein, um Aufgaben zuzuordnen, Arbeiter-Stränge aufzuwecken und eine Barriere auszuführen. Die Bedingung ist wahr, wenn die Nicht-Eintrittsfunktion außerhalb der parallelen Bereiche aufgerufen wird. Die Bedienung ist falsch, wenn die Nicht-Eintrittsfunktion innerhalb eines parallelen Bereichs aufgerufen wird, in welchem Falle die parallele Schleife von dem ausführenden Strang sequenziell ausgeführt wird.For a non-entry function such as foo (), bar (), and main core (), the compiler translates the code as if no parallel construct existed. If a non-entry function contains a parallel construct, then for each parallel construct the compiler generates a function containing the body of the parallel construct (a paged function) and then generates a conditional branch that checks whether the thread currently executing is is parent strand. In the wrong branch, the compiler inserts code that executes the loop. In true branching, the compiler inserts calls into the device runtime library to map tasks, wake worker strands, and execute a barrier. The condition is true if the non-entry function is called outside the parallel areas. The operation is false when the non-entry function is called within a parallel area, in which case the parallel loop is executed sequentially by the executing strand.
Zum Beispiel ist der übersetzte Code für die Funktion bar() nachfolgend in Tabelle 4 gezeigt.For example, the translated code for the function bar () is shown in Table 4 below.
Tabelle 4 – übersetzter Code für die Funktion bar () Table 4 - translated code for the function bar ()
signal task() und barrier() sind Funktionen in der Geräte-Laufzeitbibliothek. bar_par_frunc() ist die ausgelagerte Funktion, die dem parallelen Konstrukt in der ursprünglichen Funktion bar() entspricht.signal task () and barrier () are functions in the device runtime library. bar_par_frunc () is the paged function that corresponds to the parallel construct in the original function bar ().
In dieser Ausführungsform enthält die Geräte-Laufzeitbibliothek unter anderem die folgenden Funktionen: init(), scheduler(), signal task(), signal done(), und barrier(). Die Bibliothek realisiert auch die folgenden Funktionen für die interne Verwendung: signal(), wait() und fetch_task().In this embodiment, the device runtime library includes, among other things, the following functions: init (), scheduler (), signal task (), signal done (), and barrier (). The library also implements the following functions for internal use: signal (), wait (), and fetch_task ().
Alle Arbeiter-Stränge führen die scheduler()-Funktion aus. Die Arbeiter-Stränge durchlaufen einen Schlaf-Aufwach-Ausführen-Zyklus, bis sie angewiesen werden, zu enden.All worker strands execute the scheduler () function. The worker strands go through a sleep-wake-up-run cycle until instructed to end.
Tabelle 5 – Beispiel-Code unter Verwendung einer Marke zum Verlassen des Programms Table 5 - Example code using a mark to exit the program
Es wird eine Bool'sche Variable 'exit_flag' in dem von Blöcken gemeinsam benutzten Speicher abgelegt und auf diese kann von allen Strängen innerhalb des Strangblocks zugegriffen werden. Sie wird von dem übergeordneten Strang verwendet, um den Arbeiter-Strängen mitzuteilen, ob sie alle die Ausführung beenden sollen. Die 'exit_flag' wird in der init()-Funktion auf falsch gesetzt, und wird in der signal_done()-Funktion auf wahr gesetzt. Beide Funktionen werden von dem übergeordneten Strang aufgerufen.A Boolean variable 'exit_flag' is stored in the memory shared by blocks and can be accessed by all threads within the thread block. It is used by the parent thread to tell the worker strands if they should all finish the execution. The 'exit_flag' is set to false in the init () function, and is set to true in the signal_done () function. Both functions are called by the parent thread.
Tabelle 6 – Beispiel-Code für Änderung des Zustands der Marke für das Verlassen des Programms Table 6 - Example code for changing the state of the program exit program
Ein weiterer Teil des von Blöcken gemeinsam benutzten Speichers wird verwendet, um die aktuelle Aufgabe mitzuteilen. Die aktuelle Aufgabe wird von dem übergeordneten Strang in der signal_task()-Funktion festgelegt und wird von dem Arbeiter-Strang in der fetch_task()-Funktion abgeholt. Der von Blöcken gemeinsam benutzte Speicher enthält den Zeiger auf die ausgelagerte Funktion, die dem parallelen Konstrukt entspricht.Another part of the memory shared by blocks is used to communicate the current task. The current task is set by the parent thread in the signal_task () function and is fetched by the worker thread in the fetch_task () function. The memory shared by blocks contains the pointer to the paged function that corresponds to the parallel construct.
Tabelle 7 – Beispiel-Code zur Verwendung eines Zeigers, um eine aktuelle Aufgabe zu kennzeichnen Table 7 - Example code for using a pointer to identify a current task
Da die parallelen Bereiche der Reihe nach innerhalb eines Strangblocks ausgeführt werden, ist zu jedem Zeitpunkt nur eine einzelne Aufgabe aktiv. Wenn die parallelen Bereiche asynchron ausgeführt werden können, wird typischerweise eine kompliziertere Datenstruktur, etwa ein Stapel, eine Warteschlange oder ein Datenbaum benötigt, um die aktiven Aufgaben zu speichern.Since the parallel areas are executed sequentially within a thread block, only a single task is active at any one time. If the parallel areas can be executed asynchronously, typically a more complicated data structure, such as a stack, queue, or tree, is needed to store the active tasks.
barrier(), signal() und wait()-Funktionen werden unter Anwendung einer Hardware-Barriere realisiert.barrier (), signal (), and wait () functions are implemented using a hardware barrier.
Tabelle 8 – Beispielfunktionen für barrier(), signal() und wait() Table 8 - Example functions for barrier (), signal (), and wait ()
Der SIMT-Prozessor
Die Partitionseinheit
Die Strang-Disponiereinheit
Der Funktionsprozessor
Der Funktionsprozessor
Ähnlich zu der Verarbeitung von Nicht-Eintrittsfunktionen, die kein paralleles Konstrukt aufweisen, werden äußere Funktionen
Ein datenparalleles Programm mit Aufteilung-Vereinigung wird in ein übergeordnetes Programm und eine Gruppe aus parallelen Aufgaben unterteilt. Das übergeordnete Programm ist das Programm, das von dem übergeordneten Strang ausgeführt wird. Eine parallele Aufgabe entspricht einem parallelen Bereich, der von den Arbeiter-Strängen ausgeführt wird. Das übergeordnete Programm enthält Disponierpunkte, an denen der übergeordnete Strang eine parallele Aufgabe zuweisen bzw. bestimmen wird, die Arbeiter-Stränge aufwecken wird und darauf wartet, dass die Arbeiter-Stränge abgeschlossen sind.A data-parallel program with split-union is divided into a higher-level program and a group of parallel tasks. The parent program is the program that is executed by the parent thread. A parallel task corresponds to a parallel area performed by the worker strands. The parent program contains scheduling points where the parent thread will assign a parallel task that will wake worker strands and wait for the worker strands to complete.
Der spezielle übergeordnete Strang in der speziellen übergeordneten Gruppe wird den sequenziellen Bereich des Programms ausführen.The special parent thread in the special parent group will execute the sequential part of the program.
Alternativ kann ein Einzel–Strangverhalten in dem sequenziellen Bereich nachgebildet werden, während alle Stränge in der Gruppe den Code ausführen. Jedoch weisen Nachbildungsschemata Beschränkungen im Hinblick auf die Komplexität hinsichtlich des Leistungsverhaltens und der Erzeugung auf, so dass diese weniger brauchbar sind. Die notwendige Vorbestimmung und Synchronisierung ergibt zusätzlichen Aufwand bei der Ausführung. Ferner müssen alle Funktionen, die von dem sequenziellen Bereich und dem parallelen Bereich aufgerufen werden, unterschiedlich geklont und kontrolliert werden.Alternatively, a single-strand behavior may be replicated in the sequential region while all strands in the group execute the code. However, replication schemes have limitations in terms of performance and generation complexity, making them less useful. The necessary predetermination and synchronization results in additional effort in the execution. Furthermore, all functions called from the sequential area and the parallel area must be cloned and controlled differently.
Bei gegebener Unterteilung der Stränge und Gruppen vollführen die Arbeiter-Stränge und der übergeordnete Strang die folgenden Lebenszyklen aus: Eine Ausführungsform eines Arbeiter-Strangs durchläuft die folgenden Phasen in einem Lebenszyklus:
- 1) der Strangblock startet;
- 2) im Schlafmodus, bis das Aufwecken durch den übergeordneten Strang erfolgt;
- 3) Verlassen des Programms, wenn die Marke zum Verlassen des Programms auf wahr gesetzt ist;
- 4) Abholen und Ausführung der Aufgabe, die von dem übergeordneten Strang zugewiesen wird;
- 5) Eintreten in eine Barriere und
- 6) Zurückkehren zur Stufe 2.
- 1) the strand block starts;
- 2) in sleep mode until awakening by the parent strand;
- 3) exit the program if the mark to exit the program is true;
- 4) fetching and executing the task assigned by the parent thread;
- 5) entering a barrier and
- 6) Return to level 2.
Eine Ausführungsform eines übergeordneten Strangs durchläuft die folgenden Stufen in einem Lebenszyklus:
- 1) der Strangblock startet;
- 2) Setzen der Marke für das Verlassen des Programms auf falsch;
- 3) Ausführen des übergeordneten Programms, bis ein paralleler Bereich oder das Ende des übergeordneten Programms erreicht werden;
- 4) am Anfang eines parallelen Bereichs:
- a. Festlegen einer parallelen Aufgabe,
- b. Aufwecken der Arbeiter-Stränge,
- c. Eintreten in eine Barriere, und
- d. Fortsetzen des übergeordneten Programms (Stufe 3); und
- 5) am Ende des übergeordneten Programms:
- a. Setzen der Marke für das Verlassen eines Programms auf wahr,
- b. Aufwecken der Arbeiter-Stränge, und c. Ende.
- 1) the strand block starts;
- 2) set the mark for leaving the program to false;
- 3) execute the parent program until a parallel area or the end of the parent program is reached;
- 4) at the beginning of a parallel area:
- a. Defining a parallel task,
- b. Awakening the worker strands,
- c. Entering a barrier, and
- d. Continue the parent program (level 3); and
- 5) at the end of the parent program:
- a. Setting the mark for leaving a program to true,
- b. Awakening the worker strands, and c. The End.
Die anderen Stränge in der übergeordneten Gruppe warten im Wesentlichen auf das Ende des Programms in untätiger Weise. Das Programm wird abwechselnd von dem übergeordneten Strang und den Arbeiter-Strängen ausgeführt. Dies führt zu einer guten Auslastung des Befehls-Cache-Speichers, die besser ist als die Auslastung, die durch ein Verfahren hervorgerufen wird, in welchem sowohl der übergeordnete Strang als auch die Arbeiter-Stränge aktiv sind und unterschiedliche Code-Pfade ausführen.The other threads in the parent group are essentially waiting for the end of the program to idle. The program is executed alternately by the parent strand and the worker strands. This results in a good utilization of the instruction cache that is better than the load caused by a method in which both the parent thread and the worker threads are active and execute different code paths.
Der Fachmann auf diesem Gebiet, an den sich diese Anmeldung richtet, erkennt, dass andere und weitere Hinzufügungen, Streichungen, Ersetzungen und Modifizierungen an den beschriebenen Ausführungsformen vorgenommen werden können.Those skilled in the art to which this application pertains recognize that other and further additions, deletions, substitutions, and alterations can be made to the described embodiments.
Claims (10)
Applications Claiming Priority (4)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
US201261722661P | 2012-11-05 | 2012-11-05 | |
US61/722,661 | 2012-11-05 | ||
US13/724,359 | 2012-12-21 | ||
US13/724,359 US9747107B2 (en) | 2012-11-05 | 2012-12-21 | System and method for compiling or runtime executing a fork-join data parallel program with function calls on a single-instruction-multiple-thread processor |
Publications (1)
Publication Number | Publication Date |
---|---|
DE102013018380A1 true DE102013018380A1 (en) | 2014-05-08 |
Family
ID=50489874
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
DE201310018380 Pending DE102013018380A1 (en) | 2012-11-05 | 2013-11-04 | System for compilation or run-time execution of data-parallel program in single-instruction multi-strand processor, has strand association unit connected to partition unit and determining strand from parent group to run as parent strand |
Country Status (1)
Country | Link |
---|---|
DE (1) | DE102013018380A1 (en) |
-
2013
- 2013-11-04 DE DE201310018380 patent/DE102013018380A1/en active Pending
Similar Documents
Publication | Publication Date | Title |
---|---|---|
DE68921906T2 (en) | Method for a multiprocessor system with self-assigning processors. | |
DE69909829T2 (en) | MULTIPLE PROCESSOR FOR THREAD SOFTWARE APPLICATIONS | |
DE68925646T2 (en) | PIPELINE MULTIPROCESSOR SYSTEM | |
DE102015112202A1 (en) | Combining paths | |
DE4217012C2 (en) | Microprocessor operating with a variety of instruction streams and static interleaving | |
DE69032418T2 (en) | Private thread storage in a multi-filament digital data processing system | |
DE102018126003A1 (en) | Combine states of multiple threads in a multithreaded processor | |
DE102018126001A1 (en) | Synchronization in a multi-tile processing array | |
DE102012220267B4 (en) | Arithmetic work distribution - reference counter | |
DE112011101391T5 (en) | GPU-enabled database systems | |
DE102018126004A1 (en) | Synchronization in a multi-tile processing arrangement | |
DE102017213160B4 (en) | Compilation for node device GPU-based parallel processing | |
DE102012221502A1 (en) | A system and method for performing crafted memory access operations | |
DE112012002465T5 (en) | Graphics processor with non-blocking concurrent architecture | |
DE102013202173A1 (en) | Uniform load processing for subsets of parallel threads | |
DE102018109538A1 (en) | Techniques for fully synchronizing a thread execution | |
DE102013214756A1 (en) | METHOD AND DEVICE FOR IMPROVING THE PROCESSING CAPACITY OF A MULTI-CORE PROCESSOR | |
DE112012004728T5 (en) | Method, program and system for simulation execution | |
EP0825540B1 (en) | Pipeline processor | |
DE102013015658A1 (en) | System and method for starting callable functions | |
DE102012222391B4 (en) | Multichannel Time Slice Groups | |
DE102016221526A1 (en) | Apparatus and method for processing a plurality of tasks | |
Binet et al. | Multicore in production: Advantages and limits of the multiprocess approach in the ATLAS experiment | |
EP2386949A1 (en) | Method and device for allocating a number of sub-tasks from a task to a number of computing units in a pre-defined processor architecture | |
DE102013017514A1 (en) | WORKING FOR GRAPHIC PROCESSING UNIT BASED ON WORKING SERVICE TIMES |
Legal Events
Date | Code | Title | Description |
---|---|---|---|
R012 | Request for examination validly filed | ||
R082 | Change of representative |
Representative=s name: KRAUS & WEISERT PATENTANWAELTE PARTGMBB, DE |