OFFLOAD SERVER AND OFFLOAD CONTROL METHOD

Abstract

To automatically improve performance even when an environment in which multiple types of accelerators are mixed together is a migration target.SOLUTION: An offload server 1 comprises: a code analysis unit for analyzing whether a source code of a software program includes a loop statement that can be executed by a GPU or a many-core CPU; a processing pattern generation unit for generating a pattern for offloading the loop statement to the GPU or the many-core CPU; a performance measurement unit for mapping each loop statement to a gene pattern, deploying each pattern for offloading each loop statement to a verification environment, and measuring performance; and an executable file creation unit for setting a higher fitness for a pattern having a shorter processing time, selecting a pattern by elite selection of a genetic algorithm, and repeating the creation of a next generation pattern by processing crossover and mutation.SELECTED DRAWING: Figure 7

Description

[Technical field]

[0001] The present invention relates to an offload server that offloads software processing programs processed by a normal CPU (Central Processing Unit) so as to process them at high speed in an environment where accelerators such as GPUs (Graphics Processing Units), FPGAs (Field Programmable Gate Arrays) and many-core CPUs are mixed, and an offload control method. [Background technology]

[0002] In recent years, it has been said that Moore's Law, which states that the semiconductor integration density of CPUs doubles every 1.5 years, may be slowing down. In light of this situation, the use of devices such as FPGAs and GPUs, as well as CPUs with fewer cores, is increasing. For example, Microsoft is working to improve the efficiency of its search engines by using FPGAs. Amazon provides FPGAs and GPUs as cloud instances.

[0003] However, in order to properly utilize devices other than ordinary CPUs with few cores in a system, it is necessary to configure and write programs that take into account the characteristics of the device. This means that knowledge of OpenMP (registered trademark) (Open Multi-Processing), OpenCL (registered trademark) (Open Computing Language), and CUDA (registered trademark) is required, which presents a high skill barrier for most programmers. Note that from here on, the term (registered trademark) will be omitted for OpenMP (registered trademark), OpenCL (registered trademark), and CUDA (registered trademark).

[0004] It is expected that the number of systems that utilize accelerators other than few-core CPUs, such as GPUs, FPGAs, and many-core CPUs, will continue to increase in the future, but there are high technical barriers for programmers to make the most of them.Therefore, in order to remove such barriers and make full use of accelerators other than few-core CPUs, there is a demand for a platform that can adaptively convert and configure software in which a programmer writes processing logic to suit the environment in which it is deployed (FPGA, GPU, many-core CPU, etc.), and enable it to operate in a way that suits the environment.

[0005] Therefore, the inventor proposed in Non-Patent Document 1 an environment-adaptive software that automatically performs conversion and resource settings so that code written once can be used by GPUs, FPGAs, many-core CPUs, etc. present in the deployment environment, thereby enabling applications to run with high performance. The inventor has proposed a method of automatically offloading loop statements and function blocks of source code to an FPGA or GPU as elements of environment adaptive software in Non-Patent Documents 2, 3, and 4. [Prior art documents] [Non-patent literature]

[0006] [Non-Patent Document 1] Y. Yamato, H. Noguchi, M. Kataoka and T. Isoda, “Proposal of Environment Adaptive Software,” The 2nd International Conference on Control and Computer Vision (ICCCV 2019), pp.102-108, Jeju, June 2019. [Non-Patent Document 2] Y. Yamato, “Study of parallel processing area extraction and data transfer number reduction for automatic GPU offloading of IoT applications,” Journal of Intelligent Information Systems, Springer, DOI: 10.1007 / s10844-019-00575-8, Aug. 2019. [Non-Patent Document 3] Y. Yamato, “Proposal of Automatic FPGA Offloading for Applications Loop Statements,” The 7th Annual Conference on Engineering and Information Technology (ACEAIT 2020), pp.111-123, 2020. [Non-Patent Document 4] Y. Yamato, “Proposal of Automatic Offloading for Function Blocks of Applications,” The 8th IIAE International Conference on Industrial Application Engineering 2020 (ICIAE 2020), pp.4-11, Mar. 2020. Summary of the Invention [Problem to be solved by the invention]

[0007] In recent years, CUDA has become popular as an environment for running GPGPU (General Purpose GPU), which uses the parallel computing power of GPUs for purposes other than image processing. CUDA is an environment for GPGPUs developed by NVIDIA.

[0008] Furthermore, OpenCL has been released as a specification for handling heterogeneous devices such as FPGAs, many-core CPUs, and GPUs in the same way, and its development environment is also available. CUDA and OpenCL are programs that are extensions of the C language, such as explicitly describing the copying and release of memory data between kernels such as FPGAs and CPU hosts, and are therefore highly difficult to program.

[0009] Compared to CUDA and OpenCL, this technology allows for easier use of heterogeneous devices by specifying the location of parallel processing in the directive line, and the compiler creates executable files for GPUs, many-core CPUs, etc. according to the directive line. Specifications include OpenACC (registered trademark) (OpenACCelerators) and OpenMP, and compilers include the PGI compiler and gcc. Note that from here on, the word (registered trademark) will be omitted for OpenACC (registered trademark).

[0010] By using technical specifications such as CUDA, OpenCL, OpenACC, and OpenMP, it is possible to offload to FPGAs, GPUs, and many-core CPUs. However, even if accelerator processing itself is possible, there are still challenges in accelerating it.

[0011] For example, Intel compilers are compilers with automatic parallelization functions for multi-core CPUs. During automatic parallelization, these extract parts of loop statements in the code that can be processed in parallel and parallelize them. However, due to the effects of memory processing, etc., simply parallelizing parallelizable loop statements often does not result in the desired performance. Note that the loop statements referred to here are for / while / do-while in C / C++ source code.

[0012] Specifically, when accelerating with FPGAs or GPUs, OpenCL and CUDA engineers repeatedly tune their programs or use OpenACC compilers to search for and try appropriate parallel processing ranges. For this reason, it is difficult for programmers with poor technical skills to use FPGAs, GPUs, or many-core CPUs to speed up applications. In addition, when using automatic parallelization technology, it is necessary to carry out trial and error to search for parallel processing points.

[0013] Currently, offloading to heterogeneous devices is mainly done manually. The inventors have proposed the concept of environment-adaptive software and are considering automatic offloading. However, the offloading is only to a single GPU, and no consideration has been given to a migration environment in which various accelerators such as GPUs, FPGAs, and many-core CPUs are mixed.

[0014] Therefore, an object of the present invention is to provide an offload server and an offload control method that can automatically improve performance even in a migration destination environment in which multiple types of accelerators are mixed. [Means for solving the problem]

[0015] In order to solve the above-mentioned problems, the offload server of the present invention is characterized by comprising: a code analysis unit that analyzes whether the source code of a software program includes a loop statement that can be executed by a GPU, an FPGA, or a many-core CPU; a processing pattern creation unit that creates a pattern for offloading the loop statement to any of the GPU, the FPGA, and the many-core CPU; a performance measurement unit that maps each of the loop statements to a genetic pattern and deploys each of the patterns for offloading each of the loop statements in a verification environment to measure performance; and an executable file creation unit that sets a higher fitness for patterns with shorter processing times, selects a pattern by elite selection of a genetic algorithm, and repeatedly creates the next generation pattern by crossover and mutation processing. Other means will be described in the description of the embodiment of the invention. Effect of the Invention

[0016] According to the present invention, even if the destination environment contains a mixture of multiple types of accelerators, the performance can be automatically improved. Furthermore, if the destination node contains multiple types of accelerators, offloading to the multiple types of accelerators simultaneously can achieve faster offloading than when a single accelerator is used. [Brief description of the drawings]

[0017] [Figure 1] FIG. 2 is a functional block diagram illustrating a configuration example of an offload server. [Diagram 2] FIG. 2 is a hardware configuration diagram illustrating an example of a computer that realizes the functions of an offload server. [Figure 3A] 13 is a flowchart of a process for selecting a pattern to be offloaded to an accelerator. [Figure 3B] 13 is a flowchart of a process for selecting a pattern to be offloaded to an accelerator. [Figure 3C] 13 is a flowchart of a process for selecting a pattern to be offloaded to an accelerator. [Figure 4A] 13 is a flowchart of a process of offloading a function block to one node. [Figure 4B] 13 is a flowchart of a process of offloading a loop statement to one node. [Figure 5A] 13 is a flowchart of a process for selecting a pattern for offloading a functional block to an accelerator. [Figure 5B] 13 is a flowchart of a process for selecting a pattern for offloading a functional block to an accelerator. [Figure 6] FIG. 13 illustrates an automatic offload process of an offload server. [Figure 7] FIG. 13 is a diagram illustrating an image of a search for a control unit by a simple genetic algorithm of an offload server. [Figure 8] FIG. 13 is a diagram showing gene sequence mapping of a loop statement onto a many-core CPU. [Figure 9A] 13 is a flowchart of a process for selecting a pattern for offloading a loop statement to a many-core CPU. [Figure 9B] 13 is a flowchart of a process for selecting a pattern for offloading a loop statement to a many-core CPU. [Figure 10] FIG. 13 is a diagram showing gene sequence mapping of loop statements to a GPU. [Figure 11A] 13 is a flowchart of a process for selecting a pattern for offloading a loop statement to a GPU. [Figure 11B] 13 is a flowchart of a process for selecting a pattern for offloading a loop statement to a GPU. [Figure 12A] 11 is a flowchart of a process for selecting a pattern for offloading a loop statement to an FPGA. [Figure 12B] 11 is a flowchart of a process for selecting a pattern for offloading a loop statement to an FPGA. DETAILED DESCRIPTION OF THE PREFERRED EMBODIMENTS

[0018] Hereinafter, embodiments of the present invention will be described in detail with reference to the accompanying drawings. In order to materialize the concept of environment adaptive software, the inventor has proposed automatic GPU offloading of program loop statements, automatic FPGA offloading, and automatic offloading of program function blocks. Based on these elemental technology studies, the inventor defines the target of the destination environment, describes the basic idea to be followed even when the target is diversified, and further proposes automatic offloading to each destination environment and an automatic offloading method when the destination environment is diversified.

[0019] <Basic Concepts for Diverse Destination Environments> The various migration environments targeted in this embodiment are GPU, FPGA, and many-core CPU. GPU and FPGA have a long history as heterogeneous hardware different from CPU, and there are many cases of acceleration by manual offloading using CUDA or OpenCL, and the market is large. In addition, as for many-core CPUs, CPUs equipped with many cores of 16 cores or more have been released on the market at low prices of 1,000 to several thousand dollars in recent years, and there are cases of speeding up by parallelization using technical specifications such as OpenMP and manual tuning.

[0020] In cases where the destination environment is not simply a CPU, in order to automatically offload at high speed, we take an approach in which we measure performance on actual machines in a verification environment and combine it with techniques such as evolutionary computing to gradually find a faster offloading pattern.

[0021] This is the same as in the case of GPU offloading and the like proposed by the inventor so far. The reason is that performance is difficult to predict statically because it varies greatly not only depending on the code structure but also on the actual processing contents such as the specifications of the processing hardware, data size, loop count, etc., and therefore dynamic measurement is necessary. There are automatic parallelizing compilers on the market that find loop statements and parallelize them at the compile stage, but performance measurement is necessary because the parallelization of parallelizable loop statements alone often results in slow speeds.

[0022] The targets for offloading are loop statements and function blocks of the program. This is the same as the GPU and FPGA offloading discussed so far. As for loop statements, most of the time-consuming processing in a program is currently consumed by loops, so loop statements are considered as targets for offloading.

[0023] On the other hand, when speeding up a specific process, a function block often uses an algorithm suited to the process content and processing hardware, and therefore may be able to achieve a significant speed increase compared to parallel processing of individual loop statements, etc. In this embodiment, the speed is increased by replacing frequently used function blocks such as matrix multiplication and Fourier transform with processes (IP cores, CUDA libraries, etc.) implemented with algorithms suited to accelerators such as many-core CPUs, FPGAs, and GPUs.

[0024] FIG. 1 is a functional block diagram showing an example of the configuration of the offload server 1. As shown in FIG. The offload server 1 is a device that automatically offloads specific processing of an application to an accelerator such as a GPU, FPGA, or many-core CPU. As shown in FIG. 1, the offload server 1 includes a control unit 11, an input / output unit 12, a storage unit 13, and a verification machine 14 (accelerator verification device).

[0025] The input / output unit 12 is composed of a communication interface 950 (see Figure 2) for transmitting and receiving information between each device belonging to the cloud layer 2, network layer 3, and device layer 4, and an input / output interface 960 (see Figure 2) for transmitting and receiving information between input devices such as a touch panel and a keyboard, and output devices such as a monitor.

[0026] The storage unit 13 is composed of a hard disk, a flash memory, a RAM (Random Access Memory), etc. The storage unit 13 stores a test case database 131 and a code pattern database 133, and also temporarily stores programs (offload programs) for executing the functions of the control unit 11 and information necessary for the processing of the control unit 11, such as an intermediate language file 132.

[0027] The test case database 131 stores data on test items corresponding to the software to be verified. For example, for a database system such as MySQL, the data includes transaction tests such as TPC-C. The code pattern database 133 stores replaceable accelerator libraries or accelerator IP cores corresponding to library names. The verification machine 14 includes an accelerator such as a GPU, an FPGA, or a many-core CPU as a verification environment for the environment-adaptive software system.

[0028] The control unit 11 is an automatic offload function unit that is responsible for the overall control of the offload server 1. The control unit 11 is realized, for example, by the CPU 910 in FIG.

[0029] The control unit 11 includes a code designation unit 111, a code analysis unit 112, a processing designation unit 114, a processing pattern creation unit 115, a performance measurement unit 116, an executable file creation unit 117, a production environment deployment unit 118, a performance measurement test extraction execution unit 119, and a user provision unit 120.

[0030] Code Designation 111 The code designation unit 111 designates the input source code. Specifically, the code designation unit 111 passes the source code described in the received file to the code analysis unit 112.

[0031] Code Analysis Section 112 The code analysis unit 112 analyzes the source code of the software program and identifies loop statements such as for / do-while / while, and function blocks such as FFT libraries, matrix operations, random number generation, etc. These loop statements and function blocks can be executed by an accelerator.

[0032] <<Processing designation section 114>> The process specification unit 114 replaces each function block with a process to be offloaded to an accelerator and compiles it, and specifies a parallel process specification statement to be offloaded to an accelerator for each loop statement and compiles it. The process specification unit 114 includes an offload range extraction unit 114a and an intermediate language file output unit 114b.

[0033] The offload range extraction unit 114a identifies processes that can be offloaded to a many-core CPU, GPU, or FPGA, such as loop statements such as for / do-while / while, and functional blocks such as FFT, matrix calculation, and random number generation, and extracts the offload range. The intermediate language file output unit 114b outputs an intermediate language file 132 of the extracted offloadable processes. The intermediate language extraction is not completed once, but is repeated to search for an appropriate offload range and to perform trial and error execution for optimization.

[0034] <<Processing pattern creation unit 115>> The processing pattern creation unit 115 creates a pattern for offloading the function blocks and loop statements to the accelerator based on the extracted offload range.

[0035] When using a GPU as an accelerator for a loop statement, the processing pattern creation unit 115 creates a parallel processing pattern that excludes loop statements (repeated statements) that cause a compilation error from being offloaded, and specifies whether or not to process loop statements that do not cause a compilation error in parallel.

[0036] 《Performance measurement section 116》 The performance measurement unit 116 compiles source code of the processing pattern, places it in the verification machine 14, and executes the processing for performance measurement when offloaded to the accelerator. The performance measurement unit 116 includes a binary file placement unit 116a. The binary file placement unit 116a deploys (places) an executable file derived from an intermediate language in the verification machine 14 including a many-core CPU, a GPU, and an FPGA.

[0037] The performance measurement unit 116 executes the placed binary file, measures the performance when offloaded, and returns the performance measurement result to the processing pattern creation unit 115. In this case, the processing pattern creation unit 115 creates another offload pattern. The performance measurement unit 116 attempts to measure the performance of the created other offload pattern.

[0038] Executable File Creation Section 117 The executable file creation unit 117 selects an offload pattern with the highest processing performance from among a plurality of offload patterns based on the results of performance measurements repeated a predetermined number of times, and compiles the offload pattern with the highest processing performance to create an executable file.

[0039] 《Production Environment Configuration Section 118》 The production environment deployment unit 118 deploys the created executable file in the production environment for the user ("Deployment of final binary file in production environment"). The production environment deployment unit 118 determines a pattern that specifies the final offload area, and deploys it to the production environment for the user.

[0040] 《Performance Measurement Test Extraction Execution Unit 119》 After placing the executable file, the performance measurement test extraction execution unit 119 extracts performance test items from the test case database 131 and executes the performance tests ("Placing the final binary file in the production environment"). After arranging the executable file, the performance measurement test extraction and execution unit 119 extracts performance test items from the test case database 131 and automatically executes the extracted performance tests in order to show the performance to the user.

[0041] User provision unit 120 The user providing unit 120 presents the user with information such as price and performance based on the performance test results ("Providing information such as price and performance to user"). Test case database 131 stores data on test items corresponding to the software to be verified. The user providing unit 120 presents the user with data such as performance data obtained by implementing the test items stored in test case database 131, and data such as price information on resources (such as virtual machines and GPUs) on which the user software is deployed. The user decides whether to start using the service for which charges are required based on the presented information such as price and performance.

[0042] [Application of genetic algorithms] The offload server 1 can use a genetic algorithm for optimizing the offload. The offload server 1 is configured as follows when using a genetic algorithm. That is, the processing specification unit 114 sets the number of loop statements (repetitive statements) as the gene length based on a genetic algorithm. The processing pattern creation unit 115 maps whether or not accelerator processing is possible to the gene pattern by setting either 1 or 0 when accelerator processing is performed, and the other 0 or 1 when accelerator processing is not performed.

[0043] When a genetic algorithm is used, the processing pattern creation unit 115 prepares a specified number of gene patterns in which each gene value is randomly created to be 1 or 0. The performance measurement unit 116 compiles source code specifying a parallel processing specification statement in the accelerator according to each individual, and places the compiled code in the verification machine 14. The performance measurement unit 116 executes a performance measurement process in the verification machine 14.

[0044] Here, when using a genetic algorithm, if a gene with the same parallel processing pattern as before occurs in an intermediate generation, the performance measurement unit 116 does not compile the source code corresponding to the parallel processing pattern and does not perform performance measurement, and uses the same value as the performance measurement value. Also, for source code that causes a compilation error and source code for which performance measurement does not end within a predetermined time, the performance measurement unit 116 treats the source code as a timeout and sets the performance measurement value to a predetermined time (long time).

[0045] The executable file creation unit 117 measures the performance of all the individuals, and evaluates them so that the shorter the processing time, the higher the fitness of the individual. From all the individuals, the executable file creation unit 117 selects those with a fitness higher than a predetermined value (for example, the top n% of the total number, or the top m of the total number: n and m are natural numbers) as high-performance individuals, and performs crossover and mutation processes on the selected individuals to create the next generation of individuals. After completing processing for the specified number of generations, the executable file creation unit 117 selects the parallel processing pattern with the highest performance as a solution.

[0046] FIG. 3 is a hardware configuration diagram illustrating an example of a computer that realizes the functions of the offload server. The computer 900 has a CPU 910, a RAM 920, a ROM 930, a HDD 940, a communication interface 950, an input / output interface 960, and a media interface 970.

[0047] The CPU 910 operates and controls each unit based on a program stored in the ROM 930 or the HDD 940. The ROM 930 stores a boot program executed by the CPU 910 when the computer 900 is started, programs that depend on the hardware of the computer 900, and the like.

[0048] The HDD 940 stores programs executed by the CPU 910, 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 transmits data generated by the CPU 910 to other devices via the communication network 80.

[0049] 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 acquires data from the input devices via the input / output interface 960. The CPU 910 also outputs generated data to the output devices via the input / output interface 960.

[0050] The media 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 onto the RAM 920 via the media 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 magneto-optical recording medium such as an MO (Magneto Optical disk), a tape medium, a magnetic recording medium, a semiconductor memory, or the like.

[0051] For example, when the computer 900 functions as the offload server 1 according to this embodiment, a CPU 910 of the computer 900 executes a program loaded onto a RAM 920 to realize the functions of each unit of the offload server 1. In addition, data for each unit of the offload server 1 is stored in an HDD 940. The CPU 910 of the computer 900 reads and executes these programs from a recording medium 980, but as another example, these programs may be acquired from another device via a communication network 80.

[0052] 3A, 3B, and 3C are flowcharts of a process for selecting a pattern to be offloaded to an accelerator. Offload server 1 considers automatic offloading for three destination environments: GPU, FPGA, and many-core CPU, using two methods: function blocks and loop statements.

[0053] When the C / C++ source code is input (S200), the control unit 11 executes a series of offload pattern selection processes. The control unit 11 calls the selection processes of Fig. 5A and Fig. 5B to check the offload of the functional blocks to the many-core CPU (S201), and judges whether the desired performance has been obtained (S202). If the desired performance has been obtained (Yes), the control unit 11 selects the highest performance pattern from all measurements (S208), and ends the process of this figure. If the desired performance has not been obtained (No), the control unit 11 proceeds to step S203.

[0054] In step S203, the control unit 11 calls the selection process of Fig. 5A and Fig. 5B to check the offloading of the functional block to the GPU. Then, the control unit 11 judges whether the desired performance has been obtained (S204). If the desired performance has been obtained (Yes), the control unit 11 selects the highest performance pattern from all measurements (S208) and ends the process of this figure. If the desired performance has not been obtained (No), the control unit 11 proceeds to step S205.

[0055] In step S205, the control unit 11 calls the selection process of Figures 5A and 5B to check the offloading of the functional blocks to the FPGA. The control unit 11 then determines whether the desired performance has been obtained (S206). If the desired performance has been obtained (Yes), the control unit 11 selects the highest performance pattern from all measurements (S208) and ends the process of this figure. If the desired performance has not been obtained (No), the control unit 11 removes the part of the code that could have been offloaded the functional blocks and proceeds to step S209 of Figure 3B. By performing the processes in steps S201 to S208, the performance of the functional blocks can be automatically improved even in a migration destination environment where multiple types of accelerators are mixed. Furthermore, the control unit 11 can search for a pattern that can efficiently increase the speed by trying many-core CPU, GPU, and FPGA in that order.

[0056] In step S209, the control unit 11 calls the selection process of Fig. 9A and Fig. 9B to check offloading of loop statements to the many-core CPU. The control unit 11 then determines whether the desired performance has been obtained (S210). If the desired performance has been obtained (Yes), the control unit 11 selects the highest performance pattern from all measurements (S208) and ends the process of this figure. If the desired performance has not been obtained (No), the control unit 11 proceeds to step S211.

[0057] In step S211, the control unit 11 calls the selection process of Fig. 11A and Fig. 11B to check offloading of loop statements to the GPU. The control unit 11 then determines whether the desired performance has been obtained (S212). If the desired performance has been obtained (Yes), the control unit 11 selects the highest performance pattern from all measurements (S208) and ends the process of this figure. If the desired performance has not been obtained (No), the control unit 11 proceeds to step S213.

[0058] In step S213, the control unit 11 calls the selection process of Fig. 12A and Fig. 12B to check offloading of loop statements to FPGA. Then, the control unit 11 judges whether the desired performance has been obtained (S214). If the desired performance has been obtained (Yes), the control unit 11 selects the highest performance pattern from all measurements (S208) and ends the process of this figure. If the desired performance has not been obtained (No), the control unit 11 selects the highest performance pattern from the six verifications for each migration destination (S215).

[0059] By the processing of steps S209 to S215, even if the migration destination environment is a mixture of multiple types of accelerators, the performance of the loop statement can be automatically improved. Furthermore, the control unit 11 can search for a pattern that can be efficiently accelerated by trying many-core CPU, GPU, and FPGA in that order. The control unit can search for a pattern that can be efficiently accelerated by trying a loop statement after trying a function block as a target for offloading to an accelerator.

[0060] Then, the control unit 11 judges whether or not there are multiple candidates for offloading to one node (S216). If there is no candidate that can offload multiple nodes to one node (No), the control unit 11 ends the process in this figure. If there is a candidate that can offload multiple nodes to one node (Yes), the control unit 11 proceeds to step S217 in FIG. 3C.

[0061] In step S217 in Fig. 3C, the control unit 11 determines whether or not there is a candidate to which multiple functional blocks can be offloaded. If there is a candidate to which multiple functional blocks can be offloaded (Yes), the control unit 11 proceeds to step S218. If there is no candidate to which multiple functional blocks can be offloaded (No), the control unit 11 proceeds to step S223.

[0062] In step S218, the control unit 11 creates a combination pattern of offloading a plurality of functional blocks. The control unit 11 measures the performance of this combination pattern (S219). Then, the control unit 11 determines whether or not the combination of offloading the functional blocks can increase the speed (S220).

[0063] In step S220, if the combined offload of the functional blocks is faster, the control unit 11 removes the combined functional blocks from the source code (S221) and proceeds to the processing of step S223. If the uncombined offload of the functional blocks is faster, the control unit 11 removes one functional block that can be faster from the offload targets of the source code (S222) and proceeds to the processing of step S223.

[0064] In step S223, the control unit 11 determines whether there is a loop statement offload pattern that can be accelerated and whether it is a single accelerator or multiple accelerators. If there is no loop statement offload pattern that can be accelerated, the control unit 11 proceeds to step S227, selects the pattern with the highest performance from all measurements, and ends this selection process. Here, in an example where loop statement offload pattern B offloads 30 for statements to a GPU and loop statement offload pattern C offloads 50 for statements to a many-core CPU, a case where only B can be accelerated is described as a single accelerator, and a case where both B and C can be accelerated is described as multiple accelerators.

[0065] If the result of the determination in step S223 is that the offloading pattern of the loop statement of a single accelerator can be accelerated, the control unit 11 creates a combination of the offloading of the removed function block and the offloading of the loop statement of the single accelerator (S224).Then, the control unit 11 measures the performance of the combination pattern (S226), selects the pattern with the highest performance from all the measurements (S227), and ends this selection process.

[0066] If the determination result in step S223 is that the offloading pattern of loop statements of multiple accelerators can be accelerated, the control unit 11 creates a combination of the offloading of the removed function block and the offloading of loop statements of multiple accelerators (S225). For example, if the offloading of function blocks A to the GPU, the offloading pattern of loop statements B to the GPU, and the offloading pattern of loop statements C to the many-core CPU can be accelerated, four combinations, A+B+C, A+B, A+C, and B+C, are possible. Then, the control unit 11 measures the performance of the combination patterns (S226), selects the pattern with the best performance from all the measurements (S227), and ends this selection process. By performing the processes in steps S217 to S227, if the migration destination has multiple types of accelerators in the same node, offloading to the multiple types of accelerators simultaneously can be achieved, thereby achieving faster offloading than when a single accelerator is used.

[0067] <<Verification order of offload attempts>> Consider automated offloading across multiple destinations. Regarding performance, even if the processing is the same, it varies greatly depending on the processing hardware specifications and processing contents (data size, loop count, etc.), so it is necessary to measure it on the actual machine in the verification environment. For the three destination environment GPU, FPGA, and many-core CPU, there are two offloading methods, loop statements and function blocks, so a total of six offloading methods are possible: 3 x 2.

[0068] First, when comparing the offloading of function blocks with the offloading of loop statements, the offloading of function blocks, which offloads according to the processing content including the algorithm, can achieve faster processing speeds.

[0069] Next, when virtualization is not performed, the main price range for GPUs, FPGAs, and many-core CPUs is that many-core CPUs are more expensive than GPUs, and FPGAs are even more expensive than those. Also, the verification time for one pattern is roughly the same for many-core CPUs and GPUs, but longer for FPGAs. FPGAs currently require several hours to set up the circuit, and performance measurement takes time.

[0070] Considering these current circumstances, we propose the following order for verifying the six offloads: offload of function blocks for many-core CPUs, offload of function blocks for GPUs, offload of function blocks for FPGAs, offload of loop statements for many-core CPUs, offload of loop statements for GPUs, and offload of loop statements for FPGAs. We will try offloading in this order to search for a pattern that provides high performance.

[0071] The target code may be different between the first three and the latter three. Specifically, if the first three are able to offload a function block, the latter three loop statement offloads are attempted on the code excluding the part of the function block that was able to be offloaded.

[0072] In addition, the user can specify target performance and price in the offload trial, and if six offload patterns that are sufficiently fast and low-cost within the range specified by the user are found in the first half of the trial, subsequent trials do not need to be performed. The reason for this is that although function block offloading can offload fewer targets compared to loop statement offloading, when offloading is possible, it can achieve higher speeds.

[0073] In addition, when automatically offloading, it is desirable to search for a fast pattern as quickly and inexpensively as possible. Therefore, FPGA, which requires a long verification time, is the last to be verified, and if a pattern that satisfies the user requirements is found at an earlier stage, the offload pattern for FPGA is not verified.

[0074] There is no big difference between GPUs and many-core CPUs in terms of price or verification time, but the difference between many-core CPUs and normal CPUs is smaller than that of GPUs, which are different devices with separate memory spaces. Therefore, the verification order will be many-core CPUs first, and if a pattern that satisfies user requirements is found with many-core CPUs, offload patterns for GPUs will not be verified. Since the CPU and GPU are different devices, the calculation results may differ due to differences in rounding errors, etc., even if the offload is successful.

[0075] 《Creating high-performance patterns by combining multiple offloads》 Creating a solution pattern also requires trial and error. If only one of the six offload trials is faster than the normal CPU, the executable file for that pattern can be deployed in a production environment.

[0076] When placing executable files, it is acceptable to select not only the pattern with the highest performance, but also the pattern with the highest cost performance depending on the price of the device to be used. If multiple patterns can be made faster during testing, the solution pattern must be considered.

[0077] When offloading, it is easy to increase the speed by offloading to a GPU, FPGA, or many-core CPU on the same node, but when offloading to a different node, technical specifications such as MPI are required, and it is difficult to increase the speed due to communication overhead. Also, while servers are generally configured with many-core CPUs and GPUs, or many-core CPUs and FPGAs, configurations with many-core CPUs, GPUs, and FPGAs are rare.

[0078] If it is possible to accelerate using multiple accelerators, the first condition is that it can be offloaded to the same node. Many-core CPUs and GPUs on the same node are candidates, but GPUs and FPGAs on different nodes are not candidates.

[0079] First, let us consider function block offloading. The following rules apply to function block offloading. (a) When the same functional block can be offloaded to multiple types of accelerators, the offload destination with the highest performance improvement effect is selected. (b) If different functional blocks can be offloaded to different types of accelerators and cannot be offloaded to the same node, the functional block with the greatest performance improvement effect is selected as the offload target. (c) If different functional blocks can be offloaded to different types of accelerators and can be offloaded to the same node, those functional blocks are targeted for offloading.

[0080] As an example of rule (a), consider the case where FFT processing that takes 10 seconds on a single-core CPU takes 2 seconds when offloaded to a GPU, 1 second when offloaded to an FPGA, and 5 seconds when offloaded to a many-core CPU. In this case, it is recommended to select FPGA as the accelerator to offload to.

[0081] As an example of rule (b), consider the case where random number processing that would take 20 seconds on a single-core CPU takes 2 seconds when offloaded to a GPU, FFT processing that would take 100 seconds on a single-core CPU takes 1 second when offloaded to an FPGA, and matrix calculations that would take 50 seconds on a single-core CPU take 10 seconds when offloaded to a many-core CPU. In this case, it is advisable to select the pattern of offloading the most effective FFT processing to the FPGA.

[0082] Regarding rule (c), different functional blocks are offloaded to multiple accelerators on the same node to verify performance, and if the result is faster than single-unit offloading, it is considered the solution. As an example of rule (c), consider the case where random number processing that would take 20 seconds on a single-core CPU takes 20 seconds when offloaded to a GPU, FFT processing that would take 10 seconds on a single-core CPU takes 1 second when offloaded to an FPGA, and matrix calculations that would take 50 seconds on a single-core CPU take 10 seconds when offloaded to a many-core CPU. Furthermore, assume that there is a node equipped with both a GPU and a many-core CPU. In this case, it is recommended to select a pattern in which random number processing is offloaded to the GPU and matrix calculations are offloaded to the many-core CPU on the same node.

[0083] When offloading to multiple types of accelerators on the same node, simultaneous offloading to multiple types of accelerators does not necessarily ensure faster speeds. Therefore, the offload server 1 simultaneously offloads to multiple types of accelerators to verify performance and confirms that the speed is faster than when offloaded to a single accelerator. In this way, it is possible to search for patterns that can efficiently speed up functional blocks in an environment where multiple types of accelerators are mixed.

[0084] Note that function blocks that are frequently used processes may be offloaded with a high priority. In that case, the sample tests for measuring performance in the six preceding offloading verifications should be sample tests that mainly perform the processes with high priority.

[0085] FIG. 4A is a flowchart of a process for offloading a function block to one node. The control unit 11 of the offload server 1 determines whether the same functional block can be offloaded to multiple types of accelerators (S50). If the same functional block can be offloaded to multiple types of accelerators (Yes), the control unit 11 selects the offload destination with the highest performance improvement effect (S51) and ends the processing of Fig. 4A. If the same functional block cannot be offloaded to multiple types of accelerators (No), the control unit 11 proceeds to step S52.

[0086] In step S52, the control unit 11 determines whether different functional blocks can be offloaded to different types of accelerators. If different functional blocks can be offloaded to different types of accelerators (Yes), the control unit 11 sets these functional blocks as offload targets (S55) and ends the process of FIG. 4A. If different functional blocks cannot be offloaded to different types of accelerators (No), the control unit 11 sets the functional block with the greater performance improvement effect as the offload target (S54), and ends the process of FIG. 4A.

[0087] Next, offloading of loop statements will be considered. When function block offloading is possible, offloading of loop statements is performed on the code excluding the offloadable function block portion. If the speedup achieved by offloading loop statements is to offload to a node other than the offload destination of the function block offloading, and the performance improvement effect is small, offloading of loop statements is not performed.

[0088] For example, if an FFT process that would take 100 seconds on a single-core CPU can be offloaded and completed in 1 second on an FPGA, and it is found that the speedup can be reduced from 20 seconds to 5 seconds by offloading the remaining loop statements to a GPU, the performance improvement effect is small, so offloading to a GPU on a different node will not be performed.

[0089] The offloading of loop statements is subject to the following rules. (d) If the same loop statement can be offloaded to multiple types of accelerators, select the offload destination that provides the greatest performance improvement.

[0090] (e) If different loop statements can be offloaded to different types of accelerators but cannot be offloaded to the same node, the loop statements that will produce the greatest performance improvement are selected as the offload target.

[0091] (f) If different loop statements can be offloaded to different types of accelerators and can be offloaded to the same node, the loop statements with the highest performance improvement effect are offloaded first, and the remaining loop statements that have not been offloaded are offloaded to an offload destination with a lower performance improvement effect.

[0092] As an example of rule (d), consider the case where an ABCDE loop statement that takes 20 seconds on a single-core CPU takes 2 seconds when offloaded to a GPU, 1 second when offloaded to an FPGA, and 5 seconds when offloaded to a many-core CPU. In this case, it is better to select the pattern of offloading the ABCDE loop statement to the FPGA.

[0093] As an example of rule (e), consider the case where an ABC loop statement that takes 50 seconds on a single-core CPU takes 2 seconds when offloaded to a GPU, a BCD loop statement that takes 100 seconds on a single-core CPU takes 1 second when offloaded to an FPGA, and a CDE loop statement that takes 50 seconds on a single-core CPU takes 10 seconds when offloaded to a many-core CPU. In this case, it is better to select the pattern of offloading the BCD loop statement to the FPGA.

[0094] Regarding rule (f), different loops are offloaded to multiple types of accelerators on the same node to verify performance, and if it is faster than a single accelerator, that pattern is selected as the solution. However, restrictions are imposed so that processing for multiple loop statements defined in the same file is not split into different accelerators. This is to avoid an increase in complexity due to splitting accelerators for multiple loop statements in the same file that are closely related.

[0095] As an example of rule (f), assume that an ABC loop statement that would take 50 seconds on a single-core CPU takes 2 seconds when offloaded to a GPU, a BCD loop statement that would take 10 seconds on a single-core CPU takes 1 second when offloaded to an FPGA, and a CDE loop statement that would take 50 seconds on a single-core CPU takes 10 seconds when offloaded to a many-core CPU, and there is a node equipped with both a GPU and a many-core CPU. In this case, try to offload the ABC loop statement to the GPU and the DE loop statement to the many-core CPU on the same node.

[0096] FIG. 4B is a flowchart of a process for offloading a loop statement to one node. The control unit 11 of the offload server 1 determines whether the same loop statement can be offloaded to multiple types of accelerators (S60). If the same loop statement can be offloaded to multiple types of accelerators (Yes), the control unit 11 selects the offload destination with the highest performance improvement effect (S61) and ends the processing of Fig. 4B. If the same loop statement cannot be offloaded to multiple types of accelerators (No), the control unit 11 proceeds to step S62.

[0097] In step S62, the control unit 11 determines whether different loop statements can be offloaded to different types of accelerators. If different loop statements can be offloaded to different types of accelerators (Yes), the control unit 11 sets these loop statements as targets for offloading (S65) and ends the process of FIG. 4B. If different loop statements cannot be offloaded to different types of accelerators (No), the control unit 11 selects the loop statement with the greater performance improvement effect as the offload target (S64) and ends the processing of FIG. 4B.

[0098] When offloading to multiple types of accelerators on the same node, there is no guarantee that the processing will be faster, so it is a good idea to actually offload and verify the performance to confirm that the processing is faster than when offloading to a single accelerator. In this way, it is possible to search for patterns that can efficiently speed up loop statements in an environment where multiple types of accelerators are mixed.

[0099] In this way, the best-performing pattern is created based on the pattern that provides the highest speed among the six performance measurements. Although it is possible to offload to multiple accelerators for the same node, when combining multiple offloading patterns, multiple measurements are required in order to create the best-performing pattern by measuring whether the combination provides better performance than a single pattern.

[0100] 5A and 5B are flowcharts of a process for selecting a pattern for offloading a functional block to an accelerator. This selection process uses the method described in Non-Patent Document 4. First, the code analysis unit 112 performs an offload analysis of the C / C++ source code (S30). Specifically, the code analysis unit 112 uses a syntax analysis tool such as Clang to analyze library calls and function processing included in the code.

[0101] The offload range extraction unit 114a detects the external library call of this source code (S31), and obtains a replaceable accelerator library or an accelerator IP core from the code pattern database 133 using the library name as a key (S32). Here, the accelerator is any one of a many-core CPU, a GPU, and an FPGA. Specifically, the offload range extraction unit 114a obtains executable files, OpenCL, etc. that can be accelerated from records registered in the code pattern database 133 using the called library as a key.

[0102] Here, the offload range extraction unit 114a obtains replaceable accelerator libraries or accelerator IP cores by comparing the recognized external library calls with the code pattern database 133 using the called library name as a key. Here, the accelerator library is an executable file for many-core CPUs and GPUs. The accelerator IP core is OpenCL, etc.

[0103] In addition, when searching for libraries or IP cores that can speed up a functional block, similarity detection tools are used in addition to name matches of libraries, etc. Similarity detection tools are tools used in software engineering, such as Deckard, to detect copied code and code that has been modified after copying. Similarity detection tools look at abstract syntax tree similarities, etc., to determine the similarity of code, and can detect processes that are written similarly when calculated by a CPU, or when someone else's reference code has been copied and modified.

[0104] The processing pattern creation unit 115 replaces the processing description of the source code to be replaced with a processing description of a library for many-core CPUs or GPUs, or an IP core for FPGAs, in accordance with the accelerator to be replaced (S33). Furthermore, the processing pattern creation unit 115 offloads the processing description of the replaced accelerator library or accelerator IP core to the accelerator as a functional block to be offloaded (S34). At this time, since it is not known whether the processing that can be offloaded will actually lead to an increase in speed or whether the cost-effectiveness is sufficient, an offload pattern that will result in a higher speed is automatically extracted by trying whether to offload or not through performance measurement in a verification environment.

[0105] The processing pattern creation unit 115 writes interface processing for calling the accelerator library or the accelerator IP core (S35). This interface processing is, for example, an OpenCL API for FPGA. Next, the executable file creation unit 117 compiles the created pattern (S36). The performance measurement unit 116 measures the performance of the created pattern in the verification environment (S37). This is the first performance measurement.

[0106] After the first performance measurement, the executable file creation unit 117 creates a combination pattern for the pattern that was able to increase the speed during the first measurement (S38), and compiles the created combination pattern (S39). Then, the performance measurement unit 116 measures the performance of the created combination pattern in the verification environment (S40). This is the second performance measurement.

[0107] After the second performance measurement, the production environment deployment unit 118 selects the pattern with the highest performance from the first and second measurements (S41), and ends the processing of this flow.

[0108] As evaluated in Non-Patent Document 4, offloading of tuned function blocks for accelerators that process function blocks containing multiple loop statements results in a higher degree of speedup than offloading of individual loop statements.

[0109] "implementation" The offload server 1 grasps the program structure, such as the libraries being called and the classes and structures that are defined, through code analysis.

[0110] Next, the offload server 1 detects a GPU library, an FPGA IP core, etc. that can speed up the called library. The offload server 1 uses the called library as a key to obtain executable files and OpenCL that can be accelerated from records registered in the code pattern database 133. When a replacement function that can be accelerated is detected, the offload server 1 creates the executable file.

[0111] In the case of libraries for many-core CPUs or GPUs, the original part is deleted and a replacement is written to call the replacement library (CUDA library, etc.). In the case of FPGA IP cores, the original part of the acquired OpenCL code is deleted from the host code and then replaced in the kernel code. After completing the replacement writing, compile with gcc for many-core CPUs, PGI compiler for GPUs, and Intel Acceleration Stack for FPGAs.

[0112] Although the above describes the case of library invocation, parallel processing is also performed when similarity detection is used. Deckard is used to detect the similarity between the detected partial code such as classes and structures and the comparison code registered in the DB, and functional blocks with similarity exceeding a threshold are found, along with corresponding GPU libraries and FPGA IP cores. In particular, when the interfaces of the source code and the library or IP core to be replaced, such as arguments, return values, and types, are different, the offload server 1 confirms with the user who requested the offload whether it is OK to change the interface to match the destination library or IP core, and creates the executable file after confirmation.

[0113] Here, the offload server 1 creates an executable file that can measure the performance on the GPU or FPGA in the verification environment. In function block offloading, the offload server 1 measures the performance of each replacement function block and checks whether it can be accelerated by offloading it. If there is only one pattern that is faster than the normal CPU among the six offloading trials, that pattern is selected, but if there are multiple patterns, a solution pattern is created using the following logic.

[0114] [Automatic offload operation] The offload server 1 of this embodiment is an example in which a technology for automatically offloading user application logic to an accelerator is applied as a component technology of an environment adaptive software system.

[0115] FIG. 6 is a diagram showing the automatic offload process of the offload server 1. As shown in FIG. 6, the offload server 1 is applied to the elemental technology of an environment adaptive software system. The offload server 1 includes a control unit 11, a test case database 131, an intermediate language file 132, and a verification machine 14. The offload server 1 acquires the source code 130 that is used by the user.

[0116] A user uses a resource 15. The resource 15 is, for example, a device 151, a device 152 having a CPU and a GPU, a device 153 having a CPU and an FPGA, and a device 154 having a CPU. The offload server 1 automatically offloads function processing to an accelerator of the device 152 having a CPU and a GPU or the device 153 having a CPU and an FPGA.

[0117] The operation of each unit will be described below with reference to the step numbers in FIG. Step S11: Specify source code In step S11, the code designation unit 111 (see FIG. 1) identifies a processing function (image analysis, etc.) of the service provided to the user. Specifically, the code designation unit 111 designates the input source code.

[0118] Step S12: Source code analysis In step S12, the code analysis unit 112 (see FIG. 1) analyzes the source code of the processing function and grasps the structure of loop statements, FFT library calls, and the like.

[0119] Step S13: Extracting the off-road range In step S13, the process specification unit 114 (see FIG. 1) identifies loop statements included in the source code of the application, and compiles each loop statement by specifying a parallel processing specification statement for the accelerator. Specifically, the offload range extraction unit 114a (see FIG. 1) extracts the range of loop statements, FFT, etc. that can be offloaded to the GPU / FPGA.

[0120] Step S14: Output of intermediate language file In step S14, the process designation unit 114 (see FIG. 1) causes the intermediate language file output unit 114b to output the intermediate language file 132. The output of the intermediate language is not completed once, but is repeated to perform trial and error execution to search for an appropriate offload area and to optimize the execution.

[0121] Step S15: Processing in the event of a compilation error In step S15, the processing pattern creation unit 115 (see FIG. 1) creates a parallel processing pattern that excludes loop statements that cause a compilation error from being offloaded, and specifies whether or not to perform parallel processing for loop statements that do not cause a compilation error.

[0122] Step S21: Placing executable files In step S21, the executable file placement unit 116a (see FIG. 1) deploys an executable file derived from the intermediate language to the verification machine 14 including a many-core CPU, a GPU, and an FPGA.

[0123] Step S22: Performance measurement In step S22, the performance measurement unit 116 (see FIG. 1) executes the arranged file and measures the performance when offloaded. In order to make the offloading area more appropriate, the performance measurement result is returned to the processing pattern creation unit 115, which creates another offloading pattern. Then, the performance measurement unit 116 measures the performance of the other offloading pattern (see symbol e in FIG. 6).

[0124] As indicated by an arrow e in FIG. 6, the control unit 11 repeatedly executes the processes from step S12 to step S22.

[0125] The automatic offload function of the control unit 11 can be summarized as follows. That is, the processing specification unit 114 identifies loop statements in the source code, specifies a parallel processing specification statement for each loop statement, and compiles the loop statements. The processing pattern creation unit 115 creates a parallel processing pattern that excludes loop statements that generate compilation errors from being offloaded, and specifies whether or not to perform parallel processing for loop statements that do not generate compilation errors. The performance measurement unit 116 then compiles the source code of the corresponding parallel processing pattern, places it on the verification machine 14, and executes the performance measurement process on the verification machine 14. The executable file creation unit 117 selects a pattern with the highest processing performance from multiple parallel processing patterns based on the performance measurement results repeated a predetermined number of times, and compiles the selected pattern to create an executable file.

[0126] Step S23: Deploy to production environment In step S23, the production environment deployment unit 118 determines a pattern that specifies the final offload area, and deploys it to the production environment for the user.

[0127] Step S24: Performance measurement test and execution In step S24, after arranging the executable file, the performance measurement test extraction execution unit 119 extracts performance test items from the test case database 131 and automatically executes the extracted performance tests in order to show the performance to the user.

[0128] Step S25: Providing price, performance, etc. to users In step S25, the user providing unit 120 presents information such as price, performance, etc. based on the performance test results to the user. The user decides whether to start using the service for which a fee is charged based on the presented information such as price, performance, etc.

[0129] The above steps S11 to S25 are assumed to be performed in the background of the user's use of the service, for example, during the first day of temporary use.

[0130] As described above, when applied to the elemental technology of an environment-adaptive software system, the control unit 11 of the offload server 1 extracts an area to be offloaded from the source code of an application used by a user and outputs an intermediate language in order to offload functional processing (S11 to S15). The control unit 11 places and executes an executable file derived from the intermediate language on a verification machine 14, and verifies the offloading effect (S21 to S22). After repeating the verification and determining an appropriate offload area, the control unit 11 deploys the executable file in the production environment that is actually provided to the user, and provides it as a service (S23 to S25).

[0131] [Automatic offloading of loop statements using genetic algorithms] Automatic offloading of loop statements to a GPU or many-core CPU is a process of repeating steps S12 to S22 in FIG. 6 to obtain an offload code to be finally deployed in step S23.

[0132] GPUs generally do not guarantee latency, but are devices suited to increasing throughput through parallel processing. There are a wide variety of applications suited to GPUs. Typical applications include data encryption processing, image processing for camera video analysis, and machine learning processing for analyzing large amounts of sensor data, which involve a lot of repetitive processing. Therefore, the offload server 1 aims to increase the speed of applications by automatically offloading loop statements contained in the source code of the application to the GPU.

[0133] However, as described in the section on prior art, appropriate parallel processing is necessary to increase speed. In particular, when using a GPU, due to memory transfer between the CPU and GPU, performance is often not achieved unless the data size and number of loops are large. Also, depending on the timing of memory data transfer, the combination of individual loop statements (repeated statements) that can be accelerated in parallel may not be the fastest. For example, if there are 10 loop statements (repeated statements), and numbers 1, 5, and 10 can be accelerated compared to the CPU, the combination of numbers 1, 5, and 10 will not necessarily be the fastest.

[0134] In order to specify appropriate parallel regions, there are attempts to optimize loop statements by trial and error using the PGI compiler to determine whether they can be parallelized. However, trial and error requires a lot of work, and when providing such a service, it slows down the start of user adoption and increases costs.

[0135] Therefore, in this embodiment, an appropriate offload area is automatically extracted from a general-purpose program that is not intended to be parallelized. Therefore, the offload server 1 checks the offload of the function blocks in the order of many-core CPU → GPU → FPGA, and then checks the offload of the loop statement in the order of many-core CPU → GPU → FPGA. By checking the function blocks with high offload effect before the loop statement and checking in the order of many-core CPU → GPU → FPGA, which is the order of lowest verification cost, it is possible to efficiently search for a pattern that can be accelerated.

[0136] [Search image of control unit 11 using a simple genetic algorithm] FIG. 7 is a diagram showing an image of a search performed by the control unit 11 using the simple genetic algorithm of the offload server. A genetic algorithm is a combinatorial optimization method that mimics the evolutionary process of living organisms. The flowchart for a genetic algorithm is as follows: Initialization → Evaluation → Selection → Crossover → Mutation → Termination decision.

[0137] In this embodiment, a simple genetic algorithm that simplifies the process is used among the genetic algorithms. The simple genetic algorithm is a simplified genetic algorithm in which genes are only 1 and 0, and roulette selection, one-point crossover, and mutation reverse the value of one gene.

[0138] 《Initialization》 In the initialization, after checking whether all loop statements in the source code can be parallelized, parallelizable loop statements are mapped to the gene array. If GPU processing is used, the value is set to 1, and if not, the value is set to 0. The genes prepare a specified number of individuals M, and randomly assign 1 and 0 to each loop statement.

[0139] Specifically, the control unit 11 (see FIG. 1) acquires the source code 130 (see FIG. 6) used by the user, and checks whether loop statements can be parallelized from the code pattern 141 of the source code 130. As shown in FIG. 8, when three loop statements are detected from the code pattern 141d, one digit is randomly assigned to each loop statement, in this case three digits of 1 or 0 are randomly assigned to the three loop statements. For example, 0 is assigned when processing is performed by a single-core CPU, and 1 is assigned when output to a many-core CPU. However, at this stage, 1 or 0 is randomly assigned. The code corresponding to the gene length is three digits, and the three-digit code for gene length is 2 3 = 8 patterns, for example, 100, 110, .... In Fig. 7, the circles (◯ marks) in the code pattern 141 are shown as an image of the code.

[0140] "evaluation" In the evaluation of step S300, deployment and performance measurement are performed. That is, the performance measurement unit 116 (see FIG. 1) compiles the code corresponding to the gene, deploys it to the verification machine 14, and executes it. The performance measurement unit 116 performs benchmark performance measurement. The degree of suitability of the gene of a pattern with good performance (parallel processing pattern) is increased.

[0141] "choice" In the selection of step S301, a high-performance code pattern is selected based on the fitness. The performance measurement unit 116 (see FIG. 1) selects a specified number of genes with high fitness based on the fitness. In this embodiment, roulette selection according to the fitness and elite selection of the genes with the highest fitness are performed. FIG. 7 shows, as a search image, that the number of circles (◯) in the selected code pattern 142 has been reduced to three.

[0142] 《Crossover》 In crossover, some genes are exchanged at a certain point between selected individuals at a certain crossover rate Pc to create offspring individuals. The genes of a roulette-selected pattern (parallel processing pattern) are crossed with those of other patterns. The position of the one-point crossover is arbitrary. For example, crossover is performed at the second digit of the above three-digit code.

[0143] "mutation" Mutation involves changing the value of each gene in an individual from 0 to 1 or from 1 to 0 with a constant mutation rate Pm. In addition, in order to avoid local solutions, mutation is introduced. Note that in order to reduce the amount of calculation, mutation may not be performed. Then, as shown in step S302, each code that has been subjected to these crossovers and mutations is set as a next-generation code.

[0144] End Determination Generate next generation code patterns after crossover and mutation. In the termination judgment, the process is terminated after repeating the process for a specified number of generations T, and the gene with the highest fitness is taken as the solution. For example, performance is measured and three individuals with the fastest processing speeds are selected, for example individuals with genes 110,010,001. These three are then recombined using a genetic algorithm in the next generation to create, for example, a new pattern (parallel processing pattern) 011 (one example). At this point, a mutation is introduced into the recombined pattern, such as automatically changing 0 to 1. The above process is repeated to find the pattern with the fastest processing speed. A designated generation (for example, the 20th generation) is decided, and the pattern remaining in the final generation is designated as the final solution.

[0145] 《Deploy》 The parallel processing pattern with the highest processing performance that corresponds to the gene with the highest fitness is then redeployed into a production environment and provided to the user.

[0146] "supplementary explanation" This section explains the case where there are a significant number of loop statements (loop statements; repetitive statements) that cannot be offloaded to the GPU. For example, even if there are 200 loop statements, only about 30 can be offloaded to the GPU. Here, we will exclude those that will result in an error, and run the genetic algorithm on these 30.

[0147] OpenACC has a compiler that allows you to extract bytecode for the GPU and offload it to the GPU by specifying the directive "#pragma acc kernels". By writing a loop statement command in this "#pragma", you can determine whether the loop statement can run on the GPU.

[0148] For example, when using C / C++, the C / C++ code is analyzed to detect loop statements. When a loop statement is detected, the OpenACC parallel processing syntax "#pragma acc kernels" is used to write to the loop statement. In detail, loop statements are inserted one by one into an empty "#pragma acc kernels" and compiled. If an error occurs, the loop statement is excluded since it cannot be processed by the GPU in the first place. In this way, the remaining loop statements are detected. The one that does not produce an error is taken as the length (gene length). If there are five error-free loop statements, the gene length is 5, and if there are ten error-free loop statements, the gene length is 10. Note that parallel processing is not possible when there is a data dependency such that the previous processing is used for the next processing. The above is the preparation stage. Next, genetic algorithm processing is performed.

[0149] A code pattern is obtained with a gene length corresponding to the number of loop statements. At first, parallel processing patterns such as 10010, 01001, 00101, ... are randomly assigned. A genetic algorithm process is performed and then compiled. At this time, an error may occur even though the loop statement can be offloaded. This is the case when the loop statements are hierarchical (GPU processing is possible if either is specified). In this case, the loop statement that caused the error can be left as it is. Specifically, one method is to make it take longer to process, causing a timeout.

[0150] It is deployed on a verification machine 14 and benchmarked - for example, in the case of image processing, the image processing is benchmarked - and the shorter the processing time, the higher the fitness is evaluated. For example, the inverse of the processing time, that is, an individual whose processing time is 10 seconds, has a fitness of 1. An individual whose processing time is 100 seconds has a fitness of 0.1. An individual whose processing time is 1 second has a fitness of 10.

[0151] The offload server 1 selects the one with the highest adaptability. For example, the offload server 1 selects 3 to 5 individuals from 10 individuals and rearranges them to create a new code pattern. During the creation process, the same code pattern as before may be created. In this case, since there is no need to perform the same benchmark, the offload server 1 uses the same adaptability data as before. The offload server 1 of this embodiment stores the code pattern and its processing time in the storage unit 13.

[0152] The above describes the image of the search performed by the control unit 11 using the simple genetic algorithm.

[0153] FIG. 8 is a diagram showing gene sequence mapping of a loop statement to a many-core CPU. Code pattern 141d includes three loop statements, where one binary digit is assigned to each loop statement, and three digits of 1 or 0 are randomly assigned to the three loop statements.

[0154] The first for loop statement in the code pattern 141d is given the directive "#pragma omp parallel for". In this case, the first for loop statement is compiled so as to be offloaded to the many-core. Furthermore, the gene sequence 100 of the code pattern 141d is shown on the left side of the code pattern 141d.

[0155] 9A and 9B are flowcharts of a process for selecting a pattern for offloading a loop statement to a many-core CPU. First, the code analysis unit 112 analyzes a loop statement in the C / C++ source code (S90). The process designation unit 114 identifies the loop statement and the reference relationship in the C / C++ source code (S91).

[0156] The process designation unit 114 operates a benchmark tool to grasp the number of loops of each loop statement and to allocate each loop to a threshold value (S92), and then checks the possibility of parallel processing of each loop statement (S93).

[0157] The PGI compiler used for automatic offloading to the GPU would output an error if parallelization was not possible. However, with OpenMP compilers such as gcc, such errors are the programmer's responsibility. Therefore, in order to automate this process, the processing performed by the OpenMP directives has been simplified to only determine whether or not the loop statement should be parallelized on a many-core CPU, and a check is also performed during performance measurement to see if the final calculation result when parallelized is correct, creating a mechanism in which only patterns that produce correct calculation results remain in the evolutionary computation.

[0158] Next, the process specification unit 114 counts the number of loop statements that can be processed in parallel, sets the number as the gene length (S94), and prepares a specified number of gene sequences as initial values ​​(S95). Here, the process specification unit 114 creates a specified number of gene sequences to which 0 and 1 are randomly assigned.

[0159] The process specification unit 114 maps the C / C++ source code to genes and prepares a specified population pattern (S96). Specifically, the process specification unit 114 inserts a directive that specifies parallel processing when the gene value is 1 according to the prepared gene sequence into the C / C++ code (see, for example, the #pragma directive in FIG. 8).

[0160] The control unit 11 repeats the processes of steps S100 to S108 for the designated number of generations. In addition, in the repetition of the designated number of generations, the processing of steps S101 to S104 is further repeated the designated number of times. That is, within the repetition of the designated number of generations, the repetition of the designated number of individuals is processed in a nested state.

[0161] In the repetition of the specified number of individuals, the processing pattern creation unit 115 (see FIG. 1) compiles the C / C++ source code specified by the directive according to the gene pattern, i.e., each offload pattern, with the gcc compiler (S102). That is, the processing pattern creation unit 115 compiles the created C / C++ code with the gcc compiler on the verification machine 14 equipped with a many-core CPU. Here, a compilation error may occur when specifying parallel processing of multiple nested loop statements. In this case, it is treated the same as when the processing time during performance measurement times out.

[0162] The performance measurement unit 116 (see FIG. 1) deploys an executable file compiled from an offload pattern, i.e., the compiled offload pattern, to a verification machine 14 equipped with a many-core CPU (S103), executes the deployed executable file, and measures the benchmark performance when offloaded (S104).

[0163] Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same values ​​are used. In other words, if a gene with the same pattern as before is generated during the genetic algorithm process, the individual is not compiled or its performance is not measured, and the same measured value as before is used.

[0164] In step S105, if the repetition of the designated number of individuals has not been completed, the control unit 11 returns to the process of step S101, and if the repetition of the designated number of individuals has been completed, the control unit 11 proceeds to step S106.

[0165] Next, in step S106, the executable file creation unit 117 (see FIG. 1) evaluates individuals with shorter processing times so that they have a higher fitness, and selects individuals with higher performance. The executable file creation unit 117 (see FIG. 1) evaluates the difference between the calculation result of each individual and the calculation result of a single-core CPU, and selects an individual that falls within the acceptable range (S107). When measuring performance, the executable file creation unit 117 compares the final calculation result with the result when the original code is processed by a normal CPU to ensure that it is the same as the result when parallel processing is not performed, and if the difference is unacceptably large, the fitness of the pattern is set to 0 so that it is not selected for the next generation. The executable file creation unit 117 also checks whether the final calculation result when parallel processing is performed is correct during performance measurement, thereby making it possible to narrow down the patterns in the evolutionary computation to those that produce correct calculation results.

[0166] The executable file creation unit 117 performs crossover and mutation processes on the selected individuals to create the next generation individuals (S108). This is to perform compilation, performance measurement, fitness setting, selection, crossover, and mutation processes on the next generation individuals.

[0167] That is, after measuring the benchmark performance for all individuals, the executable file creation unit 117 sets the fitness of each gene sequence according to the benchmark processing time.The executable file creation unit 117 selects individuals to be left according to the set fitness.The executable file creation unit 117 then performs genetic algorithm processing on the selected individuals by crossover, mutation, and direct copying to create the next generation of individuals.

[0168] In step S109, if the repetition of the designated number of generations has not been completed, the control unit 11 returns to the process of step S100, and if the repetition of the designated number of generations has been completed, the control unit 11 proceeds to step S110.

[0169] In step S110, after completing the genetic algorithm process for the designated number of generations, the executable file creation unit 117 determines that the C / C++ code corresponding to the highest-performance gene sequence (the highest-performance parallel processing pattern) is a solution, and ends the process of FIG. 9B. By doing this, it is possible to search for patterns that can efficiently speed up loop statements, even when the target environment is a many-core CPU.

[0170] Like GPUs, many-core CPUs also speed up processing by making use of a large number of computing cores and parallelizing processing. The difference with GPUs is that many-core CPUs share the same memory, so there is no need to consider the overhead of data transfer between the CPU and GPU memory, which is often an issue when offloading to a GPU. In addition, the OpenMP specification is often used to parallelize program processing on many-core CPUs. OpenMP is a specification that specifies parallel processing, etc. for a program using directives such as "#pragma omp parallel for". The OpenMP programmer is responsible for parallelizing processing with OpenMP, and if a process that cannot be parallelized is parallelized, the compiler does not output an error, but rather outputs an incorrect calculation result.

[0171] Taking these factors into consideration, in this embodiment, automatic offloading of loop statements to many-core CPUs employs an evolutionary computing method that gradually increases the speed by creating multiple patterns for specifying whether or not the loop can be processed in parallel using "#pragma" in OpenMP and repeatedly measuring actual performance in a verification environment.

[0172] "implementation" The offload server 1 analyzes the C / C++ source code to find loop statements and grasp the program structure, such as variable data used in the loop statements and the processing of those variables.

[0173] Since it is necessary to eliminate loop statements that cannot be processed in parallel, the offload server 1 attempts to insert a directive to process each loop statement using the GPU, and excludes loop statements that generate errors from the genetic algorithm. Here, the number of loop statements that do not generate errors is the gene length.

[0174] Next, the offload server 1 prepares a gene sequence for the specified number of individuals as the initial value. Each gene value is created by randomly assigning 0 and 1. Depending on the prepared gene sequence, the offload server 1 inserts a directive into the C / C++ code that specifies GPU or many-core CPU processing when the gene value is 1.

[0175] The offload server 1 compiles the C / C++ code with the directives inserted using gcc, deploys the compiled executable file, and measures performance. In the performance measurement, the offload server 1 checks whether the calculation results obtained by parallel processing using, for example, the PCAST ​​function of the PGI compiler, along with the processing time, are significantly different from the original code, and if the results are unacceptable, the fitness is set to 0.

[0176] After measuring the performance of all individuals, the offload server 1 sets the fitness of each individual according to the processing time. The individuals to be left are selected according to the set fitness. The selected individuals are subjected to genetic algorithm processing by crossover, mutation, and direct copying to create the next generation of individuals.

[0177] The offload server 1 further performs instruction insertion, compilation, performance measurement, fitness setting, selection, crossover, and mutation processes on the next generation of individuals, and repeats this process for the specified number of generations. After completing the genetic algorithm process for the specified number of generations, the offload server 1 determines the C / C++ code with directives that corresponds to the gene sequence with the highest performance as the solution.

[0178] FIG. 10 is a diagram showing gene sequence mapping of loop statements to a GPU. The first for loop statement of the code pattern 141e is given "#pragma acc kernels". At this time, the first for loop statement is compiled so as to be offloaded to the GPU. Furthermore, the gene sequence 100 of the code pattern 141e is shown on the left side of the code pattern 141e.

[0179] 11A and 11B are flowcharts of a process for selecting a pattern for offloading a loop statement to a GPU. This selection process uses the method described in Non-Patent Document 2. The offload server 1 performs the following processing using an OpenACC compiler for C / C++.

[0180] The code analysis unit 112 (see FIG. 1) analyzes loop statements by code analysis of the C / C++ source code (S120). Then, the process designation unit 114 (see FIG. 1) identifies the loop statements in the C / C++ source code and their reference relationships (S121), runs a benchmark tool, grasps the number of loops of each loop statement, and assigns them according to a threshold (S122). Then, the process designation unit 114 checks the possibility of parallel processing of each loop statement (S123).

[0181] Next, the control unit 11 repeats the processes of steps S124 to S127 the number of times corresponding to the number of loop statements. The process specification unit 114 specifies parallel processing for each loop statement and compiles it (S125). As shown in Fig. 10, parallel processing in OpenACC is specified by "#pragma acc kernels". When a compilation error occurs, the process designation unit 114 deletes the designation of parallel processing and excludes the loop statement from the offload target (S126). Then, when the control unit 11 has repeated the process for all the loop statements, the process proceeds to step S128.

[0182] The process specification unit 114 counts the number of loop statements that do not cause a compilation error and sets this as the gene length (S128), and prepares gene sequences for the specified number of individuals. Here, gene sequences for the specified number of individuals are created with 0 and 1 randomly assigned as initial values. Next, the process specification unit 114 maps the C / C++ source code to the genes and prepares patterns for the specified number of individuals (S129). Specifically, the process specification unit 114 inserts a directive that specifies parallel processing when the gene value is 1, according to the prepared gene sequence, into the C / C++ code (see, for example, the #pragma directive in FIG. 10).

[0183] The control unit 11 repeats the processes of steps S130 to S137 for the designated number of generations. In addition, in the repetition of the designated number of generations, the processing of steps S131 to S134 is further repeated the designated number of times. That is, within the repetition of the designated number of generations, the repetition of the designated number of individuals is processed in a nested state.

[0184] In the repetition of the specified number of individuals, the processing pattern creation unit 115 (see FIG. 1) compiles the C / C++ source code specified by the directive according to the gene pattern, i.e., each offload pattern, with the PGI compiler (S132). That is, the processing pattern creation unit 115 compiles the created C / C++ code with the PGI compiler on the verification machine 14 equipped with a GPU. Here, a compilation error may occur when multiple nested loop statements are specified in parallel, etc. In this case, it is treated the same as when the processing time during performance measurement times out.

[0185] The performance measurement unit 116 (see FIG. 1) deploys an executable file into which the offload pattern has been compiled, i.e., the compiled offload pattern, to a verification machine 14 equipped with a CPU and a GPU (S133), executes the deployed executable file, and measures the benchmark performance when offloaded (S134).

[0186] Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same values ​​are used. In other words, if a gene with the same pattern as before is generated during the genetic algorithm process, the individual is not compiled or its performance is not measured, and the same measured value as before is used.

[0187] In step S135, if the control unit 11 has not completed the repetition of the designated number of individuals, the process returns to step S131, and if the control unit 11 has completed the repetition of the designated number of individuals, the process proceeds to step S136.

[0188] In step S136, the executable file creation unit 117 (see FIG. 1) evaluates the individuals with shorter processing times so that they have higher fitness, and selects the individuals with higher performance. The executable file creation unit 117 then performs crossover and mutation processes on the selected individuals to create the next generation of individuals (S137). This is to perform compilation, performance measurement, fitness setting, selection, crossover, and mutation processes on the next generation of individuals. That is, after measuring the benchmark performance for all individuals, the executable file creation unit 117 sets the fitness of each gene sequence according to the benchmark processing time.Then, the executable file creation unit 117 selects individuals to be left according to the set fitness.Then, the executable file creation unit 117 performs genetic algorithm processing such as crossover, mutation, and direct copying on the selected individuals to create the next generation of individuals.

[0189] In step S138, if the repetition of the designated number of generations has not been completed, the control unit 11 returns to the process of step S130, and if the repetition of the designated number of generations has been completed, the control unit 11 proceeds to step S139.

[0190] In step S139 after the genetic algorithm process for the designated number of generations is completed, the executable file creation unit 117 determines the C / C++ code corresponding to the highest-performance gene sequence (the highest-performance parallel processing pattern) as a solution, and ends the process of FIG. 11B.

[0191] The fundamental challenge is that while it is possible for a compiler to find a restriction that says this loop cannot be parallelized on a GPU, it is difficult to find a match that says this loop is suitable for parallelization on a GPU.

[0192] It is generally said that loops with a large number of loops are more suitable for GPU offloading, but it is difficult to predict how much performance will be achieved by actually offloading to the GPU without actual measurements. For this reason, it is common to manually instruct a program to offload a loop to the GPU and measure performance through trial and error.

[0193] Non-Patent Document 2 proposes to automatically discover suitable loop statements to offload to a GPU using a genetic algorithm, which is one of the evolutionary computing techniques. First, a general-purpose program for a normal CPU that does not assume GPU processing is checked for parallelizable loop statements, then the parallelizable loop statements are geneticized by assigning a value of 1 when executed by a GPU and a value of 0 when executed by a CPU, and performance verification is repeated in a verification environment to search for an appropriate region.

[0194] By narrowing down to parallelizable loop statements and then retaining and recombining parallel processing patterns that can be sped up in the form of genetic parts, the system is able to search for patterns that can be sped up efficiently from the vast number of possible parallel processing patterns.

[0195] In addition to appropriate extraction of loop statements, Non-Patent Document 2 proposes that, for variables used in nested loop statements, when offloading loop statements to a GPU, if CPU-GPU transfer is performed at a lower level of the nesting, transfer will be performed for each lower loop, which is inefficient, so for variables that are acceptable for CPU-GPU transfer at a higher level, they should be transferred collectively at the higher level. Automatic GPU offloading of loop statements is made possible by optimization using evolutionary computing techniques and reducing CPU-GPU transfers.

[0196] 12A and 12B are flowcharts of a process for selecting a pattern for offloading a loop statement to an FPGA. This selection process uses the method described in Non-Patent Document 3. First, the code analysis unit 112 analyzes the source code to be offloaded (S150), and analyzes information on loop statements and variables according to the language of the source code.

[0197] Then, when the offload range extraction unit 114a identifies loop statements and reference relationships in the source code (S151), it narrows down candidates for whether or not to attempt FPGA offloading for the identified loop statements. Arithmetic strength is one indicator of whether or not there is an offloading effect for a loop statement.

[0198] The offload range extraction unit 114a calculates the arithmetic strength of the loop statements of the application using an arithmetic strength analysis tool (S152). The arithmetic strength is an index that increases with a large number of calculations and decreases with a large number of accesses. A process with high arithmetic strength is a heavy process for the processor. Therefore, the offload range extraction unit 114a analyzes the arithmetic strength of the loop statements using the arithmetic strength analysis tool, and narrows down the offload candidates to loop statements with high arithmetic strength.

[0199] Even if a loop statement has high arithmetic strength, it is a problem if it consumes excessive FPGA resources when it is processed by an FPGA. Therefore, this article describes how to calculate the amount of resources required when processing a high arithmetic strength loop statement on an FPGA.

[0200] When compiling into an FPGA, the process involves converting from a high-level language such as OpenCL to the hardware description level such as HDL, and then performing the actual wiring process based on that. At this time, the wiring process takes a lot of time, but it only takes a few minutes to reach the intermediate state of HDL. Even at the intermediate state of HDL, the resources such as flip flops and look up tables used by the FPGA can be known. Therefore, by looking at the intermediate state of HDL, the amount of resources used can be known in a short time even before compilation is complete.

[0201] Therefore, in this embodiment, the processing pattern creation unit 115 converts the target loop statement into a high-level language such as OpenCL, and first calculates the resource amount. In addition, since the arithmetic strength and resource amount when the loop statement is offloaded are determined, the resource efficiency is set to arithmetic strength / resource amount or arithmetic strength×loop count / resource amount. Then, loop statements with high resource efficiency are further narrowed down as offload candidates.

[0202] 12A, the processing pattern creation unit 115 measures the number of loops in the loop statements of the application using a profiling tool such as gcov or gprof (S153), and narrows down the loop statements to those with high arithmetic strength and a high number of loops (S154).

[0203] The processing pattern creation unit 115 creates an OpenCL for offloading each of the narrowed down loop statements to the FPGA (S155).

[0204] Here, we will provide additional information on converting loop statements to OpenCL. In other words, when converting loop statements to a high-level language such as OpenCL, two processes are required. One is to divide the CPU processing program into the kernel (FPGA) and the host (CPU) according to the grammar of the high-level language such as OpenCL. The other is to incorporate techniques for increasing speed when dividing. In general, methods for increasing speed using FPGA include local memory cache, stream processing, multiple instantiation, loop statement unrolling, nested loop statement integration, memory interleaving, etc. Although these methods are not necessarily effective for some loop statements, they are often used as techniques for increasing speed.

[0205] Next, several loop statements with high resource efficiency are selected, and offload patterns are created using these to measure performance, the number of which is to be measured. FPGA speedup can be achieved by concentrating FPGA resources on a single process, or by distributing FPGA resources among multiple processes. Offload server 1 creates a fixed number of patterns for the selected single loop statements, and precompiles them as a preliminary step to running them on the actual FPGA.

[0206] The processing pattern creation unit 115 precompiles the created OpenCL to calculate the amount of resources to be used (S156), and narrows down loop statements with high resource efficiency (S157).

[0207] Then, the executable file creation unit 117 compiles OpenCL for offloading the narrowed-down loop statements (S158). The performance measurement unit 116 measures the performance of the compiled program (S159). This is the first performance measurement.

[0208] The processing pattern creation unit 115 creates a list of loop statements that have been measured and have higher performance than the CPU (S160), and creates an OpenCL for offloading by combining the loop statements in the list (S161). The processing pattern creation unit 115 calculates the amount of resources to be used by precompiling the combined offloaded OpenCL (S162). This is the second calculation of the resource amount. Note that the processing pattern creation unit 115 may not precompile, but may instead use the sum of the resource amounts obtained in the precompilation before the first measurement. In this way, the number of precompilations can be reduced.

[0209] The executable file creation unit 117 compiles the combined offload OpenCL (S163). The performance measurement unit 116 measures the performance of the compiled program (S164). This is the second performance measurement.

[0210] The production environment deployment unit 118 selects the pattern with the highest performance from the first and second measurements (S165), and ends the processing of this flow.

[0211] In this way, the automatic FPGA offloading of loop statements is limited to loop statements with high arithmetic intensity and loop count, which are highly resource efficient, and offload patterns are created, and high-speed patterns are searched for through actual measurements in the verification environment.

[0212] When considering offloading specific loop statements that take a long time to process to FPGAs to speed them up, it is difficult to predict which loops should be offloaded to make them faster, so we propose automatically measuring performance in a verification environment, just like with GPUs. However, since it takes several hours to compile OpenCL and run it on an actual FPGA, measuring performance by repeating it multiple times using a genetic algorithm for automatic GPU offloading would require enormous processing time and is not realistic.

[0213] Therefore, in this embodiment, the loop statements that are candidates for offloading to FPGA are narrowed down, and then performance measurement trials are performed. Specifically, first, for the discovered loop statements, loop statements with high arithmetic intensity are extracted using an arithmetic intensity analysis tool such as the ROSE framework. Furthermore, loop statements with a large number of loop counts are also extracted using a profiling tool such as gcov or gprof.

[0214] Then, loop statements with high arithmetic intensity or loop count are selected as candidates for OpenCL conversion. For offload candidate loop statements with high arithmetic intensity or loop count, the created OpenCL is precompiled to find loop statements with high resource efficiency. During the compilation process, the resources to be created, such as flip flops and look up tables, are known, so the loop statements that use fewer resources are further narrowed down.

[0215] Several loop statements were selected, and their performance was measured using these. Speeding up on an FPGA can be achieved by concentrating FPGA resources on just one process, or by distributing FPGA resources across multiple processes. Therefore, the selected single loop statements are compiled to run on an actual FPGA device and performance is measured, and for the single loop statements that could be sped up further, combination patterns are created and a second performance measurement is performed. Offload server 1 selects the fastest pattern from the multiple patterns whose performance was measured in the verification environment as the solution.

[0216] Regarding automatic FPGA offloading of loop statements, candidate loop statements are narrowed down using arithmetic strength, number of loops, and amount of resources, and then performance measurements of multiple patterns are performed in the verification environment to enable automatic offloading.

[0217] "implementation" The offload server 1 analyzes the C / C++ source code to find loop statements and to grasp the program structure, such as variable data used within the loop statements. Next, the offload server 1 executes ROSE to obtain the arithmetic strength of each loop statement, and obtains the loop count of each loop statement using gcov.

[0218] Next, the offload server 1 generates OpenCL code to offload each loop statement with high arithmetic intensity and high loop count to the FPGA. The OpenCL code is divided into the loop statement as an FPGA kernel and the rest as a CPU host program.

[0219] Next, the offload server 1 precompiles the created OpenCL code and calculates the amount of resources such as flip flops to be used. The amount of resources used is displayed as a percentage of the total resource amount. Here, the offload server 1 selects a loop statement with high arithmetic intensity, high loop count, and low resource amount (high resource efficiency) based on the arithmetic intensity, loop count, and resource amount.

[0220] Next, the offload server 1 creates a pattern to be measured using the selected loop statements as candidates, and compiles it. The offload server 1 measures the performance of each of the selected loop statements when they are offloaded, and checks whether they can be sped up. If multiple statements are sped up as a result, the offload server 1 creates an OpenCL for that combination and measures its performance to check whether it can be sped up more than a single statement. However, because FPGAs have limited resource amounts, the offload server 1 will not create an OpenCL if the amount of resources does not fall within the upper limit when offloading multiple statements. Finally, the offload server 1 selects a high-speed pattern from the multiple measurement patterns as a solution.

[0221] <<Original Application of the Present Invention and Its Effects>> In the invention described in claim 1, an offload server is characterized by comprising: a code analysis unit that analyzes whether the source code of a software program includes a functional block that can be executed by an accelerator; a processing pattern creation unit that creates a pattern for offloading the functional block to the accelerator; a performance measurement unit that deploys each of the patterns for offloading the functional block to a verification environment to measure performance; and a control unit that, if the performance measured by the performance measurement unit is desired, creates a pattern for offloading the remaining loop statements of the software program excluding the functional block, and deploys each of the patterns to the verification environment to measure performance, by the processing pattern creation unit, and, if the performance measured by the performance measurement unit is not desired, creates a pattern for offloading loop statements of the software program, by the processing pattern creation unit, and deploys each of the patterns to the verification environment to measure performance, by the performance measurement unit.

[0222] In this way, even if the destination environment contains a mixture of multiple types of accelerators, the performance can be automatically improved. Furthermore, if the destination node has multiple types of accelerators, offloading to multiple types of accelerators simultaneously can achieve faster offloading than when using a single accelerator.

[0223] In the invention described in claim 2, the accelerator is of multiple types, and the control unit selects the fastest pattern among the patterns to be offloaded to each of the accelerators.

[0224] By doing this, performance can be automatically improved even in a migration environment that contains a mixture of multiple types of accelerators.

[0225] In the invention described in claim 3, the control unit tries the loop statement after trying the functional block as a target for offloading to the accelerator, and tries a many-core CPU, GPU, and FPGA in that order as the accelerator to which the functional block and / or the loop statement is to be offloaded.This is an offload server as described in claim 1 or 2, characterized in that

[0226] In this way, it is possible to search for a pattern that can efficiently increase speed in an environment where multiple types of accelerators are mixed.

[0227] In the invention described in claim 4, the control unit is configured to, in offloading the function block or the loop statement, if a pattern that is faster than the original processing is found for a plurality of accelerators and the same function block or loop statement can be offloaded to the plurality of accelerators, select the accelerator with higher performance as the offload destination, if a different function block or a different loop statement can be offloaded to a different accelerator but cannot be offloaded to an accelerator on the same node, select the accelerator with higher performance as the offload destination, and if a different function block or a different loop statement can be offloaded to a different accelerator and can be offloaded to an accelerator on the same node, select a different accelerator as the offload destination for the function block or the loop statement.

[0228] In this way, it is possible to search for patterns that can efficiently speed up function blocks or loop statements in an environment where multiple types of accelerators are mixed.

[0229] In the invention described in claim 5, the control unit is configured to, in offloading the function block and the loop statement, if the offload destination of the loop statement is a different node from the offload destination of the function block, if the performance improvement due to offloading the loop statement is less than or equal to the performance improvement due to offloading the function block, not attempt a pattern of offloading the loop statement and attempt a pattern of offloading the function block, if the performance improvement due to offloading the loop statement exceeds the performance improvement due to offloading the function block, not attempt a pattern of offloading the function block and attempt a pattern of offloading the loop statement, and if the offload destination of the loop statement is the same node as the offload destination of the function block, attempt a pattern of offloading the function block and the loop statement.

[0230] In this way, it is possible to search for patterns that can efficiently speed up function blocks or loop statements in an environment where multiple types of accelerators are mixed.

[0231] In the invention described in claim 6, the offload server comprises a code analysis unit that analyzes whether the source code of a software program contains a loop statement that can be executed by a many-core CPU; a processing pattern creation unit that creates a pattern for offloading the loop statement to the many-core CPU; and a performance measurement unit that maps each of the loop statements to a genetic pattern, deploys each of the patterns for offloading each of the loop statements to a verification environment to measure performance, sets high performance as high fitness, and repeats the process of creating next-generation patterns through elite selection, crossover, and mutation processes of a genetic algorithm, wherein the performance measurement unit, in measuring performance in the verification environment, compares the calculation result when processed by the many-core CPU with the calculation result of the original processing, and if the error in the calculation result when processed by the many-core CPU exceeds an allowable threshold, sets the fitness of the genetic pattern to low.

[0232] By doing this, it is possible to search for patterns that can efficiently speed up loop statements, even when the target environment is a many-core CPU.

[0233] In the invention described in claim 7, an offload control method is provided which executes the steps of: a code analysis unit analyzing whether or not a source code of a software program includes a functional block that can be executed by an accelerator; a processing pattern creation unit creating a pattern for offloading the functional block to the accelerator; and a performance measurement unit deploying each of the patterns for offloading the functional block to a verification environment and measuring performance; and if the performance measured by the performance measurement unit is desired, the processing pattern creation unit creates a pattern for offloading remaining loop statements of the software program excluding the functional block; and the performance measurement unit deploys each of the patterns to the verification environment and measures performance; and if the performance measured by the performance measurement unit is not desired, the processing pattern creation unit creates a pattern for offloading loop statements of the software program; and the performance measurement unit deploys each of the patterns to the verification environment and measures performance.

[0234] In this way, even if the destination environment contains a mixture of multiple types of accelerators, the performance can be automatically improved. Furthermore, if the destination node has multiple types of accelerators, offloading to multiple types of accelerators simultaneously can achieve faster offloading than when using a single accelerator.

[0235] In the invention described in claim 8, an offloading program is provided for causing a computer to execute the steps of: analyzing whether or not source code of a software program includes a functional block that can be executed by an accelerator; creating patterns for offloading the functional blocks to the accelerator; deploying each of the patterns for offloading the functional blocks to a verification environment to measure performance; if the measured performance is desired, creating patterns for offloading the remaining loop statements of the software program excluding the functional blocks and deploying each of the patterns to the verification environment to measure performance; if the measured performance is not desired, creating patterns for offloading loop statements of the software program and deploying each of the patterns to the verification environment to measure performance.

[0236] In this way, even if the destination environment contains a mixture of multiple types of accelerators, the performance can be automatically improved. Furthermore, if the destination node has multiple types of accelerators, offloading to multiple types of accelerators simultaneously can achieve faster offloading than when using a single accelerator. [Explanation of symbols]

[0237] 1 Offload Server 2. Cloud Layer 3 Network Layer 4 Device Layer 10. Gateway 20 Network Edge 30 Data Center 11 Control section 12 Input / output section 13 Storage section 14 Test machine 15 Resources 111 Code designation section 112 Code Analysis Department 114 Processing specification section 114a Off-road range extraction section 114b Intermediate language file output section 115 Processing pattern creation unit 116 Performance measurement section 116a Executable file placement section 117 Executable File Creation Department 118 Production Environment Deployment Department 119 Performance Measurement Test Extraction Execution Department 120 User Provision Department 130 Source Code 131 Test Case Database 132 Intermediate language files 133 Code Pattern Database 141,142 Code Patterns 151~154 equipment 80 Communication Network 900 Computers 910 CPU 920 RAM 930 ROM 940 HDD 950 Communication Interface 960 Input / Output Interface 970 Media Interface 980 Recording media

Claims

1. a code analysis unit that analyzes whether a source code of a software program includes a loop statement that can be executed by a GPU, an FPGA, or a many-core CPU; a processing pattern creation unit that creates a pattern for offloading the loop statement to any one of the GPU, the FPGA, and the many-core CPU; a performance measurement unit that maps each of the loop statements to a gene pattern, and deploys each of the patterns that offload each of the loop statements to a verification environment to measure performance; an executable file creation unit which repeats the process of setting a higher fitness for a pattern with a shorter processing time, selecting a pattern by elite selection of a genetic algorithm, and creating a next-generation pattern by processing of crossover and mutation; Equipped with the performance measurement unit, in measuring the performance in the verification environment, compares a calculation result obtained by processing using any one of the GPU, the FPGA, and the many-core CPU with a calculation result of an original process, and sets a degree of conformance of the gene pattern to low if an error in the calculation result obtained by processing using any one of the GPU, the FPGA, and the many-core CPU exceeds an allowable threshold value; 13. An offload server comprising:

2. the executable file creation unit performs roulette selection according to the fitness of a gene pattern in elite selection of a genetic algorithm; The offload server according to claim 1 .

3. the execution file creation unit selects the pattern with the highest performance as a solution after completing the processing for the specified number of generations; The offload server according to claim 1 .

4. A code analysis unit analyzes whether a source code of the software program includes a loop statement that can be executed by a GPU, an FPGA, or a many-core CPU; a processing pattern creation unit creating a pattern for offloading the loop statement to any one of the GPU, the FPGA, and the many-core CPU; A performance measurement unit maps each of the loop statements to a gene pattern, and deploys each of the patterns for offloading each of the loop statements to a verification environment to measure performance; a step in which the executable file creation unit sets a higher fitness for a pattern with a shorter processing time, selects a pattern by elite selection of a genetic algorithm, and creates a next-generation pattern by crossover and mutation processing; The performance measurement unit further executes a step of comparing a calculation result obtained by processing using any one of the GPU, the FPGA, and the many-core CPU with a calculation result obtained by an original process in the measurement of performance in the verification environment, and setting a goodness of fit of the gene pattern to a low level if an error in the calculation result obtained by processing using any one of the GPU, the FPGA, and the many-core CPU exceeds an allowable threshold value. An off-load control method comprising:

Images