KR20200140210A - Idempotent kernel generateing method and apparatus - Google Patents

Abstract

The present specification discloses a method and device for generating an idempotent kernel to ensure the interruption and re-execution of a GPU kernel without disruption in an application in a GPU general-purpose computing environment. Specifically, the method for generating an idempotent kernel according to one embodiment of the present invention may comprise the steps of: identifying a first memory variable related to a risk of write after read (WAR) not statically determined in a source code related to a GPU kernel function under the control of an analysis unit; inserting a code for setting an access flag bitmap for tracking read access or write access to the first memory variable under the control of a code insertion unit; declaring a temporary pointer variable in the GPU kernel function as a local variable under the control of a variable declaration unit and initializing the local variable as a first memory variable; and inserting a code for dynamically determining whether to use the temporary pointer variable as the first memory variable or a second memory variable before the write access occurs after the read access to the temporary pointer variable under the control of a decision code insertion unit.

Description

Idempotent kernel generation method and device {IDEMPOTENT KERNEL GENERATEING METHOD AND APPARATUS}

The present specification relates to generation of an idempotent kernel, and more particularly, to a method and apparatus for preventing data hazards from occurring due to interruption of a GPU.

The graphics processing unit (GPU) does not support preemption due to hardware limitations. Unlike the central processing unit (CPU), the GPU executes threads in parallel on multiple cores, so the cost of storing and fetching the context is high for the context switch of the instruction unit.

In addition, since the GPU does not support preemptive scheduling, a priority reversal may occur. Also, if there is a kernel that has already started executing, the GPU cannot switch context to another kernel until the kernel has finished executing. Here, the priority reversal problem refers to a phenomenon in which execution of a high priority task waits for execution due to the execution of a low priority task. This is a problem that must be solved in a real-time system.

The GPU generally provides an abort command, but when the kernel is stopped with the abort command, the entire application is stopped and must be run from the beginning. One application generally executes a plurality of GPU kernels sequentially, and each kernel may have a dependency on the result of the previous execution kernel. Therefore, if the GPU kernel is stopped at any moment, the input data changed while the kernel is executing is changed, and the GPU system does not currently provide a means to recover the changed data, so that the input data has the same operation result. Initialization may require re-execution at the application level. In order to support preemptive scheduling in such a GPU environment, a technology for ensuring consistency of input data that is changed during execution of the GPU kernel is required.

An object of the present specification is to provide a method of guaranteeing interruption and re-execution of a GPU kernel without interrupting an application in a general-purpose GPU computing environment.

In addition, an object of the present specification is to provide a method capable of supporting scheduling according to priority of a GPU kernel without hardware support.

In addition, an object of the present specification is to provide a method in which data consistency is guaranteed and re-executable even if the GPU kernel is stopped at any time.

The technical problems to be achieved in the present invention are not limited to the technical problems mentioned above, and other technical problems that are not mentioned will be clearly understood by those of ordinary skill in the technical field to which the present invention belongs from the following description. I will be able to.

In the method of generating an idempotent kernel according to an embodiment of the present invention, the step of checking a first memory variable related to a risk of a write after read (WAR) that is not statically determined in a source code related to a GPU kernel function, and the first memory Inserting a code for setting an access flag bitmap for tracking read access or write access to variables, and declaring a temporary pointer variable in the GPU kernel function as a local variable and initializing it as the first memory variable And inserting a code for dynamically determining whether to use the temporary pointer variable as a first memory variable or a second memory variable after a read access to the temporary pointer variable and before a write access occurs. can do.

In addition, variables related to the risk of write after read (WAR) that are not statically determined may be variables related to conditional statements or dynamically determined indexes.

In addition, the access flag bitmap may be a set of bits indicating write access or read access to each of the regions included in the global memory.

In addition, the code for setting the access flag bitmap may cause a bit mapped to the specific area to have a value of 1 when a read access or a write access occurs to a specific area included in the global memory.

Also, the first memory variable and the second memory variable may be global memory variables.

An idempotent kernel generation apparatus according to an embodiment of the present invention includes an analysis unit for checking a first memory variable related to a risk of a write after read (WAR) that is not statically determined in a source code related to a GPU kernel function, and the first A flag code insertion unit that inserts a code for setting an access flag bitmap for tracking read access or write access to a memory variable, and a temporary pointer variable in the GPU kernel function is declared as a local variable, and the first memory variable A decision code for inserting a variable declaration unit initialized with and a code for dynamically deciding whether to use the temporary pointer variable as a first memory variable or a second memory variable after read access to the temporary pointer variable and before write access occurs It may be characterized in that it includes an insert.

In addition, variables related to the risk of write after read (WAR) that are not statically determined may be variables related to conditional statements or dynamically determined indexes.

In addition, the access flag bitmap may be a set of bits indicating write access or read access to each of the regions included in the global memory.

In addition, the code for setting the access flag bitmap may cause a bit mapped to the specific area to have a value of 1 when a read access or a write access occurs to a specific area included in the global memory.

Also, the first memory variable and the second memory variable may be global memory variables.

According to the present specification, there is an effect of guaranteeing the interruption and re-execution of the GPU kernel without interrupting an application in a general-purpose GPU computing environment.

In addition, according to the present specification, it is possible to support scheduling according to the priority of the GPU kernel without hardware support.

In addition, according to the present specification, since the execution of the application is maintained even if the execution of the kernel is stopped, when a GPU kernel of high priority occurs, the execution of the existing kernel is stopped and the kernel having a high priority can be executed.

In addition, according to the present specification, even if the GPU kernel is stopped at any time, data consistency is guaranteed and re-execution is possible.

In addition, according to the present specification, it is possible to support preemptive priority scheduling by executing an idempotent GPU kernel without hardware support and without additional context storage.

The effects that can be obtained in the present invention are not limited to the above-mentioned effects, and other effects not mentioned will be clearly understood by those of ordinary skill in the art from the following description. .

BRIEF DESCRIPTION OF THE DRAWINGS The accompanying drawings, which are included as part of the detailed description to aid in understanding of the present invention, provide embodiments of the present invention, and describe technical features of the present invention together with the detailed description.
1 shows an example of simply expressing whether or not idempotency is guaranteed according to a memory access pattern.
2 is a diagram for describing an access flag bitmap.
3 is a flowchart illustrating an operation of an idempotent kernel generating apparatus according to an embodiment of the present invention.
4 is a simple example of source code for explaining the operation of the idempotent kernel generation device.
5 is a flowchart illustrating a method of generating an idempotent kernel according to an embodiment of the present invention.
6 is a flowchart illustrating a method of generating an idempotent kernel according to another embodiment of the present invention.
7 is a diagram illustrating an idempotent kernel generation device according to an embodiment proposed in the present specification.
8 is a diagram illustrating an idempotent kernel generation device according to another embodiment proposed in the present specification.
9 is a diagram for explaining a management structure.

Hereinafter, preferred embodiments of the present invention will be described in detail with reference to the accompanying drawings. The detailed description to be disclosed hereinafter together with the accompanying drawings is intended to describe exemplary embodiments of the present invention, and is not intended to represent the only embodiments in which the present invention may be practiced. The following detailed description includes specific details to provide a thorough understanding of the present invention. However, one of ordinary skill in the art appreciates that the invention may be practiced without these specific details.

Satisfying the idempotence of the memory level in the GPU kernel means that the global memory variables passed to the kernel are guaranteed to be consistent, so that the same results can be expected even if it is stopped at any moment and then executed again. .

Hereinafter, the concept of a WAR hazard and idempotence will be described through a simple example before describing the method and apparatus for generating the idempotent kernel of the present invention.

If there is a memory access pattern in which writing to the variable occurs after reading access to the variable in the source code, the idempotence of the corresponding source code may not be guaranteed. If a variable with a WAR data hazard is initialized to an arbitrary value before initial access, the idempotence of the corresponding source code can be guaranteed because it is always executed with the same data even if the source code is rerun. .

1 shows an example of simply expressing whether or not idempotency is guaranteed according to a memory access pattern.

Referring to FIG. 1, in the case of FIG. 1(a), variables A and B have a WAR risk access pattern. The code of FIG. 1(a) that is normally executed may have a result in which values of A and B are replaced with each other. However, if the source code is interrupted after storing B in A (A=B;), when the source code is rerun, the initial A value has been changed to B, so the expected result cannot be obtained.

For example, it can be assumed that the initial values are A=1 and B=2. In the normal case, tmp stores a value of 1 (tmp=A), and A is changed to a value of 2 (A=B). Next, B is changed to a value of 1 (B=tmp). As a result, A=2 and B=1 may have result values that are substituted with each other. However, after changing the value of A to the value of B (A=B), if it is interrupted, A=2, B=1, and when rerunning the source code again, tmp stores the value of 2 (tmp =A), A keeps the value of 2 (A=B). Next, B holds the value of 2 (B=tmp). This can have erroneous results.

In the case of FIG. 1(b), the values of A and B are initialized (A=1; B=2;) before the source code is executed. In this case, even if the values of A and B are stopped at a certain point, the expected result is obtained. Can be obtained.

For example, in the case of normal, at the beginning of the code, it is initialized to A=1, B=2. Next, the variable tmp has a value of 1 (tmp=A). Next, A has a value of 2 (A=B). Next, B has a value of 1 (B=tmp). If, after A has a value of 2 (A=B), it is stopped, and when performing again, starting from initialization with A=1 and B=2 again, the problem as shown in FIG. 1(a) does not occur.

This specification proposes a method for generating an idempotent GPU kernel code that can be rerun immediately without storing the GPU kernel in a GPU environment (a method for generating an idempotent kernel). This method can support immediate preemption support.

Hereinafter, the present specification divides the idempotent kernel generation method into an idempotent kernel generation method through static analysis (hereinafter, a first embodiment) and an idempotent kernel generation method through an access flag bitmap (hereinafter, a second embodiment). Take a look.

Hereinafter, the embodiments described in the present specification are only classified for convenience of description, and some methods and/or configurations of certain embodiments may be substituted with the methods and/or configurations of other embodiments, or may be combined and applied. Yes, of course.

Embodiment 1

First, let's look at how to create an idempotent kernel through static analysis.

(Static analysis for generating idempotent GPU kernel)

The idempotent kernel generation device can perform static analysis to generate an idempotent GPU kernel.

The idempotent kernel generation device tracks the occurrence of the WAR data risk pattern (or WAR risk) for the code area executed by the GPU in the GPGPU program source code to check whether the code is idempotent or not, and can check the information of the variables that generate the WAR pattern. have.

Specifically, the idempotent kernel generation device can check GPU function information. The GPU function information may be kernel function information and/or device function information performed by GPU hardware in the GPGPU program source code.

In the parallel computing library (CUDA/OpenCL) that supports the GPGPU programming model, functions executed in the GPU are classified into kernel/device functions through specific directives. Through these indicators, the idempotent kernel generator can generate a structure for static analysis by classifying functions executed in GPU hardware and extracting information about the indicator, function name, parameter, and return type.

In the case of OpenCL, the function called by the code executed in the GPU is defined in the form of an inline function, so to track this, the idempotent kernel generator can identify device function information by tracking the function call information in the code. .

In order to track the WAR risk that may occur due to the function call relationship while GPU code is being executed, the idempotent kernel generation device first traces the WAR risk inside each function block, and then tracks the global memory variable transfer relationship between functions and WAR risk can be identified.

Next, the idempotent kernel generation device can check the global memory variable information passed to the GPU function.

Like general C and/or C++, the parameters of the GPU function may have the form of Call by Value and/or Call by Reference. Global memory variables can be transferred in the form of Call by Reference. Global memory variables are passed to GPU functions in the form of a pointer or reference, and the idempotent kernel generation device is a global memory variable through '*' and/or '&' according to the C and/or C++ specification. You can check.

(WAR pattern information generation)

Next, the idempotent kernel generation device may generate and/or check global memory variable access pattern information (WAR pattern information).

The idempotent kernel generation device can distinguish and check variable read and/or write access through an operator pattern in the source code.

For example, (1) the'=' assignment operator can be identified as a write access to the leftmost variable at the end and read access to all other variables.

(2) The '+=' compound operator can be verified as a write access after read access to the leftmost variable at the end.

(3) The'++' unary operator can be verified by reading and writing access to the target variable.

(4) When passed as a parameter when calling a function, the access pattern can be confirmed through global analysis after analysis within each function.

(5) All operators except the three operators described above can be verified by read access to variables.

And/or, the idempotent kernel generation device can check WAR pattern information through the formation of a variable access management structure.

You can check whether a WAR pattern has occurred or not by forming a management structure for each global memory variable and recording sequential variable access patterns in the code.

Since global memory variables are passed by call-by-reference (Call-by-Reference), the area pointed to by the pointer is structured (Struct). 'Or'->' type pointer access can be made. In consideration of this case, the variable that is finally accessed can be checked.

And/or, when an array-type variable is used (generally), the location of the memory area to be actually accessed may be determined through an index. The idempotent kernel generating device may determine whether a WAR pattern has occurred in consideration of the corresponding index, and generate and/or check WAR pattern information.

And/or, the idempotent kernel generating device may check alias pointer information. When the global memory variable (pointer) passed to the GPU kernel function is stored in the function's local pointer variable and used, access to the corresponding local pointer variable may need to be checked because it is actually an access to the global memory variable. The idempotent kernel generation device recognizes when a local pointer variable is declared in a function, and traces the variables used to initialize the variable, and can finally check and/or create associated global memory variable and/or WAR pattern information. .

(Idempotent GPU kernel generation)

Next, the idempotent kernel generation device can create an idempotent GPU kernel. In the first embodiment, a static idempotent kernel can be created.

The idempotent kernel generating device may perform redirection to guarantee the idempotency of the GPU kernel. The idempotent kernel generation device may redirect access to variables in which WAR risk occurs to other areas based on WAR pattern information. When a write access occurs after a read access by an initial variable access, the kernel generating device of the plane can redirect the write operation to another additionally allocated global memory area before the write access occurs. The idempotent kernel generating device may replace the variable name with respect to the redirected variable for writing so that both reading and/or writing thereafter are performed in the redirected area. The idempotent kernel generation device can maintain data consistency by reflecting the redirected area in the original after the GPU kernel is all executed and immediately before the kernel is terminated.

Second embodiment

Next, a method of creating an idempotent kernel through an access flag bitmap will be described in detail.

An idempotent kernel generating device can generate an idempotent GPU kernel in consideration of dynamic determining factors. The idempotent kernel generation device can check for amorphous idempotence in code that is not statically determined. If the access and/or location of the variable cannot be statically determined according to the conditional statement and/or the dynamically determined index in the source code, the idempotence of the code cannot be guaranteed through name conversion.

In order to generate the idempotent GPU kernel based on the dynamically determined element, the idempotent kernel generation device is an access flag bit capable of tracking read access and/or write access to the absolute position in the memory of the corresponding variable, as shown in FIG. 2. You can insert code to set up the map. As shown in FIG. 2, the apparatus for generating an idempotent kernel may map a 1-bit access flag bit to a region determined based on an index of a global memory variable and replace it with 1 when an access occurs.

FIG. 2 is a diagram for explaining an overall overview of an idempotent kernel generating apparatus according to an exemplary embodiment. Referring to FIG. 2, an idempotent kernel is statically created through the GPU Kernel Transfiler, which is a sub-item, and due to a dynamically determined element (for example, when an index is dynamically determined, it is not possible to know where the index is accessed, so WAR tracking is impossible). A part that cannot be analyzed and/or converted can be tracked through a Global Memory Access Flag Bitmap.

And/or, if a WAR memory access pattern is found through such a static and/or dynamic method, a redirection space instead of the original space in the global memory region in order to maintain the original region. Write to, and then read from here. In the case of a pattern other than WAR, the original space can be accessed as it is.

3 is a flowchart illustrating an operation of an idempotent kernel generating apparatus according to an embodiment of the present invention.

2 and 3, first, in order to check (S320) the idempotent kernel generating device the WAR hazard (pattern) of the global memory variable within the GPU kernel function that is not statically determined (S320), the GPU kernel function source The code can be statically analyzed (S310).

Next, the idempotent kernel generating device may determine whether or not there is a risk of WAR in the global memory variable within the GPU kernel function through static analysis of the GPU kernel function source code. If there is no risk of WAR, the idempotent kernel generation device may insert a scheduling hint delivery code into the GPU device driver (S350).

Next, the idempotent kernel generating device generates an access flag bitmap setting code (or a dynamic flag setting code) to track read access and/or write access to the variable when the risk of WAR is identified. Can be inserted (S330). Through the code, the value of the access flag bitmap may be changed according to read access or write access, as shown in FIG. 2.

Next, the idempotent kernel generating device may insert a redirection code for a corresponding variable according to read access or write access (S340). Through the code, the processing device can dynamically check the WAR risk based on the access flag bitmap, determine whether to use the source area of the global memory or the redirection area of the global memory, and perform the operation.

4 is a simple example of source code for explaining the operation of the idempotent kernel generation device. Fig. 4(a) shows the kernel function code without change, and Fig. 4(b) shows the kernel function code generated and/or converted by the idempotent kernel generation device.

4(b), rbit[(tid+i)/8]=!(wbit[(tid+i)/8]&(tid+i)%8)I(rbit) of FIG. 4(b) [(tid+i)/8]&(tid+i)%8); And, wbit[(tid+i)/8]=1<<(tid+i)%8; is the source code and/or operation that sets the access flag bitmap, which is the data type size of the global memory variable and 1 bit. It includes an operation that maps the flag of and also traces the presence or absence of a RAW pattern by referring to the previous write access flag for read access, so that the WAR risk according to the actual access can be checked, not the code order.

For example, an arbitrary global memory variable used in a GPU kernel function can be configured in an arbitrary type array. For such a global memory variable (array), a unit (array element) accessible by an index can be mapped to one bit of the access flag memory area. For example, when a global memory variable G called int G[10] is transferred, one bit of the rbit[] and/or wbit[] memory areas may correspond to a word indicated by one index of G.

In addition, rbit has two roles: recording the fact that a read has occurred when the kernel accesses the global memory area during operation, and if a RAW (Read after Write) pattern has occurred, it is It may be to record the presence or absence. The reason for setting the rbit by checking whether RAW has occurred is that the reason for RAW occurrence is that the original data area is always initialized and then read occurs, so idempotency (even if it is stopped at any moment during execution and then re-executed, the same result This is because it does not harm what you can expect).

The idempotent kernel generating device may determine whether or not to perform a redirection according to a dynamically determined variable access pattern.

In order to dynamically determine the redirection, an initial access can be made in the original area by declaring a temporary pointer variable (A_addr) as a local variable inside the function as shown in FIG. 4(b) and initializing it with the transferred global memory variable.

After the access flag for read and/or write is set, before the write occurs, it is possible to determine whether to use the temporary pointer variable as the original area or the redirect area depending on whether or not WAR occurs, and access after that is made in the replaced area. .

5 is a flowchart illustrating a method of generating an idempotent kernel according to an embodiment of the present invention.

The idempotent kernel generation method shown in FIG. 5 may be performed by an idempotent kernel generation device. For example, the idempotent kernel generation device may be a graphic processing unit (GPU), a central processing unit (CPU), or other computing devices.

Referring to FIG. 5, first, the idempotent kernel generation method includes generating WAR (Write After Read) pattern information from source code related to a GPU kernel function (S510), and a first memory variable based on the WAR pattern information. When the write access is confirmed after the read access that causes the risk of WAR for the write access, inserting a code for replacing the first memory variable with the second memory area variable before the code for write access (S520) may be included.

The WAR pattern information may include one or more global memory variables in which a WAR hazard may or may occur. Alternatively, the WAR pattern information may include an operator and/or source code that causes a WAR risk.

The first memory variable and the second memory variable may be global memory variables.

Generating WAR pattern information (S510) includes: checking a GPU kernel function in the source code (S511), checking a global memory variable transferred to the GPU kernel function (S512), and the GPU kernel It may include generating the WAR pattern information (S513) based on a function and the global memory variable.

First, the idempotent kernel generating device can check the GPU kernel function in the source code (S511).

For example, the source code may be General-Purpose on Graphics Processing Units (GPGPU) program source code.

The idempotent kernel generating apparatus may check (or distinguish) the GPU kernel function and the device function according to a specific indicator. For example, in the case of the OpenCL library, a kernel function can be indicated through the __kernel directive (specific indicator), and in the case of the CUDA library, the kernel function can be expressed through the __global__ directive (specific indicator).

For example, in the parallel computing library (CUDA/OpenCL) that supports the GPGPU programming model, functions executed in the GPU are classified into kernel functions and device functions through specific indicators. Accordingly, the idempotent kernel generating apparatus can classify (or confirm) a function executed by the GPU hardware into a kernel function and a device function through this indicator.

And/or, the idempotent kernel generating device may generate information about a function for static analysis. Information on the function may include indicator information, function information, parameter information, and/or return information. For example, the indicator information may be information on an indicator indicating whether a corresponding function is a kernel function or a device function. The function information may be information indicating a function name. The parameter information may be information indicating a parameter used in a function. The return information may be information indicating the return type of the function. And/or, the idempotent kernel generating device may generate a structure through information on the function.

Here, the generated structure may refer to a class that manages information of each kernel function. In this class, the function name, parameter, return value, and number of lines written in the code can be included as basic information. In addition, a structure for managing access patterns for global memory variables and alias pointers passed to kernel functions to be described later may be included. For example, the structure for managing information on a function in FIG. 9 may be a variable Func Info of Class FI.

And/or, the idempotent kernel generating apparatus may check (or distinguish) device functions by tracking function call information in the source code.

And/or, the idempotent kernel generation device tracks the WAR risk inside each function block in order to track the WAR risk that may occur due to the function call relationship while executing the GPU code, and then tracks the global memory variable transfer relationship between the functions, You can check the global WAR risk of the variable.

Next, the idempotent kernel generation device may check information on the global memory variable and/or the global memory variable transmitted to the GPU kernel function (S512).

The parameters of the GPU function may have the form of a call by value and a call by reference.

In addition, global memory variables may be transferred in the form of a call by reference. For example, the global memory variable may be transmitted in the form of a pointer or a reference to a GPU function. Therefore, the idempotent kernel generation device can check the global memory variable through the'*' and/or'&' indicators.

Next, the step of generating WAR pattern information based on the GPU kernel function and the global memory variable (S513) is a step of checking read access and write access through an operator related to the GPU kernel function and the global memory variable (S513-1). And, it may include a step (S513-2) of generating WAR pattern information by read access and write access to the global memory variable. Here, the operator related to the global memory variable may mean an operator using the global memory variable.

For example, the idempotent kernel generation device can check read access and write access as follows through operators related to global memory variables.

(1) The assignment operator (=) can be verified as a write access to the leftmost variable and a read access to all other variables.

(2) The unary operator (++) can be verified as a write access after read access to the target variable.

(3) The compound operator (+=) can be verified as a write access after read access to the leftmost variable at the end.

(4) Function calls can be verified as read access and/or write access through global analysis after analyzing each function.

(5) Other operators can be identified as read access.

In this way, the idempotent kernel generation device can check read access and write access through an operator, and generate WAR pattern information including variables and/or codes for which write access occurs after read access.

And/or, the idempotent kernel generation device can generate WAR pattern information through the formation of a variable access management structure. For example, the idempotent kernel generating device can check the WAR risk by forming a management structure for each global memory variable and recording and/or storing sequential variable access patterns in the source code. Since global memory variables are passed by call by reference, the area pointed to by the pointer is structured, and the'.' Alternatively, it can be made with a pointer access in the form of'->'. The idempotent kernel generation device can check the variables that are finally accessed in consideration of this approach. Here, the management structure for the global memory variable may be Class Param, Class Value, and Class CallF generated by inheriting Class VI shown in FIG. 9. These can basically include the variable name, data type, location in the code (Line, Token) and Class UI (use info) to manage access information. In the case of Class UI, information on where the variable is used in the code (line, token) and how (use type) can be included.

And/or, since the location of the memory area is determined through the index of the variable in the form of an array, the risk of WAR can be checked by considering the index and WAR pattern information can be generated.

And/or, if the global memory variable passed to the GPU kernel function is stored in the function's local pointer variable and used, access to the local pointer variable can be viewed as access to the global memory variable, so it is necessary to track this. There may be. Therefore, the idempotent kernel generation device can check the local pointer variable declared in the function, trace the variables used to initialize the variable to check the WAR risk of the associated global memory variable, and generate WAR pattern information. .

Next, if the write access is confirmed after a read access that causes a risk of WAR to the first memory variable based on the WAR pattern information, the idempotent kernel generation device puts the first memory variable into the second memory area before the code for write access. It may include a step (S520) of inserting a code to replace with a variable.

6 is a flowchart illustrating a method of generating an idempotent kernel according to another embodiment of the present invention.

The method for generating an idempotent kernel illustrated in FIG. 6 may be performed by an idempotent kernel generating device. For example, the idempotent kernel generating device may be a graphic processing unit (GPU), a central processing unit (CPU), or an upper and lower processing unit.

Referring to FIG. 6, the method of generating an idempotent kernel includes a step (S610) of checking a first memory variable related to a risk of a write after read (WAR) that is not statically determined in a source code related to a GPU kernel function (S610), Inserting a code for setting an access flag bitmap for tracking read access or write access to the first memory variable (S620), a temporary pointer variable in the GPU kernel function is declared as a local variable, and the first memory variable Initializing the temporary pointer variable to (S630) and inserting a code for dynamically determining whether to use the temporary pointer variable as a first memory variable or a second memory variable after read access to the temporary pointer variable and before the write access occurs (S640) ) Can be included. The first memory variable and the second memory variable may be global memory variables.

For example, variables related to a WAR (Write After Read) risk that are not statically determined may be variables related to a conditional statement or a dynamically determined index. And/or, the idempotent kernel generation device tracks the presence or absence of a RAW pattern, so that the risk of WAR according to the actual approach, not the code order, can be identified.

The access flag bitmap may be a set of bits indicating write access or read access to each of the regions included in the global memory. Each bit of the bit set may be mapped to a different area included in the global memory.

The code for setting the access flag bitmap may make the bit mapped to the specific area have a value of '1' when a read or write access occurs to a specific area included in the global memory. According to an embodiment, a bit mapped to a specific region may include a bit related to a read access and/or a bit related to a write access. For example, when a read access occurs, a bit related to read access may have a value of '1', and when a write access occurs, a bit related to write access may have a value of '1'.

The idempotent kernel generation method of the present invention has been described by dividing into the idempotent kernel generation method shown in FIG. 5 and the idempotent kernel generation method shown in FIG. 6, but all and/or some configurations of the idempotent kernel generation method of FIG. And, an idempotent kernel generation method in which all and/or some components of the idempotent kernel generation method of FIG. 6 are combined and/or substituted is also possible.

Through this, the present invention can guarantee priority scheduling in a real-time system using a GPU.

In addition, the present invention is utilized as a checkpointing technology, so that it is possible to build a fault tolerant system in a large-scale system using an increasing GPU.

In addition, the present invention may include a GPU kernel operating in a cloud computing environment using a GPU, and support immediate scheduling to enable a system to guarantee QoS of a service.

7 is a diagram illustrating an idempotent kernel generation device according to an embodiment proposed in the present specification.

Referring to FIG. 7, the idempotent kernel generation apparatus 700 may include a pattern generation unit 701 and a redirection unit 702.

The pattern generator 701 may generate write after read (WAR) pattern information from source codes related to a GPU kernel function. The idempotent kernel generation apparatus 700 may check (or distinguish) a kernel function and a device function by a GPU specific indicator.

The pattern generator 701 may check the GPU kernel function in the source code, check a global memory variable transmitted to the GPU kernel function, and generate WAR pattern information based on the GPU kernel function and the global memory variable.

The pattern generator 701 may check read access and write access through an operator related to a GPU kernel function and a global memory variable, and generate WAR pattern information by read access and write access to a global memory variable.

The redirection unit 702, when the write access is confirmed after the read access to the first memory variable in which a WAR risk occurs based on the WAR pattern information, before the code for write access, the first memory variable is converted to the second memory variable. Replace code can be inserted. The first memory variable and the second memory variable may be global memory variables.

The operation method of the idempotent kernel generating apparatus shown in FIG. 7 is the same as the method of generating the idempotent kernel described with reference to FIG. 5, and thus detailed descriptions thereof will be omitted.

8 is a diagram illustrating an idempotent kernel generation device according to another embodiment proposed in the present specification.

Referring to FIG. 8, the idempotent kernel generation apparatus 800 may include an analysis unit 801, a flag code insertion unit 802, a variable declaration unit 803, and a decision code insertion unit 804. .

The analysis unit 801 may check a first memory variable related to a risk of write after read (WAR) that is not statically determined in the source code related to the GPU kernel function.

The flag code insertion unit 802 may insert a code for setting an access flag bitmap for tracking read access or write access to the first memory variable.

The variable declaration unit 803 may declare a temporary pointer variable in the GPU kernel function as a local variable and initialize it as a first memory variable.

The decision code insertion unit 804 may insert a code for dynamically determining whether to use the temporary pointer variable as a first memory variable or a second memory variable after a read access to the temporary pointer variable and before a write access occurs. The first memory variable and the second memory variable may be global memory variables.

Variables related to a WAR (Write After Read) risk that are not statically determined may be variables related to a conditional statement or a dynamically determined index.

The access flag bitmap may be a set of bits indicating write access or read access to each of the regions included in the global memory. Each bit of the bit set may be mapped to a different area included in the global memory.

The code for setting the access flag bitmap may make the bit mapped to the specific area have a value of '1' when a read or write access occurs to a specific area included in the global memory.

The method of operating the idempotent kernel generating apparatus shown in FIG. 8 is the same as the method of generating the idempotent kernel described with reference to FIG. 6, and thus detailed descriptions thereof will be omitted.

The idempotent kernel generation device of the present invention has been described by dividing into the idempotent kernel generation device shown in FIG. 7 and the idempotent kernel generation device shown in FIG. 8, but all and/or some configurations of the idempotent kernel generation device of FIG. And, an idempotent kernel generation device in which all and/or some configurations of the idempotent kernel generation device of FIG. 8 are combined and/or substituted is also possible.

The embodiments described above are those in which components and features of the present invention are combined in a predetermined form. Each component or feature should be considered optional unless explicitly stated otherwise. Each component or feature may be implemented in a form that is not combined with other components or features. In addition, it is also possible to construct an embodiment of the present invention by combining some components and/or features. The order of operations described in the embodiments of the present invention may be changed. Some configurations or features of one embodiment may be included in other embodiments, or may be replaced with corresponding configurations or features of other embodiments. It is apparent that claims that do not have an explicit citation relationship in the claims may be combined to constitute an embodiment or may be included as a new claim by amendment after filing.

The embodiment according to the present invention may be implemented by various means, for example, hardware, firmware, software, or a combination thereof. In the case of implementation by hardware, an embodiment of the present invention provides one or more ASICs (application specific integrated circuits), DSPs (digital signal processors), DSPDs (digital signal processing devices), PLDs (programmable logic devices), and FPGAs ( field programmable gate arrays), processors, controllers, microcontrollers, microprocessors, etc.

In the case of implementation by firmware or software, an embodiment of the present invention may be implemented in the form of a module, procedure, or function that performs the functions or operations described above. The software code can be stored in a memory and driven by a processor. The memory may be located inside or outside the processor, and may exchange data with the processor through various known means.

It is obvious to those skilled in the art that the present invention can be embodied in other specific forms without departing from the essential features of the present invention. Therefore, the above detailed description should not be construed as restrictive in all respects and should be considered as illustrative. The scope of the present invention should be determined by rational interpretation of the appended claims, and all changes within the equivalent scope of the present invention are included in the scope of the present invention.

Claims (10)

In the method of generating an idempotent kernel generating device including an analysis unit, a flag code insertion unit, a variable declaration unit, and a decision code insertion unit,
Checking a first memory variable related to a risk of a write after read (WAR) that is not statically determined in a source code related to a GPU kernel function under the control of the analysis unit;
Inserting a code for setting an access flag bitmap for tracking read access or write access to the first memory variable under the control of the code insertion unit;
Declaring a temporary pointer variable in the GPU kernel function as a local variable under the control of the variable declaration unit and initializing it as the first memory variable; And
And inserting a code for dynamically determining whether to use the temporary pointer variable as a first memory variable or a second memory variable after a read access to the temporary pointer variable under the control of the decision code insertion unit and before a write access occurs. Idempotent kernel generation method, characterized in that.
The method of claim 1,
Variables related to the risk of write after read (WAR) that are not statically determined are variables related to conditional statements or dynamically determined indexes.
The method of claim 1,
The access flag bitmap is a set of bits indicating write access or read access to each of the regions included in the global memory.
The method of claim 3,
The code for setting the access flag bitmap is an idempotent kernel generation method in which a bit mapped to the specific area has a value of 1 when a read access or a write access to a specific area included in the global memory occurs.
The method of claim 1,
The first memory variable and the second memory variable are global memory variables.
An analysis unit that checks a first memory variable related to a write after read (WAR) risk that is not statically determined in the source code related to the GPU kernel function;
A flag code inserting unit inserting a code for setting an access flag bitmap for tracking read access or write access to the first memory variable;
A variable declaration unit for declaring a temporary pointer variable in the GPU kernel function as a local variable and initializing it as the first memory variable; And
And a decision code insertion unit for inserting a code for dynamically determining whether to use the temporary pointer variable as a first memory variable or a second memory variable after read access to the temporary pointer variable and before a write access occurs. Idempotent kernel generation device.
The method of claim 6,
The variables related to the risk of write after read (WAR) that are not statically determined are variables related to conditional statements or dynamically determined indexes.
The method of claim 6,
The access flag bitmap is an idempotent kernel generation device, which is a set of bits indicating write access or read access to each of the regions included in the global memory.
The method of claim 8,
The code for setting the access flag bitmap is an idempotent kernel generating device that, when a read access or a write access occurs to a specific area included in the global memory, a bit mapped to the specific area has a value of 1.
The method of claim 6,
The first memory variable and the second memory variable are global memory variables.

Images