Offload server, offload control method, and offload program

Abstract

An offload server (1) is provided with: a requested processing load analysis unit (120) that analyzes a requested processing load based on a current offload mode of data used by a user; a representative data selection unit (121) that specifies an application for which the processing load for analysis is higher, and selects representative data from request data when the application is used; a representative data selection unit (121) that specifies an application for which the processing load for analysis is higher, and selects representative data from request data when the application is used; an improvement degree calculation unit (122) that creates a new offload pattern on the basis of the selected representative data, compares the processing time and use frequency of the created new offload pattern with the processing time and use frequency of the current offload pattern, and calculates a performance improvement effect; and a reconstruction suggestion unit (123) that suggests GPU reconstruction when the performance improvement effect is equal to or greater than a predetermined threshold value.

Description

Technical Field

[0001] The present invention relates to an offloading server capable of automatically offloading functional processing to a GPU (Graphics Processing Unit), an offloading control method, and an offloading program. Background Art

[0002] The use of heterogeneous computing resources other than a CPU (Central Processing Unit) is increasing. For example, image processing has started to be performed using a server with an enhanced GPU (accelerator), or signal processing has been accelerated using an FPGA (accelerator). An FPGA is a programmable gate array whose structure can be set by a designer or the like after manufacturing, and is a type of PLD (Programmable Logic Device). In the Amazon Web Services (AWS) (registered trademark) cloud platform, GPU instances and FPGA instances are provided, and these resources can also be used on demand. Microsoft (registered trademark) uses an FPGA to improve search efficiency.

[0003] In an OpenIoT (Internet of Things) environment, by using service collaboration technology or the like, it is expected to create various types of application programs, and furthermore, by making full use of advanced hardware, it is possible to expect an improvement in the performance of the running application programs. However, for this purpose, programs and settings that match the operating hardware are required. A large amount of technical knowledge such as CUDA (Compute Unified Device Architecture) and OpenCL (Open Computing Language) is required, which is difficult. OpenCL is an open API (Application Programming Interface) that can uniformly process all computing resources (not limited to a CPU or a GPU) without being restricted by specific hardware.

[0004] In order to easily use a GPU and an FPGA in a user's IoT application program, there are the following requirements. That is, when deploying general application programs such as running image processing and encryption processing to an OpenIoT environment, it is desired that the OpenIoT platform analyzes the application logic and automatically offloads the processing to the GPU and the FPGA.

[0005] The development environment CUDA for GPGPU (General Purpose GPU), which uses the computing power of GPUs outside of image processing, is evolving. CUDA is a development environment for GPGPUs. In addition, OpenCL has emerged as a standard specification for unified processing of heterogeneous hardware such as GPUs, FPGAs, and many-core CPUs.

[0006] In CUDA and OpenCL, programs are based on extensions of the C language. However, it is necessary to describe memory copying, release, etc. between devices such as GPUs and CPUs, and the description is difficult. In fact, the number of technicians who are good at using CUDA and OpenCL is not large.

[0007] To simply perform processing using GPGPUs, there is such a technology: the position for parallel processing in response to loop statements, etc. is specified with basic instructions, and according to the instructions, the compiler converts it into device-oriented code. Technical specifications include OpenACC (Open Accelerator), etc., and compilers include the PGI compiler (registered trademark), etc. For example, in an example using OpenACC, the user specifies parallel processing, etc. to be executed with OpenACC instructions in code written in C / C++ / Fortran languages. The PGI compiler checks the parallelizability of the code, generates GPU and CPU execution binary codes, and modularizes the execution. IBM JDK (registered trademark) supports the function of offloading parallel processing in the form of Java (registered trademark) compliant lambdas to the GPU. By using these technologies, programmers do not need to consciously allocate data to GPU memory, etc.

[0008] In this way, through technologies such as OpenCL, CUDA, and OpenACC, offloading processing to GPUs and FPGAs can be performed.

[0009] However, even if the offloading processing itself can be performed, there are many problems for appropriate offloading. For example, there is a compiler such as the Intel compiler (registered trademark) that has an automatic parallelization function. During automatic parallelization, parallel processing parts such as for statements (repetitive statements) in the program are extracted. However, when operating in parallel using a GPU, due to the data exchange overhead between CPU-GPU memories, the performance is often not high. When using GPU acceleration, skilled personnel need to explore appropriate parallel processing parts using tuning under OpenCL and CUDA, the PGI compiler, etc.

[0010] Therefore, it is difficult for unskilled users to use GPUs to improve the performance of applications. Even when using automatic parallelization techniques, it takes a lot of time to trial and error optimize whether to parallelize for-loops before starting to use them.

[0011] As measures to automate the trial and error of parallel processing positions, Non-Patent Documents 1 and 2 can be cited. In Non-Patent Documents 1 and 2, by using evolutionary computing methods to repeatedly measure performance in a verification environment, loop statements suitable for GPU offloading are appropriately extracted, and variables within nested loop statements are transferred between the CPU and GPU as uniformly as possible in the uppermost loop to achieve automatic acceleration.

[0012] In addition, in FPGAs, long compilation times are required and multiple measurements cannot be performed. Therefore, a method has been proposed to reduce the arithmetic intensity of loop statements and the FPGA resource usage rate during offloading to select candidate loop statements, then OpenCL them and perform measurements to explore appropriate patterns (Non-Patent Document 2).

[0013] As for GPU offloading of loop statements, as a measure to automate the exploration of GPU processing positions for loop statements, offloading using the genetic algorithm (GA: Genetic Algorithm), which is an evolutionary computing method, has been proposed (Non-Patent Document 3).

[0014] Prior Art Documents

[0015] Non-Patent Documents

[0016] Non-Patent Document 1: Y.Yamato, T.Demizu, H.Noguchi and M.Kataoka, “Automatic GPUOffloading Technology for Open IoT Environment,” IEEE Internet of ThingsJournal, Sep. 2018.

[0017] Non-Patent Document 2: Y.Yamato, “Study of parallel processing area extractionand data transfer number reduction for automatic GPU offloading of IoTapplications,” Journal of Intelligent Information Systems, Springer, DOI: 10.1007 / s10844-019-00575-8, Aug. 2019.

[0018] Non-Patent Document 3: Y. Yamato, "Study and Evaluation of Improved Automatic GPU Offloading Method," International Journal of Parallel, Emergent and Distributed Systems, Taylor and Francis, DOI: 10.1080 / 17445760.2021.1941010, June 2021. Summary of the Invention

[0019] Problems to be Solved by the Invention

[0020] In the prior art, the offloading of heterogeneous hardware mainly relies on manual operations, and the automatic offloading methods are limited to the technologies described in Non-Patent Documents 1 and 2.

[0021] In Non-Patent Documents 1 to 3, regarding the concept of environment-adaptive software, the automatic offloading methods of GPUs and FPGAs such as loop statements are verified.

[0022] However, in the automatic offloading methods described in Non-Patent Documents 1 to 3, it is premised on performing adaptation processes such as conversion and configuration before the application starts to be used, and no reconstruction is envisaged according to changes in usage characteristics after the start of use. That is, Non-Patent Documents 1 to 3 are all technologies before the application starts to be used, and no research has been conducted on the reconstruction after the start of use.

[0023] For example, there is a way of constructing the logic to perform GPU processing on the premise that there is a lot of category classification processing in image processing before the start of use. However, in the actual request count three months after the start of use, there is a lot of object detection processing in image processing, and the logic for performing GPU processing should be changed, etc.

[0024] In this way, there is a problem that no reconstruction is envisaged according to changes in usage characteristics after the start of use.

[0025] The present invention is completed in view of such points, and its object is to reconstruct into a more appropriate logic not only before the start of use but also according to the usage characteristics after the start of use, so as to achieve efficient resource utilization in GPUs and the like with limited resource amounts.

[0026] Solutions to Solve the Problems

[0027] In order to solve the above problems, the present invention provides an offloading server that offloads specific processing of an application program to a GPU, and is characterized by having: a processing load analysis unit that analyzes the request processing load of data used by a user based on a current offloading mode; a representative data selection unit that determines an application program with a higher request processing load analyzed by the processing load analysis unit, and selects representative data from the request data when using this application program; an improvement degree calculation unit that creates a new offloading mode based on the representative data selected by the representative data selection unit, compares the processing time and usage frequency of the created new offloading mode with the processing time and usage frequency of the current offloading mode, and calculates the performance improvement effect; and a reconstruction recommendation unit that recommends GPU reconstruction when the performance improvement effect is equal to or higher than a specified threshold.

[0028] Advantages of the Invention

[0029] According to the present invention, it is possible to reconstruct into a more appropriate logic according to the usage characteristics after the start of operation, and it is possible to achieve efficient resource utilization in a GPU or the like with limited resources. Brief Description of the Drawings

[0030] Figure 1 A diagram showing an environment-adaptive software system of an offloading server including an embodiment of the present invention.

[0031] Figure 2 A functional block diagram showing a structural example of the offloading server according to the above embodiment.

[0032] Figure 3 A diagram showing the automatic offloading process of the offloading server according to the above embodiment.

[0033] Figure 4 A diagram showing an exploration image of a control unit (automatic offloading function unit) based on a simple genetic algorithm (Simple GA) of the offloading server according to the above embodiment.

[0034] Figure 5 A diagram showing an example of a conventional CPU program of a comparative example.

[0035] Figure 6 A diagram showing an example of a loop statement showing the case of transferring data from a CPU to a GPU using a simple CPU program of a comparative example.

[0036] Figure 7 A diagram showing an example of a loop statement showing the case of transferring data from a CPU to a GPU in the case of nested integration of the offloading server according to the above embodiment.

[0037] Figure 8A diagram showing an example of a loop statement for transferring data from the CPU to the GPU in the case of transmission integration of the offloading server according to the above-described embodiment.

[0038] Figure 9 A diagram showing an example of a loop statement for transferring data from the CPU to the GPU in the case of transmission integration of the offloading server according to the above-described embodiment and using a temporary area.

[0039] Figure 10A A flowchart for explaining the outline of the operations implemented by the offloading server according to the above-described embodiment.

[0040] Figure 10B A flowchart for explaining the outline of the operations implemented by the offloading server according to the above-described embodiment.

[0041] Figure 11 A flowchart showing the reconstruction after the offloading server according to the above-described embodiment starts operation.

[0042] Figure 12 A detailed flowchart showing the analysis process of commercial request data records of the offloading server according to the above-described embodiment.

[0043] Figure 13 A detailed flowchart showing the improvement degree calculation process of the offloading server according to the above-described embodiment.

[0044] Figure 14 A hardware structure diagram showing an example of a computer that implements the functions of the offloading server according to the above-described embodiment. Detailed implementation mode

[0045] Next, the offloading server 1 and the like in the mode for implementing the present invention (hereinafter referred to as "the present embodiment") will be described.

[0046] (Background description)

[0047] The application programs to be offloaded are diverse. In addition, in application programs with a large amount of calculation and time consumption, such as image analysis for image processing and machine learning processing for analyzing sensor data, the repeated processing of loop statements takes a long time. Therefore, it is considered to aim at accelerating by automatically offloading loop statements to the GPU.

[0048] First, the basic problems in the case of automatically offloading loop statements to the GPU are as follows. That is, currently, even if the compiler can detect the limitation that "this loop statement cannot be parallelized on the GPU", it is difficult to detect the suitability that "this loop statement is suitable for parallel processing on the GPU". In addition, in the case of automatically offloading loop statements to the GPU, it is generally considered that loops with a high computational density such as a large number of loop iterations are more suitable. However, in fact, it is difficult to predict to what extent the performance will be improved without actual testing. Therefore, the instruction to offload loop statements to the GPU is manually performed, and trial and error is carried out on the performance measurement.

[0049] Based on the above problems, Non-Patent Document 1 proposed to automatically discover loop statements suitable for offloading to the GPU using a genetic algorithm (GA: Genetic Algorithm). That is, in Non-Patent Document 1, for a general program not envisioned for parallelization, first, the check for parallelizable loop statements is performed. Then, for the group of parallelizable loop statements, they are gene-ized in such a way that the value is 1 when executed on the GPU and 0 when executed on the CPU, and performance verification is repeatedly attempted in a verification environment to explore an appropriate area. Based on reducing to parallelizable loop statements, in the form of gene parts, the parallel processing patterns that can be accelerated are maintained and recombined, so as to efficiently explore the accelerable patterns based on the huge parallel processing patterns that can be obtained.

[0050] In Non-Patent Document 1, for variables used in nested loop statements, when offloading the loop statements to the GPU, data transfer between the CPU and the GPU is performed. However, when performing CPU-GPU transfer at a lower level in the nesting, transfer is required every time in the lower-level loop, which is inefficient.

[0051] Non-Patent Document 2 proposed to perform the transfer at a higher level for variables for which there is no problem with CPU-GPU transfer at a higher level. Most loops with a large number of loop iterations that consume processing time are nested, so this method shows a certain effect in acceleration based on reducing the number of transfers.

[0052] In Non-Patent Documents 1 and 2, it was confirmed that in fact, even medium-scale application programs with loop statements greater than 100 are automatically accelerated. Considering practicality, further acceleration is required.

[0053] In Non-Patent Document 3, based on Non-Patent Documents 1 and 2, it was confirmed that more practical improvements such as being able to further reduce CPU-GPU transfer and expand to loop types that can be specified for GPU processing can be achieved.

[0054] (Basic idea)

[0055] The basic idea of the present invention will be described.

[0056] First, the environment-adaptive software is described.

[0057] [Environment-adaptive software]

[0058] The environment-adaptive software executed by the offloading server of the present invention has the following characteristics. That is, through the execution of the environment-adaptive software, the offloading server automatically performs conversions, resource settings, configuration decisions, etc. in such a way that the already described program code can utilize the GPUs, FPGAs, multi-core CPUs, etc. of the environment at the configuration destination, enabling the application program to operate with high performance. As elements of the environment-adaptive software, there are methods of automatically offloading the loop statements and functional blocks of the code to GPUs and FPGAs, and methods of appropriately allocating the processing resource amounts of GPUs and the like.

[0059] In the case of GPUs, the genetic algorithm is used to perform combinations on loop statements that can be processed by GPUs, and measurements on a scale of 1000 times are carried out to explore the optimal mode. In the case of FPGAs, since one compilation takes more than 6 hours, in the offloading of loop statements, the loop statements with high arithmetic intensity, high loop count, and high resource efficiency are reduced, and an offloading mode is created for measurement to perform fast mode exploration.

[0060] [Reconfiguration after the application program starts to operate]

[0061] An explanation is given for the reconfiguration after the application program starts to operate.

[0062] In the automatic offloading methods described in Non-Patent Documents 1 and 2, it is premised on performing adaptation processes such as conversion and configuration before the application program starts to operate.

[0063] It should be noted that, except for special uses such as the circuit reconfiguration of artificial satellites, in the application of using GPUs for the acceleration of application servers, looking at commercial clouds (such as GPU instances of Amazon Web Services (AWS) (registered trademark)), there are no examples of reconfiguring GPU logic according to usage characteristics during the operation of the application program ("during operation" is a way of "after starting to operate"). This is because it is difficult to reconfigure GPU logic during operation.

[0064] The present invention reconfigures the GPU offloading logic (hereinafter referred to as GPU logic) according to changes in usage characteristics, etc. after the application program starts to operate. For example, the GPU logic is reconfigured according to usage characteristics during operation.

[0065] The present invention first offloads the normal CPU-oriented program to the GPU and starts to operate (<adaptation processes such as conversion and configuration before starting to operate>).

[0066] Next, the request characteristics are analyzed, and the GPU logic is changed in another program (<reconfiguration after starting to operate>).

[0067] (Embodiment)

[0068] Figure 1 FIG. for showing an environment-adaptive software system including the offloading server 1 of the present embodiment.

[0069] The environment-adaptive software system of the present embodiment is characterized in that, in addition to the structure of the existing environment-adaptive software, it further includes an offloading server 1. The offloading server 1 is an offloading server that offloads specific processing of an application program to an accelerator. In addition, the offloading server 1 is communicably connected to each device located in the three layers of the cloud layer 2, the network layer 3, and the device layer 4. A data center 30 is provided in the cloud layer 2, a network edge 20 is provided in the network layer 3, and a gateway 10 is provided in the device layer 4.

[0070] Therefore, in the environment-adaptive software system including the offloading server 1 of the present embodiment, in each of the device layer, the network layer, and the cloud layer, efficiency is achieved by appropriately performing function configuration and processing offloading. It mainly aims at the efficiency of function configuration by configuring functions at appropriate positions in the three layers and performing processing, and the efficiency achieved by offloading function processing such as image analysis to heterogeneous hardware such as GPUs and FPGAs.

[0071] Hereinafter, a structural example of the offloading server 1 of the present embodiment when performing offloading processing in the background of service utilization for users in the environment-adaptive software system will be described.

[0072] It is assumed that when providing a service, the service is provided to a user in the form of a trial or the like on the first day, and offloading processing such as image analysis is performed in the background. After the second day, the image analysis is offloaded to a GPU, and the service can be provided at a reasonable price.

[0073] Figure 2 Functional block diagram showing a structural example of the offloading server 1 of the embodiment of the present invention.

[0074] The offloading server 1 is a device that executes environment-adaptive software processing. As one mode of this environment-adaptive software, the offloading server 1 automatically offloads specific processing of an application program to an accelerator (<automatic offloading>).

[0075] In addition, the offloading server 1 can be connected to an emulator.

[0076] As Figure 2 shown, the offloading server 1 is configured to include a control unit 11, an input / output unit 12, a storage unit 13, and a verification machine 14 (Verification machine) (device for accelerator verification).

[0077] The control unit 11 is responsible for the overall control of the offloading server 1 and executes reconstruction control after the application starts running. In a commercial environment, another OpenACC is started for reconstruction (described later). Furthermore, when the control unit 11 processes other servers (not shown) of OpenACC after GPU reconstruction and when the GPU performs reconstruction, it stops the OpenACC before reconstruction, newly starts other servers, and switches the routing to other servers for new requests.

[0078] The input / output unit 12 is composed of a communication interface for transmitting and receiving information between various devices, etc., and an input / output interface for transmitting and receiving information between input devices such as touch panels and keyboards and output devices such as monitors.

[0079] The storage unit 13 is composed of a hard disk, a flash memory, a RAM (Random Access Memory), etc.

[0080] In this storage unit 13, a code mode DB 131, a device resource DB 132, and a test case DB (Test case database) 133 are stored, and a program (offloading program) for executing each function of the control unit 11 and information required for the processing of the control unit 11 (for example, an intermediate language file (Intermediate file) 134) are temporarily stored.

[0081] Performance test items are stored in the test case DB 133. The test case DB 133 stores information for testing the performance of an application program for measuring acceleration. For example, if it is a deep learning application program for image analysis processing, it is a sample image and a test item for executing it.

[0082] The verification machine 14 has a CPU (Central Processing Unit), a GPU, and an FPGA (accelerator) as a verification environment for environment-adaptive software.

[0083] In addition, in Figure 2 the offloading server 1 is configured to have the verification machine 14, but the verification machine 14 can also be outside the offloading server 1.

[0084] The control unit 11 is an automatic offloading function unit (Automatic Offloading function) that undertakes the overall control of the offloading server 1. The control unit 11 is realized by, for example, a CPU (not shown) expanding and executing a program (offloading program) stored in the storage unit 13 in the RAM.

[0085] The control unit 11 includes: an application code specifying unit (Specify application code) 111, an application code analyzing unit (Analyze application code) 112, a data transfer specifying unit 113, a parallel processing specifying unit 114, a parallel processing mode creating unit 115, a performance measuring unit 116, an execution file creating unit 117, a formal environment configuring unit (Deployfinal binary files to production environment) 118, a performance measurement test extraction and execution unit (Extract performance test cases and run automatically) 119, a request processing load analyzing unit 120 (processing load analyzing unit), a representative data selecting unit 121, an improvement degree calculating unit 122, a refactoring suggestion unit 123, and a user providing unit (Provide price and performance to a user to judge) 124.

[0086] <Application code specifying unit 111>

[0087] The application code specifying unit 111 specifies the input application code. Specifically, the application code specifying unit 111 determines the processing functions (such as image analysis) of the services provided to the user.

[0088] <Application code analyzing unit 112>

[0089] The application code analyzing unit 112 analyzes the source code of the processing functions to grasp structures such as loop statements and FFT library calls.

[0090] <Data transfer specifying unit 113>

[0091] The data transfer specifying unit 113 analyzes the reference relationships of the variables used in the loop statements of the application program, and specifies data transfer for the data that can be transferred outside the loop. This data transfer specification uses an explicit specification line (#pragma acc data copyin / copyout / copy(a[…]) where the variable a) that explicitly specifies data transfer outside the loop.

[0092] The data transfer specifying unit 113 specifies data transfer using a specifying line that clearly specifies data transfer from the CPU to the GPU (#pragma acc data copyin(a[…])), a specifying line that clearly specifies data transfer from the GPU to the CPU (#pragma acc data copyout(a[…])), and a specifying line that uniformly and clearly specifies round-trip data copying when data transfer from the CPU to the GPU and data transfer from the GPU to the CPU occur simultaneously (#pragma acc datacopy(a[…])).

[0093] When variables defined on the CPU program side and variables referenced on the GPU program side overlap, the data transfer specifying unit 113 instructs data transfer from the CPU to the GPU, and sets the position for specifying data transfer to the outermost loop in a loop statement for GPU processing or a loop statement at a higher level that does not include the setting or definition of the corresponding variable. Further, when variables set on the GPU program side and variables referenced on the CPU program side overlap, the data transfer specifying unit 113 instructs data transfer from the GPU to the CPU, and sets the position for specifying data transfer to the outermost loop in a loop statement for GPU processing or a loop statement at a higher level that does not include the reference, setting, or definition of the corresponding variable.

[0094] Here, based on the result of code analysis, the data transfer specifying unit 113 specifies GPU processing of a loop statement using at least one selected from the kernels directive, parallel loop directive, and parallel loop vector directive of OpenACC.

[0095] The kernels directive of OpenACC is used for single loops and tightly nested loops.

[0096] The parallel loop directive of OpenACC is used for non-tightly nested loops.

[0097] The parallel loop vector directive of OpenACC is used for loops that cannot be parallelized but can be vectorized.

[0098] <Parallel processing specifying unit 114>

[0099] The parallel processing specifying unit 114 determines the loop statements (repetitive statements) of the application program, and for each determined loop statement, creates multiple offloading processing modes specified by OpenCL for pipeline processing and parallel processing in the GPU for compilation.

[0100] The parallel processing specifying unit 114 includes an offloadable area extraction unit (Extract offloadable area) 114a and an intermediate language file output unit (Output intermediate file) 114b.

[0101] The offloadable area extraction unit 114a determines the processing that can be offloaded to the GPU for loop statements, FFT (Fast Fourier Transformation), etc., and extracts the intermediate language corresponding to the offload processing.

[0102] The intermediate language file output unit 114b outputs the extracted intermediate language file 134. The intermediate language extraction does not end in one go and is repeated for exploring appropriate offload areas and for trial optimization of execution.

[0103] <Parallel processing mode creation unit 115>

[0104] The parallel processing mode creation unit 115 creates such a parallel processing mode (GPU processing mode) that excludes loop statements (repetitive statements) that cause compilation errors from the offload targets, and specifies whether to perform parallel processing (GPU processing) for repetitive statements that do not cause compilation errors.

[0105] <Performance measurement unit 116>

[0106] The performance measurement unit 116 compiles the application of the created parallel processing (GPU processing) mode, configures it in the verification machine 14, and performs the performance measurement processing when offloading to the GPU.

[0107] The performance measurement unit 116 includes a binary file deployment unit (Deploy binary files) 116a. The binary file deployment unit 116a deploys (configures) the execution files exported from the intermediate language in the verification machine 14 having a GPU and an FPGA.

[0108] The performance measurement unit 116 executes the configured binary files, measures the performance during offloading, and returns the performance measurement result to the offloadable area extraction unit 114a. In this case, the offloadable area extraction unit 114a extracts another parallel processing mode, and the intermediate language file output unit 114b performs trial performance measurement based on the extracted intermediate language (refer to the reference numeral aa described later Figure 3 ).

[0109] <Execution file creation unit 117>

[0110] The execution file creation unit 117 selects the parallel processing mode with the highest processing performance from multiple PLD processing modes based on the performance measurement results repeated a specified number of times, and compiles and creates an execution file for the PLD processing mode with the highest processing performance.

[0111] <Formal environment configuration unit 118>

[0112] The formal environment configuration unit 118 configures the created execution file in the formal environment for users ("configuration of the final binary file in the formal environment"). The formal environment configuration unit 118 determines the mode specifying the final uninstallation area and deploys it in the formal environment for users.

[0113] <Performance measurement test extraction and execution unit 119>

[0114] After the execution file is configured, the performance measurement test extraction and execution unit 119 extracts performance test items from the test case DB 133 and executes performance tests.

[0115] After the execution file is configured, in order to display the performance to the user, the performance measurement test extraction and execution unit 119 extracts performance test items from the test case DB 133 and automatically executes the extracted performance tests.

[0116] <Request processing load analysis unit 120>

[0117] The request processing load analysis unit 120 analyzes the request processing load based on the current uninstallation mode of the commercial representative data (the data actually used by users).

[0118] The request processing load analysis unit 120 calculates the actual processing time and the total number of uses from the usage records of each application program in a fixed period (a long time; for example, one month, etc.). Specifically, the request processing load analysis unit 120 calculates the improvement factor based on the test records of the assumed usage data before operation as (actual processing time when only CPU is processed) / (actual processing time when GPU offloading is performed). Here, the sum of the values obtained by multiplying the actual processing time by the improvement factor is used as the total processing time for comparison. The request processing load analysis unit 120 compares the total actual processing time among all application programs. Then, the request processing load analysis unit 120 reorders according to the total actual processing time and determines multiple application programs with the highest processing time load.

[0119] The request processing load analysis unit 120 obtains the request data of the application programs with the highest load in a fixed period, arranges the data size at fixed intervals, and makes a frequency distribution.

[0120] <Representative data selection unit 121>

[0121] The representative data selection unit 121 selects representative data from the request data used by the application program when the processing load analyzed by the request processing load analysis unit 120 is high among the application programs. Specifically, the representative data selection unit 121 selects any one of the actual request data that conforms to the mode of the data size frequency distribution analyzed by the request processing load analysis unit 120 and designates it as the representative data.

[0122] <Improvement degree calculation unit 122>

[0123] The improvement degree calculation unit 122 determines a new offloading mode (a new offloading mode discovered in the verification environment) by executing the application program code analysis unit 112, the parallel processing designation unit 114, the parallel processing mode creation unit 115, the performance measurement unit 116, and the execution file creation unit 117 based on the representative data selected by the representative data selection unit 121, and compares the processing time and usage frequency of the determined new offloading mode with those of the current offloading mode to calculate the performance improvement effect.

[0124] The improvement degree calculation unit 122 measures the processing time of the current offloading mode and multiple new offloading modes, and calculates the performance improvement effect based on the commercial usage frequency according to (actual processing reduction time in the verification environment) × (usage frequency in the commercial environment). That is, the improvement degree calculation unit 122 calculates (actual processing reduction time in the verification environment) × (usage frequency in the commercial environment) for the current offloading mode under the commercial representative data: formula (1), and calculates (actual processing reduction time in the verification environment) × (usage frequency in the commercial environment) for multiple new offloading modes: formula (2). In addition, the calculation of formula (1) and formula (2) can also be performed by the reconstruction suggestion unit 123.

[0125] Here, in the case of offloading a normal CPU-oriented program to the GPU and starting operation (adaptation processing such as conversion and configuration before starting operation), and in the case of analyzing the request characteristics and changing the GPU logic in another program (reconstruction after starting operation), specific examples of the parameters of the improvement degree calculation unit 122 are as follows.

[0126] The improvement degree calculation unit 122 extracts, through measurement in the verification environment, offloading modes that accelerate the test cases of the commercial representative data in multiple application programs with high loads. Specifically, the improvement degree calculation unit 122 counts the number of for statements by syntax parsing in the application programs with high loads, and discovers and removes for statements that cannot be processed by the GPU through the compiler function.

[0127] The improvement degree calculation unit 122 sets the remaining number of for statements as the gene length, creates a fixed number of gene patterns, makes 1 correspond to GPU execution and 0 correspond to non - execution, and appends OpenACC instructions. The instructions are both GPU calculations such as ￥#pragma acc kernels and data processing such as ￥#pragma acc data copy.

[0128] The improvement degree calculation unit 122 compiles the OpenACC file corresponding to the gene pattern using the verification environment machine, measures the performance through test cases of commercial representative data, and sets a higher fitness for faster patterns.

[0129] The improvement degree calculation unit 122 performs processes such as crossover according to the performance fitness to create the next - generation gene pattern. The improvement degree calculation unit repeats the genetic algorithm process for a certain number of generations and takes the final highest - performance pattern as the solution.

[0130] <Restructuring recommendation unit 123>

[0131] The restructuring recommendation unit 123 determines the restructuring recommendation based on whether the performance improvement effect of the new offloading mode is above the specified threshold of the current offloading mode. Specifically, the restructuring recommendation unit 123 confirms whether "the calculation result of formula (2) / the calculation result of formula (1)" calculated by the improvement degree calculation unit 122 is above the specified threshold. If "the calculation result of formula (2) / the calculation result of formula (1)" is above the specified threshold, it recommends restructuring; if it is less than the specified threshold, it does nothing (no restructuring recommendation).

[0132] <User providing unit 124>

[0133] The user providing unit 124 provides the user with information such as price - performance based on the performance test results ("provides the user with price - performance and other information"). The test data for automatically performing tests is stored in the test case DB133, and this test measures the performance of the application. The user providing unit 124 prompts the user with the results of the test data for which the test case DB133 has been executed and the overall price of the system determined by the unit prices of the resources (virtual machines, FPGA instances, GPU instances, etc.) provided by the system. The user judges whether to start paying for the service based on the prompted price - performance and other information.

[0134] [Application of genetic algorithm]

[0135] The offloading server 1 can use the genetic algorithm in the optimization of offloading. The structure of the offloading server 1 in the case of using the genetic algorithm is as described below.

[0136] That is, the parallel processing specifying unit 114 sets the number of loop statements (repetitive statements) without compilation errors as the gene length based on a genetic algorithm. The parallel processing mode creation unit 115 maps the availability of accelerator processing to the gene pattern by setting the case of performing accelerator processing as either 1 or 0, and the case of not performing it as the other of 0 or 1.

[0137] The parallel processing mode creation unit 115 prepares gene patterns of a specified number of individuals with each value of the gene randomly created as 1 or 0. The performance measurement unit 116 compiles the application program code specifying the parallel processing statement in the accelerator according to each individual and configures it in the verification machine 14. The performance measurement unit 116 performs performance measurement processing in the verification machine 14.

[0138] Here, when genes of the same parallel processing mode are generated in the intermediate generation, the performance measurement unit 116 does not compile the application program code corresponding to the parallel processing mode and perform performance measurement, but uses the same value as the performance measurement value.

[0139] In addition, for the application program code that generates compilation errors and the application program code whose performance measurement does not end within the specified time, the performance measurement unit 116 sets the performance measurement value to the specified time (long time) as a timeout process.

[0140] The execution file creation unit 117 performs performance measurement on all individuals and evaluates them in such a way that the shorter the processing time of an individual, the higher the fitness. The execution file creation unit 117 selects individuals with a fitness higher than a specified value (for example, the top n% of all numbers, or the top m of all numbers, where n and m are natural numbers) from all individuals as high-performance individuals, and performs crossover and mutation processing on the selected individuals to create the next generation of individuals. After the processing of the specified number of generations is completed, the execution file creation unit 117 selects the parallel processing mode with the highest performance as the solution.

[0141] Next, the automatic unloading operation of the offloading server 1 configured as described above will be described.

[0142] The offloading server 1 of this embodiment is characterized by performing reconstruction after the FPGA logic starts operation. The <automatic offloading process> executed by the offloading server 1 as a form of environment-adaptive software is the same in the reconstruction before and after the start of operation. That is, the difference is that Figure 3 The automatic offloading process of the offloading server 1 shown is the same in the reconstruction before and after the start of operation, but the data processed before the start of operation is assumed usage data, while in the reconstruction after the start of operation, the data processed is the data used in actual commercial use (commercial representative data).

[0143] [Automatic offloading operation]

[0144] The offloading server 1 of the present embodiment is an example of an offloading server that automatically offloads function processing to a GPU or the like.

[0145] Figure 3 It is a diagram showing the automatic offloading process of the offloading server 1. Figure 3 The <automatic offloading process> is the same in the reconstruction before and after the start of operation.

[0146] As Figure 3 As shown, the offloading server 1 is applied to the basic technology of environment-adaptive software. The offloading server 1 includes: a control unit (automatic offloading function unit) 11 that executes environment-adaptive software processing, a code mode DB 131, a device resource DB 132, a test case DB 133, an intermediate language file 134, and a verification machine 14.

[0147] The offloading server 1 acquires the application program code (Application code) 130 used by the user.

[0148] The user uses OpenIoT Resources 15 as a commercial environment. The OpenIoT Resources 15 utilize, for example, various devices (Device) 151, a device 152 having a CPU-GPU, a device 153 having a CPU-FPGA, and a device 154 having a CPU. The offloading server 1 automatically offloads the processing function to the accelerators of the device 152 having a CPU-GPU and the device 153 having a CPU-FPGA.

[0149] The offloading server 1 executes environment-adaptive software processing by coordinating the platform functions composed of the code mode DB 131, the device resource DB 132, and the test case DB 133 based on the environment adaptation function of the commercial environment and the verification environment provided by the operator.

[0150] Hereinafter, with reference to Figure 3 the step numbers, the operations of each part of the automatic offloading operation (before the start of operation) will be described.

[0151] [Automatic offloading operation (before the start of operation)]

[0152] First, the conversion of the code required before the start of operation of the application program, the adjustment of the resource amount, the adjustment of the configuration position, and the verification are performed.

[0153] <Step S11: Specify application code: Application code specification>

[0154] In step S11, the application code specifying unit 111 (refer to Figure 2)Determine the processing function (such as image analysis) of the service provided by the user. Specifically, the application code specifying unit 111 specifies the input application code.

[0155] <Step S12: Analyze application code: Analysis of application code>

[0156] In step S12, the application code analysis unit 112 (refer to Figure 2 ) analyzes the source code of the processing function and grasps the structure such as loop statements and FFT library calls.

[0157] <Step S21: Extract offloadable area: Extraction of offloadable area>

[0158] In step S21, the parallel processing specifying unit 114 (refer to Figure 2 ) determines the loop statements (repetitive statements) of the application, and for each repetitive statement, specifies parallel processing or pipeline processing in the GPU and compiles it using a high-level synthesis tool. Specifically, the offload range extraction unit 114a (refer to Figure 2 ) determines the processing that can be offloaded to the FPGA, such as loop statements, and extracts OpenCL as the intermediate language corresponding to the offload processing.

[0159] <Step S22: Output intermediate file: Output of intermediate language file>

[0160] In step S22, the intermediate language file output unit 114b (refer to Figure 2 ) outputs the intermediate language file 134. The intermediate language extraction does not end at once. In order to explore an appropriate offload area and perform trial optimizations, it is repeated.

[0161] <Step S23: Compile error: Creation of parallel processing mode>

[0162] In step S23, the parallel processing mode creation unit 115 (refer to Figure 2 ) creates such a parallel processing mode (GPU processing mode) that excludes the loop statements with compilation errors from the offload targets, and for the repetitive statements without compilation errors, specifies whether to perform GPU processing.

[0163] <Step S31: Deploy binary files: Deployment of executable files>

[0164] In step S31, the binary file deployment unit 116a (refer to Figure 2)Deploy the executable file for intermediate language export on the verification machine 14 with a GPU. The binary file configuration unit 116a starts the configured file, executes the envisioned test cases, and measures the performance during uninstallation.

[0165] <Step S32: Measure performances: Performance measurement for retrieving an appropriate mode>

[0166] In step S32, the performance measurement unit 116 (refer to Figure 2 ) executes the configured file and measures the performance during uninstallation.

[0167] To make the uninstalled area more appropriate, the performance measurement result is returned to the uninstall range extraction unit 114a, and the uninstall range extraction unit 114a extracts another mode. Then, the intermediate language file output unit 114b performs a trial performance measurement based on the extracted intermediate language (refer to Figure 3 the reference numeral aa in the figure). The performance measurement unit 116 repeats the performance measurement in the verification environment and determines the code mode for final deployment.

[0168] As Figure 3 shown by the reference numeral aa in the figure, the control unit 11 repeatedly executes the above steps S12 to S23. The automatic uninstall function of the control unit 11 is summarized as follows. That is, the parallel processing designation unit 114 determines the loop statements (repetitive statements) of the application program, and for each repetitive statement, specifies parallel processing or pipeline processing in the GPU using OpenCL and compiles it using a high-level synthesis tool. Then, the parallel processing mode creation unit 115 creates such a parallel processing mode, excludes the loop statements with compilation errors from the uninstallation target, and for the loop statements without compilation errors, specifies whether to perform parallel processing. Then, the binary file configuration unit 116a compiles the application program corresponding to the parallel processing mode, configures it in the verification machine 14, and the performance measurement unit 116 performs the performance measurement processing using the verification machine 14. The executable file creation unit 117 selects the mode with the highest processing performance from multiple parallel processing modes based on the performance measurement results repeated a specified number of times, compiles the selected mode, and creates an executable file.

[0169] <Step S41: Determination of resource size>

[0170] The control unit 11 determines the resource size (refer to Figure 3 the reference numeral bb in the figure).

[0171] <Step S51: Selection of an appropriate configuration location>

[0172] The control unit 11 selects an appropriate configuration location with reference to the device resource DB132.

[0173] <Step S61: Deploy final binary files to production environment: Configuration of the production environment for the final files>

[0174] In Step S61, the formal environment configuration department 118 determines the mode of the designated final uninstallation area and deploys it in the formal environment facing users.

[0175] <Step S62: Extract performance test cases and run automatically: Extraction of test cases and confirmation of normality>

[0176] In Step S62, after the execution file is configured, the performance measurement test extraction and execution department 119 extracts performance test items from the test case DB 133 to display the performance to the user and automatically executes the extracted performance tests.

[0177] <Step S63: Provide price and performance to a user to judge: User prompt for judging the price and performance at the start of use>

[0178] In Step S63, the user provision department 124 prompts the user with information such as price and performance based on the performance test results. Based on the prompted price, performance and other information, the user judges to start paying to use the service.

[0179] The above Steps S11 to S63 are assumed to be carried out in the background of the user's service use, for example, within the first day of starting use, etc. In addition, the processing carried out in the background to reduce costs can be targeted only at GPU and FPGA offloading.

[0180] As described above, when the control unit (automatic offloading function unit) 11 of the offloading server 1 is applied to the basic technology of the environment adaptive software, in order to perform offloading function processing, it extracts the offloading area from the source code of the application program used by the user and outputs the intermediate language (Steps S11 to S23). The control unit 11 configures and executes the execution file for intermediate language export in the verification machine 14 to verify the offloading effect (Steps S31 to S32). After repeating the verification and determining the appropriate offloading area, the control unit 11 deploys the execution file in the formal environment actually provided to the user and provides it as a service (Steps S41 to S63).

[0181] In addition, in the above, the processing flow of code conversion, resource amount adjustment, and configuration position adjustment required for unified environment adaptation has been described, but it is not limited to this, and it is also possible to extract only the processing that one wants to perform. For example, in the case of only wanting to perform code conversion for the GPU, only the parts required for the environment adaptation function and verification environment in the above steps S11 to S63 need to be used.

[0182] <GPU Automatic Unloading>

[0183] The above code analysis uses a syntax parsing tool such as Clang to analyze the application code. Code analysis needs to analyze the device assumed to be unloaded, so it is difficult to generalize. Among them, it is possible to grasp the structure of the code such as loop statements and variable reference relationships, or to grasp the functional blocks that perform FFT processing as functional blocks, call libraries that perform FFT processing, etc. The judgment of functional blocks is difficult to be automatically judged by the unloading server. This can also be grasped by similarity determination using a similar code detection tool such as Deckard. Here, Clang is a tool for C / C++, and a tool that matches the parsed language needs to be selected.

[0184] In addition, in the case of the process of unloading an application, in the GPU, FPGA, IoT GW, etc., research that matches the unloading destination is required respectively. Generally, regarding performance, it is difficult to automatically find the setting that becomes the highest performance at once. Therefore, try the unloading mode by repeatedly measuring performance multiple times in the verification environment to find the mode that can be accelerated.

[0185] [GPU Automatic Unloading Using Genetic Algorithm]

[0186] An explanation is given for the automatic acceleration using the genetic algorithm. The automatic acceleration using the genetic algorithm is used in the automatic unloading operation (before starting operation).

[0187] GPU automatic unloading is used to repeat for the GPU Figure 3 Steps S21 to S23, S31 to S32, and finally obtain the unloading code deployed in steps S61 to S63.

[0188] The GPU is a device that generally does not guarantee latency but is suitable for improving throughput by using parallel processing. There are various applications for IoT operations. Representative ones include encryption processing of IoT data, image processing for camera image analysis, machine learning processing for analyzing a large amount of sensor data, etc., and they have a lot of repetitive processing. Therefore, acceleration is achieved by automatically unloading the repetitive statements of the application to the GPU.

[0189] However, as described in the prior art, appropriate parallel processing is required during acceleration. Especially when using a GPU, for the memory transfer between the CPU and the GPU, if the data size and the number of loops are small, the performance is often not high. In addition, there are cases where the combination of each loop statement (repetitive statement) that can be parallel-accelerated is not the fastest due to the timing of memory data transfer, etc. For example, when 3 out of 10 for statements (repetitive statements), namely the 1st, 5th, and 10th, can be accelerated compared to the CPU, the combination of the 1st, 5th, and 10th is not necessarily the fastest, etc.

[0190] In order to specify an appropriate parallel region, there has been an attempt to use the PGI compiler to optimize by trial and error on whether a for statement can be parallelized. However, in order to spend a lot of work in the trial and error, when provided as a service, there are problems such as a delay in the start of user use and an increase in cost.

[0191] Therefore, in the present embodiment, an appropriate offloading region is automatically extracted from a general program that does not assume parallelization. Therefore, it is realized that first, the check of the for statements that can be parallelized is performed, and then the genetic algorithm is used for the group of for statements that can be parallelized to repeatedly perform performance verification in a verification environment to explore an appropriate region. On the basis of reducing to the for statements that can be parallelized, in the form of a part of a gene, the parallel processing patterns that can be accelerated are maintained and recombined, so that the patterns that can be accelerated can be efficiently explored from the huge number of parallel processing patterns that can be obtained.

[0192] [Exploration image of the control unit (automatic offloading function unit) 11 based on the basic genetic algorithm]

[0193] Figure 4 It is a diagram showing the exploration by the control unit (automatic offloading function unit) 11 based on the basic genetic algorithm. Figure 4 It shows the exploration image of the process and the mapping of the gene sequence of the for statement.

[0194] The genetic algorithm is one of the combinatorial optimization methods that mimics the biological evolution process. The flowchart of the genetic algorithm is initialization → evaluation → selection → crossover → mutation → end determination.

[0195] In the present embodiment, the basic genetic algorithm that simplifies the process is used in the genetic algorithm. The basic genetic algorithm is a simplified genetic algorithm in which the gene is only set to 1 or 0, and the roulette wheel selection, one-point crossover, and mutation change the value of the gene at one position to the opposite value, etc.

[0196] <Initialization>

[0197] In initialization, after checking the parallelizability of all for statements in the application code, the parallelizable for statements are mapped to gene sequences. The case of performing GPU processing is set to 1, and the case of not performing GPU processing is set to 0. The number of individuals M specified in gene preparation is randomly assigned 1 or 0 in one for statement.

[0198] Specifically, the control unit (automatic unloading function unit) 11 (refer to Figure 2 ) obtains the application code 130 used by the user (refer to Figure 3 ), and as Figure 4 shown, checks the parallelizability of for statements according to the code patterns 141 of the application code 130. As Figure 4 shown, in the case where 5 for statements are found in the code pattern 141 (refer to the reference numeral a in Figure 4 ), a random 1-bit is assigned to each for statement. Here, 1 or 0 of 5 bits are randomly assigned to 5 for statements. For example, the case of using CPU processing is set to 0, and the case of handing over to GPU is set to 1. Among them, 1 or 0 is randomly assigned at this stage.

[0199] The code corresponding to the gene length is 5 bits, and the code of the gene length of 5 bits is 2 5 = 32 patterns, such as 10001, 10010,.... In addition, in Figure 4 , the circle mark (○ mark) in the code pattern 141 shows the appearance of the code.

[0200] <Evaluation>

[0201] In evaluation, deployment and performance measurement are performed (refer to the reference numeral b in Figure 4 ). That is, the performance measurement unit 116 (refer to Figure 2 ) compiles the code equivalent to the gene and deploys it on the verification machine 14 for execution. The performance measurement unit 116 performs benchmark test performance measurement. The fitness of the gene in the mode with good performance (parallel processing mode) is improved.

[0202] <Selection>

[0203] In selection, based on fitness, high-performance code patterns are selected (Select high performance codepatterns) (refer to the reference numeral c in Figure 4 ). The performance measurement unit 116 (refer to Figure 2) Based on the fitness, a specified number of individuals are selected for genes with high fitness. In this embodiment, roulette wheel selection corresponding to the fitness and elitist selection of the gene with the highest fitness are performed.

[0204] In Figure 4 , reducing the circle marks (○ marks) among the selected code patterns 142 to 3 shows the appearance of exploration.

[0205] <Crossover>

[0206] In crossover, at a fixed crossover rate Pc, a part of the genes are exchanged between the selected individuals at a certain point to create offspring individuals.

[0207] Perform gene crossover between a certain pattern (parallel processing pattern) selected by roulette wheel selection and other patterns. The position of one-point crossover is arbitrary. For example, crossover occurs at the 3rd bit among the above 5-bit codes.

[0208] <Mutation>

[0209] In mutation, at a fixed mutation rate Pm, each value of the genes of an individual is changed from 0 to 1 or from 1 to 0.

[0210] In addition, in order to avoid local solutions, mutation is introduced. Alternatively, it can be a method where mutation is not performed to reduce the computational amount.

[0211] <End determination>

[0212] As Figure 4 shown, generation of the next-generation code patterns after crossover and mutation (Generate nextgeneration code patterns after crossover&mutation) (refer to reference numeral d in FIG. 4) is performed.

[0213] In end determination, after repeating the specified number of generations T times, the process ends, and the gene with the highest fitness is taken as the solution.

[0214] For example, performance measurement is performed, and 3 with high speed 10010, 01001, 00101 are selected. The next generation recombines these 3 through the genetic algorithm. For example, a new pattern (parallel processing pattern) 10101 (an example) is created. At this time, mutations such as changing 0 to 1 are arbitrarily inserted into the recombined patterns. The above operations are repeated to find the pattern with the highest speed. A specified generation (for example, 20 generations) is determined, and the remaining patterns in the last generation are taken as the final solution.

[0215] <Deployment (Configuration)>

[0216] It is redeployed in the formal environment in the parallel processing mode with the highest processing performance corresponding to the gene with the highest adaptation degree and provided to users.

[0217] <Supplementary Explanation>

[0218] Explain the case where there are a considerable number of for statements (loop statements; repetition statements) that cannot be unloaded to the GPU. For example, even if there are 200 for statements, about 30 for statements can be unloaded to the GPU. Here, exclude the incorrect for statements and perform a genetic algorithm on these 30.

[0219] In OpenACC, there is a compiler that can specify with the instruction #pragma acc kernels, extract bytecodes for the GPU, and perform GPU offloading through execution. By writing the command for the for statement in this #pragma, it is possible to determine whether the for statement is running on the GPU.

[0220] For example, in the case of using C / C++, analyze the C / C++ code to find for statements. When a for statement is found, write the for statement using #pragma acc kernels, which is the syntax for parallel processing in OpenACC. Specifically, add the for statements one by one to an empty #pragma acc kernels and compile. If there is an error, the for statement itself cannot be processed by the GPU, so it is excluded. In this way, find the remaining for statements. Then, use the for statements without errors as the length (gene length). If there are 5 for statements without errors, the gene length is 5; if there are 10 for statements without errors, the gene length is 10. In addition, there are cases where parallel processing cannot be performed due to data dependencies such as using the previous processing for the next processing.

[0221] The above is the preparation stage. Next, perform genetic algorithm processing.

[0222] Obtain a code pattern with a gene length corresponding to the number of for statements. First, randomly assign parallel processing modes 10010, 01001, 00101,.... Perform genetic algorithm processing and compile. At this time, although the for statements can be unloaded, errors sometimes occur. This is the case where the for statements are in a hierarchical structure (GPU processing can be performed if any one is specified). In this case, the incorrect for statements can also remain. Specifically, there is a method that causes the processing time to become longer and time out.

[0223] Deployed in the verification machine 14, benchmark testing is performed. For example, if it is image processing, benchmark testing is performed through this image processing. The shorter the processing time, the higher the evaluation of the fitness. For example, set it to the reciprocal of the processing time. Set the fitness of an individual with a processing time of 10 seconds to 1, the fitness of an individual with a processing time of 100 seconds to 0.1, and the fitness of an individual with a processing time of 1 second to 10.

[0224] Select individuals with high fitness. For example, select 3 to 5 from 10, and recombine these to create a new code pattern. At this time, during the creation process, sometimes the same code pattern as before is obtained. In this case, since the same benchmark test is not required, the same data as before is used. In this embodiment, the code pattern and its processing time are stored in the storage unit 13.

[0225] Above, the exploration of the control unit (automatic unloading function unit) 11 based on the basic genetic algorithm has been described. Next, the unified processing method of data transmission will be described.

[0226] [Unified Processing Method of Data Transmission]

[0227] <Basic Idea>

[0228] In order to reduce CPU-GPU transmission, in addition to transmitting the variables of the nested loop as much as possible at the upper level, the present invention unifies the transmission timing of multiple variables, and further reduces such transmissions automatically transmitted by the compiler.

[0229] When reducing transmission, not only the nested units, but also the variables with consistent transmission timing to the GPU are unified for transmission. For example, if variables such as the results of GPU processing are processed by the CPU and then processed by the GPU again, etc., it is also possible to handle variables defined by the CPU used in multiple loop statements by uniformly sending them to the GPU before the start of GPU processing and returning them to the CPU after all GPU processing is completed.

[0230] In order to grasp the reference relationship between loops and variables during code analysis, for variables defined by multiple files according to their results, GPU processing and CPU processing are not nested. For variables separated between CPU processing and GPU processing, the data copy statement of OpenACC is used to specify the unified output specification.

[0231] For variables that do not need to be transmitted at the timing of loop statement processing and are uniformly transmitted before the start of GPU processing, use data present to clarify that transmission is not required.

[0232] When transferring data between the CPU and GPU, a temporary area is created (#pragma acc declare create). After the data is stored in the temporary area, the transfer is indicated by synchronizing the temporary area (#pragma acc update).

[0233] <Comparative Example>

[0234] First, the comparative example will be described.

[0235] The comparative examples are a conventional CPU program (refer to Figure 5 ), simply using the GPU (refer to Figure 6 ), and unified nesting (Non-Patent Document 2) (refer to Figure 7 ). In addition, the <1> to <4> etc. at the beginning of the loop statements in the following descriptions and figures are marked for ease of explanation (the same applies to other figures and their descriptions).

[0236] Figure 5 The loop statement of the shown conventional CPU program is described on the CPU program side.

[0237] In the <1> loop [for (i = 0; i < 10; i++)] {

[0238] }

[0239] Among them,

[0240] there is a <2> loop [for (j = 0; j < 20; j++)] {

[0241] }. Figure 5 The reference numeral e in the figure is the setting of the variables a and b in the above <2> loop.

[0242] In addition, next is

[0243] <3> loop [for (k = 0; k < 30; k++)] {

[0244] }

[0245] and

[0246] <4> loop [for (l = 0; l < 40; l++)] {

[0247] }. Figure 5 The reference numeral f in the figure is the setting of the variables c and d in the above <3> loop. Figure 5 The reference numeral g in the figure is the setting of the variables e and f in the above <4> loop.

[0248] Figure 5 The shown conventional CPU program is executed by the CPU (without using the GPU).

[0249] Figure 6 It represents a loop statement diagram for the case of solely using a GPU to transfer data from the CPU to the GPU for the conventional CPU program shown in Figure 5 . There are two types of data transfers: data transfer from the CPU to the GPU and data transfer from the GPU to the CPU. Hereinafter, the data transfer from the CPU to the GPU will be taken as an example.

[0250] Figure 6 The loop statement that solely uses the GPU shown in

[0251] is described on the CPU program side. In <1> loop [for(i = 0; i < 10; i++)] {

[0252] }

[0253] among them,

[0254] there exists <2> loop [for(j = 0; j < 20; j++)] {

[0255] }.

[0256] Furthermore, as shown by the reference numeral h in Figure 6 in <1> loop [for(i = 0; i < 10; i++)] {

[0257] }, at the upper part, a processable part that can be parallel-processed for a for statement, etc. based on the PGI compiler is specified by the OpenACC instruction #pragma acc kernels (parallel processing specification statement).

[0258] As shown by the dashed box including the reference numeral h in Figure 6 , through #pragma acc kernels, data is transferred from the CPU to the GPU. Here, a and b are transferred at this timing, so the transfer is performed 10 times.

[0259] In addition, as shown by the reference numeral i in Figure 6 in <3> loop [for(k = 0; k < 30; k++)] {

[0260] }, at the upper part, a processable part that can be parallel-processed for a for statement, etc. based on the PGI compiler is specified by the OpenACC instruction #pragma acc kernels.

[0261] As shown by the dashed box including the reference numeral i in Figure 6 , through #pragma acc kernels, c and d are transferred at this timing.

[0262] Here, in <4> loop [for(l = 0; l < 40; l++)] {

[0263] At the upper part of}, #pragma acc kernels is not specified. Since the efficiency of GPU processing for this loop is poor, GPU processing is not performed.

[0264] Figure 7 A diagram of a loop statement showing data transfer from the CPU to the GPU and from the GPU to the CPU based on unified nesting (Non-Patent Document 2).

[0265] In Figure 7 In the loop statement shown, at the position indicated by the reference numeral j in Figure 7 insert a data transfer instruction line from the CPU to the GPU, which here is the #pragma acc data copyin(a, b) of the copyin clauses for variables a and b.

[0266] The above #pragma acc data copyin(a, b) is specified for the uppermost loop that does not include the setting and definition of variable a (here, the upper part of the <1> loop [for(i = 0; i < 10; i++)]{

[0267] }).

[0268] At the timing indicated by the dotted box including the reference numeral j of Figure 7 transfer a and b, so one transfer occurs.

[0269] In addition, in Figure 7 In the loop statement shown, at the position indicated by the reference numeral k in Figure 7 insert a data transfer instruction line from the GPU to the CPU, which here is the #pragma acc data copyout(a, b) of the copyout clauses for variables a and b.

[0270] The above #pragma acc data copyout(a, b) is specified at the lower part of the <1> loop [for(i = 0; i < 10; i++)]{

[0271] }.

[0272] In this way, in the data transfer from the CPU to the GPU, by inserting the #pragma acc data copyin(a, b) of the copyin clause for variable a at the above position, the data transfer can be clearly indicated. Thus, it is possible to perform data transfer as uniformly as possible in the upper loop and avoid inefficient transfers such as transferring data each time in each loop as in the loop statement that simply uses the GPU as indicated by the reference numeral h in Figure 6}.

[0273] <Embodiment>

[0274] Next, this embodiment will be described.

[0275] 《Using data present to identify variables that do not need to be transferred》

[0276] In this embodiment, for variables defined by multiple files, GPU processing and CPU processing are not nested. For variables where CPU processing and GPU processing are separate, the specification of unified output is specified using the data copy statement of OpenACC. At the same time, unified transmission is performed, and at this time, data present is used to identify variables that do not need to be transferred.

[0277] Figure 8 A diagram showing a loop statement for unifying data transmission between CPU and GPU based on this embodiment. Figure 8 Compared with the comparative example Figure 7 of unified nesting.

[0278] In Figure 8 the loop statement shown, at the position indicated by the reference numeral l in Figure 8 insert an instruction line for data transfer from CPU to GPU, which is the #pragma acc datacopyin(a, b, c, d) of the copyin section for variables a, b, c, and d here.

[0279] The above #pragma acc data copyin(a, b, c, d) is specified for the uppermost loop that does not include the setting and definition of variable a (here it is the upper part of the <1> loop [for(i = 0; i < 10; i++)]{

[0280] }).

[0281] In this way, for variables defined by multiple files, GPU processing and CPU processing are not nested. For variables where CPU processing and GPU processing are separate, the specification of unified output is specified using the OpenACC data copy statement #pragma acc datacopyin(a, b, c, d).

[0282] In Figure 8 at the timing indicated by the dotted box containing the reference numeral l, a, b, c, and d are transferred, so there is 1 transfer.

[0283] Then, regarding the unified transmission using the above #pragma acc data copyin(a, b, c, d), for variables that do not need to be transferred at this time, use to identify in Figure 8The timing indicated by the double-dashed line frame including the figure mark m has been specified by the data present statement #pragma acc data present(a, b) of the GPU variable.

[0284] For variables that are transferred collectively using the above #pragma acc data copyin(a, b, c, d) and do not need to be transferred at this time, use the Figure 8 The timing indicated by the double-dashed line frame including the figure mark n has been specified by the data present statement #pragma acc data present (c, d) of the GPU existence variable.

[0285] exist <1> , <3> When the GPU processing is completed in the loop, insert the #pragma acc datacopyout(a, b, c, d) of the copyout statement section of variables a, b, c, d into the data transfer instruction line from GPU to CPU. Figure 8 of <3> The position of the symbol o where the loop ends.

[0286] By specifying unified transfer, variables that can be transferred uniformly are transferred uniformly, and variables that have been transferred and do not need to be transferred are displayed using data present, thereby reducing transfers and achieving further efficiency of the offload unit. However, even if OpenACC instructs transfer, the compiler may automatically determine and transfer depending on the compiler. Compiler-based automatic transfer refers to the phenomenon that, unlike OpenACC instructions, transfers between the CPU and GPU that were originally not required are automatically transferred due to reliance on the compiler.

[0287] 《Temporary regional storage of data》

[0288] Figure 9 It is a diagram showing a loop statement of unified transfer during data transfer between CPU and GPU according to the present embodiment. Figure 9 and Figure 8 The unified nesting and the variables that do not need to be transferred clearly correspond.

[0289] exist Figure 9 In the loop statement shown, Figure 9At the position indicated by the reference numeral p, the OpenACC declare create statement #pragma acc declare create for creating a temporary area is specified during the data transfer between the CPU and the GPU. Thus, a temporary area is created (#pragma acc declare create) during the data transfer between the CPU and the GPU, and the data is stored in the temporary area.

[0290] In addition, at Figure 9 the position indicated by the reference numeral q, the OpenACC declare create statement #pragma acc update for synchronizing the temporary area is used to indicate the transfer.

[0291] In this way, a temporary area is created, and the parameters are initialized in the temporary area for CPU-GPU transfer, thereby cutting off unnecessary CPU-GPU transfers. It is possible to reduce the undesired low-performance transfers under the indication of OpenACC.

[0292] [GPU offloading process]

[0293] Through the above unified processing method for data transfer, it is possible to extract appropriate loop statements during offloading and avoid inefficient data transfer.

[0294] However, even when using the above unified processing method for data transfer, there are programs that are not offloaded to the GPU. In effective GPU offloading, the number of loop times for the offloading process is relatively large.

[0295] Therefore, in this embodiment, as a preparatory stage for exploring the formal offloading process, an analysis tool is used to investigate the number of loop times. When using the analysis tool, it is possible to investigate the execution times of each line. Therefore, for example, a program with a loop of more than 50 million times can be assigned in advance as an object for exploring the offloading process, etc. The following is a specific description (partially repeated with Figure 3 the content described).

[0296] In this embodiment, first, the application program of the exploration offloading processing unit is analyzed to grasp loop statements such as for, do, and while. Next, sampling processing is performed, and an analysis tool is used to investigate the number of loop times of each loop statement. By determining whether there is a loop above a fixed value, it is decided whether to formally explore the offloading processing unit.

[0297] In the case of deciding to formally conduct the exploration, the processing of the genetic algorithm is started (refer to Figure 4)。In the initialization step, after checking the parallelizability of all loop statements in the application code, the cases where GPU processing is to be performed on parallelizable loop statements are set to 1, and the cases where GPU processing is not to be performed are set to 0, which are mapped to gene sequences. The gene prepares a specified number of individuals, and 1s and 0s are randomly assigned to each value of the gene.

[0298] Among them, in the code equivalent to the gene, based on the variable data reference relationship within the loop statement specified for GPU processing, explicit instructions for data transfer (#pragma acc data copyin / copyout / copy) are appended.

[0299] In the evaluation step, the code equivalent to the gene is compiled, deployed and executed on a verification machine, and benchmark test performance measurement is performed. The fitness of the gene of the mode with good performance improvement is increased. As described above, the code equivalent to the gene is inserted with parallel processing instruction lines (for example, refer to Figure 5 the reference numeral e) and data transfer instruction lines (for example, refer to Figure 5 the reference numeral g, Figure 6 the reference numeral h, Figure 7 the reference numeral j).

[0300] In the selection step, based on the fitness, genes with high fitness corresponding to the specified number of individuals are selected. In this embodiment, roulette wheel selection corresponding to the fitness and elite selection of the gene with the highest fitness are performed. In the crossover step, at a fixed crossover rate Pc, a part of the genes are exchanged at a certain point between the selected individuals to create offspring individuals. In the mutation step, each value of the gene of the individual is changed from 0 to 1 or from 1 to 0 at a fixed mutation rate Pm.

[0301] Until the mutation step ends, when creating the genes of the next generation of the specified number of individuals, similar to the initialization step, explicit instructions for data transfer are appended, and the evaluation, selection, crossover, and mutation steps are repeated.

[0302] Finally, in the end determination step, after repeating the specified number of generations, the process ends, and the gene with the highest fitness is used as the solution. The code mode with the highest performance corresponding to the gene with the highest fitness is redeployed in the formal environment and provided to the user.

[0303] Hereinafter, the implementation of the offloading server 1 will be described. This implementation is used to confirm the effectiveness of this embodiment.

[0304] [Implementation]

[0305] An implementation of automatically offloading a C / C++ application using a general PGI compiler will be described.

[0306] In this implementation, for the purpose of confirming the effectiveness of GPU automatic offloading, the target application is set to an application in C / C++ language, and the existing PGI compiler is used for GPU processing for illustration.

[0307] The C / C++ language has leading popularity in the development of OSS (Open Source Software) and proprietary software, and many applications are developed in C / C++ language. To confirm the offloading of applications used by general users, general applications of OSS such as encryption processing and image processing are used.

[0308] GPU processing is performed through the PGI compiler. The PGI compiler is a compiler for C / C++ / Fortran that interprets OpenACC. In this embodiment, the parallel processing section such as the for statement is specified by the OpenACC instruction #pragma acc kernels (parallel processing specification statement). Thereby, bytecode for the GPU is extracted, and GPU offloading can be performed by executing it. Furthermore, an error is prompted when there is a dependency between the data within the specified for statement and parallel processing is not possible, or in the case of multiple hierarchical structures with different nested for statements, etc. At the same time, instructions for displaying data transfer can be given through instructions such as #pragma acc data copyin / copyout / copy.

[0309] Corresponding to the specification in the above #pragma acc kernels (parallel processing specification statement), an instruction for displaying data transfer is given by inserting the #pragma acc data copyout(a[…]) of the OpenACC copyin clause to the above position.

[0310] <Summary of the actions in the implementation>

[0311] Describe the summary of the actions in the implementation.

[0312] The implementation performs the following processing.

[0313] Before starting the processing of the following Figure 10A -B process, prepare an accelerated C / C++ application and a benchmark tool for measuring its performance.

[0314] In the implementation, when there is a commission for using a C / C++ application, first, the code of the C / C++ application is parsed to find the for statement and understand the program structure such as the variable data used within the for statement. In the syntax parsing, a syntax parsing library such as LLVM / Clang is used.

[0315] In implementation, first, obtain the expectation of whether the application has GPU offloading effect, perform a benchmark test, and grasp the number of loops of the for statement mastered by the above syntax parsing. In grasping the number of loops, gcov of GNU coverage is used, etc. As analysis tools, "GNU profiler (gprof)" and "GNU coverage (gcov)" are known. Since both can investigate the execution times of each line, either can be used. The execution times can be targeted at, for example, only applications with a loop count of 10 million or more, but this value can be changed.

[0316] General-purpose applications for CPUs are not implemented with parallelization in mind. Therefore, first, it is necessary to exclude for statements that cannot perform GPU processing by themselves. Therefore, the #pragma acc kernels instruction for parallel processing is inserted into each for statement one by one, and a determination is made as to whether an error occurs during compilation. Regarding compilation errors, there are several types: cases where an external routine is called in a for statement, cases where different hierarchical structures are repeatedly specified in nested for statements, cases where there is a process of exiting midway through a for statement via break, etc., and cases where there is a data dependency for the data of the for statement. Different applications have different types of compilation errors, and there are also other cases. Compilation errors are excluded from the processing targets, and the #pragma instruction is not inserted.

[0317] Compilation errors are difficult to handle automatically, and in addition, even if handled, the effects often cannot be shown. In the case of external routine calls, sometimes #pragma acc routine can be used to avoid them, but many external calls are libraries. Even if GPU processing is included, their calls become bottlenecks and the performance is not high. By trying each for statement one by one, no compilation errors will occur for nested errors. In addition, in the case of exiting midway via break, etc., parallel processing requires fixing the number of loops and the program needs to be changed. In the case of having data dependencies, parallel processing cannot be performed in the first place.

[0318] Here, when the number of loop statements for which no error occurs even with parallel processing is a, a is the gene length. 1 of the gene corresponds to having a parallel processing instruction, and 0 corresponds to having no parallel processing instruction. The application code is mapped to a gene of length a.

[0319] Next, as an initial value, prepare a gene sequence specifying the number of individuals. As Figure 4As described, each value of the gene is randomly assigned 0 and 1 for creation. According to the prepared gene sequence, when the value of the gene is 1, the instructions for specifying GPU processing, namely #pragma acc kernels, #pragma acc parallel loop, and #pragma acc parallel loop vector, are inserted into the C / C++ code. The reason for not parallelizing in a single loop or the like is that if it is the same processing, the kernel performs better with the PGI compiler. At this stage, the part to be processed by the GPU is determined in the code corresponding to a certain gene.

[0320] Use the PGI compiler on a machine with a GPU to compile the C / C++ code with the instructions for parallel processing and data transfer inserted. Deploy the compiled executable file and measure the performance with a benchmark tool.

[0321] After measuring the benchmark performance for all individuals, set the fitness of each gene sequence according to the benchmark processing time. Based on the set fitness, select the remaining individuals. Perform genetic algorithm processing of crossover processing, mutation processing, and direct replication processing on the selected individuals to create a new generation of individual groups.

[0322] For the individuals of the new generation, perform instruction insertion, compilation, performance measurement, fitness setting, selection, crossover, and mutation processing. Here, when a gene with the same pattern as before is generated in the genetic algorithm processing, compile the individual, do not perform performance measurement, and use the same measurement value as before.

[0323] After the genetic algorithm processing for the specified number of generations ends, use the C / C++ code with instructions corresponding to the gene sequence with the highest performance as the solution.

[0324] Among them, the number of individuals, the number of generations, the crossover rate, the mutation rate, the fitness setting, and the selection method are parameters of the genetic algorithm and are specified separately. The proposed technology can automate the GPU offloading that previously required the time and skills of professional technicians by automating the above processing.

[0325] Figure 10A -B is a flowchart for explaining the action outline of the above implementation, Figure 10A and Figure 10B are connected by connectors.

[0326] Perform the following processing using the OpenACC compiler for C / C++.

[0327] <Code analysis>

[0328] In step S101, the application code analysis unit 112 (refer to Figure 2)Perform code parsing of C / C++ application programs.

[0329] <Determination of loop statements>

[0330] In step S102, the parallel processing specifying unit 114 (refer to Figure 2 ) determines the loop statements and reference relationships of the C / C++ application program.

[0331] <Parallel processing possibility of loop statements>

[0332] In step S103, the parallel processing specifying unit 114 checks the GPU processing possibility (#pragma acc kernels) of each loop statement.

[0333] <Repetition of loop statements>

[0334] The control unit (automatic unloading function unit) 11 repeats the number of times of the loop statement for the processing of steps S105 - S116 between the loop start of step S104 and the loop end of step S117.

[0335] <Repetition the number of times of the loop statement (Part 1)>

[0336] The control unit (automatic unloading function unit) 11 repeats the number of times of the loop statement for the processing of steps S106 - S107 between the loop start of step S105 and the loop end of step S108.

[0337] In step S106, the parallel processing specifying unit 114 compiles each loop statement by specifying GPU processing with OpenACC (#pragma acc kernels).

[0338] In step S107, when there is an error, the parallel processing specifying unit 114 checks the GPU processing possibility with the next directive sentence (#pragma acc parallel loop).

[0339] <Repetition the number of times of the loop statement (Part 2)>

[0340] The control unit (automatic unloading function unit) 11 repeats the number of times of the loop statement for the processing of steps S110 - S111 between the loop start of step S109 and the loop end of step S112.

[0341] In step S110, the parallel processing specifying unit 114 compiles each loop statement by specifying GPU processing with OpenACC (#pragma acc parallel loop).

[0342] In step S111, when there is an error in the parallel processing designated section 114, the next directive sentence is used to check the GPU processing possibility (#pragma acc parallel loop vector).

[0343] <Repeat the loop number of times (the third one)>

[0344] The control section (auto-unload function section) 11 repeats the processing of steps S114 - S115 the number of times of the loop statement between the loop start of step S113 and the loop end of step S116.

[0345] In step S114, the parallel processing designated section 114 compiles each loop statement by specifying GPU processing with OpenACC (#pragma acc parallel loop vector).

[0346] In step S115, when there is an error, the parallel processing designated section 114 removes the GPU processing directive sentence from the loop statement.

[0347] <Count of the For statement>

[0348] In step S118, the parallel processing designated section 114 counts the number of for statements without compilation errors as the gene length.

[0349] <Prepare the mode of the specified number of individuals>

[0350] Next, as an initial value, the parallel processing designated section 114 prepares gene sequences of the specified number of individuals. Here, 0 and 1 are randomly assigned to create them.

[0351] In step S119, the parallel processing designated section 114 maps the C / C++ application code to the genes and prepares the mode of the specified number of individuals.

[0352] According to the prepared gene sequences, when the value of the gene is 1, insert the instruction for the specified parallel processing into the C / C++ code (for example, refer to Figure 4 the #pragma instruction).

[0353] The control section (auto-unload function section) 11 repeats the processing of steps S121 - S128 the specified number of generations between the loop start of step S120 and the loop end of step S129.

[0354] In addition, during the above-mentioned specified number of generations, the processing of steps S122 - S125 is also repeated the specified number of individuals between the loop start of step S121 and the loop end of step S126. That is, during the specified number of generations, the specified number of individuals is repeated in a nested state for processing.

[0355] <Data transfer specification>

[0356] In step S122, based on the variable reference relationship, the data transfer specification unit 113 performs data transfer specification using explicitly indicated lines (#pragma acc data copy / copyin / copyout / present and #pragam acc declarecreate, #pragma acc update).

[0357] <Compilation>

[0358] In step S123, the parallel processing mode creation unit 115 (refer to Figure 2 ) compiles the C / C++ code specified by the instruction with the PGI compiler according to the gene pattern. That is, the parallel processing mode creation unit 115 compiles the created C / C++ code with the PGI compiler on the verification machine 14 equipped with a GPU.

[0359] Here, in the case of making multiple parallel specifications for nested for statements, etc., compilation errors sometimes occur. In this case, the same processing as in the case of the processing time timeout during performance measurement is performed.

[0360] In step S124, the performance measurement unit 116 (refer to Figure 2 ) deploys the execution file on the verification machine 14 equipped with a CPU-GPU.

[0361] In step S125, the performance measurement unit 116 executes the deployed execution file (binary file) and measures the benchmark performance at the time of uninstallation.

[0362] Here, in the intermediate generation, genes of the same pattern as before are not measured, and the same value is used. That is, when genes of the same pattern as before are generated in the genetic algorithm process, the individual is not compiled and performance-measured, but the same measured value as before is used.

[0363] In step S127, the execution file creation unit 117 (refer to Figure 2 ) evaluates in such a way that the shorter the processing time of an individual, the higher the fitness, and selects an individual with high performance.

[0364] In step S128, the execution file creation unit 117 performs crossover and mutation processing on the selected individuals to create the next generation of individuals. The execution file creation unit 117 performs compilation, performance measurement, fitness setting, selection, crossover, and mutation processing on the next generation of individuals.

[0365] That is, after measuring the benchmark performance of all individuals, the fitness of each gene sequence is set according to the benchmark processing time. Based on the set fitness, the remaining individuals are selected. The execution file creation unit 117 performs genetic algorithm processing such as crossover processing, mutation processing, and direct copy processing on the selected individuals to create a new generation of individual groups.

[0366] In step S130, after the genetic algorithm processing of the specified generation ends, the execution file creation unit 117 uses the C / C++ code corresponding to the gene sequence with the highest performance (the highest performance parallel processing mode) as the solution.

[0367] <Parameters of the genetic algorithm>

[0368] The above-mentioned number of individuals, number of generations, crossover rate, mutation rate, fitness setting, and selection method are parameters of the genetic algorithm. The parameters of the genetic algorithm can also be set as follows, for example.

[0369] The parameters and conditions of the basic genetic algorithm to be executed can be set as follows, for example.

[0370] Gene length: Number of parallel loop statements

[0371] Number of individuals M: Less than or equal to the gene length

[0372] Number of generations T: Less than or equal to the gene length

[0373] Fitness: (Processing time) (-1 / 2)

[0374] With this setting, the shorter the benchmark processing time, the higher the fitness. In addition, by making the fitness the (-1 / 2) power of the processing time, it is possible to prevent the fitness of specific individuals with short processing times from being too high and the exploration range from becoming narrow. In addition, when the performance measurement does not end within a fixed time, it is timed out, and the fitness is calculated as a time (long time) such as 1000 seconds for the processing time. This timeout time can also be changed according to the performance measurement characteristics.

[0375] Selection: Roulette wheel selection

[0376] Among them, the gene with the highest fitness in this generation is neither crossed nor mutated, and elite preservation for the next generation is performed simultaneously.

[0377] Crossover rate Pc: 0.9

[0378] Mutation rate Pm: 0.05

[0379] <Cost-performance ratio>

[0380] The cost-performance ratio of the automatic unloading function is described.

[0381] When only looking at the price of the hardware of GPU boards such as NVIDIA Tesla, the price of a machine equipped with a GPU is about twice that of a machine with only a CPU. However, generally in the costs of a data center or the like, the hardware and system development costs are less than one-third, the operating costs such as electricity costs and maintenance / operation systems exceed one-third, and other costs such as service customization are around one-third. In this embodiment, it is possible to improve the performance of processing that takes time in application programs that perform operations such as encryption processing and image processing by more than twice. Therefore, even if the server hardware price itself doubles, a sufficient cost effect can be expected.

[0382] In this embodiment, gcov, gprof, etc. are used to pre-determine application programs with many loops and long execution times and perform offloading attempts. Thereby, it is possible to find application programs that can be effectively accelerated.

[0383] <Time until starting to use the official service>

[0384] Describe the time until starting to use the official service.

[0385] When setting the time from compilation to one performance measurement to be about 3 minutes, in a genetic algorithm with a population size of 20 and 20 generations, the exploration of solutions takes at most about 20 hours. However, since the compilation and measurement of gene patterns that are the same as before are omitted, it ends within 8 hours. In the case of multiple clouds, virtual hosts, and network services, it actually takes about half a day to start using the service. In this embodiment, for example, automatic offloading within half a day can be performed. Therefore, if it is automatic offloading within half a day and a trial can be made from the beginning, a sufficient improvement in user satisfaction can be expected.

[0386] In order to explore the offloading part in a short time, consider parallelizing the population size by multiple verification machines to perform performance measurement. Adjusting the timeout time according to the application program also helps to shorten the time. For example, consider a timeout when the time taken for the offloading process is twice the execution time of the CPU. In addition, the larger the population size and the number of generations, the higher the possibility of finding a high-performance solution. However, when setting each parameter to the maximum, it is necessary to perform the compilation and performance benchmark tests for the number of population size × number of generations. Therefore, it takes time until starting to use the official service. In this embodiment, the genetic algorithm is performed with a smaller population size and number of generations, and a wide range of exploration is performed by setting the crossover rate Pc to a relatively high value of 0.9, so as to quickly find a solution with a certain level of performance.

[0387] [Expansion of the directive sentence]

[0388] In this embodiment, in order to increase the applicable applications, the directive sentences are extended. Specifically, as the directive sentences specifying GPU processing, in addition to the kernels directive sentence, the parallel loop directive sentence and the parallelloop vector directive sentence are also extended.

[0389] Under the OpenACC standard, kernels is used for single loop and tightly nested loop. In addition, parallel loop is used for loops that also include non-tightly nested loops. parallel loop vector is used for loops that cannot be parallelized but can be vectorized. Here, a tightly nested loop means a simple loop where, through nested loops, for example, when two loops incrementing i and j become nested, the lower loop processes using i and j and the upper loop does not perform processing. In addition, there are the following differences in the implementation of PGI compilers, etc.: the determination of parallelization for kernels is performed by the compiler, and the determination of parallelization for parallel is performed by the programmer. In addition, in Non-Patent Document 2, for simple loops, in kernels such as loops that cannot be non-tightly nested loops or parallelized, the loop statements that produce errors are excluded from the target, so the application range is narrow.

[0390] Therefore, in this embodiment, kernels is used in single and tightly nested loops, and parallel loop is used in non-tightly nested loops. In addition, parallel loop vector is used in loops that cannot be parallelized but can be vectorized.

[0391] Here, by setting it as the parallel directive sentence, there is a risk that the reliability will further decrease in the case of kernels. However, it is envisaged to sample and test the final offloading program, check the difference from the CPU result, display the result to the user, and have the user confirm it. Since the CPU and GPU are inherently different in hardware, there are differences in the number of significant digits and rounding errors, etc., and even for kernels alone, it is necessary to check the difference from the CPU result.

[0392] The above describes [Automatic offloading action (before starting operation)]. The automatic offloading action (before starting operation) is summarized as follows.

[0393] Figure 3 Steps S11 and S12: Code analysis

[0394] Figure 3 Steps S21 to S23: Extracting the detachable part

[0395] Figure 3 Steps S31 to S32: Exploring the appropriate detachable part

[0396] Figure 3 Step S41: Adjusting the resource amount

[0397] Figure 3 Step S51: Adjusting the configuration position

[0398] Figure 3 Steps S61 to S63: Executing file configuration and action verification

[0399] The above steps S11 to S63 are for the conversion of code, adjustment of resource amount, adjustment of configuration position, and verification required before the application starts to be used.

[0400] [Automatic uninstallation action (after starting to be used)]

[0401] Next, with reference to the Figure 3 step numbers mentioned above, the actions of each part of the automatic uninstallation action (after starting to be used) will be described.

[0402] As described above, the uninstallation server 1 performs reconstruction after the GPU logic starts to be used.

[0403] In the reconstruction after starting to be used, after the application starts to be used, the usage characteristics, etc. are analyzed, and the necessary reconstruction is performed. The objects of reconstruction are the same as those before starting to be used, namely code conversion, adjustment of resource amount, and adjustment of configuration position.

[0404] <Basic guidelines for FPGA reconstruction during operation>

[0405] First, the basic guidelines for FPGA reconstruction during operation will be described.

[0406] It is possible to Figure 3 by the method (for details, refer to the Figure 11 flowchart described later), for the application program specified by the user, automatically unload the loop statement part suitable for the GPU to the GPU. After unloading to the commercial environment used by the user, the actual performance and price in the commercial environment are confirmed, and the user starts to use the application program. Among them, the test cases for performance optimization (the items for measuring performance when comparing performances in multiple unloading modes) use the assumed usage data before starting to be used specified by the user, which may deviate significantly from the data actually used after starting to be used.

[0407] Therefore, hereinafter, in cases where the usage mode after the start of operation is different from the original assumption, and offloading another logic to the FPGA improves performance, etc., research has been conducted on suppressing the user impact and reconstructing the GPU logic. Reconstruction sometimes changes to loop statement offloading even within the same application program, and sometimes changes to offloading of a different application program.

[0408] There are two types of GPU reconstruction: dynamic reconstruction and static reconstruction. Dynamic reconstruction is a technology for changing the circuit structure while the GPU is operating, and the interruption time for rewriting is in the order of milliseconds. On the other hand, static reconstruction is a technology for changing the circuit structure after stopping the GPU, and the interruption time is about 1 second. Which method of dynamic reconstruction or static reconstruction to adopt can be selected according to the degree of user impact of the interruption time, and the reconstruction method provided by the GPU manufacturer can be chosen. However, in both methods, there is an interruption time and a test for action confirmation for rewriting to another logic, so restrictions such as not reconstructing frequently and only making suggestions when the effect is above a threshold are set.

[0409] The research on reconstruction starts with the analysis of the request trend at fixed time intervals (such as one month, etc.). By analyzing the request trend, it is determined whether there are application programs with a higher or equal processing load compared to the application program currently being offloaded. Next, for requests with a high processing load, instead of using the assumed data, the data actually used for commercial purposes (the data actually used by users) is used, and optimization attempts for GPU offloading are made in the verification environment ( Figure 3 ).

[0410] Based on whether the calculation results of the processing time and usage frequency exceed the threshold or are below the threshold, it is determined whether the improvement effect of the new offloading mode found through verification is sufficiently high compared to the current offloading mode. When the calculation results of the processing time and usage frequency are higher than the threshold, the user is advised to reconstruct. After the user agrees, the commercial environment is reconstructed, but the reconstruction is carried out while suppressing the user impact as much as possible. In addition, when the calculation results of the processing time and usage frequency are below the threshold, the user is not advised to reconstruct.

[0411] <Flowchart showing reconstruction after the start of operation>

[0412] Figure 11 This is a flowchart showing the reconstruction after the start of operation of offloading server 1. It is an example applied to GPU reconstruction.

[0413] In step S71, the request processing load analysis unit 120 analyzes the request processing load of the data actually used by the user. In addition, the detailed process of recording and analyzing commercial request data is described later in Figure 12 .

[0414] In step S72, the representative data selection unit 121 extracts the offloading mode that accelerates the test cases of commercial representative data through verification environment measurement for the applications with higher loads among multiple loads. Specifically, the representative data selection unit 121 selects any one data from the actual request data corresponding to the mode of the data size frequency distribution analyzed by the request processing load analysis unit 120 and designates it as the representative data.

[0415] In the above steps S71 - S72, the applications with higher loads are selected.

[0416] In step S73, the improvement degree calculation unit 122 extracts the offloading mode that accelerates the test cases of commercial representative data through verification environment measurement among the applications with higher loads among multiple loads.

[0417] Specifically, in step S73, the improvement degree calculation unit 122 determines a new offloading mode (the new offloading mode found in the verification environment) by executing the application code analysis unit 112, the parallel processing specification unit 114, the parallel processing mode creation unit 115, the performance measurement unit 116, and the execution file creation unit 117 based on the representative data selected by the representative data selection unit 121, and compares the processing time and usage frequency of the determined new offloading mode with those of the current offloading mode to calculate the performance improvement effect.

[0418] That is, the improvement degree calculation unit 122 measures the processing time of the current offloading mode and the multiple new offloading modes extracted, and obtains the performance improvement effect based on the commercial usage frequency. Specifically, the improvement degree calculation unit 122 calculates in the test cases under commercial representative data (the actual processing reduction time in the verification environment) × (the usage frequency in the commercial environment) for the current offloading mode: Equation (1). Then, the improvement degree calculation unit 122 calculates (the actual processing reduction time in the verification environment) × (the usage frequency in the commercial environment) for the multiple new offloading modes: Equation (2). In addition, the detailed process of the extraction process of commercial representative data is described later in Figure 13 below.

[0419] In this way, in step S73, for the applications that are performing GPU offloading, the case of not being offloaded is calculated by multiplying by the improvement degree coefficient, so as to correct it to only CPU processing for comparison. In addition, when selecting representative data, sometimes the average data size has a large difference from the actual used data, so the mode of the data size is used.

[0420] In step S74, the improvement degree calculation unit 122 measures the processing time of the current offloading mode and the multiple new offloading modes extracted, and obtains the performance improvement effect based on the commercial usage frequency.

[0421] In step S75, the reconstruction recommendation unit 123 determines the reconstruction recommendation by checking whether the performance improvement effect of the new offloading mode is above a specified threshold of the current offloading mode. Specifically, the reconstruction recommendation unit 123 confirms whether "the calculation result of formula (2) / the calculation result of formula (1)" calculated by the improvement degree calculation unit 122 is above the specified threshold. If "the calculation result of formula (2) / the calculation result of formula (1)" is above the specified threshold, it recommends reconstruction; if it is less than the specified threshold, it does nothing (no reconstruction recommendation is made).

[0422] In step S76, the reconstruction recommendation unit 123 makes a recommendation to the subscribed user to prompt the execution of FPGA reconstruction and obtains an answer of whether GPU reconstruction can be performed or not from the subscribed user.

[0423] In step S77, the control unit 11 performs reconstruction by starting another OpenACC in the commercial environment, and ends the processing of this flow. Specifically, the control unit 11 stops the old OpenACC and starts the new OpenACC. Then, the control unit 11 rebuilds the server that processes the new OpenACC and changes the routing in such a way that the new processing flow goes to the new server.

[0424] Figure 12 It is a detailed flowchart of the analysis process for commercial request data and is a Figure 11 subroutine of step S71.

[0425] When called by the Figure 11 subroutine of step S71, in step S81, the request processing load analysis unit 120 calculates the actual processing time and the total number of uses based on the usage records of each application program over a fixed period (a long time; for example, one month, etc.). Among the application programs that are performing GPU offloading, the processing time in the case where offloading is not performed is temporarily calculated. Based on the test records under the assumed usage data before the start of operation, the improvement degree coefficient is obtained as (the actual processing time when only CPU processing is performed) / (the actual processing time when GPU offloading is performed). The request processing load analysis unit 120 totals the values obtained by multiplying the actual processing time by the improvement degree coefficient as the total processing time for comparison.

[0426] In step S82, the request processing load analysis unit 120 compares the total actual processing time among all application programs.

[0427] In step S83, the request processing load analysis unit 120 sorts in order of the total actual processing time in sequence and determines multiple application programs with the highest processing time load.

[0428] In step S84, the request processing load analysis unit 120 obtains request data of the application programs with higher loads in a fixed time (short time; 12 hours, etc.), arranges the data sizes at fixed intervals, creates a frequency distribution, and returns Figure 11 to step S71.

[0429] Figure 13 The detailed flowchart of the improvement degree calculation process is a Figure 11 subroutine of step S73.

[0430] When called by the Figure 11 subroutine of step S73, in step S91, the improvement degree calculation unit 122 extracts the offloading mode that accelerates the test cases for commercial representative data among multiple application programs with higher loads through verification environment measurement. Specifically, the improvement degree calculation unit 122 counts the number of for statements in the application programs with higher loads by using syntax parsing, finds the for statements that cannot be processed by the GPU through the compiler function, and removes them.

[0431] In step S92, the improvement degree calculation unit 122 sets the number of remaining for statements as the gene length, creates a fixed number of gene patterns, with 1 corresponding to GPU execution and 0 corresponding to non-execution, and appends OpenACC instructions.

[0432] In step S93, the improvement degree calculation unit 122 compiles the OpenACC file corresponding to the gene pattern with the verification environment machine, measures the performance with the test cases of commercial representative data, and sets a higher fitness for the faster pattern.

[0433] In step S94, the improvement degree calculation unit 122 performs processes such as crossover according to the performance fitness, creates the next-generation gene pattern. The improvement degree calculation unit repeats the genetic algorithm process for a certain number of generations, takes the final highest-performance pattern as the solution, and returns Figure 11 to step S73.

[0434] [Implementation supplement]

[0435] Supplementary description of the implementation.

[0436] Before starting the operation, the GPU is offloaded in the same manner as the previous implementation tool (Non-Patent Document 3).

[0437] Regarding the reconstruction during operation ( "during operation" is a way of "after starting the operation"), it is implemented in such a way that the processing steps of the flowchart showing Figures 11 to 13 the reconstruction after starting the operation act in sequence.

[0438] Hereinafter, supplements are made to the content determined during the implementation, etc.Figures 11 to 13 The process will be described as follows.

[0439] First, in the load analysis of step S71 described above, Figure 11 the request data for a fixed time (long time) is analyzed to determine the application with the highest load, and the actual request data of this application for a fixed time (short time) is obtained.

[0440] Here, the administrator can arbitrarily set the number of applications with the highest load and the fixed time, so there is a degree of freedom. Among them, the above-mentioned long time is assumed to be a relatively long period of more than one month. In addition, the above-mentioned short time is assumed to be a relatively short period such as 12 hours. By analyzing the request data, the actual processing time and the number of uses of the application are totaled and obtained using the time command of Linux (registered trademark). Since the actual time consumed by the application appears in the log, the time command can calculate the value obtained by summing the log times and the time.

[0441] · Before starting operation

[0442] In the selection of representative data (before starting operation) of step S72 described above, Figure 11 the actual request data of the application with the highest load for a fixed time (short time) is arranged at regular intervals of a fixed size to make a frequency distribution. At this time, the number of levels of the frequency distribution is determined by Sturges' rule. Sturges' rule is a formula that determines the number of levels to be 1 + log2n as appropriate when the number of uses of the application is n (n is an arbitrary natural number). In order to use Sturges' rule, the number of levels needs to be determined. After selecting the modal class group, a representative data is selected from the modal class group. In addition, when selecting the representative data, the data with the value closest to the center of the class group is selected as the representative data.

[0443] · After starting operation

[0444] In the selection of representative data (after starting operation) of step S72 described above, Figure 11 using the selected representative data, in the same processing as before starting operation, GPU offloading is performed on the application with the highest load. The test case for performance measurement does not use the assumed usage data, but uses commercial representative data, which is different from before starting operation.

[0445] In the improvement degree calculation of step S73 described above, Figure 11 it is necessary to observe the improvement effect when reconstructing the commercial environment in the new offloading mode. Since it is before suggesting the user to reconstruct, it has to be verified using the verification environment server. The improvement degree calculation unit 122 measures the improvement for one-time processing using commercial representative data and calculates the overall improvement degree using the commercial usage frequency. Then,Figure 11 In the calculation of the improvement degree in step S74, the effects obtained under the condition of reconstructing the commercial environment are compared.

[0446] In the above-mentioned Figure 11 In the reconstruction recommendation in step S75, when the "calculation result of formula (2) / calculation result of formula (1)" is less than the specified threshold, the reconstruction recommendation unit 123 does not recommend reconstruction to the user. Since frequent reconstruction recommendations cause inconvenience to the user, by setting the threshold for effect improvement to a value sufficiently larger than 1 time, the frequent occurrence of reconstruction recommendations can be suppressed, and the cases of truly effective reconstruction can be retained. The reconstruction threshold is an implementation that can be variably set. For example, 1.5 is set.

[0447] In the above-mentioned Figure 11 In the reconstruction recommendation in step S76, the reconstruction recommendation unit 123 provides information on price changes or improvement effects and recommends reconstruction to the user. The information on price changes or improvement effects is the price that changes through reconstruction and the information on how many times the improvement effect is in the verification environment even if the price does not change. Thus, the subscribing user can judge whether reconstruction is better.

[0448] In the above-mentioned Figure 11 In the reconstruction execution in step S77, the control unit 11 stops the pre-reconstruction OpenACC through the GPU processing machine (offloading server 1) and starts the post-reconstruction OpenACC as the implementation of the reconstruction. Stopping and starting OpenACC requires an interruption time of about 1 second. If you want to further reduce the interruption time, you can also adopt the method of newly starting another machine (other server) that processes the post-reconstruction OpenACC and switching the routing of new requests to the new machine. In the method of restarting another machine, 2 times the GPU resources are used, but this resource is only temporarily increased before and after the switch.

[0449] The above describes the automatic offloading operation (after starting operation).

[0450] [Embodiment]

[0451] Describe the embodiment.

[0452] To confirm the effectiveness of GPU reconstruction, the following <Implementation Utilization Tool> is used.

[0453] The target application program is an application program in C / C++ language, and the GPU uses NVIDIA Tesla T4 * 2.

[0454] The machine for compilation is Supermicro SYS-1029GP-TR (CPU: Intel Xeon Silver 4210R, RAM: 128GB).

[0455] GPU processing uses the PGI compiler 19.10 and CUDA toolkit 10.1. The PGI compiler is a C / C++ / Fortran compiler that interprets OpenACC. The PGI compiler generates and executes GPU-oriented bytecode by specifying loop statements such as for statements with OpenACC directives such as #pragma acc kernels and #pragma acc parallel loop. By generating and executing GPU-oriented bytecode, the PGI compiler can perform GPU offloading. At the same time, the PGI compiler can indicate explicit data transfer and no data transfer through directives such as #pragma acc data copyin / copyout / copy and #pragma acc data present.

[0456] In the syntax parsing of the C / C++ language, the syntax parsing library of LLVM / Clang 6.0 (python binding of libClang) is used.

[0457] In GPU offloading, for statements that cannot use GPU processing are found through the functions of the PGI compiler and removed from the genetic algorithm. Here, loops with a small number of loop iterations (e.g., less than 1000 times) are likely to have poor effects, so they are excluded, and the gene length can also be shortened. In the analysis of the number of loop iterations, gcov of the profiler can be used.

[0458] [Hardware Structure]

[0459] The offloading server 1 of this embodiment is implemented by a computer 900 with a structure as shown, for example, Figure 14 In addition, Figure 2 The verification machine 14 shown is outside the offloading server 1.

[0460] Figure 14 It is a hardware structure diagram showing an example of the computer 900 that implements the functions of the offloading server 1.

[0461] The computer 900 has a CPU 910, a RAM 920, a ROM 930, an HDD 940, a communication interface (I / F: Interface) 950, an input / output interface (I / F) 960, and a media interface (I / F) 970.

[0462] The CPU 910 operates based on programs stored in the ROM 930 or HDD 940 and controls each part. The ROM 930 stores the boot program executed by the CPU 910 when the computer 900 starts up, programs dependent on the hardware of the computer 900, etc.

[0463] The HDD 940 stores programs executed by the CPU 910 and data used by such programs, etc. The communication interface 950 receives data from other devices via the communication network 80 and sends it to the CPU 910, and sends the data generated by the CPU 910 to other devices via the communication network 80.

[0464] The CPU 910 controls output devices such as a display and a printer, and input devices such as a keyboard and a mouse via the input / output interface 960. The CPU 910 obtains data from the input devices via the input / output interface 960. In addition, the CPU 910 outputs the generated data to the output devices via the input / output interface 960.

[0465] The medium interface 970 reads a program or data stored in the recording medium 980 and provides it to the CPU 910 via the RAM 920. The CPU 910 loads the program from the recording medium 980 to the RAM 920 via the medium interface 970 and executes the loaded program. The recording medium 980 is, for example, an optical recording medium such as a DVD (Digital Versatile Disc), a PD (Phase - change rewritable Disk), a magneto - optical recording medium such as an MO (Magneto - Optical disk), a tape medium, a magnetic recording medium, or a semiconductor memory, etc.

[0466] For example, when the computer 900 functions as the offloading server 1 of the present embodiment, the CPU 910 of the computer 900 realizes the functions of each part of the offloading server 1 by executing the program loaded onto the RAM 920. In addition, data of each part of the offloading server 1 is stored in the HDD 940. The CPU 910 of the computer 900 reads and executes these programs from the recording medium 980. As other examples, these programs can also be obtained from other devices via the communication network 80.

[0467] [Effect]

[0468] As described above, the offloading server 1 of the present embodiment offloads specific processing of an application to a GPU, and has: a request processing load analysis unit 120 that analyzes the request processing load based on the current offloading mode of the data used by the user; a representative data selection unit 121 that determines the application with the highest processing load analyzed by the request processing load analysis unit 120, and selects representative data from the request data used by the application; an improvement degree calculation unit 122 that creates a new offloading mode (the new offloading mode found in the verification environment) based on the representative data selected by the representative data selection unit 121, compares the processing time and usage frequency of the created new offloading mode with those of the current offloading mode, and calculates the performance improvement effect; an improvement degree calculation unit 122 that compares the processing time and usage frequency of the new offloading mode (the new offloading mode found in the verification environment) with those of the current offloading mode based on the representative data selected by the representative data selection unit 121, and calculates the performance improvement effect; and a reconstruction recommendation unit 123 that recommends GPU reconstruction when the performance improvement effect is equal to or greater than a specified threshold value.

[0469] In this way, the offloading server 1 is reconstructed into a more appropriate GPU logic not only according to the usage characteristics before the start of operation but also according to the usage characteristics after the start of operation, thereby achieving efficient resource utilization in a GPU with a limited amount of resources.

[0470] Specifically, since the offloading server 1 can reconstruct the GPU logic according to changes in usage characteristics after the application starts to operate, when the data actually used after the start of operation may deviate significantly, for example, when the usage mode after the start of operation is different from the initial assumption, or when offloading another logic to the GPU results in higher performance, the GPU logic can be reconstructed with less impact on the user.

[0471] In addition, even in the same application, GPU reconstruction sometimes changes to different loop statement offloading and sometimes changes to offloading of different applications. The reconstruction targets include various aspects such as the GPU, FPGA offloading logic, resource amount, and configuration location.

[0472] Consider the cost of applying the present embodiment. A GPU board such as NVIDIA Tesla (registered trademark) costs about 200,000 yen per piece. Therefore, when only looking at the hardware price, the price of a machine equipped with a GPU is usually twice that of a machine with only a CPU. However, even though the hardware price itself doubles, for this technology that can improve the performance of loop statement processing that has taken time for CPU processing by more than three times and then reconstructs it into an appropriate logic, it can be evaluated as cost-effective.

[0473] In the offloading server 1 of the present embodiment, it is characterized by having: an application code analysis unit 112 that analyzes the source code of an application; a data transfer designation unit 113 that analyzes the reference relationship of variables used in the loop statements of the application, and for data that can be transferred outside the loop, performs data transfer designation using a clearly specified line that clearly specifies data transfer outside the loop; a parallel processing designation unit 114 that determines the loop statements of the application, and for each determined loop statement, designates a parallel processing designation statement in the accelerator for compilation; a parallel processing mode creation unit 115 that creates a parallel processing mode that excludes duplicate statements with compilation errors from the offloading target, and for loop statements without compilation errors, designates whether to perform parallel processing; a performance measurement unit 116 that compiles the application in the parallel processing mode, configures it in an accelerator verification device, and executes performance measurement processing when offloading to the GPU; and an execution file creation unit 117 that selects the parallel processing mode with the highest processing performance from multiple parallel processing modes based on the performance measurement results of the performance measurement processing, compiles the parallel processing mode with the highest processing performance, creates an execution file, and an improvement degree calculation unit 122 determines a new offloading mode by executing the application code analysis unit 112, the parallel processing designation unit 114, the parallel processing mode creation unit 115, the performance measurement unit 116, and the execution file creation unit 117 based on the representative data selected by the representative data selection unit 121, compares the processing time and usage frequency of the determined new offloading mode with the processing time and usage frequency of the current offloading mode, and calculates the performance improvement effect.

[0474] By doing so, the improvement degree calculation unit 122 of the offloading server 1 reconstructs the new offloading mode into a more appropriate GPU logic according to the usage characteristics after the start of operation by executing the application code analysis unit 112, the parallel processing designation unit 114, the parallel processing mode creation unit 115, the performance measurement unit 116, and the execution file creation unit 117 based on the representative data, so as to achieve efficient resource utilization in a GPU with limited resources.

[0475] In the offloading server 1 of the present embodiment, it is characterized in that the request processing load analysis unit 120 obtains an improvement degree coefficient by (actual processing time in the case of only CPU processing) / (actual processing time in the case of GPU offloading) based on the test records in the assumed usage data before the start of operation, and takes the sum of the values obtained by multiplying the actual processing time by the improvement degree coefficient as the total processing time for comparison.

[0476] In this way, when the offloading server 1 selects an application with a higher load, for an application undergoing GPU offloading, it can multiply by the improvement factor to calculate the situation when it is not offloaded, and thus correct it to CPU-only processing for comparison. In the case of non-offloading, correction is performed by multiplying by the improvement factor, so the actual processing time can be calculated more precisely.

[0477] In the offloading server 1 of the present embodiment, it is characterized in that the request processing load analysis unit 120 acquires request data of an application with a higher load during a specified period, arranges the data sizes at fixed intervals and creates a frequency distribution, and the representative data selection unit 121 selects any one data from the actual request data corresponding to the mode of the frequency distribution as the representative data.

[0478] In this way, when the offloading server 1 selects representative data, the average of the data sizes sometimes differs greatly from the actually used data. By using the mode of the data sizes, more appropriate representative data can be selected.

[0479] In the offloading server 1 of the present embodiment, it is characterized in that the improvement degree calculation unit 122 measures the processing times of the current offloading mode and multiple new offloading modes, and calculates the performance improvement effect based on the commercial usage frequency according to (actual processing reduction time in the verification environment) × (commercial environment usage frequency).

[0480] In this way, the offloading server 1 calculates (actual processing reduction time in the verification environment) × (commercial environment usage frequency) in multiple new offloading modes, and calculates (actual processing reduction time in the verification environment) × (commercial environment usage frequency) in the current offloading mode. By dividing the former by the latter and using the two parameters of processing time and usage frequency, a more accurate performance improvement degree effect can be obtained.

[0481] In the offloading server 1 of the present embodiment, it is characterized in that it has: another server that processes OpenACC after GPU reconstruction; and a control unit 11 that, when performing GPU reconstruction, stops the OpenACC before reconstruction, newly starts the other server, and switches the routing to the other server for new requests.

[0482] In this way, the control unit 11 starts the reconstructed OpenACC in the other server for GPU reconstruction, so OpenACC can be stopped and started without an interruption time.

[0483] The present invention is an offloading program for causing a computer to function as the above-mentioned offloading server.

[0484] In this way, using an ordinary computer, the various functions of the above-mentioned offloading server 1 can be realized.

[0485] In addition, in the offloading server 1 of the present embodiment, the request processing load analysis unit 120 calculates the actual processing time and the total number of uses based on the usage records of each application program during a specified period. In this way, the offloading server 1 can analyze the request processing load of the data actually used by the user based on the usage records of each application program during the specified period.

[0486] In addition, among the processes described in the above embodiments, all or part of the processes described as being automatically performed can be performed manually, or all or part of the processes described as being performed manually can be automatically performed by known methods. Additionally, regarding the information including the processing order, control order, specific names, various data, and parameters shown in the above text and drawings, unless otherwise specified, they can be arbitrarily changed.

[0487] In addition, each structural element of each illustrated device is a functional concept, and it is not necessarily physically configured as shown. That is, the specific manner of dispersion and integration of each device is not limited to the illustration, and all or part of it can be dispersed and integrated functionally or physically in any unit according to various loads, usage conditions, etc.

[0488] In addition, regarding the above-described various structures, functions, processing units, processing methods, etc., part or all of them can also be implemented by hardware such as designing in an integrated circuit. In addition, the above-described various structures, functions, etc. can also be implemented by software for interpreting and executing programs that enable a processor to implement each function. Information such as the programs, tables, and files for implementing each function can be stored in a recording device such as a memory, hard disk, SSD (Solid State Drive), or in a recording medium such as an IC (Integrated Circuit) card, SD (Secure Digital) card, or optical disc.

[0489] In addition, in the present embodiment, as long as GPU processing can be offloaded, it can be arbitrary.

[0490] In addition, in the present embodiment, as a repetitive statement (loop statement), a for statement is exemplified, but it also includes while statements and do-while statements other than the for statement. Among them, the for statement that specifies the continuation condition of the loop is more suitable.

[0491] Description of Reference Numerals

[0492] 1: Offloading server

[0493] 11: Control unit

[0494] 12: Input / output unit

[0495] 13: Storage unit

[0496] 14: Verification machine (accelerator verification device)

[0497] 15: OpenIoT resources

[0498] 111: Application code specifying unit

[0499] 112: Application code analysis unit

[0500] 113: Data transfer specifying unit

[0501] 114: Parallel processing specifying unit

[0502] 114a: Unloading range extraction unit

[0503] 114b: Intermediate language file output unit

[0504] 115: Parallel processing mode creation unit

[0505] 116: Performance measurement unit

[0506] 116a: Binary file configuration unit

[0507] 117: Execution file creation unit

[0508] 118: Formal environment configuration unit

[0509] 119: Performance measurement test extraction execution unit

[0510] 120: Request processing load analysis unit (processing load analysis unit)

[0511] 121: Representative data selection unit

[0512] 122: Improvement degree calculation unit

[0513] 123: Refactoring suggestion unit

[0514] 124: User provision unit

[0515] 130: Application code

[0516] 131: Code pattern DB

[0517] 132: Device resource DB

[0518] 133: Test case DB

[0519] 134: Intermediate language file

[0520] 151: Various devices

[0521] 152: Device with CPU-GPU

[0522] 153: Device with CPU-FPGA

[0523] 154: Device with CPU

Claims

1. An offloading server that offloads specific processing of an application to a GPU (Graphics Processing Unit), characterized in that, comprising: a processing load analysis unit that analyzes a request processing load based on a current offloading mode of data used by a user; a representative data selection unit that determines an application program with a higher request processing load analyzed by the processing load analysis unit, and selects representative data from request data used when the application program is used; an improvement degree calculation unit that creates a new offloading mode based on the representative data selected by the representative data selection unit, compares the processing time and usage frequency of the created new offloading mode with the processing time and usage frequency of the current offloading mode, and calculates a performance improvement effect; and a reconstruction recommendation unit that recommends GPU reconstruction when the performance improvement effect is equal to or greater than a specified threshold.

2. The offloading server according to claim 1, wherein comprising: an application program code analysis unit that analyzes the source code of an application program; a data transfer specification unit that analyzes the reference relationship of variables used in a loop statement of the application program, and performs data transfer specification using a line that clearly specifies data transfer outside the loop for data that can be transferred outside the loop, where the line that clearly specifies data transfer outside the loop clearly specifies data transfer outside the loop; a parallel processing specification unit that determines loop statements of the application program, and for each of the determined loop statements, specifies a parallel processing specification statement in the GPU for compilation; a parallel processing mode creation unit that creates a parallel processing mode that excludes loop statements with compilation errors from offloading targets, and for loop statements without compilation errors, specifies whether to perform parallel processing; a performance measurement unit that compiles the application program of the parallel processing mode, configures it in an accelerator verification device, and performs performance measurement processing when offloading to the GPU; and an execution file creation unit that selects a parallel processing mode with the highest processing performance from multiple parallel processing modes based on the performance measurement result of the performance measurement processing, compiles the parallel processing mode with the highest processing performance, and creates an execution file, the improvement degree calculation unit determines a new offloading mode by executing the application program code analysis unit, the parallel processing specification unit, the parallel processing mode creation unit, the performance measurement unit, and the execution file creation unit based on the representative data selected by the representative data selection unit, compares the processing time and usage frequency of the determined new offloading mode with the processing time and usage frequency of the current offloading mode, and calculates a performance improvement effect.

3. The offloading server according to claim 1, wherein the processing load analysis unit obtains an improvement degree coefficient by (actual processing time in the case of only CPU processing) / (actual processing time in the case of GPU offloading) based on a test record in assumed usage data before the start of operation, and uses the sum of the values obtained by multiplying the actual processing time by the improvement degree coefficient for comparison.

4. The offloading server according to claim 1, wherein the processing load analysis unit obtains request data of an application program with a higher load during a specified period, arranges the data sizes at fixed intervals, and creates a frequency distribution. The representative data selection unit selects any one of the actual request data corresponding to the mode of the frequency distribution as the representative data.

5. The offloading server according to claim 1, wherein: The improvement degree calculation unit measures the processing times of the current offloading mode and the multiple new offloading modes, and calculates the performance improvement effect based on the commercial usage frequency according to (actual processing reduction time in the verification environment) × (commercial environment usage frequency).

6. The offloading server according to claim 1, characterized in that, It has: Another server that processes OpenACC (Open Accelerator) after GPU reconstruction; and A control unit that, when performing GPU reconstruction, stops the OpenACC before reconstruction, newly starts the other server, and switches the routing to the other server for new requests.

7. An offloading control method, which is an offloading control method for offloading specific processing of an application program to an offloading server of a GPU (Graphics Processing Unit), characterized in that: The offloading server performs: A processing load analysis step of analyzing the request processing load based on the current offloading mode of the data used by the user; A representative data selection step of determining the application program with the highest request processing load among the analyzed ones, and selecting representative data from the request data when the application program is used; An improvement degree calculation step of creating a new offloading mode based on the selected representative data, comparing the processing time and usage frequency of the created new offloading mode with the processing time and usage frequency of the current offloading mode, and calculating the performance improvement effect; And A reconstruction recommendation step of recommending GPU reconstruction when the performance improvement effect is equal to or higher than a specified threshold.

8. An offloading program for causing a computer to function as the offloading server according to any one of claims 1 to 6.

Images