KR101980999B1 - Method and apparatus for bypassing thread group level cache - Google Patents

Abstract

The present invention relates to a method and apparatus for bypassing cache with a thread group level. According to one embodiment of the present invention, the method performed by the apparatus comprises the steps of: extracting memory information and cache information for a memory and cache provided in the apparatus; determining at least one switching point for each thread group using cache by using the extracted memory information and cache information; and bypassing the cache used by the thread group according to the determined switching point.

Description

[0001] METHOD AND APPARATUS FOR BYPASSING THREAD GROUP LEVEL CACHE [

BACKGROUND OF THE INVENTION 1. Field of the Invention The present invention relates to a cache bypassing technique, and more particularly, to a thread-group level cache bypassing method and apparatus.

Modern graphics processing units (GPUs) have widely adopted cache hierarchies as shown in Table 1 below. A relatively small cache of graphics processing units can easily become congested with thousands of threads.

Modern graphics processing units (GPUs) employ a multi-tiered memory architecture to reduce long global memory access latency. The graphics processing unit has an L1 / L2 cache managed by hardware. However, due to the limited size of cache and memory resources, bottlenecks occur and performance degradation occurs.

1 is a diagram for explaining a general cache bypassing method.

As shown in FIG. 1, the cache bypassing design scheme allows certain instructions or some thread groups to bypass the cache to address these bottlenecks. First, an instruction-level cache bypassing method 110 and a thread-level cache bypassing method 120 are used in the cache bypassing method. Here, the thread group may be referred to as " Warp " or " Tread Group ".

The thread group level cache bypassing method 120 is better than the instruction level cache bypassing method 110 in GPU computing. Because the thread group level cache bypassing method 120 can use more cache and memory resources than the instruction level cache bypassing method 110, Thus, cache bypassing techniques are being explored to allow other data, threads, and more cache and memory resources to be used. Thread-group-level cache bypassing, which can eliminate memory resource concentration, performs better in graphics processing units (GPUs), where thousands of threads execute concurrently in parallel. The thread-group-level cache bypassing method uses a cache for thread group 0 (121) for Warp 0, thread group 1 (122) for Warp 1, and some threads group Warp 2 < / RTI > 123 is to bypass the cache.

2 is a diagram for explaining a difference in execution time between thread groups.

As shown in FIG. 2, if the entire program is cached or bypassed collectively, there will be a difference in execution time between the thread groups. Even if you do the same task, the workload, the thread group that uses the cache will run faster. On the other hand, a group bypassing the cache takes a relatively long execution time. A thread block is a unit that is scheduled to a streaming multiprocessor (SM).

Korean Registered Patent No. 10-1662363 (registered on September 27, 2016)

In a general thread group level cache bypassing method, a fast-running thread group 202, 203, 204 must stall until a slow thread group 201 is completed, despite completion of execution. The scheduled thread block TB waits until the slow thread group 201 is completed and the execution is started after the slow thread group 201 is completed. That is, the execution of the next thread block (TB) can not be started before execution of all the thread groups of the reserved thread block (TB) is completed. For this reason, general thread group level cache bypassing techniques require improvements to achieve optimal performance improvements.

Embodiments of the present invention reduce the execution time gap between thread groups that occurs when a general thread group level cache bypassing technique is applied to a graphics processing apparatus, And to provide a method and apparatus for cache bypassing at a group level.

Embodiments of the present invention are based on the assumption that when a general thread group level cache bypassing technique is applied, a performance loss caused by a difference in execution time between a thread group using a cache and a thread group not using a cache Group-level cache bypassing that can achieve better performance than normal thread group-level cache bypassing by minimizing the execution time gap between thread groups by switching groups that use the cache at specific program points Method, and apparatus.

According to an embodiment of the present invention, there is provided a cache bypass bypassing method performed by a cache bypassing apparatus, comprising: determining at least one switching point for each thread group using a cache; And bypassing a cache used by the thread group according to the determined at least one switching point.

The method may further comprise extracting memory access information that occurs when accessing the provided memory and cache bypassing information that occurs when bypassing the cache, wherein determining the switching point comprises: At least one switching point may be determined for each thread group using the extracted memory access information and the extracted cache bypassing information.

The method may further comprise searching for at least one switchable point for each thread group.

The step of retrieving the switchable point may set at least one switchable point by searching for at least one point where the data cached by the thread group is not reused.

The step of searching for the switchable point may search for at least one switchable point according to the number of threads in the thread block divided by the total number of thread groups divided by the number of thread groups simultaneously using the cache.

The determining the switching point may determine at least one switching point for each thread group among the at least one switchable point searched.

Wherein determining the switching point comprises: calculating an execution time gap between a thread group using a cache and a thread group not using a cache among the searched at least one switchable point, One switching point can be determined as at least one switching point.

The step of determining the switching point may determine the at least one switchable point at which the difference in distance between the switching points among the at least one switchable point searched is minimized as the at least one switching point.

The method comprising: modifying the program code based on the determined switching point; And generating a cache bypassing program by compiling the modified program code, wherein bypassing the cache further comprises: caching the cache used by the thread group according to the generated cache bypassing program, You can pass it.

The modifying the program code may modify the program code using a first cache operator for bypassing the thread group and a second cache operator for caching of the thread group based on the determined switching point .

According to another embodiment of the present invention, there is provided a cache management method comprising: a switching point determining unit determining at least one switching point for each thread group using a cache; And a cache bypassing unit for bypassing a cache used by the thread group according to the determined at least one switching point.

Wherein the apparatus further comprises an information extracting unit for extracting memory access information generated when accessing the memory and cache bypassing information generated when bypassing the cache, At least one switching point may be determined for each thread group using the information and the extracted cache bypassing information.

The apparatus may further include a switching point searching unit for searching at least one switchable point for each thread group.

The switching point searching unit may search at least one point where data cached by the thread group is not reused to set the at least one switchable point.

The switching point searching unit may search at least one switchable point according to the number of the thread groups divided by the total number of the thread groups simultaneously by the number of the thread groups using the cache.

The switching point determination unit may determine at least one switching point for each thread group among the searched at least one switchable point.

Wherein the switching point determiner calculates an execution time difference between a thread group that uses a cache and a thread group that does not use a cache among the searched at least one switchable point and performs at least one switching operation in which the calculated execution time gap is minimized The possible points can be determined as at least one switching point.

The switching point determination unit may determine at least one switchable point at which the difference in distance between the switching points among the at least one switchable points searched is minimized as the at least one switching point.

The apparatus comprising: a code modifying unit for modifying the program code based on the determined switching point; And a program compiler for generating a cache bypassing program by compiling the modified program code, wherein the cache bypassing unit is configured to bypass the cache used by the thread group according to the generated cache bypassing program .

The code modifying unit may modify the program code using a first cache operator for bypassing the thread group and a second cache operator for caching the thread group based on the determined switching point.

Embodiments of the present invention solve the performance degradation phenomenon by reducing the execution time gap between thread groups that occurs when a general thread group level cache bypassing technique is applied to a graphics processing apparatus, thereby achieving better performance.

Embodiments of the present invention are based on the assumption that when a general thread group level cache bypassing technique is applied, a performance loss caused by a difference in execution time between a thread group using a cache and a thread group not using a cache To achieve a better performance than normal thread group level cache bypassing by minimizing the execution time gap between thread groups by switching groups that use the cache at specific program points.

1 is a diagram for explaining a general cache bypassing method.
2 is a diagram for explaining a difference in execution time between thread groups.
3 is a block diagram illustrating a configuration of a thread bypass leveling apparatus according to an embodiment of the present invention.
4 is a diagram for explaining a general thread group level cache bypassing method.
5 is a diagram for explaining a switching point in a thread bypassing method of a thread group level according to an embodiment of the present invention.
FIG. 6 is a view for explaining a plurality of switching points in a thread bypassing method of a thread group level according to an embodiment of the present invention.
FIG. 7 is a diagram for explaining a search for a switchable point in a thread bypassing method of a thread group level according to an embodiment of the present invention.
8 is a diagram for explaining switching point determination in a thread bypass level cache-level method according to an embodiment of the present invention.
9 to 12 are diagrams for explaining a code modification procedure in a thread bypass level cache-level bypassing method according to an embodiment of the present invention.
FIG. 13 is a flow chart for explaining a flow of a thread-group-level cache bypassing method according to an embodiment of the present invention.
14 is a view showing execution time according to a cache unused, a cache use, and a general thread group level cache bypass passing method.
FIG. 15 is a diagram illustrating an execution time according to a general cache bypass passing method and a cache group bypassing method according to an embodiment of the present invention.
16 is a diagram illustrating an execution time according to a conventional cache bypass transfer method that is not balanced and a balanced cache bypass transfer method according to an embodiment of the present invention.

While the invention is susceptible to various modifications and alternative forms, specific embodiments thereof are shown by way of example in the drawings and will herein be described in detail.

It is to be understood, however, that the invention is not to be limited to the specific embodiments, but includes all modifications, equivalents, and alternatives falling within the spirit and scope of the invention.

The terms first, second, etc. may be used to describe various components, but the components should not be limited by the terms. The terms are used only for the purpose of distinguishing one component from another. For example, without departing from the scope of the present invention, the first component may be referred to as a second component, and similarly, the second component may also be referred to as a first component. And / or < / RTI > includes any combination of a plurality of related listed items or any of a plurality of related listed items.

It is to be understood that when an element is referred to as being "connected" or "connected" to another element, it may be directly connected or connected to the other element, . On the other hand, when an element is referred to as being "directly connected" or "directly connected" to another element, it should be understood that there are no other elements in between.

The terminology used in this application is used only to describe a specific embodiment and is not intended to limit the invention. The singular expressions include plural expressions unless the context clearly dictates otherwise. In the present application, the terms "comprises" or "having" and the like are used to specify that there is a feature, a number, a step, an operation, an element, a component or a combination thereof described in the specification, But do not preclude the presence or addition of one or more other features, integers, steps, operations, elements, components, or combinations thereof.

Unless defined otherwise, all terms used herein, including technical or scientific terms, have the same meaning as commonly understood by one of ordinary skill in the art to which this invention belongs. Terms such as those defined in commonly used dictionaries should be interpreted as having a meaning consistent with the meaning in the context of the relevant art and are to be interpreted in an ideal or overly formal sense unless explicitly defined in the present application Do not.

Hereinafter, preferred embodiments of the present invention will be described in detail with reference to the accompanying drawings. In order to facilitate the understanding of the present invention, the same reference numerals are used for the same constituent elements in the drawings and redundant explanations for the same constituent elements are omitted.

3 is a block diagram illustrating a configuration of a thread bypass leveling apparatus according to an embodiment of the present invention.

3, a thread group level cache bypassing apparatus 300 according to an embodiment of the present invention includes an information extracting unit 310, a switching point determining unit 330, and a cache bypassing unit 360, . Here, the cache bypassing apparatus 300 according to another embodiment of the present invention may further include a switching point searching unit 320. In addition, the cache bypassing apparatus 300 according to another embodiment of the present invention may further include a code correction unit 340 and a program compiler 350.

However, not all illustrated components are required. The cache bypassing apparatus 300 of the thread group level may be implemented by a larger number of components than the illustrated components, and the thread group level cache bypassing apparatus 300 may be implemented by fewer components.

The concrete configuration and operation of each component of the cache bypassing apparatus 300 of the thread group level in FIG. 1 will be described below.

The information extracting unit 310 extracts memory information and cache information for a memory and a cache provided in the graphics processing apparatus. The information extracting unit 310 may extract the memory access information generated when accessing the provided memory and the cache bypassing information generated when the provided cache is bypassed, as memory information and cache information.

The switching point determination unit 330 determines at least one switching point for each thread group using the cache by using the extracted memory information and cache information.

The cache bypassing unit 360 bypasses the cache used by the thread group according to the at least one switching point determined by the switching point determining unit 330. [

Meanwhile, the switching point searching unit 320 searches at least one switchable point for each thread group. The switching point searching unit 320 can search at least one point where data cached by the thread group is not reused and set the at least one switchable point. The switching point searching unit 320 may search for at least one switchable point according to the number of threads in the thread block divided by the total number of thread groups and the number of thread groups simultaneously using the cache.

When a switchable point is found, the switching point determining unit 330 can determine at least one switching point for each thread group among the detected at least one switchable point. The switching point determining unit 330 calculates an execution time difference between a thread group using a cache and a thread group not using a cache among at least one switchable point searched by the switching point searching unit 320. [ The switching point determination unit 330 may determine at least one switchable point at which the calculated execution time gap is minimized as the at least one switching point.

The switching point determining unit 330 may determine at least one switchable point at which the difference in distance between the switching points in the at least one switchable point detected by the switching point searching unit 320 is minimized as at least one switching point .

On the other hand, the code correcting unit 340 modifies the program code of the graphics processing apparatus based on the switching point determined by the switching point determining unit 330. [ Based on the switching point determined by the switching point determination unit 330, the code correction unit 340 performs a graphic processing (not shown) using the first cache operator for bypassing the thread group and the second cache operator for caching the thread group The program code of the device can be modified.

The program compiler 350 compiles the program code of the modified graphics processing unit in the code correction unit 340 to generate a cache bypassing program.

The cache bypassing unit 360 may perform cache bypassing at the thread group level for bypassing the cache used by the thread group according to the cache bypassing program generated by the program compiler 350. [

4 is a diagram for explaining a general thread group level cache bypassing method.

A common thread group level cache bypassing method divides the entire thread group into a thread group that uses the cache and a thread group that does not use the cache, and executes the thread group. A thread group may include at least one thread, and is not limited to a particular thread group. For example, at least one thread is op0; op1; op2; op3 < / RTI >

As shown in FIG. 4, thread group 0 410 indicated by Warp 0 and thread group 1 420 indicated by Warp 1 are cached thread groups using a cache. On the other hand, thread group 2 (430), indicated by Warp 2, and thread group 3 (440), indicated by Warp 3, are cache-bypassed thread groups without using cache.

Here, the execution of the thread group 0 410 and the thread group 1 420 using the cache is completed in advance compared to the thread group 2 430 and the thread group 3 440. Accordingly, the thread group 0 410 and the thread group 1 420 using the cache must wait until the execution of the thread group 2 430 and the thread group 3 440 is completed.

5 is a diagram for explaining a switching point in a thread bypassing method of a thread group level according to an embodiment of the present invention.

The method of bypassing a thread group level according to an embodiment of the present invention divides the entire thread group into a thread group that uses a cache and a thread group that does not use a cache, and executes a thread group, similar to FIG.

However, after the switching point, the cached thread group using the cache is bypassed, and the bypassing thread group that does not use the cache uses the cache.

As shown in FIG. 5, thread group 0 410 indicated by Warp 0 and thread group 1 420 indicated by Warp 1 use a cache, and cache-bypassed after the switching point It becomes a thread group.

On the other hand, thread group 2 (430) indicated by Warp 2 and thread group 3 (440) indicated by Warp 3 are bypassed and do not use the cache, but become a cached thread group using the cache after the switching point .

Here, the thread group 0 (410) and the thread group 1 (420) using the cache are almost similar to the thread group 2 (430) and the thread group 3 (440). This is because thread group 0 410, thread group 1 420, thread group 2 430, and thread group 3 440 both use the cache equally once. Therefore, it is not necessary to wait until the execution of another thread group is completed in both the thread group that uses the cache first or the thread group that uses the cache later. Accordingly, the thread bypass method of the thread group level according to an embodiment of the present invention can reduce the execution time difference between the thread groups.

FIG. 6 is a view for explaining a plurality of switching points in a thread bypassing method of a thread group level according to an embodiment of the present invention.

The thread bypass method of the thread group level according to an embodiment of the present invention switches each thread group individually on the basis of the switching point # 1 and the switching point # 2. For example, a thread-group-level cache bypassing method cache-bypasses a thread group using a cache so as not to use the cache after switching point # 1. As another example, the thread-group level cache bypassing method caches non-cached thread groups to use the cache after switching point # 1.

The individual cache bypass transfer method of each thread group will be described with reference to FIG. FIG. 6 shows a case where a thread block is composed of six thread groups and at the same time, the number of suitable thread groups using the cache is two.

Thread group 0 (410), which is represented by the first group Warp 0, and thread group 1 (420), which is represented by Warp 1, uses the cache before the switching point # 1. Thread group 0 410 and thread group 1 420 are switched at switching point # 1 by cache bypassing device 300. After switching point # 1, thread group 0 410 and thread group 1 420 are bypassed to not use the cache.

Thread group 2 430 indicated by Warp 2 and thread group 3 440 indicated by Warp 3 do not use the cache before switching point # 1. Thread group 2 430 and thread group 3 440 are switched at switching point # 1 by cache bypassing device 300. After switching point # 1, thread group 2 430 and thread group 3 440 are cached to use the cache. The thread group 2 430 and the thread group 3 440 are switched again at the switching point # 2 by the cache bypassing apparatus 300. After switching point # 2, thread group 2 430 and thread group 3 440 are bypassed so as not to reuse the cache.

Thread group 4 (450), represented by Warp 4, and thread group 5 (460), indicated by Warp 5, do not use a cache prior to switching point # 2. Thread group 4 450 and thread group 5 460 are switched at switching point # 2 by cache bypassing device 300. After switching point # 2, thread group 4 450 and thread group 5 460 are cached to use the cache.

FIG. 7 is a diagram for explaining a search for a switchable point in a thread bypassing method of a thread group level according to an embodiment of the present invention.

As shown in FIG. 7, the thread block 700 may be composed of a plurality of thread groups. The thread block 700 includes loads A (load A) to load H (load H), and a plurality of thread groups may each include a plurality of loads.

Here, the cache bypassing apparatus 300 searches for at least one switchable point with respect to the thread block 700. At this point, the switchable point represents the point where the cached data is no longer reused. That is, the cache bypassing apparatus 300 can search for at least one point at which data cached by the thread group is not reused, and set at least one point that is searched for as at least one switchable point.

In the example shown in Figure 7, the cache bypassing device 300 has three switchable points # 1 (Switchable point # 1), Switchable point # 2 (Switchable point # 2), and switchable point # 3 (switchable point # 3).

8 is a diagram for explaining switching point determination in a thread bypass level cache-level method according to an embodiment of the present invention.

As shown in FIG. 8, a thread block may be composed of eight thread groups (warp). Assuming that there are two proper thread groups that can use the cache at the same time, it consists of a total of four thread groups. Then, three switching points are required.

Thus, the cache bypassing apparatus 300 according to an embodiment of the present invention determines three switching points among the switchable points. At this time, the cache bypassing apparatus 300 can determine the switching point so that the distance between the switching points is the closest.

In the example shown in FIG. 8, the cache bypassing apparatus 300 can determine the switchable point # 2, the switchable point # 4, and the switchable point # 6 as the switching points among the switchable points # 1 to # 6. That is, the cache bypassing apparatus 300 may switch thread groups at the determined switching points # 2, # 4 and # 6.

9 to 12 are diagrams for explaining a code modification procedure in a thread bypass level cache-level bypassing method according to an embodiment of the present invention.

Hereinafter, the code modification process in the thread bypass level cache-level bypassing method will be described with reference to FIGS. 9 to 12. FIG.

As shown in FIG. 9, the cache bypassing apparatus 300 divides dedicated basic blocks (BB) in a thread group level cache bypassing method. Here, the dedicated basic block may be at the thread group level. For example, the cache bypassing apparatus 300 divides a dedicated basic block BB0 900 into a dedicated basic block BB0_0 910 and a dedicated basic block BB0_1 920 based on a switching point.

Here, if the switching point is the N switching points, the cache bypassing apparatus 300 allocates the dedicated basic block BB # to the dedicated basic block BB0 # _0, the dedicated basic block BB0 # , And a dedicated basic block BB0 # _N.

As shown in FIG. 10, the cache bypassing apparatus 300 may duplicate partitioned dedicated basic blocks and rename the duplicated dedicated basic blocks. By way of example, the cache bypassing device 300 may duplicate the partitioned dedicated basic block BB0_0 910 and reassign it to the dedicated basic blocks BB0_0_0 911 and BB0_0_1 912. In addition, the cache bypassing apparatus 300 can duplicate the divided dedicated basic block BB0_1 920 and reassign the dedicated basic blocks BB0_1_0 921 and BB0_1_1 922.

The cache bypassing device 300 may add cache operators to the dedicated basic block and the replicated dedicated basic block. For example, the cache bypassing apparatus 300 adds the bypassing operator " cg " to the instruction in the dedicated basic block BB0_0_0 911 to obtain ld.global.cg # 1, ld.global.cg # 2, ld.global. It can be modified like cg # 3. Conversely, the cache bypassing device 300 adds the caching operator " ca " to the instruction in the replicated dedicated basic block BB0_0_1 912 to obtain ld.global.ca # 1, ld.global.ca # 2, ld.global .ca # 3 and so on. The same can be applied to the dedicated basic block BB0_1_0 921 and the replicated dedicated basic block BB0_1_1 922.

As shown in FIG. 11, the cache bypassing apparatus 300 generates a control flow for each thread group. By way of example, the cache by-passing device 300 may determine that {BB # _0_1}, {BB # _1_0}, ... , {BB # _N - 0}.

For example, the cache bypassing apparatus 300 may execute a dedicated basic block BB0_0_1 912 to which a caching operator is added, and generate a control flow to execute a dedicated basic block BB0_1_0 921 to which a bypassing operator is added. In addition, the cache bypassing apparatus 300 may execute a dedicated basic block BB0_0_0 911 to which a bypassing operator is added, and generate a control flow to execute a dedicated basic block BB0_1_1 922 to which a caching operator is added. That is, the cache bypassing apparatus 300 can generate a control flow such that each thread group alternates with each other based on the switching point, and does not use the cache or use the cache. This allows the cache to be used fairly for each thread group.

As shown in FIG. 12, the cache bypassing apparatus 300 collects dedicated basic blocks. The cache bypassing device 300 inserts control flow instructions.

In this manner, the cache bypassing apparatus 300 can divide the dedicated basic block based on the switching point, and modify the parallel thread execution (PTX) code using the cache operators.

FIG. 13 is a flow chart for explaining a flow of a thread-group-level cache bypassing method according to an embodiment of the present invention.

The thread bypassing method of the thread group level according to an embodiment of the present invention includes the following steps in order to use an even cache for each thread group in the thread bypassing of the thread group level.

In step S101, the cache bypassing apparatus 300 of the thread group level extracts memory access information and cache bypassing information. The cache bypassing apparatus 300 extracts memory access information and cache bypassing information through an information extracting unit 310 created based on the GPU simulator. Here, the cache bypassing information may include a thread group ID for bypassing the cache.

In step S102, the thread-group-level cache bypassing device 300 searches for at least one switchable point for each thread group. The cache bypassing apparatus 300 searches all the switching points that do not drop the cache hit rate based on the extracted memory access information through the switching point searching unit 320. [ The cache bypassing apparatus 300 can reduce the execution time gap between the thread groups due to the difference in the memory access delay time which occurs in the cache bypassing at the general thread group level by using all of the thread groups alternately using the cache in a fairly alternate manner.

To this end, the cache bypassing device 300 modifies the program to switch the thread group that uses the cache at a particular point in the program. At this time, since the cache hit ratio may be lowered depending on the point at which the thread group using the cache is switched, the cache bypassing apparatus 300 searches for a switchable point where the cache hit ratio does not fall.

In step S103, the cache bypassing apparatus 300 of the thread group level determines at least one switching point for each thread group. The cache bypassing apparatus 300 determines the switching point as a point at which the data cached by the thread group that previously used the cache is no longer reused. The cache hit ratio is not lowered even when the cache bypassing apparatus 300 switches the thread group using the cache because the cached data at the corresponding switching point is not reused any more.

The cache bypassing apparatus 300 determines, through the switching point determining unit 330, a switching point to actually switch among all switchable points that do not drop the cache hit ratio. Switching should be minimized because no thread group will use the cache during the time difference to reach the switching point between thread groups in the switching process. Thus, the cache bypassing apparatus 300 can determine the switching point by the number of switching times among the switchable points. That is, the cache bypassing apparatus 300 can determine a minimum number of switching points.

The number of switching times can be determined by general thread group level cache bypassing. The cache bypassing apparatus 300 can calculate the total number of thread groups by increasing the value obtained by dividing the cache by the number of thread groups to be used. The cache bypassing apparatus 300 can determine the switching point by the number of switching times among the switchable points. At this time, the cache bypassing apparatus 300 may determine that the distance between the switching points is maximally similar or the difference between the respective switching points is the minimum.

In step S104, the cache bypassing apparatus 300 of the thread group level modifies the program code of the graphics processing apparatus. As an example, the thread-group level cache bypassing device 300 may modify the Parallel Thread Execution (PTX) code. When the switching point is determined in step S103, the cache bypassing apparatus 300 at the thread group level can modify the PTX code based on the information about the switching point to be applied to an actual program. The PTX code represents representative of the code written in the intermediate language of the GPU computing program.

Cache bypassing device 300 specifies which thread group will use the cache in a particular program area using the first cache operator and the second cache operator. Here, the first cache operator represents a cg bypassing operator for bypassing a thread group. The second cache operator represents a ca caching operator for caching of the thread group.

For example, if the program area 0 is expressed as BB0, the cache bypassing apparatus 300 attaches a first cache operator (cg bypassing operator) to the load instruction in BB0_0, which bypasses the cache using BB0. The cache bypassing apparatus 300 attaches a second cache operator (ca caching operator) to the load instruction in BB0_1 using the cache. Thereafter, the cache bypassing apparatus 300 at the thread group level can create the control flow to perform BB0_1 for the thread group to use the cache and BB0_0 to the thread group for bypassing the cache.

In step S105, the thread-group-level cache bypassing device 300 compiles the modified program code to generate a cache bypassing program. The thread-group level cache bypassing device 300 performs the compilation step. When the PTX code is generated, the GPU compiler (for example, NVCC, NVIDIA's CUDA Compiler) compiles the PTX code to generate an executable GPU program.

In step S106, the thread-group-level cache bypassing device 300 bypasses the cache used by the thread group according to the cache bypassing program.

Meanwhile, experimental results and experimental results comparing a general thread group level cache bypass passing method and a thread group level cache bypass passing method according to an embodiment of the present invention will be described.

First, we will explain the experiment to see how the general thread group level cache bypassing method causes the execution time gap between the cache-use group and the non-cache use group.

14 is a view showing execution time according to a cache unused, a cache use, and a general thread group level cache bypass passing method.

The experiment is to measure normalized execution time between thread groups in the first kernel of the ATAX application of a benchmark (Polybench / gpu) that performs mathematical operations such as matrix and linear algebra in the GPU simulator based on cycle accuracy (GPGPUSim) As shown in FIG. 14.

If the kernel does not use the L1 cache as a workload where all thread groups perform the same task, L1 Bypass (1401) has little runtime gap between the thread groups. Also, the L1 cache 1402, which uses the L1 cache due to the concentration of the L1 cache and memory resource requests, is rather degraded in performance. The WLB 1403, which is a case where thread group level cache bypassing is applied to efficiently use the L1 cache, shows a performance improvement as in the rightmost area in FIG. However, the WLB 1403 can see that the execution gap between the group using the cache and the group using the cache is about 26.53%.

FIG. 15 is a diagram illustrating an execution time according to a general cache bypass passing method and a cache group bypassing method according to an embodiment of the present invention.

The result shown in FIG. 15 can be obtained by applying the thread-group-level cache bypassing method according to an embodiment of the present invention which can reduce the execution time gap between the thread groups. The bar graph in FIG. 15 shows (1) L1 Cache Disable 1501 in which L1 cache is not used, (2) L1 Cache Enable 1502 in case of using L1 cache, and (3) thread group level cache bypassing (Cache Bypassing 1503), (4) Applying the method according to an embodiment of the present invention, Balanced Cache Bypassing 1504 measures the IPC (instruction per cycle) of a specific kernel of each application to be. The Balanced Cache Bypassing 1504 in the case of applying the method according to the embodiment of the present invention shows a higher performance improvement than the L1 Cache Disable 1501 and the L1 Cache Enable 1502. [ Balanced Cache Bypassing (1504), which is the case of applying the method according to the embodiment of the present invention, shows an average improvement of about 6% over Warp-level Cache Bypassing (1503).

16 is a diagram illustrating an execution time according to a conventional cache bypass transfer method that is not balanced and a balanced cache bypass transfer method according to an embodiment of the present invention.

Unbalanced Cache Bypassing 1601, which is an unbalanced general cache bypassing method, can see that the execution gap between a group using a cache and a group using no cache is about 26.53%, as described in FIG.

As shown in FIG. 16, Balanced Cache Bypassing 1602, which is a balanced cache bypassing method according to an embodiment of the present invention, includes a case where there is no execution gap between a group using a cache and a group using no cache Can be seen. In addition, Balanced Cache Bypassing 1602, which is a balanced cache bypassing method according to an embodiment of the present invention, reduces the computation time because each thread group is fairly cached, and unbalanced cache bypassing 1601, which is a general cache bypassing method, The performance is improved by 14.27%.

The thread-group-level cache bypassing method described above can be implemented as computer-readable code on a computer-readable recording medium.

A thread-group level cache bypassing method is a computer-readable storage medium comprising instructions executable by a processor, the instructions causing the processor to perform the steps of: extracting memory information and cache information for a memory and a cache, Determining at least one switching point for each thread group using the cache using the extracted memory information and cache information, and bypassing the cache used by the thread group according to the determined at least one switching point And a computer-readable storage medium configured to perform the method.

The computer-readable recording medium includes all kinds of recording media storing data that can be decoded by a computer system. For example, there may be a ROM (Read Only Memory), a RAM (Random Access Memory), a magnetic tape, a magnetic disk, a flash memory, an optical data storage device and the like. The computer-readable recording medium may also be distributed and executed in a computer system connected to a computer network and stored and executed as a code that can be read in a distributed manner.

It will be apparent to those skilled in the art that various modifications and variations can be made in the present invention without departing from the spirit or scope of the invention as defined by the appended claims. It will be understood by those skilled in the art that various changes in form and details may be made therein without departing from the spirit and scope of the invention.

In particular, the described features may be implemented within digital electronic circuitry, or computer hardware, firmware, or combinations thereof. The features may be implemented in a computer program product embodied in a storage device in a machine-readable storage device, for example, for execution by a programmable processor. And the features may be performed by a programmable processor executing a program of instructions for performing the functions of the described embodiments by operating on input data and generating an output. The described features include at least one programmable processor, at least one input device, and at least one output device, coupled to receive data and directives from a data storage system and to transmit data and directives to a data storage system, Such as a computer-readable recording medium. A computer program includes a set of directives that can be used directly or indirectly within a computer to perform a particular operation on a given result. A computer program may be written in any form of programming language including compiled or interpreted languages and may be implemented as a module, element, subroutine, or other unit suitable for use in other computer environments, or as a standalone program Can be used.

Suitable processors for execution of the program of instructions include, for example, both general purpose and special purpose microprocessors, and one of multiple processors of a single processor or other type of computer. Also, storage devices suitable for implementing the computer program instructions and data embodying the described features may be embodied in a computer-readable medium, such as, for example, semiconductor memory devices such as EPROM, EEPROM, and flash memory devices, Devices, magneto-optical disks, and non-volatile memory including CD-ROM and DVD-ROM disks. The processor and memory may be integrated within ASICs (application-specific integrated circuits) or added by ASICs.

While the invention has been shown and described with reference to certain preferred embodiments thereof, it will be understood by those of ordinary skill in the art that various changes in form and details may be made therein without departing from the spirit and scope of the invention as defined by the appended claims. It will be apparent to one skilled in the art to which the present invention pertains.

The combination of the above-described embodiments is not limited to the above-described embodiments, and various combinations and combinations of the above-described embodiments as well as the implementation and / or the necessity may be provided.

In the above-described embodiments, the methods are described on the basis of a flowchart as a series of steps or blocks, but the present invention is not limited to the order of steps, and some steps may occur in different orders or in a different order than the steps described above have. It will also be understood by those skilled in the art that the steps depicted in the flowchart illustrations are not exclusive and that other steps may be included or that one or more steps in the flowchart may be deleted without affecting the scope of the invention You will understand.

The foregoing embodiments include examples of various aspects. While it is not possible to describe every possible combination for expressing various aspects, one of ordinary skill in the art will recognize that other combinations are possible. Accordingly, it is intended that the invention include all alternatives, modifications and variations that fall within the scope of the following claims.

It will be apparent to those skilled in the art that various modifications and variations can be made in the present invention without departing from the spirit or scope of the inventions as defined by the following claims It will be understood that various modifications and changes may be made thereto without departing from the spirit and scope of the invention.

300: cache bypassing device
310: Information extracting unit
320: Switching point searching unit
330: Switching point determining unit
340: code counting
350: Program Compiler
360: Cache bypassing section

Claims (20)

CLAIMS 1. A method for cache bypassing at a thread group level performed by a cache bypassing device,
Determining at least one switching point for each thread group using the cache; And
Bypassing a cache used by the thread group according to the determined at least one switching point. The method according to claim 1,
Further comprising extracting memory access information generated when accessing the provided memory and cache bypassing information generated when the cache is bypassed,
Wherein the step of determining the switching point determines at least one switching point for each thread group using the extracted memory access information and the extracted cache bypassing information. The method according to claim 1,
Further comprising: retrieving at least one switchable point for each thread group. The method of claim 3,
Wherein the step of searching for the switchable point comprises:
Wherein at least one switchable point is searched for at least one point at which the thread cached data is not reused. The method of claim 3,
Wherein the step of searching for the switchable point comprises:
A thread group-level cache bypassing method for searching at least one switchable point according to the number of threads in the thread block divided by the total number of thread groups divided by the number of thread groups simultaneously using the cache. The method of claim 3,
Wherein determining the switching point comprises:
And determining at least one switching point for each thread group among the searched at least one switchable point. The method of claim 3,
Wherein determining the switching point comprises:
Calculating an execution time difference between a thread group using a cache and a thread group not using a cache among the searched at least one switchable point and calculating at least one switchable point at which the calculated execution time gap is minimized, A method for cache bypassing at a thread group level that determines a switching point. The method of claim 3,
Wherein determining the switching point comprises:
And determining at least one switchable point at which a difference in distance between switching points in the retrieved at least one switchable point is minimized as at least one switching point. The method according to claim 1,
Modifying the program code based on the determined switching point; And
Further comprising the step of generating a cache bypassing program by compiling the modified program code,
Wherein the bypassing of the cache bypasses a cache used by the thread group according to the generated cache bypassing program. 10. The method of claim 9,
Wherein the modifying the program code comprises:
And modifying the program code using a first cache operator for bypassing the thread group and a second cache operator for caching the thread group based on the determined switching point. A switching point determining unit for determining at least one switching point for each thread group using the cache; And
And a cache bypassing unit for bypassing a cache used by the thread group according to the determined at least one switching point. 12. The method of claim 11,
Further comprising an information extracting unit for extracting memory access information generated when accessing the provided memory and cache bypassing information generated when the cache is bypassed,
Wherein the switching point determination unit determines at least one switching point for each thread group using the extracted memory access information and the extracted cache bypassing information. 12. The method of claim 11,
And a switching point searching unit for searching at least one switchable point for each thread group. 14. The method of claim 13,
The switching point searching unit,
Wherein at least one switchable point is located at least one point at which the thread cached data is not reused to set the at least one switchable point. 14. The method of claim 13,
The switching point searching unit,
A thread group-level cache bypassing apparatus for searching at least one switchable point according to the number of threads in the thread block divided by the total number of thread groups divided by the number of thread groups simultaneously using the cache. 14. The method of claim 13,
The switching point determination unit may determine,
And determines at least one switching point for each thread group among the searched at least one switchable point. 14. The method of claim 13,
The switching point determination unit may determine,
Calculating an execution time difference between a thread group using a cache and a thread group not using a cache among the searched at least one switchable point and calculating at least one switchable point at which the calculated execution time gap is minimized, A thread group level cache bypassing device that determines the switching point. 14. The method of claim 13,
The switching point determination unit may determine,
And determines at least one switchable point at which the difference in distance between switching points in the retrieved at least one switchable point is minimized as at least one switching point. 12. The method of claim 11,
A code correcting unit for correcting the program code based on the determined switching point; And
Further comprising a program compiler for compiling the modified program code to generate a cache bypassing program,
Wherein the cache bypassing unit bypasses a cache used by the thread group according to the generated cache bypassing program. 20. The method of claim 19,
The code modifying unit,
And modifies the program code using a first cache operator for bypassing the thread group and a second cache operator for caching the thread group based on the determined switching point.

Images