System and method for compiling or running a data-parallel program with split-union and function calls in a single-instruction multi-strand processor

Abstract

A system for compiling or running a data-parallel program with split-union using function calls, including: a SIMT processor (100) with a pipeline control unit (108); a partition unit configured to divide strand groups of the SIMT processor (100) into a parent group and at least one worker group prior to compilation or runtime execution of the data-parallel program; and a string allocation unit that is connected to and configured with the partition unit to determine one of the strings from the parent group to be executed as the parent string and to determine all strings in the at least one worker group to be executed as worker strings before compiling or running the data-parallel program, where the parent strand of the parent strand group operates on the SIMT processor during a sequential phase of the data-parallel program, the worker strands of the worker group operate during a parallel phase of the data-parallel program on the SIMT processor, The higher-level strand waits until the parallel phase is completed. Other strands of the parent group remain inactive during the data-parallel program, and The system configures a shared memory (110) of the SIMT processor to identify a running task.

Description

TECHNICAL AREA

[0001] This application relates generally to parallel processors and in particular to a system for compiling or running a data-parallel program with split-union with function calls in a single-instruction multi-strand (SIMT) processor. BACKGROUND

[0002] As experts in this field know, applications or programs can be executed in parallel to improve their performance. Data-parallel programs execute the same process on different data simultaneously. Task-parallel programs execute different processes on the same data simultaneously. Static parallel programs have a degree of parallelism that can be determined before execution. In contrast, the parallelism achievable by dynamic parallel programs can only be determined during execution. Regardless of whether the program is data-parallel, task-parallel, static, or dynamic parallel, it can be executed in a pipeline or parallel processing line, which is common for graphical programs.

[0003] A SIMT processor is particularly adept at executing programs running in parallel with data. A control unit within the SIMT processor creates groups of execution strands and schedules them for execution, during which all strands in the group execute the same instruction simultaneously. In a specific processor, each group, or "bulge" or "chain," has 32 strands, corresponding to 32 execution pipelines or paths within the SIMT processor.

[0004] A split-merge data-parallel program begins with a main program that has only one strand. At this stage, the program is in a sequential phase or sequential region. At some point during the execution of the main program, the main or "master" strand encounters a sequence of parallel phases or regions. Each parallel region has an independent data set and can be executed by multiple strands concurrently. The number of concurrent tasks in each parallel region is determined when the parallel region begins and does not change during the parallel region. When a parallel region is encountered, the main strand splits into a group of strands (called worker strands) to execute the parallel regions in parallel. The program then enters the parallel region.When a worker strand encounters a new parallel section, the new parallel section is serialized; that is, the parallel section is executed by the arriving worker strand itself. The main strand waits until the parallel section is complete. Upon exiting the parallel section, the worker strands merge with the main strand, which then resumes execution of the main program, at which point the program enters a sequential section.

[0005] Table 1 below provides an example of a data-parallel program with split-union.

[0006] For the purpose of understanding Table 1 and the rest of this disclosure, the terms "foo" and "bar" are arbitrary names of functions. Therefore, any function may be used in place of "foo" or "bar".

[0007] The data-parallel split-branch model is frequently used in parallel programming. For example, the OpenMP standard uses this model as a fundamental strand execution model. The OpenACC standard uses this model for the worker strands in a group referred to as a "workgroup." US 2007 / 0 294 512 A1 discloses a system for selecting a processor unit in a parallel system with many processors. US 2009 / 0 013 323 A1 discloses a processor with synchronizers that synchronize different strands. US 2011 / 0 265 068 A1 describes a system that identifies parts of an application that can be executed in parallel. An object of the invention is to improve the logical grouping of strands. OVERVIEW

[0008] One aspect provides a system for compiling or running a data-parallel program with partitioning and joining, using function calls according to claim 1. In one embodiment, the system includes, among other things: (1) a partition unit configured to divide strand groups into a parent group and at least one worker group, and (2) a strand assignment unit connected to the partition unit and configured to assign or determine only a single strand from the parent group for execution and all strands in the at least one worker group for execution. SHORT DESCRIPTION

[0009] Reference is now made to the following descriptions in conjunction with the accompanying drawings, in which: Fig. 1 a block view of a SIMT processor configured to contain or execute a system or procedure for compiling or running a data-parallel program with split branching and function calls; Fig. 2 a block view of an embodiment of a system for compiling or running a data-parallel program with split branching and function calls; and Fig. 3 is a flowchart of an embodiment of a method for compiling or running a data-parallel program with split branching and function calls. DETAILED DESCRIPTION

[0010] In a SIMT processor, multiple execution strands are divided into groups. All strands within a group execute the same instruction simultaneously. In graphics processing units (GPUs), such as those commercially available from Nvidia, Santa Clara, California, which are a type of SIMT processor, these groups are referred to as "bulges" or "chains," and they execute in blocks.

[0011] A pipeline control unit of the SIMT processor creates, manages, schedules, executes, and provides a mechanism to synchronize groups. Nvidia GPUs provide a bar.sync instruction for synchronizing groups. Nvidia GPUs further support the execution of a "divergent" conditional branch by a group; some strands of the group must take the branch (because the branch condition preselection determines "true"), and the other strands must skip to the next instruction (because the branch condition preselection determines "false"). The pipeline control unit monitors active strands in the group. It executes one of the paths (branch taken or branch not taken) first and then executes the other path; the corresponding strands are activated in each path.

[0012] It is recognized here that, while all strands within a GPU strand block start at the same program address, the pipeline control unit would benefit from a software mechanism that divides and disposes of the strands into a main strand and worker strands so that they can be executed in the split-union model.

[0013] It is further recognized herein that certain embodiments of the software mechanism should manage and synchronize the strands as groups, since the pipeline control unit manages the strands as groups.

[0014] It is further recognized herein that, since the main program is a single-strand program in the split-unification model, certain implementations of the software mechanism should achieve a semantics of sequential domain without introducing side effects. Examples of instructions that cause side effects are those that use shared resources, such as a read or write instruction for shared memory, or any code operation that can invoke a shared exception handler (for example, division).

[0015] It is further recognized herein that certain embodiments of the software mechanism should support functions that can be called within the sequential domain and within the parallel domain. Such functions may themselves contain parallel constructs. It is further recognized herein that certain embodiments of the software mechanism should support function calls that can split and merge parallel domains.

[0016] Furthermore, it is recognized here that certain implementations of the software mechanism should support external functions, that is, functions that are not compiled by the same compiler as the program. Examples include mathematical functions in existing GPU math libraries and system functions such as malloc, free, and print. In certain implementations, both the parent strand in the sequential areas and the worker strands in the parallel area should be able to call an external function.

[0017] Consequently, various embodiments of a system and a method for compiling and executing data-parallel programs with partitioning-union and with function calls in a SIMT processor, such as a GPU, are described herein.

[0018] Before describing certain embodiments of the system and the method, a SIMT processor is described which is configured to include or execute a system or method for compiling or running data-parallel programs with split-union using function calls.

[0019] Fig. Figure 1 is a block view of a SIMT processor 100. The SIMT processor 100 contains several strand processors or cores 106, which are divided into strand groups 104 or "bulges" or "chains." The SIMT processor 100 contains J strand groups 104-1 to 104-J, each K of which has cores 106-1 to 106-K. In certain embodiments, the strand groups 104-1 to 104-J may be further divided into one or more strand blocks 102. Certain embodiments include thirty-two cores 106 per strand group 104. Other embodiments may contain fewer than four cores in a strand group and up to several tens of thousands of cores. Certain embodiments divide the cores 106 into a single strand group 104, while other embodiments have hundreds or even thousands of strand groups 104.Alternative embodiments of the SIMT processor 100 can divide the cores 106 exclusively into the strand groups 104, whereby the division at the strand block level is not present.

[0020] The SIMT processor 100 further comprises a pipeline control unit 108, a block-shared memory 110, and an array of local memories 112-1 to 112-J, which are assigned to the strand groups 104-1 to 104-J. 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 operate in parallel. The strand groups 104-1 to 104-J communicate with the block-shared memory 110 via a memory bus 116. The strand groups 104-1 to 104-J communicate with the local memories 112-1 to 112-J via local buses 118-1 to 118-J. For example, a string group 104-J uses the local memory 112-J by means of communication via a local bus 118-J.Certain embodiments of the SIMT processor 100 assign a shared portion of the block-shared memory 110 to each 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. Some embodiments include strand groups 104 that use only the local memory 112. Many other embodiments include strand groups 104 that achieve a balanced use of the local memory 112 and the block-shared memory 110.

[0021] The embodiment from Fig. 1 contains a parent strand group 104-1. Each of the remaining strand groups 104-2 through 104-J is considered a "worker" strand group. The parent strand group 104-1 contains numerous kernels, one of which is a parent kernel 106-1, which ultimately executes a parent strand. Programs executed in the SIMT processor 100 are structured as a sequence of kernels. Typically, each kernel finishes executing before the next kernel begins. In certain embodiments, the SIMT processor 100 can execute multiple kernels in parallel, depending on the size of the kernels. Each kernel is structured as a hierarchy of strands to be executed in the kernels 106.

[0022] Following the description of a SIMT processor in which the system or method as introduced herein can be contained or executed, various embodiments of the system and the method will now be described.

[0023] One embodiment of the system introduced herein comprises a compiler and a device runtime library. The device runtime library implements the branch and group management function. The compiler translates a split-union, data-parallel program into a main branch program and a group of outer functions, each corresponding to a parallel construct. The compiled code makes calls to functions of the device runtime library to execute the branch and group management.

[0024] Table 2 below shows an example program to illustrate the compiler translation and the implementation of the device runtime.

[0025] The execution of the `main()` program from Table 2 begins with the single parent strand. The parent strand calls the function `foo()`, which has a body that is transparent to this compiler and is translated by it. The parent strand then calls a function `ext()`, that is, an external or outside function with a body that is invisible to this compiler. Calls to outside functions are translated as they are, without any special handling by the compiler. The parent strand then encounters the first parallel section. Worker strands execute the parallel section while the parent strand waits for them to finish. Within the parallel section, each worker strand calls the functions `foo()` and `bar()`. The function `bar()` contains another parallel section; however, `bar()` is already within a parallel section.Since bar() is already within a parallel section, the parallel section within bar() is executed sequentially by each worker strand.

[0026] After the first parallel section, the parent strand encounters a second parallel section. Within this second parallel section, each worker strand calls the external function `ext()`. After the second parallel section, the parent strand calls the function `bar()`. Within `bar()`, the parent strand encounters a third parallel section, which is then processed by the worker strands.

[0027] The `main()` function is known as an entry function because it is the point at which the program starts. Functions such as `foo()` and `bar()` are non-entry functions.

[0028] For an entry function, the compiler first creates a cloned copy called `main_core()`. This cloned copy is then processed as a non-entry function, as described below. For the `main()` function, the compiler generates code shown in Table 3 below, where `groupID()` returns the ID of the thread group containing a command-executing thread. `threadID()` returns the ID of the thread. `init()`, `signal_done()`, and `scheduler()` are functions in the device runtime library.

[0029] When a GPU strand block starts, all strands within the block execute `main()`; however, they take different paths. Strand 0 is the parent strand and executes `init()`, `main_core()`, and `signal_done()`. Other strands in group 0 go directly to the end of the `main()` function and wait there. The strands in the remaining groups execute `scheduler()`.

[0030] For a non-entry function like `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, the compiler creates a function for each parallel construct that contains the body of the parallel construct (an outsourced function) and then creates a conditional branch that checks whether the currently executing strand is the parent strand. In the false branch, the compiler inserts code that executes the loop. In the true branch, the compiler inserts calls to the device runtime library to assign tasks, wake up worker strands, and execute a barrier. The condition is true if the non-entry function is called outside of the parallel strands.The operation is incorrect if the non-entry function is called within a parallel section, in which case the parallel loop is executed sequentially by the executing strand.

[0031] For example, the translated code for the function bar() is shown below in Table 4.

[0032] `signal_task()` and `barrier()` are functions in the device runtime library. `bar_par_frunc()` is the outsourced function that corresponds to the parallel construct in the original `bar()` function.

[0033] 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().

[0034] All worker strands execute the scheduler() function. The worker strands go through a sleep-wake-execute cycle until they are instructed to stop.

[0035] A Boolean variable 'exit_flag' is stored in the memory shared by blocks, and this variable can be accessed by all strands within the strand block. It is used by the parent strand to inform the worker strands whether they should all terminate execution. 'exit_flag' is set to false in the `init()` function and set to true in the `signal_done()` function. Both functions are called by the parent strand.

[0036] Another portion of the memory shared by blocks is used to communicate the current task. The current task is set by the parent branch in the `signal_task()` function and retrieved by the worker branch in the `fetch_task()` function. The memory shared by blocks contains the pointer to the swapped-out function corresponding to the parallel construct.

[0037] Since the parallel sections are executed sequentially within a string block, only a single task is active at any given time. If the parallel sections can be executed asynchronously, a more complex data structure, such as a stack, a queue, or a data tree, is typically required to store the active tasks.

[0038] The barrier(), signal() and wait() functions are implemented using a hardware barrier.

[0039] Fig. Figure 2 is a block view of an embodiment of a system 200 for compiling or running a data-parallel program using partition-union with function calls. The program 210 contains entry functions 212, non-entry functions 214, and outer functions 216. The system 200 comprises a partition unit 202, a strand allocation unit 204, a strand scheduling unit 206, a function processor 208, a device runtime library 218, and a SIMT processor 100. Fig. 1.

[0040] The SIMT processor 100 comprises the pipeline control unit 108, the data bus 114, the local buses 118-1 and 118-2, and the shared memory 110. Fig. 1. In the embodiment from Fig. Figure 2 shows the SIMT processor 100 as having a single strand block containing two strand groups: the master strand group 104-1 and the worker strand group 104-2. Strand groups 104-1 and 104-2 each contain strands 106.

[0041] The partition unit 202 designates or identifies strand group 104-1 as the superior strand group and the remaining strand groups as worker strand groups. In the embodiment shown Fig. Figure 2 shows a single worker strand group 104-2. In alternative embodiments, many worker strand groups can be used. The strand assignment unit 204 designates a parent strand 106-1 of the parent strand group 104-1. All other strands in the parent strand group 104-1 are inactive. The strand assignment unit 204 also designates each of the strands 106 in the worker strand group 104-2 as a worker strand.

[0042] The strand scheduling unit 206 translates the program 210 such that the pipeline control unit 108 appropriately controls the execution of the parent strand 106-1 and the various worker strands in the worker strand group 104-2. The strand scheduling unit 206 translates the program 210 such that when the execution of the parent strand begins, a program exit marker is reset. The strand scheduling unit 206 schedules the parent strand 106-1 so that it executes until a parallel section or the end of program 210 is reached. When a parallel section of program 210 is reached, the strand scheduling unit 206 assigns a parallel task, and the worker strands in the worker strand group 104-2 begin execution.The strand scheduling unit 206 further sets a barrier for each of the worker strands such that if the barrier is entered, the parent strand 106-2 continues its execution. When the end of program 210 is reached, the exit marker is activated, causing all worker strands to cease execution.

[0043] The function processor 208 operates on the functions of the program 210. The processing of the entry functions 212 involves creating a clone copy of an entry function, which is then processed as a non-entry function. The original entry function is processed such that the parent strand 106-1 will process the clone copy in addition to other calls, and the worker strands execute cycles of sleep mode, wake-up, fetch, and execution of the parallel task specified by the strand scheduling unit 206.

[0044] The Function Processor 208 translates the non-entry functions 214 in two ways. If no parallel construct is present in a non-entry function, the function is simply processed as is. If a parallel construct is present, an outer function containing the body of the parallel construct is created. The Function Processor 208 then creates a branching condition that either executes the parallel construct sequentially or uses the Device Runtime Library 218 to assign a task, wake up worker strands, and execute a barrier, as described previously. The wake-up and sleep functions are implemented using hardware barrier functions of the Device Runtime Library 218. Strands at barriers are not disposed of for execution by the hardware, so they do not waste work cycles.In the parent strand group 104-1, only the parent strand 106-1 participates in the barriers. This is because the hardware barrier is group-based. A group is considered to be at a barrier when every strand within the group is at the barrier.

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

[0046] A data-parallel program with split-union is divided into a parent program and a group of parallel tasks. The parent program is the program executed by the parent strand. A parallel task corresponds to a parallel section executed by the worker strands. The parent program contains scheduling points at which the parent strand will assign or determine a parallel task, which will wake up the worker strands and wait for them to complete.

[0047] The special parent strand in the special parent group will execute the sequential part of the program.

[0048] Alternatively, single-strand behavior can be emulated in the sequential area, while all strands in the group execute the code. However, emulation schemes have limitations regarding complexity in terms of performance and generation, making them less practical. The necessary predetermination and synchronization add extra overhead during execution. Furthermore, all functions called by the sequential and parallel areas must be cloned and controlled differently.

[0049] Given the subdivision of strands and groups, the worker strands and the superior strand go through the following life cycles: One embodiment of a worker strand goes through the following phases in a life cycle: 1) the strand block starts; 2) in sleep mode until it is woken up by the parent strand; 3) Exiting the program when the exit marker is set to true; 4) Retrieving and executing the task assigned by the parent strand; 5) Entering a barrier and 6) Return to level 2.

[0050] An embodiment of a higher-level string goes through the following stages in a life cycle: 1) the strand block starts; 2) Setting the exit marker to false; 3) Execution of 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. Continuation of the overarching program (level 3); and 5) at the end of the overarching program: a. Setting the exit marker to true, b. Awakening the worker strands, and c. End.

[0051] The other strands in the parent group essentially wait idly for the program to finish. The program is executed alternately by the parent strand and the worker strands. This results in good utilization of the instruction cache memory, which is better than the utilization caused by a method in which both the parent strand and the worker strands are active and execute different code paths.

[0052] Fig.Figure 3 is a flowchart of an embodiment of a method for compiling or executing a data-parallel program using split-merge with function calls. The method begins in a start step 310. In step 320, strand groups within a strand block are split into a parent strand group and at least one worker strand group. In step 330, a single strand from the parent strand group is designated as the parent strand. The remaining strands of the parent group are essentially inactive during execution. Furthermore, in step 330, all strands in the at least one worker strand group are designated as worker strands. The method ends in a finish step 340.

[0053] The person skilled in the art in this field, to whom this application is addressed, recognizes that other and further additions, deletions, replacements and modifications can be made to the described embodiments.

Claims

[1] A system for compiling or running a data-parallel program with partitioning and unifying with function calls, including: a SIMT processor (100) with a pipeline control unit (108); a partition unit configured to divide strand groups of the SIMT processor (100) into a parent group and at least one worker group prior to compilation or runtime execution of the data-parallel program; and a string allocation unit that is connected to and configured with the partition unit to determine one of the strings from the parent group to be executed as the parent string and to determine all strings in the at least one worker group to be executed as worker strings before compiling or running the data-parallel program, where the parent strand of the parent strand group operates on the SIMT processor during a sequential phase of the data-parallel program, the worker strands of the worker group operate during a parallel phase of the data-parallel program on the SIMT processor, The higher-level strand waits until the parallel phase is completed. Other strands of the parent group remain inactive during the data-parallel program, and The system configures a shared memory (110) of the SIMT processor to identify a running task. [2] The system according to claim 1, further comprising a strand scheduling unit connected to the strand allocation unit and configured to cause the pipeline control unit of a single-instruction multi-strand processor to schedule the execution of the higher-level strand as follows: A marker is set to exit a program to an initial state when the parent branch begins execution; The higher-level strand is executed until it reaches a parallel section or the end of the program; Upon reaching the parallel area, a parallel task is defined, a first barrier is entered, a second barrier is entered, and the execution of the higher-level strand continues; and Upon reaching the end, the marker for exiting a program is set to a second state, a barrier is entered, and the parent strand terminates execution. [3] The system according to claim 1 or 2, further comprising a strand scheduling unit connected to the strand allocation unit and configured to cause a pipeline control unit of a single-instruction multi-strand processor to schedule the execution of the worker strands as follows: The worker lines enter a first barrier; The worker strands begin performing a parallel task, as determined by the parent strand, when the parent strand reaches a parallel area and enters a barrier; The worker strands enter a second barrier upon completion of the parallel task; and The worker strands end when a flag indicating the exit of a program to a second state is set and the parent strand enters a barrier. [4] The system according to one of claims 1-3, further comprising a strand scheduling unit connected to the strand allocation unit and configured to use a barrier function of the pipeline control unit of a single-instruction multi-strand processor for controlling the execution and termination of the worker strands. [5] The system according to one of claims 1-4, further comprising a function processor connected to the string allocation unit and configured to generate a cloned copy of an entry function and to process the entry function as a non-entry function. [6] The system according to any one of claims 1-5, further comprising a function processor connected to the string allocation unit and configured to translate a non-entry function comprising a parallel construct, by: Generation of a function containing a field of the parallel construct; and Insertion of calls that target a device runtime library when the function is executed in the parent branch. [7] The system according to any one of claims 1-6, further comprising a function processor which translates calls to external functions unchanged. [8] The system according to one of claims 1-7, further comprising a function processor connected to the strand allocation unit and configured to use a device runtime library which provides functions that can be called by compiled user code and internal functions. [9] The system according to any one of claims 1-8, further comprising a function processor connected to the strand allocation unit and configured to use an exit marker stored in a shared memory and used to enable the parent strand to communicate to the worker strands when to stop their execution. [10] The system according to any one 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