KR102267500B1 - Idempotent kernel generateing method and apparatus - Google Patents

Abstract

The present specification discloses a method and apparatus for generating an idempotent kernel. Specifically, the method for generating an idempotent kernel according to an embodiment of the present invention includes the steps of checking 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 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 inserting 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 it as a first memory variable; and inserting a code for dynamically determining whether the temporary pointer variable is 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 inserting unit and before a write access occurs. do.

Description

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 occurrence of a data hazard due to the interruption of a GPU.

The Graphics Processing Unit (GPU) does not support the preemption function due to hardware limitations. Unlike the central processing unit (CPU), the GPU executes threads in parallel on multiple cores, so the cost of saving and retrieving a context is high for a context switch in an instruction unit.

Also, since the GPU does not support preemptive scheduling, a priority inversion phenomenon may occur. Also, if there is a kernel that has already started to be executed, the GPU cannot perform context switching to another kernel until execution of the kernel is completed. Here, the priority reversal problem refers to a phenomenon in which a high-priority task is waiting for execution due to the execution of a low-priority task and execution is delayed. This is a problem that must be solved in real-time systems.

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

An object of the present specification is to provide a method for ensuring that the GPU kernel is stopped and re-executed without application interruption in a GPU general-purpose computing environment.

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

In addition, an object of the present specification is to provide a method that guarantees data consistency and can be re-executed 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 not mentioned will be clearly understood by those of ordinary skill in the art to which the present invention belongs from the following description. will be able

The method for generating an idempotent kernel according to an embodiment of the present invention includes: identifying a first memory variable related to a write after read (WAR) risk that is not statically determined in a source code related to a GPU kernel function; Inserting a code for setting an access flag bitmap for tracking read access or write access to variables; Declaring a temporary pointer variable in the GPU kernel function as a local variable, and initializing it as the first memory variable and inserting code that dynamically determines whether the temporary pointer variable is 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, the variables related to the write after read (WAR) risk that are not statically determined may be variables related to a conditional statement or a dynamically determined index.

Also, 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.

Also, 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 apparatus for generating an idempotent kernel according to an embodiment of the present invention includes an analysis unit that identifies a first memory variable related to a write after read (WAR) risk that is not statically determined in a source code related to a GPU kernel function, and the first A flag code insertion unit for inserting 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 and a decision code for inserting code for dynamically determining whether the temporary pointer variable is a first memory variable or a second memory variable after a read access to the temporary pointer variable and before a write access occurs It may be characterized in that it includes an insert.

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

Also, 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.

Also, 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 ensuring that the GPU kernel is stopped and re-executed without stopping the application in the GPU general-purpose computing environment.

In addition, according to the present specification, there is an effect that scheduling according to the priority of the GPU kernel can be supported without hardware support.

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

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

In addition, according to the present specification, there is an effect that preemptive priority scheduling can be supported by performing an idempotent GPU kernel without additional context storage without hardware support.

The effects obtainable 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 to which the present invention pertains from the description below. .

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

Hereinafter, preferred embodiments according to the present invention will be described in detail with reference to the accompanying drawings. DETAILED DESCRIPTION The detailed description set forth below in conjunction with the appended 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 in order to provide a thorough understanding of the present invention. However, it will be apparent to one skilled in the art that the present invention may be practiced without these specific details.

Satisfying memory level idempotence in the GPU kernel means that the same result can be expected even if it is interrupted and restarted at any moment by ensuring consistency for the global memory variable passed to the kernel. .

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

If there is a memory access pattern in which a write to a variable occurs after a read access to the variable exists in the source code, the idempotence of the source code may not be guaranteed. If a variable with a WAR data hazard is initialized to a random value before initial access, even if the source code is re-executed, it is always performed with the same data, so the idempotence of the source code can be guaranteed. .

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 approach pattern. The normally executed code of FIG. 1( a ) may have a result in which the values of A and B are substituted with each other. However, if the source code is interrupted after storing B in A (A=B;), the expected result cannot be obtained because the initial value of A is changed to B when the source code is re-executed.

For example, it may be assumed that the initial values are A=1 and B=2. Normally, 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 substituted for each other. However, after changing the value of A to the value of B (A=B), if it is stopped, A=2, B=1, and when the source code is executed again, tmp stores the value of 2 (tmp =A), A maintains the value of 2 (A=B). Next, B maintains a value of 2 (B=tmp). This can lead to 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, the values of A and B show the expected results even if they are stopped at any time. can be obtained

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

The present specification proposes a method for generating an idempotent GPU kernel code that can be immediately re-executed without storing the GPU kernel in a GPU environment (a method for generating an idempotent kernel). The method can support immediate preemption support.

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

Hereinafter, the embodiments described in the present specification are only divided for convenience of description, and some methods and/or components of one embodiment may be substituted or applied in combination with methods and/or components of other embodiments. of course there is

first embodiment

First, a detailed look at how to create an idempotent kernel through static analysis.

(Static analysis to create idempotent GPU kernels)

The idempotent kernel generator may perform static analysis for generating an idempotent GPU kernel.

The idempotent kernel generation device traces the occurrence of a WAR data risk pattern (or WAR risk) in the code area executed on the GPU within the GPGPU program source code to check whether the code is idempotent, and to check the information of the variables that the WAR pattern occurs. have.

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

The parallel computing library (CUDA/OpenCL) that supports the GPGPU programming model classifies functions performed on the GPU into kernel/device functions through specific directives. Through these directives, the idempotent kernel generator can create a structure for static analysis by classifying functions performed in GPU hardware and extracting information about the indicators, function names, parameters, and return types.

In the case of OpenCL, a function called by the code executed on the GPU is defined in the form of an inline function. In order to track this, the idempotent kernel generator tracks the function call information in the code and distinguishes the device function information. .

The idempotent kernel generator first tracks the WAR risk inside each function block and then traces the global memory variable transfer relationship between functions in order to track the WAR risk that can occur due to the function call relationship during GPU code execution. WAR risk can be identified.

Next, the idempotent kernel generator may check the global memory variable information passed to the GPU function.

The parameters of the GPU function may have the form of a call by value and/or a call by reference as in general C and/or C++. A global memory variable may be passed in the form of a call by reference. Global memory variables are passed to GPU functions in the form of pointers or references, and the idempotent kernel generator uses ' * ' and/or ' & ' directives according to the C and/or C++ specification to determine global memory variables. can be checked.

(Generate WAR pattern information)

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

The idempotent kernel generator can be checked by distinguishing variable read and/or write access through operator patterns in the source code.

For example, (1) ' = ' assignment operator can be confirmed as write access to the last variable on the left, and read access to all other variables.

(2) The ' += ' compound operator can be checked as a read access and then a write access to the leftmost variable.

(3) The '++' unary operator can be identified as a read-and-write access to the target variable.

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

(5) All operators except for the above three operators can be identified as read access to variables.

And/or, the idempotent kernel generator may check WAR pattern information by forming a variable access management structure.

By forming a management structure for each global memory variable and recording sequential variable access patterns in the code, it is possible to check whether a WAR pattern occurs.

Since global memory variables are passed by call-by-reference, the area pointed to by the pointer is structured so that ' . ' or '->' type of pointer access can be made. Considering this case, it is possible to check the finally accessed variable.

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

And/or, the apparatus for generating an idempotent kernel may check alias pointer information. When using the global memory variable (pointer) passed to the GPU kernel function by storing it in the function's local pointer variable, it may be necessary to check the access to the local pointer variable because it is an access to the actual global memory variable. The idempotent kernel generator can recognize when a local pointer variable is declared within a function, track the variables used to initialize the variable, and finally check and/or create the associated global memory variable and/or WAR pattern information. .

(creating an idempotent GPU kernel)

Next, the idempotent kernel generating apparatus may generate an idempotent GPU kernel. In the first embodiment, a static idempotent kernel may be generated.

The idempotent kernel generator may perform redirection to ensure idempotency of the GPU kernel. The idempotent kernel generator may redirect access to a variable in which a WAR risk occurs to another area based on WAR pattern information. When a write access occurs after a read access as an initial variable access, the kernel-generated device may induce (redirection) to perform a write operation in another global memory area that is additionally allocated before the write access occurs. The idempotent kernel generator may substitute a variable name for a redirected variable for write so that both reads and/or writes thereafter are performed in the redirected area. The idempotent kernel generator can maintain data consistency by reflecting the redirected area to the original just before the kernel is terminated after the GPU kernel is all executed.

second embodiment

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

The idempotent kernel generator may generate an idempotent GPU kernel in consideration of a dynamic determinant. An idempotent kernel generator can check for non-deterministic idempotency in code that is not statically determined. If the access and/or position of the variable cannot be statically determined according to a conditional statement and/or a dynamically determined index in the source code, idempotence of the code cannot be guaranteed through name conversion.

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

2 is a diagram for explaining an overall outline of an apparatus for generating an idempotent kernel according to an embodiment. Referring to Figure 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 track the WAR because it is not known where it is accessed) Static It can be traced through the Global Memory Access Flag Bitmap for parts that cannot be analyzed and/or converted to

And/or, if a WAR memory access pattern is found through such a static and/or dynamic method, a redirection space, not the original space, in the global memory region to maintain the original region writes to, and then reads can be performed from there. In the case of a non-WAR pattern, access to the original space can be made as it is in the initial state.

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

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

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

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

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

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

4(b), rbit[(tid+i)/8]=!(wbit[(tid+i)/8]&(tid+i)%8)I(rbit) in 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 to set 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 flags of , and by tracking the existence of RAW patterns by referring to the previous write access flag for read access, you can check the WAR risk according to the actual access, not the code order.

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

In addition, the role of rbit is twofold. If the kernel accesses the global memory area during operation, it records the fact that a read has occurred, and if a RAW (Read after Write) pattern occurs, it is determined through wbit. It may be to record the presence or absence. The reason for setting the rbit by checking whether RAW occurs is idempotency (even if it is interrupted and re-executed at any moment during execution, the result is the same as the read occurs after the original data area is always initialized when RAW occurs). Because it does not harm what can be expected).

The idempotent kernel generator may determine whether or not redirection exists according to a dynamically determined variable access pattern.

In order to dynamically determine the redirection, the 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 global memory variable passed.

After the access flag for read and/or write is set, before write occurs, it is decided whether to use the temporary pointer variable as the original area or the redirection area according to the occurrence of WAR, and then access can be made in the replaced area. .

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

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

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

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

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

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

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

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

The apparatus for generating an idempotent kernel may identify (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 directive), and in the case of the CUDA library, the kernel function can be indicated through the __global__ directive (specific directive).

For example, in the parallel computing library (CUDA/OpenCL) that supports the GPGPU programming model, functions performed on the GPU are divided into kernel functions and device functions through specific directives. Accordingly, the idempotent kernel generating apparatus may classify (or check) a function performed in GPU hardware into a kernel function and a device function through such an indicator.

And/or, the idempotent kernel generating device may generate information about a function for static analysis. The information about 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 the 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 a return type of a function. And/or, the idempotent kernel generating apparatus may generate a structure through information on the function.

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

And/or, the apparatus for generating an idempotent kernel may identify (or identify) a device function by tracing function call information in the source code.

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

Next, the apparatus for generating an idempotent kernel may check information about 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, the global memory variable may be transferred in the form of a call by reference. For example, a global memory variable may be passed to a GPU function in the form of a pointer or a reference. Accordingly, the idempotent kernel generator may identify global memory variables through '*' and/or '&' directives.

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 the read access and the write access through the operator related to the GPU kernel function and the global memory variable (S513-1) and generating WAR pattern information by read access and write access to the global memory variable (S513-2). Here, the operator related to the global memory variable may mean an operator using the global memory variable.

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

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

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

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

(4) The call of a function can be confirmed as read access and/or write access through global analysis after analysis of each function.

(5) Other operators can be checked with read access.

The idempotent kernel generator may check read access and write access through the operator as described above, and may generate WAR pattern information including variables and/or codes in which write access occurs after read access.

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

And/or, since the location of the memory area is determined through the index of the array-type variable, the WAR risk may be checked and WAR pattern information may be generated in consideration of the corresponding index.

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

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

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

The idempotent kernel generating method illustrated in FIG. 6 may be performed by an idempotent kernel generating apparatus. For example, the idempotent kernel generating device may be a graphic processing unit (GPU), a central processing unit (CPU), or an out-of-state arithmetic unit.

Referring to FIG. 6, the idempotent kernel generation method includes a step (S610) of checking 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 (S610), Inserting a code for setting an access flag bitmap for tracking read access or write access to the first memory variable (S620), declares a temporary pointer variable in the GPU kernel function as a local variable, and the first memory variable (S630) and inserting a code for dynamically determining whether to use the temporary pointer variable as the first memory variable or the second memory variable after read access to the temporary pointer variable and before the write access occurs (S640) ) may be included. The first memory variable and the second memory variable may be global memory variables.

For example, variables related to a write after read (WAR) risk that are not statically determined may be variables related to a conditional statement or a dynamically determined index. And/or, the idempotent kernel generator can check the WAR risk according to the actual access rather than the code order by tracking the existence of RAW patterns.

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 region included in the global memory.

The code for setting the access flag bitmap may cause the bit mapped to the specific area to have a value of '1' when a read access or write access occurs to a specific area included in the global memory. According to an embodiment, the bit mapped to the 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 a read access may have a value of '1', and when a write access occurs, a bit related to a write access may have a value of '1'.

The method for generating an idempotent kernel of the present invention has been described as being divided into the method for generating an idempotent kernel shown in FIG. 5 and the method for generating an idempotent kernel shown in FIG. 6, but all and/or part of the method for generating the idempotent kernel shown in FIG. And, an idempotent kernel generating method in which all and/or some configurations of the idempotent kernel generating method of FIG. 6 are combined and/or substituted is also possible.

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

In addition, the present invention can be utilized as a checkpointing technology to enable the establishment of a fault tolerant system in a large-scale system using GPUs, which are gradually increasing.

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

7 is a diagram illustrating an apparatus for generating an idempotent kernel according to an embodiment proposed in the present specification.

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

The pattern generator 701 may generate write after read (WAR) pattern information from source code related to the GPU kernel function. The idempotent kernel generating apparatus 700 may identify (or distinguish) the kernel function and the device function according to the GPU specific indicator.

The pattern generator 701 may check the GPU kernel function from the source code, check the 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 the GPU kernel function and global memory variable, and generate WAR pattern information by read access and write access to the global memory variable.

The redirection unit 702 converts the first memory variable to the second memory variable before the code for the write access when the write access is confirmed after the read access where the WAR risk to the first memory variable occurs based on the WAR pattern information. Substitution code can be inserted. The first memory variable and the second memory variable may be global memory variables.

Since the operation method of the apparatus for generating an idempotent kernel shown in FIG. 7 is the same as the method for generating an idempotent kernel described with reference to FIG. 5 , a detailed description other than that is omitted.

8 is a diagram illustrating an apparatus for generating an idempotent kernel according to another embodiment proposed in the present specification.

Referring to FIG. 8 , the idempotent kernel generating 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 analyzer 801 may identify 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.

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 inserting unit 804 may insert a code for dynamically determining whether to use the temporary pointer variable as the first memory variable or the 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 write after read (WAR) 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 region included in the global memory.

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

Since the operation method of the apparatus for generating an idempotent kernel shown in FIG. 8 is the same as the method for generating an idempotent kernel described with reference to FIG. 6 , a detailed description other than that is omitted.

The device for generating an idempotent kernel of the present invention has been described as being divided into the device for generating an idempotent kernel shown in FIG. 7 and the device for generating an idempotent kernel shown in FIG. 8, but all and/or part of the device for generating the idempotent kernel shown in FIG. And, an idempotent kernel generating device in which all and/or some configurations of the idempotent kernel generating device of FIG. 8 are combined and/or substituted is also possible,

The embodiments described above are those in which elements 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 configure an embodiment of the present invention by combining some elements and/or features. The order of operations described in embodiments of the present invention may be changed. Some features or features of one embodiment may be included in another embodiment, or may be replaced with corresponding features or features of another embodiment. It is obvious that claims that are not explicitly cited in the claims can be combined to form an embodiment or included as a new claim by amendment after filing.

Embodiments 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 application specific integrated circuits (ASICs), digital signal processors (DSPs), digital signal processing devices (DSPDs), programmable logic devices (PLDs), FPGAs ( field programmable gate arrays), a processor, a controller, a microcontroller, a microprocessor, and the like.

In the case of implementation by firmware or software, an embodiment of the present invention may be implemented in the form of modules, procedures, functions, etc. that perform the functions or operations described above. The software code may be stored in the memory and driven by the processor. The memory may be located inside or outside the processor, and may transmit/receive data to and from the processor by various known means.

It is apparent to those skilled in the art that the present invention may be embodied in other specific forms without departing from the essential characteristics of the present invention. Accordingly, the above detailed description should not be construed as restrictive in all respects but as exemplary. The scope of the present invention should be determined by a reasonable interpretation of the appended claims, and all modifications within the equivalent scope of the present invention are included in the scope of the present invention.

Claims (10)

delete delete delete delete delete an analysis unit that identifies 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 insertion unit for 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
After read access to the temporary pointer variable, before the write access occurs, it characterized in that it comprises a decision code insertion unit for inserting a code for dynamically determining whether the temporary pointer variable as a first memory variable or a second memory variable An idempotent kernel generator.
7. The method of claim 6,
The variable related to the write after read (WAR) risk that is not statically determined is an idempotent kernel generator that is a variable related to a conditional statement or a dynamically determined index.
7. The method of claim 6,
The access flag bitmap is an idempotent kernel generating apparatus that is a set of bits indicating a write access or a read access to each of the regions included in the global memory.
9. The method of claim 8,
The code for setting the access flag bitmap causes 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.
7. The method of claim 6,
wherein the first memory variable and the second memory variable are global memory variables.

Images