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

Abstract

The system (2) has a partition unit (202) designed to divide groups in a parent group and a working group. A strand association unit (204) is connected to the partition unit. The strand association unit determines a strand from the parent group to run as a parent strand. The strand association unit determines all of the strands in the working group to run as workers strands. The strand association unit is connected to a strand disposition unit (206). A function processor (208) is coupled to the strand association unit for generating and processing an ingress function as a non-entry function.

Description

CROSS-REFERENCE TO RELATED APPLICATION

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.

TECHNICAL AREA

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.

BACKGROUND

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.

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.

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.

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

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".

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".

OVERVIEW

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.

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.

SHORT DESCRIPTION

Reference is now made to the following descriptions taken in conjunction with the accompanying drawings, in which:

1 FIG. 4 is a block diagram of a SIMT processor configured to include or execute a system or method for compiling or executing runtime a data-parallel program with split-branch and with function calls;

2 Figure 12 is a block diagram of one embodiment of a compile-time or run-time execution system of a split-branch data-parallel program with function calls; and

3 Figure 3 is a flowchart of one embodiment of a method for compilation or run-time execution of a data-parallel program with split-branch and with function calls.

DETAILED DESCRIPTION

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.

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.

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.

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.

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).

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.

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.

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.

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.

1 is a block diagram of a SIMT processor 100 , The SIMT processor 100 contains several strand processors or cores 106 that are in strand groups 104 or "bulges or chains" are divided. The SIMT processor 100 contains J string groups 104-1 to 104-J , of which each K cores 106-1 to 106-K having. In certain embodiments, the strand groups 104-1 to 104-J further into one or more strand blocks 102 be divided. Certain embodiments include thirty-two cores 106 per strand group 104 , Other embodiments may include less than four cores in a strand group and up to several tens of thousands of cores. Certain embodiments share the cores 106 into a single strand group 104 while other embodiments include hundreds or even thousands of strand groups 104 exhibit. Alternative embodiments of the SIMT processor 100 can the cores 106 exclusively in the strand groups 104 divide, where the division on the level of strand blocks is not present.

The SIMT processor 100 further includes a pipeline control unit 108 , a memory shared by blocks 110 and an array of local stores 112-1 to 112-J that are the strand groups 104-1 to 104-J assigned. The pipeline control unit 108 distributes tasks to the various strand groups 104-1 to 104-J via a data bus 114 , The cores 106 in a strand group 106-j work parallel to each other. The strand groups 104-1 to 104-J communicate via a memory bus 116 with memory shared by blocks 110 , The strand groups 104-1 to 104-J communicate accordingly with the local memories 112-1 to 112-J via local buses 118-1 to 118-J , For example, a strand group uses 104-J the local store 112-J by means of communication via a local bus 118-J , Certain embodiments of the SIMT processor 100 have a shared portion of the memory shared by blocks 110 every strand block 102 and allow access to shared portions of the block-shared memory 110 for all strand groups 104 within a strand block 102 , Certain embodiments include strand groups 104 that exclusively uses the local store 112 to use. Many other embodiments include strand groups 104 that balance the local store 112 and the memory shared by blocks 110 accomplish.

The embodiment of 1 contains a parent strand group 104-1 , Each of the remaining strand groups 104-2 to 104-J is considered a "worker" strand group. The parent strand group 104-1 contains numerous cores, one of which is a superordinate core 106-1 who eventually performs a parent strand. Programs included in the SIMT processor 100 are executed as a sequence of kernels. Typically, each kernel ends its execution before the next kernel begins. In certain embodiments, the SIMT processor 100 run multiple kernels in parallel, depending on the size of the kernel. Each kernel is built as a hierarchy of strands that are in the cores 106 are to be executed.

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.

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.

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

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.

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.

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.

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

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 ().

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.

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 () 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 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 ().

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

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

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

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 (), 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 ()

2 Figure 13 is a block diagram of one embodiment of a system 200 for compilation or runtime execution of a data-parallel program under splitting-union with function calls. The one program 210 contains entry functions 212 , Non-entry functions 214 and external functions 216 , The system 200 includes a partition unit 202 , a strand allocation unit 204 , a string scheduler 206 , a functional processor 208 , a device runtime library 218 and a SIMT processor 100 out 1 ,

The SIMT processor 100 includes the pipeline control unit 108 , the data bus 114 , the local buses 118-1 and 118-2 and the shared memory 110 out 1 , In the embodiment of 2 is the SIMT processor 100 shown to have a single strand block containing two strand groups: the parent strand group 104-1 and the workers strand group 104-2 , The strand groups 104-1 and 104-2 each contain strands 106 ,

The partition unit 202 determines or assigns the strand group 104-1 as the parent strand group and the remaining strand groups as worker strand groups. In the embodiment of 2 is a single worker strand group 104-2 shown. In alternative embodiments, many worker strand groups can be used. The strand allocation unit 204 determines a parent strand 106-1 the parent strand group 104-1 , All other strands in the parent strand group 104-1 are idle. The strand allocation unit 204 also determines each of the strands 106 in the workers strand group 104-2 as workers strands out.

The string scheduler 206 translates the program 210 such that the pipeline control unit 108 the execution of the parent strand 106-1 and the various worker strands in the workers strand group 104-2 controlled in a suitable manner. The string scheduler 206 translates the program 210 such that when the parent thread execution begins, a program exit flag is reset. The string scheduler 206 schedules the parent thread 106-1 so that it runs until a parallel area or the end of the program 210 are reached. If a parallel area of the program 210 is reached, sets the strand scheduler 206 a parallel task and the worker strands in the worker strand group 104-2 start with the execution. The string scheduler 206 Further establishes a barrier for each of the worker strands so that when the barrier is entered, the parent strand 106-2 continues its execution. When the end of the program 210 is reached, the program exit flag is activated, causing all worker strands to finish execution.

The function processor 208 operates on the functions of the program 210 , The processing of the entry functions 212 involves generating a clone copy of an entry function, which is then processed as a non-entry function. The original entry function is processed so that the parent strand 106-1 the clone copy will process in addition to other calls, and the worker strands will execute sleep mode, wakeup, retrieve, and parallel task cycles performed by the strand dispatcher 206 is fixed.

The function processor 208 translates the non-entry functions 214 in two ways. If there is no parallel construct in a non-entry function, then the function is simply processed as it is. When a parallel construct exists, an outer function containing the body of the parallel construct is created. The function processor 208 then creates a branch condition that either executes the parallel construct sequentially or the device runtime library 218 used to assign a task to wake up worker strands and perform a barrier as previously described. The wakeup and sleep functions are performed using hardware barrier functions of the device runtime library 218 realized. Strands on barriers are not scheduled for execution by the hardware, so they do not waste work cycles. In the parent strand group 104-1 takes only the parent strand 106-1 participate in the barriers. This is because the hardware barrier is group based. A group is considered to be at a barrier when each strand within the group is at the barrier.

Similar to the processing of non-entry functions that do not have a parallel construct, they become outer functions 216 from the function processor 208 processed as they are.

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.

The special parent thread in the special parent group will execute the sequential part of the program.

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.

• 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.

Given the subdivision of the strands and groups, the worker strands and parent strand perform the following life cycles: An embodiment of a worker strand goes through the following phases in a life cycle:

• 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.

• 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.

One embodiment of a parent thread goes through the following stages in a lifecycle:

• 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.

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.

3 FIG. 10 is a flowchart of one embodiment of a method for compiling or running-time execution of a data-parallel program under partition-union with function calls. The process starts in a starting step 310 , In one step 320 strand groups within a strand block are split into a parent strand group and at least one worker strand group. In one step 330 a single thread from the parent thread group is designated or designated as the parent thread. The remaining strands of the parent group are essentially idle during execution. Further, in step 330 all strands in the at least one worker strand group are designated or determined as worker strands. The procedure ends in a final step 340 ,

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)

A system for compiling or running a data-parallel program with split-union with function calls, with: a partition unit configured to divide groups into a parent group and at least one worker group; and a strand allocation unit connected to the partition unit and configured to designate only one thread from the parent group for execution as a parent thread and to determine all threads in the at least one worker group for execution as worker threads. The system of claim 1, further comprising a thread scheduler coupled to the thread allocation unit and configured to cause a single-instruction multi-strand processor pipeline control unit to schedule the execution of the parent thread as follows: a program exit flag is set to a first state when the parent strand begins execution; the parent thread is executed until it reaches a parallel area or an end of the program; upon reaching the parallel area, a parallel task is set, entered a first barrier, entered a second barrier, and continued to execute the parent strand; and on reaching the end, the flag is set to a second state to exit a program, entered a barrier, and the parent thread terminates execution. The system of claim 1 or 2, further comprising a strand scheduler coupled to the strand allocation unit and configured to cause a single-instruction multi-strand processor pipeline control unit to execute the worker threads as follows to schedule: the worker strands enter a first barrier; the worker strands begin execution of a parallel task specified by the parent thread when the parent thread reaches a parallel region and the parent strand enters a barrier; the worker strands enter a second barrier at the completion of the parallel task; and the worker strands end when a program exit flag is set to a second state and the parent strand enters a barrier. The system of claim 1, further comprising a strand scheduler coupled to and configured with the strand mapper, a barrier function of a single-instruction multi-strand pipeline controller to control execution and termination of the pipeline To use worker strands. The system of any one of claims 1-4, further comprising a function processor coupled to the strand allocation unit and configured to generate a cloned copy of an entry function and to process the entry function as a non-entry function. The system of any one of claims 1-5, further comprising a function processor coupled to the strand mapper and configured to translate a non-ingress function having a parallel construct by: Creating a function containing a body of the parallel construct; and Inserting calls that target a device runtime library when the function runs in the parent thread. The system of any one of claims 1-6, further comprising a function processor coupled to the thread allocation unit and configured to translate calls from outer functions as they are. The system of any one of claims 1-7, further comprising a function processor coupled to the thread mapper and configured to use a device runtime library that provides functions that can be invoked by compiled user code and internal functions , The system of any of claims 1-8, further comprising a functional processor coupled to the strand mapper and configured to use an egress tag stored in a shared memory and used to facilitate the parent thread to tell the worker strands when to finish their execution. The system of any of claims 1-9, wherein the system is configured to configure a shared memory of a single-instruction multi-processor to identify a current task.

Images