1. Field of the Invention
The present invention relates to a computer-implemented divergence analysis method for a pointer-based program.
2. Description of the Related Art
In recent years, GPU-CPU heterogeneous architectures have been increasingly adopted in high performance computing due to their capabilities of providing high computational throughput.
Modern GPUs employ a single instruction multiple data (SIMD) execution model that groups parallel threads executing the same instruction in lock-step. The group is called a warp in the NVIDIA GPU terminology and a wavefront in the AMD GPU terminology. However, a divergence may occur when threads that execute in lock-step are forced to take different program paths because of branches in a code. When branch divergence occurs, some threads have to wait or be idle while other threads are executing, and the divergent branches will be executed serially on GPUs. It can be seen that divergence causes serious performance degradation of GPUs.
Although much research focuses on the divergence issue, the divergence analysis for pointer-based programs has not been thoroughly looked into.
In one embodiment of the present invention, a computer-implemented divergence analysis method for a pointer-based program comprises generating an intermediate representation of a pointer-based program, wherein the intermediate representation is in gated single static assignment form, wherein, in one embodiment, the intermediate representation is in gated single static assignment and memory static assignment forms; providing a control flow graph of the intermediate representation; selecting a variable of the intermediate representation as an analysis candidate; using the analysis candidate as a traced variable and a root node; determining a definition site of the trace variable according to a corresponding use-define chain and the control flow graph; defining a node for each definition site variable of the definition site if the definition site of the traced variable does not comprise a predetermined built-in function; defining an edge by using each definition site variable of the definition site and the traced variable if the definition site of the traced variable does not comprise the predetermined built-in function; using each definition site variable of the definition site as a traced variable; repeating the steps of determining a definition site, defining a node, defining an edge, and using each definition site to obtain a divergence relation graph; transforming the divergence relation graph into a directed acyclic graph; identifying at least one divergent node of the directed acyclic graph; and determining whether the analysis candidate is divergent or not according to the at least one divergent node and the directed acyclic graph.
The objectives and advantages of the present invention will become apparent upon reading the following description and upon referring to the accompanying drawings in which:
OpenCL (Open Computing Language) is a programming framework for programs operating on heterogeneous platforms including CPUs, GPUs, DSPs and other hardware accelerators. Although the present invention employs OpenCL, the present invention is not limited to employing OpenCL.
Table 1 shows two fragments of an exemplary OpenCL program according to one embodiment of the present invention.
The program fragment A comprises an OpenCL built-in function, get_global_id( ), which returns different work-item IDs on different threads. Thus, the results of the conditions in if statements of some threads may be different from those of others, and consequently, divergence may happen. The OpenCL platform includes other built-in functions, such as get_local_id( ), atomic_add( ), atomic_sub( ), atomic_xchg( ), atomic_inc( ), atomic_dec( ), atomic_cmpxchg( ), atomic_min( ), atomic_max( ), atomic_and( ), atomic_or( ), and atomic_xor( ), that may also cause divergence.
The x values of the program fragment B may be different between different threads since different tids cause different iteration numbers. As can be seen from the above examples, a variable is divergent if the value of the variable is computed from a built-in function, for example get_global_id( ); the variable that is data-dependent on a divergent variable is divergent; or the variable that is control-dependent on a divergent variable is divergent. Similarly, a pointer is divergent if the pointer points to a divergent variable.
In the present disclosure, the variable can be a character, integer, float, double or pointer variable.
The methods of at least some embodiments are based on two extended SSA (Static Single Assignment) forms: the memory SSA and gated SSA. Referring to
The pointer-based program can be converted to an IR in memory SSA form, as shown in
The gated SSA form enhances the φ function by applying control condition operands. The IR of
The γ function adds a condition for multiple definitions to reach a confluence node that has no incoming back edges like the conditions in “if-then-else” statements. For instance, x3=β(P, x1, x2) represents x3=x1 if P, and x3=x2 if P.
The μ function, which only exists at loop headers, selects initial and loop-carried values. For instance, x2=μ(x0, x3) represents x2=x0 in the first iteration of a loop and x2=x3 after the first iteration of the loop.
The η function determines the value of a variable at the exit of the end of a loop. For instance, x4=η(P, x3) represents that x3 is the value assigned in the last iteration of a loop, and P is a condition to exit the loop.
In at least some embodiments, a revised gated SSA intermediate representation of
Table 2 shows a statement or program classification according to one embodiment of the present invention.
In Table 2, p and q are one-level pointers; r is a two-level pointer; i is a scalar variable; and A is an array of scalar functions or variables.
A statement of an Updated by Function uses an OpenCL built-in function to return a value. A statement of a Points-to Location includes a pointer pointing to one or more locations. A statement of a Pointers Aliased includes an alias of pointers with an offset, or includes a one-level-higher pointer pointing to a one-level-lower pointer. A statement of Indirect Store and Indirect Load accesses pointers indirectly when a pointer is dereferenced by * or →. In at least some embodiments, the location naming is defined as (base, offset), and the offset can be composed of an affine style.
A compiler, such as Open64, is used to compile the OpenCL program to generate a WHIRL intermediate representation, an intermediate representation in SSA form, and a CFG (control flow graph). However, the present invention is not limited to using Open64. In some embodiments, the intermediate representation in SSA form could be an intermediate representation in memory SSA. In some embodiments, the intermediate representation in SSA form could be an intermediate representation in gated SSA. In some embodiments, the intermediate representation in SSA form could be an intermediate representation in memory and gated SSA. In some embodiments, the intermediate representation in SSA form uses γ function to replace η function after a loop and uses φ function to replace μ function before a loop.
Referring to
Referring to
Referring to
Referring to
Referring to
Referring to
Referring to
Referring to
Referring to lines 32 to 36 of
p4 is selected as the analysis candidate and is passed to the function Build_DRG. Nodes for p2, p3 and c1 are added and passed to the function Build_DRG because the definition site or statement of p4 includes a γ grating function whose definition site variables or operands are p2, p3 and c1 at BB5. The node x4 is added and passed to the function Build_DRG because x4 is in the RHS of the p4 statement at BB8.
After p2 is passed to the function Build_DRG, a statement including a definition site of p2 is determined. According to the definition site of p2 at BB3, the statement of the BB3 is classified as Pointer Aliased. Therefore, the definition site variable or aliased pointer b1 is added and passed to the function Build_DRG. Since b1 is the parameter of an analyzed function, which is defined in a host code, the definition site of b1 can be found and this branch is terminated.
After p3 is passed to the function Build_DRG, a statement including a definition site of p3 is determined. According to the definition site of p3, statement at BB4 is a Points-to location. The base address a1 and the offsets tid1, which are definition site variables and array indices, are added and passed to the function Build_DRG. Since the definition site of a1 can be found, this branch is terminated.
After the function Build_DRG receives the offset tid1, a statement, including a definition site of the offset tid1 at BB1, is determined. Because the definition site of the offset tid1 includes an update function get_global_id(0), which is an OpenCL C built-in function listed in the above Table 2, the function BuildDRG_Callee is called. The function BuildDRG_Callee sets the offset tid1 to be divergent and returns.
After the function Build_DRG receives x4, a statement at BB8, including definition site of x4, is found. The definition site of x4 includes a γ grating function. Thus, the definition site variables or operands x2, x3, and i2 are added and passed to the function Build_DRG.
Because x2 is defined by a φ function at BB6, only one node x1 is added and passed to the function Build_DRG because x3 has been visited.
The definition site of x3 only includes a variable, x2, which has been visited. Thus, the analysis is terminated.
The analysis of x1 is terminated because a final definition is found.
The analyses of i1, i2, and i3 are similar to those of x1, x2, and x3, and therefore, redundant explanations are omitted. The final DRG is built as shown in
Referring to
Referring to
In one embodiment, the present invention comprises a hardware computer. The computer comprises a processor for executing each or every step of the methods of the above embodiments and a hardware memory for storing codes to cause the processor to execute each or every step of the methods of the above embodiments.
The data structures and code described in this detailed description are typically stored on a non-transitory computer-readable storage medium, which may be any device or medium that can store code and/or data for use by a computer system. The non-transitory computer-readable storage medium includes, but is not limited to, volatile memory, non-volatile memory, magnetic and optical storage devices such as disk drives, magnetic tape, CDs (compact discs), DVDs (digital versatile discs or digital video discs), or other media capable of storing code and/or data now known or later developed.
The methods and processes described in the detailed description section can be embodied as code and/or data, which can be stored in a non-transitory computer-readable storage medium as described above. When a computer system reads and executes the code and/or data stored on the non-transitory computer-readable storage medium, the computer system performs the methods and processes embodied as data structures and code stored within the non-transitory computer-readable storage medium. Furthermore, the methods and processes described below can be included in hardware modules. For example, the hardware modules can include, but are not limited to, application-specific integrated circuit (ASIC) chips, field-programmable gate arrays (FPGAs), and other programmable-logic devices now known or later developed. When the hardware modules are activated, the hardware modules perform the methods and processes included within the hardware modules.
The above-described embodiments of the present invention are intended to be illustrative only. Those skilled in the art may devise numerous alternative embodiments without departing from the scope of the following claims.
Number | Name | Date | Kind |
---|---|---|---|
7155704 | Williams | Dec 2006 | B2 |
7284219 | Manaker, Jr. | Oct 2007 | B1 |
7477255 | Lindholm et al. | Jan 2009 | B1 |
7619628 | White | Nov 2009 | B2 |
7834881 | Liu et al. | Nov 2010 | B2 |
7996798 | Srinivasan | Aug 2011 | B2 |
8161464 | Archambault | Apr 2012 | B2 |
8302085 | Schellekens | Oct 2012 | B2 |
8381203 | Beylin | Feb 2013 | B1 |
8402444 | Ball | Mar 2013 | B2 |
8881293 | Brucker | Nov 2014 | B1 |
20070240137 | Archambault | Oct 2007 | A1 |
20080091697 | Cui | Apr 2008 | A1 |
20080184208 | Sreedhar | Jul 2008 | A1 |
20080295058 | Srinivasan | Nov 2008 | A1 |
20090217248 | Bently | Aug 2009 | A1 |
Number | Date | Country |
---|---|---|
200709055 | Mar 2007 | TW |
201342212 | Oct 2013 | TW |
Entry |
---|
Notice of allowance and Search Report dated Jun. 29, 2015 from the Taiwan counterpart application 103122208. |
English translation of the Search Report dated Jun. 29, 2015 from the Taiwan counterpart application 103122208. |
English abstract translation of TW 201342212. |
English abstract translation of TW 200709055. |
Number | Date | Country | |
---|---|---|---|
20150143349 A1 | May 2015 | US |