OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM

Abstract

An offload server (1) comprises: a performance measurement unit (116) that compiles a parallel processing pattern application, puts the compiled application in an accelerator verification device, and executes performance measurement processing of when the compiled application is offloaded to an accelerator; an evaluation value setting unit (116c) that, on the basis of the power usage amount and the processing time required at the time of offloading measured by the performance measurement unit (116), sets an evaluation value which includes the processing time and the power usage amount, and which increases as the processing time and the power usage amount decrease; and an execution file creation unit (117) that, on the basis of the processing time and power usage amount measurement results, selects a parallel processing pattern with the highest evaluation value, from among a plurality of parallel processing patterns, and compiles the parallel processing pattern with the highest evaluation value to create an execution file.

Description

[Technical field]

[0001] The present invention relates to an offload server, an offload control method, and an offload program that automatically offloads functional processing to an accelerator such as a GPU (Graphics Processing Unit) or an FPGA (Field Programmable Gate Array). [Background technology]

[0002] The use of heterogeneous computing resources other than CPUs (Central Processing Units) is increasing. For example, image processing is being performed on a server with enhanced GPUs (accelerators), and signal processing is being accelerated with FPGAs (accelerators). FPGAs are programmable gate arrays whose configuration can be set by designers after manufacturing, and are a type of PLD (Programmable Logic Device). Amazon Web Services (AWS) (registered trademark) provides GPU instances and FPGA instances, and these resources can be used on demand. Microsoft (registered trademark) is using FPGAs to make searches more efficient.

[0003] In the OpenIoT (Internet of Things) environment, it is expected that a wide variety of applications will be created using service integration technologies, and by utilizing more advanced hardware, it is expected that the performance of the applications will be improved. However, in order to do so, programming and settings that are tailored to the hardware on which they will be run are required. For example, CUDA (Compute Unified Device Architecture), OpenCL (Open Computing Language) (Trademark. The same applies below.) It requires a lot of technical knowledge, and the hurdles are high. OpenCL is an open API (Application Programming Interface) that can handle all computing resources (not limited to CPUs and GPUs) in a unified manner without being tied to specific hardware.

[0004] The following is required to make it easy to use GPUs and FPGAs in user applications: When deploying general-purpose applications such as image processing and encryption processing to the OpenIoT environment, it is desirable for the OpenIoT platform to analyze the application logic and automatically offload processing to the GPU or FPGA.

[0005] CUDA, a development environment for GPGPUs (General Purpose GPUs), which use the computing power of GPUs for purposes other than image processing, is being developed. CUDA is a development environment for GPGPUs. OpenCL has also emerged as a standard for unifying the handling of heterogeneous hardware such as GPUs, FPGAs, and many-core CPUs.

[0006] CUDA and OpenCL are programmed using an extension of the C language. However, it is difficult to write code that requires memory copy and release between devices such as GPUs and the CPU. In fact, there are not many engineers who can use CUDA and OpenCL well.

[0007] To simplify GPGPU, there is a technology that specifies the parts of loop statements, etc. that should be parallelized on a directive basis, and the compiler converts them into device-oriented code according to the directive. Technical specifications include OpenACC (Open Accelerator), and compilers include PGI Compiler (registered trademark). For example, in an example using OpenACC, the user specifies parallel processing for code written in C / C++ / Fortran languages ​​using OpenACC directives. The PGI Compiler checks the parallelizability of the code, generates executable binaries for GPU and CPU, and modularizes it. IBM JDK (registered trademark) supports a function that offloads parallel processing specifications that follow the lambda format of Java (registered trademark) to the GPU. By using these technologies, the programmer does not need to be aware of data allocation to GPU memory, etc. In this way, technologies such as OpenCL, CUDA, and OpenACC make it possible to offload processing to GPUs and FPGAs.

[0008] However, even if offload processing itself is possible, there are many challenges in offloading appropriately. For example, there are compilers with automatic parallelization functions, such as the Intel Compiler (registered trademark). When performing automatic parallelization, parallel processing parts such as for statements (loop statements) in a program are extracted. However, when running in parallel using a GPU, performance is often not achieved due to the overhead of data exchange between the CPU and GPU memory. When using a GPU to speed up processing, it is necessary for skilled personnel to tune with OpenCL or CUDA, or to search for appropriate parallel processing parts using the PGI compiler, etc. For this reason, it is difficult for unskilled users to use GPUs to improve the performance of their applications, and even when using automatic parallelization technology, it takes a long time before they can start using it, as it requires trial and error tuning to decide whether or not to parallelize for statements.

[0009] Non-Patent Documents 1 and 2 are examples of efforts to automate the trial and error process for parallel processing. Non-Patent Documents 1 and 2 propose environment adaptive software that aims to run applications with high performance and low power consumption by automatically performing conversion and resource settings so that code written once can use GPUs, FPGAs, many-core CPUs, etc. present in the deployment environment. Additionally, Non-Patent Documents 1 and 2 propose a method to automatically offload loop statements in application code to GPUs as an element of the environment adaptive software and evaluate the performance improvement.

[0010] Non-Patent Document 3 proposes a method for automatically offloading loop statements in application code to an FPGA as an element of environment adaptive software, and evaluates the performance improvement. Non-Patent Document 4 proposes an automatic offloading method for loop statements in application code as an element of environment adaptive software to a mixed environment of GPUs and FPGAs, and evaluates the performance improvement. [Prior art documents] [Non-patent literature]

[0011] [Non-Patent Document 1] Y. Yamato, T. Demizu, H. Noguchi and M. Kataoka, “Automatic GPU Offloading Technology for Open IoT Environment,” IEEE Internet of Things Journal, DOI: 10.1109 / JIOT.2018.2872545, Sep. 2018 (Electronic Publishing). Vol.6, Issue 2, pp.2369-2378, Apr. 2019. (Print Publishing) [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. (Electronic Publishing) Vol.54, No.3, pp.567-584, May 2020. (Print Publishing) [Non-Patent Document 3] Y. Yamato, “Automatic Offloading Method of Loop Statements of Software to FPGA,” International Journal of Parallel, Emergent and Distributed Systems, Taylor & Francis, DOI: 10.1080 / 17445760.2021.1916020, Apr. 2021. [Non-Patent Document 4] Y. Yamato, “Proposal of Automatic Offloading Method in Mixed Offloading Destination Environment,” 2020 Eighth International Symposium on Computing and Networking Workshops (CANDAR 2020), pp.460-464, Nov. 2020. Summary of the Invention [Problem to be solved by the invention]

[0012] Non-Patent Documents 1 and 2 propose a method using evolutionary computing to automate the search for parallel processing locations when offloading processing to a GPU, etc., but only the shortening of processing time was evaluated, and the reduction in power consumption was not evaluated. In addition, the automatic offloading to FPGA in Non-Patent Document 3 and the offloading to a mixed environment in Non-Patent Document 4 were not evaluated in terms of reduction in power consumption. That is, Non-Patent Documents 1 to 4 only evaluate the reduction in processing time during automatic offloading, and do not evaluate power consumption. Therefore, there is a problem that the performance and power consumption at the migration destination are not necessarily appropriate.

[0013] The present invention has been made in view of these points, and an object of the present invention is to improve performance while reducing power consumption when automatically offloading to an offload device such as a GPU or FPGA. [Means for solving the problem]

[0014] In order to solve the above-mentioned problems, an offload server that offloads specific processing of an application to a GPU includes an application code analysis unit that analyzes the source code of the application, a data transfer specification unit that specifies, among variables that require transfer between a CPU and the GPU based on the results of the code analysis, that variables for which CPU processing and GPU processing are not mutually referenced or updated and for which the results of the GPU processing are simply returned to the CPU, to be collectively transferred before and after the start and end of the GPU processing, a parallel processing specification unit that specifies loop statements of the application, and compiles each of the specified loop statements by specifying a parallel processing specification statement in the GPU, and a loop statement for which a compilation error occurs is excluded from offloading, and specifies a parallel processing specification statement for the GPU that specifies a loop statement for which a compilation error occurs, and specifies a loop statement for which a compilation error does not occur. The offload server includes a parallel processing pattern creation unit that creates a parallel processing pattern that specifies whether or not to perform parallel processing for a loop statement; a performance measurement unit that compiles the application of the parallel processing pattern, places it on an accelerator verification device, and executes processing for measuring performance when offloaded to the GPU; an evaluation value setting unit that sets an evaluation value that includes processing time and power consumption based on the processing time and power consumption required during offloading measured by the performance measurement unit, the evaluation value being higher as the processing time and power consumption are lower; and an executable file creation unit that selects the parallel processing pattern with the highest evaluation value from a plurality of the parallel processing patterns based on the measurement results of the processing time and the power consumption, and compiles the parallel processing pattern with the highest evaluation value to create an executable file. Effect of the Invention

[0015] According to the present invention, when automatic offloading is performed to an offload device such as a GPU or FPGA, it is possible to reduce power consumption while improving performance. [Brief description of the drawings]

[0016] [Figure 1]2 is a functional block diagram showing a configuration example of an offload server according to the first embodiment of the present invention; FIG. [Diagram 2] FIG. 11 is a diagram showing an automatic offload process using GA in the offload server according to the first embodiment. [Diagram 3] FIG. 13 is a diagram showing a search image of a control unit (automatic offload function unit) by Simple GA of the offload server according to the first embodiment. [Figure 4] FIG. 13 is a diagram illustrating an example of a normal CPU program of the comparative example. [Diagram 5] 13 is a diagram illustrating an example of a loop statement when data is transferred from a CPU to a GPU using a simple CPU program of a comparative example. FIG. [Figure 6] FIG. 11 is a diagram showing an example of a loop statement for transferring data from a CPU to a GPU when the offload server according to the first embodiment is nested and integrated. [Figure 7] FIG. 11 is a diagram showing an example of a loop statement for data transfer from a CPU to a GPU in the case where transfer integration of the offload server according to the first embodiment is performed. [Figure 8] FIG. 11 is a diagram showing an example of a loop statement when data is transferred from a CPU to a GPU in a case where transfer integration of the offload server according to the first embodiment is performed and a temporary area is used. [Figure 9A] 11 is a flowchart illustrating an overview of an operation of implementing the offload server according to the first embodiment. [Figure 9B] 11 is a flowchart illustrating an overview of an operation of implementing the offload server according to the first embodiment. [Figure 10] FIG. 13 is a diagram showing power consumption (Watt) and processing time when the Himeno benchmark is offloaded to a GPU of the offload server according to the first embodiment. [Figure 11] FIG. 11 is a functional block diagram showing a configuration example of an offload server according to a second embodiment of the present invention. [Figure 12] 11 is a flowchart illustrating an overview of an operation of implementing the offload server according to the second embodiment. [Figure 13]13 is a flowchart showing a performance / power consumption measuring process of a performance measuring unit in the offload server according to the second embodiment. [Figure 14] FIG. 11 is a diagram for explaining an outline of the implementation operation of the offload server according to the second embodiment. [Figure 15] FIG. 11 is a diagram for explaining a flow from the C code of the offload server to a search for an OpenCL final solution according to the second embodiment. [Figure 16] FIG. 11 is a diagram showing the power consumption (Watt) and processing time when MRI-Q is offloaded to the FPGA of the offload server according to the second embodiment. [Figure 17] FIG. 2 is a hardware configuration diagram showing an example of a computer that realizes the functions of an offload server according to each embodiment of the present invention. DETAILED DESCRIPTION OF THE PREFERRED EMBODIMENTS

[0017] An offload server in an embodiment for carrying out the present invention (hereinafter, referred to as "the present embodiment") will be described below with reference to the drawings. (Explanation of the principle) Currently, it is difficult for a compiler to determine whether a loop statement is suitable for parallel processing on a GPU. It is difficult to predict how much performance and power consumption will be achieved by offloading to a GPU without actually measuring it. For this reason, instructions to offload this loop statement to a GPU are manually given, and measurements are conducted through trial and error. The present invention automatically finds suitable loop statements to offload to the GPU using a genetic algorithm (GA), which is an evolutionary computing technique. In other words, it geneticizes parallelizable loop statements by assigning a value of 1 when the loop is executed by the GPU and a value of 0 when the loop is executed by the CPU, and iteratively measures the loop in a verification environment to search for suitable patterns.

[0018] Here, a pattern that can be processed in a short time by measurement is regarded as a gene with high fitness. At this time, power consumption is also measured, and a new process is added to give a high fitness to a pattern with low power consumption. For example, (processing time) -1 / 2×(Electricity usage) -1 / 2 The setting is made so that the shorter the processing time and the lower the power consumption, the higher the compatibility of the gene pattern. Regarding GPU offloading of loop statements, automatic speedup and power saving are achieved by using an evolutionary computation method that includes power consumption in the fitness factor, as described in detail in the first embodiment, and by reducing CPU-GPU transfers.

[0019] (First embodiment) Next, the offload server 1 and the like in an embodiment of the present invention (hereinafter, referred to as "the present embodiment") will be described.

[0020] [Automatic GPU offloading of loop statements] 1 is a functional block diagram showing an example of the configuration of an offload server 1 according to a first embodiment of the present invention. This embodiment is an example applied to GPU automatic offloading of loop statements. The offload server 1 is a device that automatically offloads specific processing of an application to an accelerator. 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).

[0021] The input / output unit 12 is composed of a communication interface for transmitting and receiving information between each device, etc., and an input / output interface for transmitting and receiving information between input devices such as a touch panel and a keyboard, and output devices such as a monitor.

[0022] The storage unit 13 is configured with a hard disk, a flash memory, a RAM (Random Access Memory), or the like. This memory unit 13 stores a test case database (DB) 131, and also temporarily stores programs (offload programs) for executing each function of the control unit 11, and information necessary for processing by the control unit 11 (e.g., an intermediate language file (Intermediate file) 132).

[0023] The test case DB 131 stores performance test items. The test case DB 131 stores information for performing tests to measure the performance of applications to be accelerated. For example, in the case of a deep learning application for image analysis processing, the test items are sample images and the test items for executing the images. The verification machine 14 includes a CPU (Central Processing Unit), a GPU, and an FPGA (accelerator) as a verification environment for the environment adaptive software.

[0024] The control unit 11 is an automatic offloading function that controls the entire offload server 1. The control unit 11 is realized, for example, by a CPU (not shown) expanding a program (offload program) stored in the storage unit 13 into a RAM and executing the program.

[0025] The control unit 11 includes an application code specification unit (Specify application code) 111, an application code analysis unit (Analyze application code) 112, a data transfer specification unit 113, a parallel processing specification unit 114, a parallel processing pattern creation unit 115, a performance measurement unit 116, an executable file creation unit 117, a production environment deployment unit (Deploy final binary files to production environment) 118, a performance measurement test extraction and execution unit (Extract performance test cases and run automatically) 119, and a user provision unit (Provide price and performance to a user to judge) 120.

[0026] <Application code designation section 111> The application code designation unit 111 designates the input application code. Specifically, the application code designation unit 111 identifies the processing function (image analysis, etc.) of the service provided to the user.

[0027] <Application Code Analysis Unit 112> The application code analysis unit 112 analyzes the source code of the processing function and grasps the structure of the use of a specific library, such as a loop statement or an FFT library call.

[0028] <Data transfer designation unit 113> Based on the results of the code analysis, the data transfer specification unit 113 specifies that, among the variables that need to be transferred between the CPU and GPU, variables that are not mutually referenced or updated between the CPU processing and the GPU processing and for which the results of GPU processing are simply returned to the CPU are to be transferred in a lump before the start and after the end of GPU processing. Here, the variables that need to be transferred between the CPU and GPU are variables that are defined in multiple files or multiple loops based on the results of code analysis.

[0029] Here, the case of a GPU will be described as an example. In the case of a GPU, the OpenACC syntax is used for specification, whereas in the case of an FPGA, the OpenCL syntax is used for specification. The data transfer specification unit 113 specifies the specification of transferring data collectively before the start and after the end of GPU processing by using the data copy of OpenACC.

[0030] The data transfer specification unit 113 adds a directive indicating that the variables to be processed by the GPU are not required to be transferred in a batch when the variables have already been transferred to the GPU side.

[0031] The data transfer specification unit 113 explicitly indicates, using data present of OpenACC, that variables that are transferred in a batch before the start of GPU processing and that do not need to be transferred at the timing of loop statement processing do not need to be transferred.

[0032] When transferring data between the CPU and GPU, the data transfer specification unit 113 creates a temporary area on the GPU side (#pragma acc declare create), stores the data in the temporary area, and then synchronizes the temporary area (#pragma acc update) to instruct the variable transfer.

[0033] Based on the result of the code analysis, the data transfer specification unit 113 specifies GPU processing for the loop statement using at least one selected from the group consisting of a kernels directive, a parallel loop directive, and a parallel loop vector directive of OpenACC.

[0034] The OpenACC kernels directive is used for single loops and tightly nested loops. The OpenACC parallel loop directive is used for non-tightly nested loops. The OpenACC parallel loop vector directive is used for loops that cannot be parallelized but can be vectorized.

[0035] <Parallel processing designation section 114> The parallel processing specification unit 114 identifies loop statements (repetitive statements) in an application, and for each repetitive statement, specifies the processing to be performed in the GPU using an OpenACC directive, and compiles the specified loop statements. The parallel processing designation unit 114 includes an offloadable area extraction unit (Extract offloadable area) 114a and an intermediate language file output unit (Output intermediate file) 114b.

[0036] The offload range extraction unit 114a identifies processes that can be offloaded to the GPU, such as loop statements, and extracts an intermediate language corresponding to the offload process. Here, the intermediate language is an OpenACC language file (a C language extension file in which processes are specified in the OpenACC grammar) in the case of a GPU, and an OpenCL language file (a C language extension file in which processes are specified in the OpenCL grammar) in the case of an FPGA.

[0037] The intermediate language file output unit 114b outputs the extracted intermediate language file 132. The intermediate language extraction is not a one-time process, but is repeated to try and optimize the execution for an appropriate offload area search.

[0038] <Parallel processing pattern creation unit 115> The parallel processing pattern creation unit 115 creates a parallel processing pattern that excludes loop statements (repetitive statements) that cause a compilation error from being offloaded, and specifies whether or not to process repetitive statements that do not cause a compilation error in parallel.

[0039] <Performance measurement section 116> The performance measurement unit 116 compiles an application of a parallel processing pattern, places it on the verification machine 14, and executes a process for measuring performance when offloaded to the GPU. The performance measurement unit 116 includes a binary file deployment unit (Deploy binary files) 116a, a power usage measurement unit 116b (performance measurement unit), and an evaluation value setting unit 116c. Note that, although the evaluation value setting unit 116c is configured to be included in the performance measurement unit 116, it may be a separate independent functional unit.

[0040] The performance measurement unit 116 executes the placed binary file, measures the performance when offloaded, and returns the performance measurement result to the offload range extraction unit 114a. In this case, the offload range extraction unit 114a extracts another parallel processing pattern, and the intermediate language file output unit 114b attempts to measure the performance based on the extracted intermediate language (see symbol a in FIG. 2 described later).

[0041] The binary file placement unit 116a deploys (places) an executable file derived from the intermediate language on the verification machine 14 equipped with a GPU.

[0042] The power consumption measurement unit 116b measures the processing time and power consumption required during offloading. In a machine equipped with a GPU, the GPU power can be measured using the nvidia-smi command of the NVIDIA (registered trademark) tool, and the CPU power can be measured using the s-tui command. In a server equipped with an FPGA, the power consumption of the entire server can be measured using the ipmitool command of the IPMI (Intelligent Platform Management Interface).

[0043] The evaluation value setting unit 116c sets an evaluation value including the processing time and power consumption required during offloading, based on the processing time and power consumption measured by the performance measuring unit 116 and the power consumption measuring unit 116b, and the evaluation value is set to a higher value as the processing time and power consumption are lower. -1 / 2 ×(Electricity usage) -1 / 2 The lower the processing time and power consumption, the higher the evaluation value and the higher the degree of compatibility. Also, when the priorities for high performance and low power consumption are different, (processing time) -1 / 2 , (power usage) -1 / 2 Weighting may be given to either of the above.

[0044] <Executable File Creation Section 117> The executable file creation unit 117 selects the parallel processing pattern with the highest evaluation value from the multiple parallel processing patterns based on the results of the measurement of processing time and power consumption repeated a predetermined number of times, and compiles the parallel processing pattern with the highest evaluation value to create an executable file.

[0045] <Production Environment Configuration Department 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.

[0046] <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 DB 131 and executes the performance tests ("Placement of final binary file in production environment"). After arranging the executable file, the performance measurement test extraction and execution unit 119 extracts performance test items from the test case DB 131 and automatically executes the extracted performance tests in order to show the performance to the user.

[0047] <User provision unit 120> The user provision 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 the user"). The test case DB 131 stores data for automatically conducting tests to measure the performance of applications. The user provision unit 120 presents the user with the results of executing the test data in the test case DB 131 and the price of the entire system, which is determined from the unit prices of the resources used in the system (virtual machines, FPGA instances, GPU instances, etc.). The user decides whether to start charging for the service based on the presented information such as price and performance.

[0048] [Application of genetic algorithms] The offload server 1 can use evolutionary computing techniques such as GA to optimize the offload. The offload server 1 is configured as follows when GA is used. That is, the parallel processing specification unit 114 sets the number of loop statements (repetitive statements) that do not cause a compilation error based on a genetic algorithm as the gene length. The parallel 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.

[0049] The parallel processing pattern creation unit 115 prepares a specified number of gene patterns by randomly creating each gene value as 1 or 0, and the performance measurement unit 116 compiles application code specifying a parallel processing specification statement in the GPU according to each individual, and places it in the verification machine 14. The performance measurement unit 116 executes a process for performance measurement in the verification machine 14.

[0050] Here, 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 application code corresponding to that parallel processing pattern or measure its performance, but uses the same value as the performance measurement value. Furthermore, for application code in which a compilation error occurs and application code for which performance measurement does not end within a predetermined time, the performance measuring unit 116 treats the application code as timeout and sets the performance measurement value to a predetermined time (long time).

[0051] 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. From all the individuals, the executable file creation unit 117 selects those with high fitness as high-performance individuals, and performs crossover and mutation processes on the selected individuals to create the next generation of individuals. The above selection can be performed using a method such as roulette selection, which selects probabilistically according to the fitness ratio. 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.

[0052] The automatic offload operation of the offload server 1 configured as above will now be described. [Automatic offload operation] The offload server 1 of this embodiment is an example in which GPU automatic offloading of user application logic is applied as an elemental technology of environment adaptive software. FIG. 2 is a diagram showing the automatic offload process using the GA of the offload server 1. As shown in FIG. 2, the offload server 1 is applied to the elemental technology of the environment adaptive software. The offload server 1 has a control unit (automatic offload function unit) 11, a test case DB 131, an intermediate language file 132, and a verification machine 14. The offload server 1 acquires an application code 130 that is used by the user.

[0053] The offload server 1 automatically offloads function processing to accelerators such as a device 152 having a CPU-GPU and a device 153 having a CPU-FPGA.

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

[0055] <Step S12: Analyze application code> In step S12, the application code analysis unit 112 (see FIG. 1) analyzes the source code of the processing function and grasps the structure of the use of a specific library, such as a loop statement or an FFT library call.

[0056] <Step S13: Extract offloadable area> In step S13, the parallel processing specification unit 114 (see FIG. 1) identifies loop statements (repetitive statements) in the application, and compiles each repetitive statement by specifying GPU processing using OpenACC. Specifically, the offload range extraction unit 114a (see FIG. 1) identifies processing that can be offloaded to the GPU, such as loop statements, and extracts an intermediate language corresponding to the offload processing.

[0057] <Step S14: Output intermediate file> In step S14, the intermediate language file output unit 114b (see FIG. 1) outputs the intermediate language file 132. The intermediate language extraction is not a one-time process, but is repeated to try and optimize the execution for an appropriate offload area search.

[0058] <Step S15: Compile error> In step S15, the parallel processing pattern creation unit 115 (see FIG. 1) creates a parallel processing pattern that excludes loop statements that produce compilation errors from being offloaded, and specifies whether or not to process repetitive statements that do not produce compilation errors in parallel.

[0059] <Step S21: Deploy binary files> In step S21, the binary file placement unit 116a (see FIG. 1) deploys an executable file derived from the intermediate language to the verification machine 14 equipped with a GPU. The binary file placement unit 116a starts the placed file, executes assumed test cases, and measures the performance when offloaded.

[0060] <Step S22: Measure performance> In step S22, the performance measurement unit 116 (see FIG. 1) executes the arranged file and measures the performance and power consumption when offloaded. In order to make the offload area more appropriate, the performance measurement result is returned to the offload range extraction unit 114a, which extracts another pattern. Then, the intermediate language file output unit 114b attempts performance measurement based on the extracted intermediate language (see symbol a in FIG. 2). The performance measurement unit 116 repeats the performance measurement in the verification environment and finally determines the code pattern to be deployed.

[0061] As indicated by the symbol a in FIG. 2, the control unit 11 compiles each iterative statement by specifying GPU processing using OpenACC.

[0062] <Step S23: Deploy final binary files 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.

[0063] <Step S24: Extract performance test cases and run automatically> In step S24, after arranging the executable file, the performance measurement test extraction execution unit 119 extracts performance test items from the test case DB 131 and automatically executes the extracted performance tests in order to show the performance to the user.

[0064] <Step S25: Provide price and performance to a user to judge> 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.

[0065] 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.

[0066] As described above, when applied to the elemental technology of the environment adaptive software, the control unit (automatic offload function unit) 11 of the offload server 1 extracts an area to be offloaded from the source code of the application used by the user and outputs an intermediate language in order to offload functional processing (steps S11 to S15). The control unit 11 places and executes an executable file derived from the intermediate language on the verification machine 14, and verifies the offloading effect (steps 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 (steps S23 to S25).

[0067] In the above, a process flow for performing all the processes required for environment adaptation, such as code conversion, resource amount adjustment, and placement location adjustment, has been described, but it is not limited to this, and it is also possible to extract only the processes you want to perform. For example, if you want to perform only code conversion for a GPU, you can use only the necessary parts of steps S11 to S21, such as the environment adaptation function and the verification environment.

[0068] [Automatic GPU offloading using GA (genetic algorithm)] The GPU automatic offload is a process in which steps S12 to S22 in FIG. 2 are repeated for the GPU to obtain the offload code to be finally deployed in step S23.

[0069] GPUs generally do not guarantee latency, but are devices suited to increasing throughput through parallel processing. There is a wide variety of applications that run on IoT. Typical applications include encryption processing of IoT data, 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 aim is to speed up applications by automatically offloading repetitive statements to the GPU.

[0070] 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 (repetitive statements) that can be accelerated in parallel may not be the fastest. For example, if there are 10 "for" statements (repetitive 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.

[0071] In order to specify appropriate parallel regions, there have been attempts to optimize parallel execution of for statements through trial and error using the PGI compiler. However, trial and error requires a lot of work, and when providing such a service, it slows down the start of user use and increases costs.

[0072] Therefore, in this embodiment, an appropriate offload area is automatically extracted from a general-purpose program that does not assume parallelization. For this purpose, parallelizable for statements are checked first, and then performance verification trials are repeated in a verification environment using GA for the parallelizable for statements to search for an appropriate area. By narrowing down to parallelizable for statements and then retaining and recombining parallel processing patterns that can be accelerated in the form of gene parts, it is possible to efficiently search for patterns that can be accelerated from a huge number of possible parallel processing patterns.

[0073] [Search image of control unit (automatic offload function unit) 11 using Simple GA] 3 is a diagram showing a search image of the control unit (automatic offload function unit) 11 by the Simple GA. FIG. 3 shows a search image of processing and gene sequence mapping of a for statement. GA is a combinatorial optimization method that mimics the evolutionary process of living organisms. The GA flowchart is as follows: Initialization → Evaluation → Selection → Crossover → Mutation → Termination decision. In this embodiment, we use Simple GA, which is a simplified version of GA. Simple GA is a simplified GA that uses only 1 and 0 as genes, roulette wheel selection, one-point crossover, and mutation by reversing the value of one gene.

[0074] <Initialization> In the initialization, after checking whether all for statements in the application code can be parallelized, parallelizable for statements are mapped to gene arrays. If GPU processing is used, the value is set to 1, and if not, the value is set to 0. A specified number of individuals M is prepared for genes, and 1 or 0 is randomly assigned to each for statement. Specifically, the control unit (automatic offload function unit) 11 (see Fig. 1) acquires the application code 130 (see Fig. 2) used by the user, and checks whether or not for statements can be executed in parallel from the code patterns 141 of the application code 130, as shown in Fig. 3. As shown in Fig. 3, if five for statements are found from the code pattern 141 (see symbol b in Fig. 3), one digit is randomly assigned to each for statement, in this case five digits of 1 or 0 are randomly assigned to the five for statements. For example, 0 is assigned when processing is done by the CPU, and 1 is assigned when output to the GPU. However, at this stage, 1 or 0 is randomly assigned. The code corresponding to the gene length is 5 digits, and the 5-digit gene length code is 2 5 =32 patterns, for example, 10001, 10010, .... In Fig. 3, the circles (○ marks) in the code pattern 141 are shown as an image of the code.

[0075] <Evaluation> In the evaluation, deployment and performance measurement are performed (see symbol c in FIG. 3). That is, the performance measurement unit 116 (see FIG. 1) compiles the code corresponding to the genes, deploys it on the verification machine 14, and executes it. The performance measurement unit 116 performs benchmark performance measurement. It increases the fitness of genes of patterns with good performance and power consumption (parallel processing patterns).

[0076] Here, a new process is added in which patterns that can be processed in a short time by measurement are regarded as genes with high fitness, while power consumption is also measured and low-power patterns are also regarded as having high fitness. For example, an evaluation value shown in formula (1) is introduced, and based on this evaluation value, the shorter the processing time and the lower the power consumption, the higher the fitness of the gene pattern. For example, (processing time) -1 / 2 is 0.1, (power consumption) -1 / 2 is 0.1, the score is 0.1 x 0.1 = 0.01. If other scores are greater than this 0.01, the fitness with the higher score is used.

[0077]

number

[0078] <Select> In the selection, high performance and low power code patterns are selected based on the fitness (see symbol d in FIG. 3). 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. 3 shows, as a search image, that the number of circles (◯) in Select code patterns 142 has been reduced to three.

[0079] <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 third digit of the above five-digit code.

[0080] <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. However, in order to reduce the amount of calculation, mutation may not be performed.

[0081] <End Judgment> As shown in Figure 3, generate next generation code patterns after crossover & mutation (see symbol e in Figure 3). 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 the three fastest ones are selected: 10010, 01001, and 00101. These three are then recombined using GA in the next generation, for example by crossing the first and second patterns to create a new pattern (parallel processing pattern) 11011. At this point, a mutation is automatically applied to the recombined pattern, such as changing 0 to 1. The above is repeated to find the fastest pattern. A designated generation (for example, the 20th generation) is decided, and the pattern remaining in the final generation is designated as the final solution.

[0082] <Deployment> 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.

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

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

[0085] For example, if C / C++ is used, the C / C++ code is analyzed to find for statements. When a for statement is found, it writes to the for statement using the OpenACC parallel processing syntax #pragma acc kernels, #prama acc parallel loop, or #prama acc parallel loop vector. In detail, for statements are inserted one by one into #pragma acc kernels, #prama acc parallel loop, or #prama acc parallel loop vector, and compiled. If an error occurs, the for statement cannot be processed by the GPU in the first place, so it is excluded.

[0086] In this way, the remaining for statements are found. The one that does not produce errors is then taken as the length (gene length). If there are five error-free for statements, the gene length is 5, and if there are 10 error-free for statements, the gene length is 10. Note that parallel processing is not possible when there is data dependency, such as the previous processing being used for the next processing. The above is the preparation stage. Next, GA processing will be performed.

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

[0088] Deploy it on the verification machine 14 and benchmark it - for example, if it is image processing, benchmark it with that image processing, and evaluate the fitness higher as the processing time is shorter. For example, the processing time is raised to the power of -1 / 2, so that a processing time of 1 second is 1, a processing time of 100 seconds is 0.1, and a processing time of 0.01 seconds is 10. Those with high adaptability are selected, for example, 3 to 5 are selected from 10, and a new code pattern is created by rearranging them. At this time, the same code pattern as before may be created during the creation. In that case, since there is no need to perform the same benchmark, the same data as before is used. In this embodiment, the code pattern and its processing time are stored in the storage unit 13. Above, we have explained the search image of the control unit (automatic offload function unit) 11 by the Simple GA. Next, we will explain the batch processing method of data transfer.

[0089] [Batch processing method for data transfer] <Basic Concept> In order to reduce CPU-GPU transfers, in addition to transferring nested loop variables at as high a level as possible, the present invention consolidates the timing of many variable transfers and further reduces transfers that are automatically transferred by the compiler. To reduce transfers, variables that can be transferred to the GPU at the same time are transferred not only on a nesting basis, but also in a batch. For example, unless the variables are ones that process the results of GPU processing in the CPU and then process them again on the GPU, it is possible to send variables defined on the CPU that are used in multiple loop statements to the GPU in a batch before GPU processing begins, and then return them to the CPU after all GPU processing is complete.

[0090] During code analysis, the reference relationships between loops and variables are identified, and from the results, for variables defined in multiple files, the GPU and CPU processing is not nested, and for variables whose CPU and GPU processing can be separated, a specification is made to transfer them together using the OpenACC data copy statement. The data is transferred in a batch before GPU processing begins, and variables that do not need to be transferred at the time of loop statement processing are explicitly indicated as not needing to be transferred using data present. When transferring data between the CPU and GPU, a temporary area is created (#pragma acc declare create), the data is stored in the temporary area, and the transfer is then instructed by synchronizing the temporary area (#pragma acc update).

[0091] <Comparative Example> First, a comparative example will be described. The comparative examples are a normal CPU program (see FIG. 4), a simple GPU usage (see FIG. 5), and nested bundling (Non-Patent Document 2) (see FIG. 6). <1> ~ <4> etc. are added for the convenience of explanation (the same applies to other drawings and their explanations). The loop statement of the normal CPU program shown in Figure 4 is written in the CPU program, <1> Loop [for(i=0; i<10; i++)] { } Among them, <2> Loop [for(j=0; j<20; j++)] { The symbol f in Figure 4 is <2> This is the setting of variables a and b in the loop. Also, <3> Loop [for(k=0; k<30; k++)] { } and, <4> Loop [for(l=0; l<40; l++)] { } The symbol g in FIG. 4 is the <3> The setting of variables c and d in the loop. The symbol h in Figure 4 is <4> This is the setting of variables e and f in the loop. The normal CPU program shown in Figure 4 is executed by the CPU (without using the GPU).

[0092] Figure 5 is a diagram showing a loop statement when the normal CPU program shown in Figure 4 is used to transfer data from the CPU to the GPU using a simple GPU. There are two types of data transfer: data transfer from the CPU to the GPU, and data transfer from the GPU to the CPU. Below, we will take data transfer from the CPU to the GPU as an example. The simple GPU-utilizing loop statement shown in Figure 5 is written in the CPU program. <1> Loop [for(i=0; i<10; i++)] { } Among them, <2> Loop [for(j=0; j<20; j++)] { There is. Furthermore, as shown in FIG. 5 (i), <1> Loop [for(i=0; i<10; i++)] { Above the}, the OpenACC directive #pragma acc kernels (parallel processing specification statement) is used to specify parallel processing parts such as for statements by the PGI compiler. As shown in the dashed box containing the symbol i in Figure 5, data is transferred from the CPU to the GPU by #pragma acc kernels. In this case, a and b are transferred at this timing, so they are transferred 10 times.

[0093] Also, as shown by reference j in FIG. <3> Loop [for(k=0; k<30; k++)] { Above the}, the OpenACC directive #pragma acc kernels is used to specify parallel processing parts such as for statements using the PGI compiler. As shown in the dashed box including the symbol j in Figure 5, c and d are transferred at this timing by #pragma acc kernels.

[0094] Where: <4> Loop [for(l=0; l<40; l++)] { Do not specify #pragma acc kernels above the}. This loop is not GPU processed because it is inefficient to use GPU.

[0095] FIG. 6 is a diagram showing loop statements in the case of data transfer from the CPU to the GPU and from the GPU to the CPU by nested consolidation (Non-Patent Document 2). In the loop statement shown in Fig. 6, a data transfer directive from the CPU to the GPU, in this case #pragma acc data copyin(a,b) of the copyin clause for variables a and b, is inserted at the position indicated by symbol k in Fig. 6. Note that parentheses () are added to copyin(a,b) in this specification for notational reasons. The same notation is used for copyout(a,b) and datacopyin(a,b,c,d) described later. The above #pragma acc data copyin(a,b) applies to the top-level loop (here, <1> Loop [for(i=0; i<10; i++)] { }) Since a and b are transferred at the timing shown in the dashed-dotted box including the symbol k in FIG. 6, one transfer occurs.

[0096] In addition, in the loop statement shown in FIG. 6, a data transfer instruction line from the GPU to the CPU, in this case #pragma acc data copyout(a,b) of the copyout clause for variables a and b, is inserted at the position indicated by symbol l in FIG. 6. The above #pragma acc data copyout(a,b) is <1> Loop [for(i=0; i<10; i++)] { } is specified at the bottom.

[0097] In this way, when transferring data from the CPU to the GPU, the data transfer is explicitly specified by inserting #pragma acc data copyin(a,b) in the copyin clause of variable a at the above-mentioned position. This allows data transfer to be performed collectively in a loop as high up as possible, and avoids inefficient data transfers for each loop, as in the simple GPU-utilizing loop statement shown in Figure 5.

[0098] <Embodiment> Next, the present embodiment will be described. 《Indicate variables that do not need to be transferred using data present》 In this embodiment, for variables defined in multiple files, GPU processing and CPU processing are not nested, and variables that can be separated into CPU processing and GPU processing are specified to be transferred in a lump using the data copy statement of OpenACC. In addition, variables that are transferred in a lump and do not need to be transferred at that time are explicitly indicated using data present.

[0099] Fig. 7 is a diagram showing a loop statement by transfer bundling during CPU-GPU data transfer in this embodiment, which corresponds to the nest bundling in Fig. 6 of the comparative example. In the loop statement shown in FIG. 7, a data transfer instruction line from the CPU to the GPU, in this case #pragma acc datacopyin(a,b,c,d) in the copyin clause for variables a,b,c,d, is inserted at the position indicated by symbol m in FIG. 7. The above #pragma acc data copyin(a,b,c,d) applies to the top-level loop (here, <1> Loop [for(i=0; i<10; i++)] { })

[0100] In this way, for variables defined in multiple files, where GPU and CPU processing is not nested and variables can be separated into CPU and GPU processing, the OpenACC data copy statement #pragma acc data copyin(a,b,c,d) is used to specify that the variables be transferred in a batch. Since a, b, c, and d are transferred at the timing shown in the dashed-dotted box including the symbol m in FIG. 7, one transfer occurs.

[0101] Then, the variables are transferred together using the above #pragma acc data copyin(a,b,c,d), and variables that do not need to be transferred at that time are specified using the data present statement #pragma acc data present (a,b), which explicitly indicates that the variables are already in the GPU at the time shown in the dashed-dotted-line frame including the symbol n in Figure 7.

[0102] The variables are transferred together using the above #pragma acc data copyin(a,b,c,d), and variables that do not need to be transferred at that time are specified using the data present statement #pragma acc data present(c,d), which explicitly indicates that the variables are already in the GPU at the time shown in the dashed-dotted-line frame including the symbol o in Figure 7. <1> , <3> When the loop is processed by the GPU and the GPU processing is completed, the data transfer instruction from the GPU to the CPU, in this case #pragma acc datacopyout(a,b,c,d) in the copyout section of variables a,b,c,d, is changed to <3> Insert at position p where the loop ended.

[0103] By specifying the batch transfer, variables that can be batch transferred are transferred in a batch, and variables that have already been transferred and do not need to be transferred are explicitly indicated using data present, reducing the transfer and making the offloading method more efficient. However, even if OpenACC instructs transfer, depending on the compiler, the compiler may automatically determine and transfer. Automatic transfer by the compiler is an event that, unlike OpenACC instructions, is automatically transferred depending on the compiler, even if transfer between the CPU and GPU is not actually necessary.

[0104] Temporary data storage 8 is a diagram showing a loop statement by bundling transfer during CPU-GPU data transfer in this embodiment. Fig. 8 corresponds to the nest bundling and explicit specification of variables not requiring transfer in Fig. 7. In the loop statement shown in Figure 8, the OpenACC declare create statement #pragma acc declare create, which creates a temporary area during CPU-GPU data transfer, is specified at the position indicated by symbol q in Figure 8. As a result, a temporary area is created (#pragma acc declare create) during CPU-GPU data transfer, and the data is stored in the temporary area.

[0105] In addition, the transfer is instructed by specifying the OpenACC declare create statement #pragma acc update to synchronize the temporary area at the position indicated by the symbol r in Figure 8.

[0106] In this way, by creating a temporary area, initializing parameters in the temporary area, and using it for CPU-GPU transfer, unnecessary CPU-GPU transfer is blocked. This makes it possible to reduce transfers that are unintended by OpenACC instructions but degrade performance.

[0107] [GPU offload processing] The data transfer batch processing technique described above makes it possible to extract loop statements suitable for offloading and to avoid inefficient data transfer. However, even if the above data transfer batch processing method is used, there are some programs that are not suitable for GPU offloading. Effective GPU offloading requires that the number of loops of the offloaded process is large.

[0108] Therefore, in this embodiment, a profiling tool is used to investigate the number of loops as a preliminary step to a full-scale offload processing search. By using a profiling tool, the number of executions of each line can be investigated, so that it is possible to pre-allocate programs, for example, programs with loops of 50 million or more times as targets for offload processing search. A specific description is given below (some of the content overlaps with that described in FIG. 2).

[0109] In this embodiment, first, an application for searching an offload processor is analyzed, and loop statements such as for, do, while, etc. are identified. Next, a sample process is executed, and the number of loops of each loop statement is investigated using a profiling tool. Depending on whether or not there are loops that are equal to or greater than a certain value, it is determined whether or not to perform a full-scale offload processor search.

[0110] When it is decided to conduct a full-scale search, the GA process begins (see Figure 2). In the initialization step, after checking whether all loop statements in the application code can be processed in parallel, parallelizable loop statements are mapped to the gene array as 1 if they can be processed by GPU, and 0 if they cannot. A specified number of genes are prepared, and each gene value is randomly assigned a value of 1 or 0.

[0111] Here, in the code corresponding to the gene, an explicit instruction for data transfer (#pragma acc data copyin / copyout / copy) is added based on the variable data reference relationship within the loop statement specified to be processed by GPU.

[0112] In the evaluation step, the code corresponding to the genes is compiled, deployed to a verification machine, and executed to measure the benchmark performance. The fitness of genes with good performance patterns is increased. As described above, parallel processing instruction lines (e.g., see symbol f in Figure 4) and data transfer instruction lines (e.g., see symbol h in Figure 4, symbol i in Figure 5, and symbol k in Figure 6) are inserted into the code corresponding to the genes.

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

[0114] When the mutation step is completed and the specified number of genes of the next generation are created, an explicit instruction for data transfer is added as in the initialization step, and the evaluation, selection, crossover, and mutation steps are repeated.

[0115] Finally, in the termination determination step, the process is terminated after the specified number of generations and repetitions, and the gene with the highest fitness is taken as the solution. The code pattern with the highest performance corresponding to the gene with the highest fitness is redeployed to the production environment and provided to the user.

[0116] The following describes the implementation of the offload server 1. This implementation is intended to confirm the effectiveness of this embodiment. [implementation] We explain the implementation of automatic offloading of C / C++ applications using the general-purpose PGI compiler. In this implementation, since the purpose is to confirm the effectiveness of GPU automatic offloading, the target application is a C / C++ language application, and the GPU processing itself is explained using the conventional PGI compiler.

[0117] C / C++ is one of the most popular languages ​​for developing OSS (Open Source Software) and proprietary software, and many applications are developed in C / C++. In order to check the offloading of applications used by general users, we use general-purpose OSS applications such as encryption and image processing.

[0118] GPU processing is performed by the PGI compiler. The PGI compiler is a compiler for C / C++ / Fortran that interprets OpenACC. In this embodiment, parallel processing units such as for statements are specified by the OpenACC directive #pragma acc kernels (parallel processing specification statement). This allows bytecode for GPUs to be extracted and executed to enable GPU offloading. Furthermore, an error is issued when there is a dependency between data in a for statement and processing cannot be performed in parallel, or when multiple levels of nested for statements with different levels are specified. In addition, explicit data transfer instructions are possible using directives such as #pragma acc data copyin / copyout / copy.

[0119] In accordance with the above #pragma acc kernels (parallel processing specification), an explicit data transfer instruction is given by inserting #pragma acc data copyout(a[…]) in the OpenACC copyin clause at the position mentioned above.

[0120] <Implementation Overview> The outline of the implementation will be explained. The implementation does the following: Before starting the process of the flow in Figures 9A-B below, prepare the C / C++ application to be accelerated and a benchmark tool to measure its performance.

[0121] In the implementation, when a request to use a C / C++ application is received, the system first analyzes the code of the C / C++ application to find for statements and to understand the program structure, such as the variable data used in the for statements. For syntax analysis, the LLVM / Clang syntax analysis library is used.

[0122] In the implementation, first, a benchmark is run to obtain an estimate of whether the application will benefit from GPU offloading, and the number of loops in the for statement identified by the syntax analysis above is determined. To determine the number of loops, tools such as gcov from GNU Coverage are used. Known profiling tools include "GNU Profiler (gprof)" and "GNU Coverage (gcov)." Either can be used, as both can investigate the number of times each line is executed. The number of executions can be set to only target applications with a loop count of 10 million or more, for example, but this value can be changed.

[0123] General-purpose applications for CPUs are not implemented with parallelization in mind. Therefore, it is necessary to first eliminate for statements that cannot be processed by the GPU. Therefore, for each for statement, we try to insert the #pragma acc kernels, #prama acc parallel loop, or #prama acc parallel loop vector directives for GPU processing, and determine whether an error occurs during compilation. There are several types of compilation errors. These include when an external routine is called in the for statement, when different hierarchical levels are specified in a nested for statement, when there is a process that exits the for statement midway using break, and when there is data dependency in the data in the for statement. There are many types of compilation errors depending on the application, and there may be other errors besides those mentioned above, but compilation errors are not processed and the #pragma directive is not inserted.

[0124] Compilation errors are difficult to deal with automatically, and even if they are dealt with, they are often ineffective. In the case of external routine calls, they can sometimes be avoided by using #pragma acc routine, but many external calls are libraries, and even if they are included in the GPU processing, the calls become a bottleneck and performance is not achieved. Since the for statements are tried one by one, no compilation errors occur for nesting errors. Also, if you exit midway using break, etc., the number of loops must be fixed for parallel processing, and program modification is required. If there is data dependency, parallel processing itself is not possible in the first place.

[0125] Here, if the number of loop statements that do not cause an error even when processed in parallel is a, a becomes the gene length. The gene 1 corresponds to the presence of a parallel processing directive, and 0 corresponds to the absence of a parallel processing directive, and the application code is mapped to the gene of length a.

[0126] Next, a gene array for the specified number of individuals is prepared as initial values. As explained in Figure 3, each gene value is created by randomly assigning 0 and 1. Depending on the prepared gene array, when the gene value is 1, the directives \#pragma acc kernels, \#pragma acc parallel loop, and \#pragma acc parallel loop vector that specify GPU processing are inserted into the C / C++ code. The reason that single loops etc. are not made parallel is that for the same processing, kernels have better performance from the PGI compiler's perspective. At this stage, the part of the code corresponding to a certain gene that will be processed by the GPU is decided.

[0127] The C / C++ code with parallel processing and data transfer directives inserted is compiled with the PGI compiler on a machine with a GPU. The compiled executable is deployed and the performance and power usage are measured with a benchmark tool.

[0128] After measuring the benchmark performance for all individuals, the fitness of each gene sequence is set according to the benchmark processing time and power consumption. Individuals to be left are selected according to the set fitness. The selected individuals are subjected to GA processing including crossover, mutation, and direct copying to create the next generation population.

[0129] The next generation of individuals are subjected to directive insertion, compilation, performance measurement, fitness setting, selection, crossover, and mutation processes. If a gene with the same pattern as before is generated during the GA process, compilation and performance measurement are not performed for that individual, and the same measurement value as before is used.

[0130] After the GA process is completed for the specified number of generations, the C / C++ code with directives that corresponds to the gene sequence with the best performance is taken as the solution.

[0131] Among these, the number of individuals, number of generations, crossover rate, mutation rate, fitness setting, and selection method are GA parameters and are specified separately. By automating the above processes, the proposed technology makes it possible to automate GPU offloading, which previously required the time and skills of a specialist engineer.

[0132] 9A-B are flowcharts outlining the operation of the above-described implementation, and FIG. 9A and FIG. 9B are connected by a connector. The following process is performed using the OpenACC compiler for C / C++.

[0133] <Code Analysis> In step S101, the application code analysis unit 112 (see FIG. 1) performs code analysis of a C / C++ application.

[0134] <Loop statement identification> In step S102, the parallel processing specification unit 114 (see FIG. 1) identifies loop statements and reference relationships in the C / C++ application.

[0135] <Parallel processing of loop statements> In step S103, the parallel processing specification unit 114 checks whether each loop statement can be processed by the GPU (#pragma acc kernels).

[0136] <Repeat loop statement> The control unit (automatic offload function unit) 11 repeats the processes of steps S105 to S116 between the loop start point of step S104 and the loop end point of step S117 for the number of loop statements.

[0137] <Repeat the number of loops (part 1)> The control unit (automatic offload function unit) 11 repeats the processes of steps S106-S107 between the loop start point of step S105 and the loop end point of step S108 for the number of loop statements. In step S106, the parallel processing specification unit 114 specifies GPU processing (#pragma acc kernels) in OpenACC for each loop statement and compiles it. In step S107, when an error occurs, the parallel processing specification unit 114 checks the possibility of GPU processing with the next directive (#pragma acc parallel loop).

[0138] <Repeat the number of loops (part 2)> The control unit (automatic offload function unit) 11 repeats the processing of steps S110-S111 between the loop start point of step S109 and the loop end point of step S112 for the number of loop statements. In step S110, the parallel processing specification unit 114 specifies GPU processing (#pragma acc parallel loop) in OpenACC for each loop statement and compiles it. In step S111, when an error occurs, the parallel processing specification unit 114 checks the possibility of GPU processing with the following directive (#pragma acc parallel loop vector).

[0139] <Repeat the number of loops (part 3)> The control unit (automatic offloading function unit) 11 repeats the processes of steps S114 - S115 for the number of loop statements between the loop start of step S113 and the loop end of step S116. In step S114, the parallel processing specifying unit 114 specifies GPU processing (#pragma acc parallel loop vector) for each loop statement and compiles it with OpenACC. In step S115, when there is an error, the parallel processing specifying unit 114 removes the GPU processing instruction statement from the loop statement.

[0140] <Count of for loop statements> In step S118, the parallel processing specifying unit 114 counts the number of for loop statements without compilation errors and uses it as the gene length.

[0141] <Prepare specified number of individual patterns> Next, as an initial value, the parallel processing specifying unit 114 prepares a gene sequence of the specified number of individuals. Here, it is created by randomly assigning 0 and 1. In step S119, the parallel processing specifying unit 114 maps the C / C++ application code to genes and prepares the specified number of individual patterns. According to the prepared gene sequence, when the value of the gene is 1, a directive specifying parallel processing is inserted into the C / C++ code (see #pragma directive in Figure 3 for example).

[0142] The control unit (automatic offloading function unit) 11 repeats the processes of steps S121 - S130 for the specified number of generations between the loop start of step S120 and the loop end of step S131. Also, in the above repetition of the specified number of generations, the processes of steps S122 - S125 are further repeated for the specified number of individuals between the loop start of step S121 and the loop end of step S126. That is, within the repetition of the specified number of generations, the repetition of the specified number of individuals is processed in a nested state.

[0143] <Specify data transfer> In step S122, the data transfer specification unit 113 specifies a data transfer using explicit directive lines (#pragma acc data copy / copyin / copyout / present and #pragma acc declarecreate, #pragma acc update) based on the variable reference relationship.

[0144] <Compile> In step S123, the parallel processing pattern creation unit 115 (see FIG. 1) compiles the C / C++ code with directives specified according to the gene pattern using a PGI compiler. That is, the parallel processing pattern creation unit 115 compiles the created C / C++ code using a PGI compiler on the verification machine 14 equipped with a GPU. Here, a compilation error may occur when multiple nested for statements are specified in parallel, etc. In this case, it is treated the same as when the processing time during performance measurement times out.

[0145] In step S124, the performance measurement unit 116 (see FIG. 1) deploys the executable file to the verification machine 14 equipped with a CPU-GPU. In step S125, the performance measurement unit 116 executes the arranged binary file and measures the benchmark performance when offloaded.

[0146] Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same value is used. In other words, if a gene with the same pattern as before is generated during the GA process, compilation or performance measurement is not performed for that individual, and the same measured value as before is used.

[0147] In step S127, the power consumption measuring unit 116b (see FIG. 1) measures the processing time and the amount of power consumption.

[0148] In step S128, the evaluation value setting unit 116c (see FIG. 1) sets an evaluation value based on the measured processing time and power consumption.

[0149] In step S129, the executable file creation unit 117 (see FIG. 1) evaluates the individual so that the higher the evaluation value, the higher the fitness, and selects the individual with the highest performance. The executable file creation unit 117 selects a pattern with a short time and low power consumption as a solution from the multiple measured patterns.

[0150] In step S130, the executable file creation unit 117 performs crossover and mutation processes on the selected individuals to create the next generation individuals. The executable file creation unit 117 performs compilation, performance measurement, fitness setting, selection, crossover, and mutation processes on the next generation individuals. That is, after benchmark performance measurement for all individuals, the fitness of each gene sequence is set according to the benchmark processing time. Individuals to be left are selected according to the set fitness. The executable file creation unit 117 performs GA processing such as crossover, mutation, and direct copy processing on the selected individuals to create the next generation population.

[0151] In step S132, the executable file creation unit 117 determines, after the GA process for the designated number of generations is completed, the C / C++ code corresponding to the gene sequence with the highest performance (parallel processing pattern with the highest performance) as a solution.

[0152] <GAのパラメータ> The above-mentioned number of individuals, number of generations, crossover rate, mutation rate, fitness setting, and selection method are parameters of the GA. The parameters of the GA may be set, for example, as follows. The parameters and conditions for the Simple GA to be executed can be, for example, as follows: Gene length: Number of parallelizable loop statements Number of individuals M: gene length or less Number of generations T: gene length or less Conformity: (Processing time) (-1 / 2) ×(Electricity usage) (-1 / 2)

[0153] With this setting, the shorter the benchmark processing time, the higher the fitness. Also, by making the fitness include the (-1 / 2)th power of the processing time, it is possible to prevent the fitness of a specific individual with a short processing time from becoming too high, thereby narrowing the search range. Also, if the performance measurement does not end within a certain time, a timeout is set and the fitness is calculated assuming a processing time of 1000 seconds or the like (long time). This timeout time can be changed according to the performance measurement characteristics. Selection: Roulette selection However, elite preservation is also performed in which the genes with the highest fitness in one generation are preserved in the next generation without undergoing crossover or mutation. Crossover rate Pc:0.9 Mutation rate Pm: 0.05

[0154] <Cost performance> This section describes the cost-performance of the automatic offload function. Looking only at the hardware price of a GPU board such as NVIDIA Tesla, the price of a machine equipped with a GPU is about twice that of a machine with only a normal CPU. However, in general, the cost of a data center, etc., is 1 / 3 or less for hardware and system development, more than 1 / 3 for electricity charges and operational costs such as maintenance and operation systems, and about 1 / 3 for other costs such as service orders. In this embodiment, the performance of time-consuming processes in applications such as encryption processing and image processing can be improved by more than 2 times. Therefore, even if the server hardware price itself doubles, sufficient cost effectiveness can be expected.

[0155] In this embodiment, applications that have many loops and take a long time to execute are identified in advance using gcov, gprof, etc., and offloading is attempted. This makes it possible to find applications that can be efficiently accelerated.

[0156] <Time until start of production service> We will explain the time until the actual service can be used. If one performance measurement from compilation takes about three minutes, a GA with 20 individuals and 20 generations will take up to 20 hours to find a solution, but since compilation and measurement of the same genetic pattern as before are omitted, it will be completed in less than eight hours. In reality, many cloud, hosting, and network services take about half a day to start using the service. In this embodiment, automatic offloading is possible within half a day, for example. Therefore, if automatic offloading within half a day is possible, and trial use can be performed at first, it is expected that user satisfaction will be sufficiently increased.

[0157] In order to search for the offload part in a shorter time, it is possible to measure the performance in parallel for the number of individuals using multiple verification machines. Adjusting the timeout time according to the application can also shorten the time. For example, if the offload process takes twice as long as the execution time on the CPU, a timeout can be set. In addition, the more individuals and generations there are, the higher the possibility of finding a high-performance solution. However, when maximizing each parameter, compilation and performance benchmarking must be performed for the number of individuals x the number of generations. This takes time before the actual service can be used. In this embodiment, the GA is performed with a small number of individuals and generations, but by setting the crossover rate Pc to a high value of 0.9 and searching a wide range, a solution with a certain level of performance can be found quickly.

[0158] [Expand directive] In this embodiment, directives are expanded to increase the number of applicable applications. Specifically, directives specifying GPU processing are expanded to include the parallel loop directive and the parallel loop vector directive in addition to the kernels directive. In the OpenACC standard, kernels are used for single loops and tightly nested loops. Parallel loops are used for loops including non-tightly nested loops. Parallel loop vector is used for loops that cannot be parallelized but can be vectorized. Here, a tightly nested loop is a simple loop in which, for example, when two loops that increment i and j are nested in a nested loop, processing using i and j is performed in the lower loop but not in the upper loop. Also, in implementations such as the PGI compiler, the difference between kernels and parallel is that the compiler decides whether to parallelize, while the programmer decides whether to parallelize.

[0159] Therefore, in this embodiment, kernels are used for single, tightly nested loops, parallel loops are used for non-tightly nested loops, and parallel loop vectors are used for loops that cannot be parallelized but can be vectorized. There is a concern that using the parallel directive here will result in less reliable results than with kernels. However, it is assumed that sample tests will be performed on the final offload program, the difference in results with the CPU will be checked, and the results will be shown to the user for confirmation. In the first place, CPUs and GPUs have different hardware, so there are differences in the number of significant digits and rounding errors, and so even just with kernels, it is necessary to check the difference in results with the CPU.

[0160] [evaluation] Explain the evaluation. In the [automatic GPU offloading of loop statements] of this embodiment, when determining the evaluation value of the measurement pattern, a method is added to the existing implementation tool so that the lower the power consumption, the higher the evaluation value is, and offloading is performed to confirm that low power consumption can be achieved.

[0161] <Evaluation target> The evaluation target is the Himeno benchmark for fluid calculation in the [GPU automatic offloading of loop statements] of this embodiment, and the MRI-Q benchmark for MRI (Magnetic Resonance Imaging) image processing in the [FPGA automatic offloading of loop statements] of the second embodiment described later.

[0162] The Himeno benchmark is a 3L benchmark software for measuring the performance of incompressible fluid analysis, and solves the Poisson equation using the Jacobi iteration method. The Himeno benchmark is also available in C and Fortran, but in order to measure power consumption, we decided to use Python, which takes a certain amount of calculation time, and wrote the processing logic in Python. Data is calculated on a large (maximum) 512 x 256 x 256 grid. CPU processing is handled by Python's Numpy, and GPU processing is handled via the Cupy library, which offloads the Numpy Interface to the GPU. Incidentally, MRI-Q will be described later in the evaluation of the second embodiment.

[0163] <Evaluation method> The code of the target application is input, and loop offloading recognized by Clang etc. is attempted for the destination GPU or FPGA to determine the offloading pattern. During this process, the processing time and power consumption are measured. For the final offloading pattern, the change in power consumption over time is obtained, and the power reduction compared to processing entirely on the CPU is confirmed. In the "GPU automatic offloading of loop statements" of this embodiment, an appropriate pattern is selected by GA. In the "FPGA automatic offloading of loop statements" of the second embodiment described later, GA is not performed, and the number of measurement patterns is narrowed down to four patterns using arithmetic strength, etc. Offload target loop statement: Himeno benchmark 13 Pattern compatibility: The evaluation value shown in formula (1), that is, (processing time) -1 / 2 ×(Electricity usage) -1 / 2 As shown in equation (1), the lower the processing time and power consumption, the higher the evaluation value and the higher the degree of compatibility.

[0164] <Evaluation environment> In this embodiment, the [GPU automatic offload of loop statements] uses a GeForce RTX 2080 Ti. The power consumption is measured by measuring the GPU power with NVIDIA's nvidia-smi (registered trademark) and the CPU power with s-tui (registered trademark). In the second embodiment, the [FPGA automatic offload of loop statements] described later uses an Intel PAC with Intel Arria10 GXFPGA (registered trademark). The power consumption is measured by measuring the power consumption of the entire server using ipmitool (registered trademark) of the IPMI (Intelligent Platform Management Interface) of the Dell (registered trademark) server.

[0165] <Results and Discussion> FIG. 10 shows the power consumption (Watt) and processing time when the Himeno benchmark was offloaded to the GPU. The symbol s in FIG. 10 shows a comparison of the power consumption Watts during each processing time for "all CPU processing" on the left side of FIG. 10 and "CPU and GPU processing" on the right side of FIG. In the Himeno benchmark, the processing time for "CPU and GPU processing" on the right side of Figure 10 is reduced from 153 seconds to 19 seconds compared to "all CPU processing" on the left side of Figure 10, but the power consumption Watts has increased from a maximum of about 26.9 W for "all CPU processing" to a maximum of about 116.2 W for "CPU and GPU processing." As a result, the Watt sec for "CPU and GPU processing" is 2071 Watt sec, roughly half the 4077 Watt sec for "all CPU processing."

[0166] We also confirmed that multiple applications could be powered down. In this embodiment, the power consumption (Watt) increases with [GPU automatic offloading of loop statements], but the overall processing time is shortened, resulting in low power consumption overall.

[0167] As described above, in the [GPU automatic offloading of loop statements] of this embodiment, an evolutionary computation method that includes power consumption in the fitness level, automatic speedup by reducing CPU-GPU transfer, and low power consumption by evaluating power consumption are realized. In particular, when measuring in a verification environment during automatic GPU offloading, power consumption is acquired in addition to processing time, and a pattern with a short time and low power consumption is considered to have a high fitness level, and low power consumption is incorporated into the automatic code conversion. As described in the evaluation of FIG. 10, low power consumption was confirmed through the automatic offloading of existing applications, and its effectiveness was confirmed.

[0168] Next, the offload server 1A and the like in the second embodiment of the present invention will be described. The second embodiment is an example of application to automatic FPGA offloading of loop statements. In this embodiment, an example in which the present invention is applied to an FPGA (Field Programmable Gate Array) as a PLD (Programmable Logic Device) will be described, but the present invention is applicable to programmable logic devices in general.

[0169] (Explanation of the principle) Since it is difficult to predict which loops will be offloaded to FPGAs to speed up the process, we propose to automatically measure in a verification environment, just like with GPUs. However, since it takes several hours to compile OpenCL and run it on an actual FPGA, it is not possible to repeatedly measure using GA with GPU automatic offloading because the processing time would be enormous. Therefore, we narrow down the loop statements that are candidates for offloading to FPGAs and then perform measurements. Specifically, we use an arithmetic strength analysis tool such as ROSE (registered trademark) to extract loop statements with high arithmetic strength from the discovered loop statements. Furthermore, we use a profiling tool such as gcov (registered trademark) to extract loop statements with a large number of loop counts.

[0170] Candidates for loop statements with high arithmetic intensity or loop counts are converted to OpenCL. When converting to OpenCL, the CPU processing program is divided into the kernel (FPGA) and host (CPU) according to the OpenCL grammar. The OpenCL created for each candidate loop statement is precompiled to find loop statements with high resource efficiency. Since the resources to be created are known during the compilation process, the loop statements are further narrowed down to those that use a sufficiently small amount of resources. Several candidate loop statements remain, and these are used to measure performance and power consumption. The selected single loop statements are compiled and measured, and for those single loop statements that can be further accelerated, combination patterns are created and a second measurement is performed. From the multiple patterns measured, the pattern that is short in time and consumes the least power is selected as the solution.

[0171] Regarding FPGA offloading of loop statements, we first narrow down the options using arithmetic strength, etc., then perform measurements and improve the evaluation value of low-power patterns to automatically increase speed and reduce power consumption.

[0172] Second embodiment 11 is a functional block diagram showing a configuration example of an offload server 1A according to a second embodiment of the present invention. In describing this embodiment, the same components as those in FIG. 1 are given the same reference numerals and descriptions of overlapping parts are omitted. The offload server 1A is a device that automatically offloads specific processing of an application to an accelerator. In addition, the offload server 1A can be connected to an emulator. As shown in FIG. 11, the offload server 1A includes a control unit 21, an input / output unit 12, a storage unit 13, and a verification machine 14 (accelerator verification device).

[0173] The control unit 21 is an Automatic Offloading function that controls the entire offload server 1A. The control unit 21 is realized, for example, by a CPU (not shown) expanding and executing a program (offloading program) stored in the storage unit 13 in the RAM.

[0174] The control unit 21 includes an application code specifying unit (Specify application code) 111, an application code analyzing unit (Analyze application code) 112, a PLD processing specifying unit 213, an arithmetic intensity calculating unit 214, a PLD processing pattern creating unit 215, a performance measuring unit 116, an execution file creating unit 117, a production environment deployment unit (Deploy final binary files to production environment) 118, a performance measurement test extraction execution unit (Extract performance test cases and run automatically) 119, and a user providing unit (Provide price and performance to a user to judge) 120.

[0175] <PLD processing specifying unit 213> The PLD processing specifying unit 213 identifies the loop statements (repetition statements) of the application, and for each identified loop statement, creates and compiles a plurality of offloading processing patterns that specify pipeline processing and parallel processing in the PLD using OpenCL. The PLD processing specifying unit 213 includes an offloading range extraction unit (Extract offload able area) 213a and an intermediate language file output unit (Output intermediate file) 213b.

[0176] The offloading range extraction unit 213a identifies processing that can be offloaded to the FPGA, such as loop statements and FFT, and extracts intermediate language corresponding to the offloading processing.

[0177] The intermediate language file output unit 213b outputs the extracted intermediate language file 132. The extraction of the intermediate language is not completed in one go, and for appropriate offload area exploration, it is repeated for trial execution and optimization.

[0178] <Arithmetic intensity calculation unit 214> The arithmetic intensity calculation unit 214 calculates the arithmetic intensity of the loop statements of the application using an arithmetic intensity analysis tool such as the ROSE framework (registered trademark). The arithmetic intensity is the value obtained by dividing the number of floating point operations (floating point number, FN) executed during the operation of the program by the number of bytes accessed to the main memory (FN operations / memory access). The arithmetic intensity is an index that increases when the number of calculations is large and decreases when the number of accesses is large. A process with a high arithmetic intensity is a heavy process for the processor. Therefore, an arithmetic intensity analysis tool is used to analyze the arithmetic intensity of the loop statements. The PLD processing pattern creation unit 215 narrows down the loop statements with high arithmetic intensity as offload candidates.

[0179] An example of the calculation of the arithmetic intensity will be described. Assume that 10 floating point calculation processes (10 FLOP) are performed in one loop, and the data used in the loop is 2 bytes. When data of the same size is used for each loop, 10 / 2 = 5 [FLOP / byte] becomes the arithmetic intensity. Note that since the loop count is not considered in the arithmetic intensity, in this embodiment, in addition to the arithmetic intensity, the loop count is also considered for narrowing down.

[0180] <PLD processing pattern creation unit 215> The PLD processing pattern creation unit 215 narrows down the loop statements with an arithmetic intensity higher than a predetermined threshold (hereinafter, appropriately referred to as high arithmetic intensity) calculated by the arithmetic intensity calculation unit 214 as offload candidates, and creates a PLD processing pattern. In addition, as a basic operation, the PLD processing pattern creation unit 215 creates a PLD processing pattern that excludes loop statements (repetitive statements) that cause compilation errors from being offloaded, and specifies whether or not to perform PLD processing on repetitive statements that do not cause compilation errors.

[0181] -Loop count measurement function As a loop count measurement function, the PLD processing pattern creation unit 215 uses a profiling tool to measure the number of loops in the loop statements of an application, and narrows down the loop statements that have high arithmetic strength and a loop count greater than a predetermined number (hereinafter referred to as "high loop count" as appropriate). To grasp the loop count, GNU Coverage's gcov or the like is used. Known profiling tools include "GNU Profiler (gprof)" and "GNU Coverage (gcov)". Either may be used, as both can investigate the number of times each loop is executed.

[0182] In addition, since the number of loops is not particularly visible in arithmetic strength analysis, a profiling tool is used to measure the number of loops in order to detect loops with a large number of loops and high loads. Here, the level of arithmetic strength indicates whether a process is suitable for offloading to an FPGA, and the number of loops x arithmetic strength indicates whether the load associated with offloading to an FPGA is high.

[0183] · OpenCL (intermediate language) creation function The PLD processing pattern creation unit 215, as an OpenCL creation function, creates OpenCL for offloading each of the narrowed-down loop statements to the FPGA (OpenCLization). That is, the PLD processing pattern creation unit 215 compiles OpenCL for offloading the narrowed-down loop statements. In addition, the PLD processing pattern creation unit 215 creates a list of loop statements that have been measured and have higher performance than the CPU, and creates OpenCL for offloading by combining the loop statements in the list.

[0184] We will discuss OpenCL conversion. The PLD processing pattern creation unit 215 converts the loop statement into a high-level language such as OpenCL. First, the CPU processing program is divided into a kernel (FPGA) and a host (CPU) according to the grammar of the high-level language such as OpenCL. For example, if one of ten "for" statements is to be processed by FPGA, that one is extracted as a kernel program and written according to the grammar of OpenCL. An example of the grammar of OpenCL will be described later.

[0185] Furthermore, when dividing, techniques for further speedup can be incorporated. In general, techniques for speedup using FPGAs include local memory cache, stream processing, multiple instantiation, loop unrolling, merging nested loops, memory interleaving, etc. These techniques are often used to speed up a program, although they may not be absolutely effective for some loops.

[0186] A kernel created according to the OpenCL C language syntax is executed on a device (e.g., FPGA) by a host (e.g., CPU) program created using the OpenCL C language runtime API. The part where the kernel function hello() is called from the host side is a call to clEnqueueTask(), one of the OpenCL runtime APIs. The basic flow of OpenCL initialization, execution, and termination described in the host code is the following steps 1 to 13. Among these steps 1 to 13, steps 1 to 10 are the procedure (preparation) until the kernel function hello() is called from the host side, and step 11 is the execution of the kernel.

[0187] 1. Platform Specific The platform on which OpenCL runs is identified using the function clGetPlatformIDs(), which provides platform identification functionality defined in the OpenCL runtime API.

[0188] 2. Device Identification The device such as the GPU to be used on the platform is identified using the function clGetDeviceIDs() that provides a device identification function defined in the OpenCL runtime API.

[0189] 3. Creating a context An OpenCL context, which serves as the execution environment for running OpenCL, is created using the function clCreateContext(), which provides a context creation function defined in the OpenCL runtime API.

[0190] 4. Create a command queue A command queue is created in preparation for controlling the device using the clCreateCommandQueue() function, which provides command queue creation functionality defined in the OpenCL runtime API. In OpenCL, the host issues commands to the device (such as issuing kernel execution commands and memory copy commands between the host and device) through the command queue.

[0191] 5. Creating a memory object A memory object that allows the host to reference the memory object is created using the function clCreateBuffer(), which provides the functionality to allocate memory on the device defined in the OpenCL runtime API.

[0192] 6. Loading the kernel file The execution of the kernel executed on the device is controlled by the host program. For this reason, the host program must first load the kernel program. A kernel program can be binary data created by the OpenCL compiler or source code written in the OpenCL C language. This kernel file is loaded (description omitted). Note that the OpenCL runtime API is not used when loading the kernel file.

[0193] 7. Creating a program object In OpenCL, a kernel program is known as a program project. This procedure is called program object creation. Create a program object that allows the host to reference the memory object using clCreateProgramWithSource(), a function defined in the OpenCL runtime API that provides the program object creation function. To create a program object from a compiled binary sequence of a kernel program, use clCreateProgramWithBinary().

[0194] 8. Build The program object registered as source code is built using the OpenCL C compiler and linker. Build a program object using the function clBuildProgram(), which executes the build by the OpenCL C compiler and linker defined in the OpenCL runtime API. Note that if you create a program object from a compiled binary string using clCreateProgramWithBinary(), this compilation step is not necessary.

[0195] 9. Creating a kernel object Create a kernel object using the function clCreateKernel(), which provides the kernel object creation function defined in the OpenCL runtime API. Since one kernel object corresponds to one kernel function, specify the name of the kernel function (hello) when creating the kernel object. Also, if multiple kernel functions are written as one program object, one kernel object corresponds one-to-one to one kernel function, so call clCreateKernel() multiple times.

[0196] 10. Kernel arguments setting The kernel arguments are set using the clSetKernel() function, which provides the ability to provide arguments to the kernel defined in the OpenCL runtime API (pass values ​​to arguments held by the kernel function). Now that steps 1 to 10 above are complete, we move on to step 11, where the kernel is executed on the device from the host side.

[0197] 11. Kernel execution Kernel execution (submission to the command queue) acts on the device, and therefore is a queuing function for the command queue. Queue a command to execute the kernel hello on the device using the function clEnqueueTask(), which provides the kernel execution functionality defined in the OpenCL runtime API. After the command to execute the kernel hello is queued, it will be executed on an executable compute unit on the device.

[0198] 12. Reading from a memory object Data is copied from the device memory area to the host memory area using the function clEnqueueReadBuffer(), which provides the functionality for copying data from device memory to host memory defined in the OpenCL runtime API. Data is also copied from the host memory area to the device memory area using the function clEnqueueWrightBuffer(), which provides the functionality for copying data from the host to host memory. Note that these functions act on the device, so the copy command is queued in the command queue once before data copying begins.

[0199] 13. Releasing Objects Finally, release the various objects that you have created so far. Above, we have explained device execution of kernels written in accordance with the OpenCL C language.

[0200] -Resource amount calculation function The PLD processing pattern creation unit 215, as a resource amount calculation function, precompiles the created OpenCL and calculates the amount of resources to be used ("first resource amount calculation"). The PLD processing pattern creation unit 215 calculates resource efficiency based on the calculated arithmetic strength and resource amount, and selects c loop statements whose resource efficiency is higher than a predetermined value from each loop statement based on the calculated resource efficiency. The PLD processing pattern creation unit 215 calculates the amount of resources to be used by precompiling with the combined offload OpenCL ("second resource amount calculation"). Here, it is also possible to use the sum of the resource amounts in precompilation before the first measurement without precompiling.

[0201] <Performance measurement section 116> The performance measurement unit 116 compiles the created application of the PLD processing pattern, places it on the verification machine 14, and executes the processing for performance measurement when offloaded to the PLD.

[0202] The performance measurement unit 116 executes the placed binary file, measures the performance when offloaded, and returns the performance measurement result to the offload range extraction unit 213a. In this case, the offload range extraction unit 213a extracts another PLD processing pattern, and the intermediate language file output unit 213b attempts to measure the performance based on the extracted intermediate language (see symbol a in FIG. 2).

[0203] The performance measurement unit 116 includes a binary file deployment unit (Deploy binary files) 116a, a power usage measurement unit 116b, and an evaluation value setting unit 116c. Note that, although the evaluation value setting unit 116c is configured to be included in the performance measurement unit 116, it may be a separate independent functional unit.

[0204] The binary file placement unit 116a deploys (places) an executable file derived from the intermediate language on the verification machine 14 equipped with a GPU.

[0205] The power consumption measuring unit 116b measures the processing time and power consumption required during FPGA offloading.

[0206] The evaluation value setting unit 116c sets an evaluation value that includes the processing time and power consumption based on the processing time and power consumption required during FPGA offloading measured by the performance measuring unit 116 and the power consumption measuring unit 116b, and that is higher as the processing time and power consumption are lower.

[0207] A specific example of performance measurement will be described. The PLD processing pattern creation unit 215 narrows down loop statements with high resource efficiency, and compiles OpenCL that offloads the loop statements narrowed down by the executable file creation unit 117. The performance measurement unit 116 measures the performance of the compiled program ("first performance measurement").

[0208] Then, the PLD processing pattern creation unit 215 creates a list of loop statements that have been measured and have higher performance than the CPU. The PLD processing pattern creation unit 215 creates an OpenCL for offloading by combining the loop statements in the list. The PLD processing pattern creation unit 215 calculates the amount of resources to be used by precompiling the combined offload OpenCL. Note that, instead of precompiling, the sum of the resource amounts in precompiling before the first measurement may be used. The executable file creation unit 117 compiles the combined offload OpenCL, and the performance measurement unit 116 measures the performance of the compiled program ("second performance measurement").

[0209] <Executable File Creation Section 117> The executable file creation unit 117 selects the PLD processing pattern with the highest evaluation value from the multiple PLD processing patterns based on the results of the measurement of processing time and power consumption repeated a predetermined number of times, and compiles the PLD processing pattern with the highest evaluation value to create an executable file.

[0210] The automatic offload operation of the offload server 1A configured as above will now be described. [Automatic offload operation] The offload server 1A of this embodiment is an example in which an elemental technology of environment adaptive software is applied to FPGA automatic offloading of user application logic. The automatic offload processing of the offload server 1A shown in FIG. 2 will be described. 2, the offload server 1A is applied to the elemental technology of the environment adaptive software. The offload server 1A has a control unit (automatic offload function unit) 11, a test case DB 131, an intermediate language file 132, and a verification machine 14. The offload server 1 acquires an application code 130 that is used by the user.

[0211] A user uses, for example, various devices 151, a device having a CPU-GPU 152, a device having a CPU-FPGA 153, and a device having a CPU 154. The offload server 1 automatically offloads function processing to accelerators of the device having a CPU-GPU 152 and the device having a CPU-FPGA 153.

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

[0213] <Step S12: Analyze application code> In step S12, the application code analysis unit 112 (see FIG. 11) analyzes the source code of the processing function and grasps the structure of the use of a specific library, such as a loop statement or an FFT library call.

[0214] <Step S13: Extract offload able area> In step S13, the PLD processing specification unit 213 (see FIG. 11) identifies loop statements (repetitive statements) in the application, and specifies parallel processing or pipeline processing in the FPGA for each repetitive statement, and compiles it with a high-level synthesis tool. Specifically, the offload range extraction unit 213a (see FIG. 11) identifies processing that can be offloaded to the FPGA, such as loop statements, and extracts OpenCL as an intermediate language corresponding to the offload processing.

[0215] <Step S14: Output intermediate file> In step S14, the intermediate language file output unit 213b (see FIG. 11) outputs the intermediate language file 132. The intermediate language extraction is not a one-time process, but is repeated to perform trial and error optimization to find an appropriate offload area.

[0216] <Step S15: Compile error> In step S15, the PLD processing pattern creation unit 215 (see FIG. 11) creates a PLD processing pattern that excludes loop statements that cause a compilation error from being offloaded, and specifies whether or not to perform FPGA processing for repetitive statements that do not cause a compilation error.

[0217] <Step S21: Deploy binary files> In step S21, the binary file placement unit 116a (see FIG. 11) deploys an executable file derived from the intermediate language to the verification machine 14 equipped with an FPGA. The binary file placement unit 116a starts the placed file, executes assumed test cases, and measures the performance when offloaded.

[0218] <Step S22: Measure performance> In step S22, the performance measurement unit 116 (see FIG. 11) executes the arranged file and measures the performance and power consumption when offloaded. In order to make the offload area more appropriate, the performance measurement result is returned to the offload range extraction unit 213a, which extracts another pattern. Then, the intermediate language file output unit 213b attempts performance measurement based on the extracted intermediate language (see symbol a in FIG. 2). The performance measurement unit 116 repeats the performance and power consumption measurements in the verification environment, and finally determines the code pattern to be deployed.

[0219] As indicated by the reference character a in FIG. 2, the control unit 21 repeatedly executes the above steps S12 to S22. The automatic offload function of the control unit 21 can be summarized as follows. That is, the PLD processing designation unit 213 identifies loop statements (repeated statements) of the application, and for each repeated statement, specifies parallel processing or pipeline processing in FPGA in OpenCL (intermediate language), and compiles it with a high-level synthesis tool. Then, the PLD processing pattern creation unit 215 creates a PLD processing pattern that excludes loop statements in which a compilation error occurs from the offload target, and specifies whether or not to perform PLD processing for loop statements in which no compilation error occurs. Then, the binary file placement unit 116a compiles the application of the corresponding PLD processing pattern, places it in the verification machine 14, and the performance measurement unit 116 executes the performance measurement processing on the verification machine 14. The execution file creation unit 117 selects the highest evaluation value (e.g., evaluation value=(processing time)) from among a plurality of PLD processing patterns based on the performance and power consumption measurement results repeated a predetermined number of times. -1 / 2 ×(Electricity usage) -1 / 2 It then selects the pattern with the highest match rate (the one with the highest match rate) and compiles the selected pattern to create an executable file.

[0220] <Step S23: Deploy final binary files 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.

[0221] <Step S24: Extract performance test cases and run automatically> In step S24, after arranging the executable file, the performance measurement test extraction execution unit 119 extracts performance test items from the test case DB 131 and automatically executes the extracted performance tests in order to show the performance to the user.

[0222] <Step S25: Provide price and performance to a user to judge> 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.

[0223] The above steps S21 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. In addition, the processing performed in the background to reduce costs may be limited to GPU / FPGA offloading.

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

[0225] In the above, a process flow for performing all the processes required for environment adaptation, such as code conversion, resource amount adjustment, and placement location adjustment, has been described, but it is not limited to this, and it is also possible to extract only the processes you want to perform. For example, if you want to perform only code conversion for FPGA, you can use only the necessary parts such as the environment adaptation function and the verification environment from step S21 to step S21 above.

[0226] [FPGA automatic offload] The code analysis described above uses a syntax analysis tool such as Clang to analyze application code. Code analysis requires analysis that assumes the device to be offloaded, so it is difficult to generalize. However, it is possible to grasp the code structure, such as loop statements and variable reference relationships, and to determine that a function block performs FFT processing, or that a library that performs FFT processing is being called. It is difficult for the offload server to automatically determine the function block. This can also be grasped by using a similar code detection tool such as Deckard to determine the degree of similarity. Here, Clang is a tool for C / C++, but it is necessary to select a tool that matches the language to be analyzed.

[0227] In addition, when offloading application processing, consideration must be given to the offload destination for each, such as GPU, FPGA, and IoT GW. In general, it is difficult to automatically discover the settings that maximize performance in one go. For this reason, offload patterns are tried out by repeating performance measurements several times in a verification environment to find a pattern that can increase speed.

[0228] Below, we explain a method for offloading loop statements in application software to FPGAs. [flowchart] FIG. 12 is a flowchart for explaining the outline of the operation of the offload server 1A. In step S201, the application code analysis unit 112 analyzes the source code of the application to be offloaded. The application code analysis unit 112 analyzes information on loop statements and variables according to the language of the source code.

[0229] In step S202, the PLD processing specification unit 213 identifies the loop statements and reference relationships of the application.

[0230] Next, the PLD processing pattern creation unit 215 performs processing to narrow 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. In step S203, the arithmetic strength calculation unit 214 calculates the arithmetic strength of the loop statement of the application using an arithmetic strength analysis tool. The arithmetic strength is an index that increases with a large number of calculations and decreases with a large number of accesses, and processing with high arithmetic strength is heavy processing for the processor. Therefore, the arithmetic strength of the loop statement is analyzed using the arithmetic strength analysis tool, and high-density loop statements are narrowed down to offload candidates. Therefore, the arithmetic strength of the loop statement is analyzed using the arithmetic strength analysis tool, and high-density loop statements are narrowed down to offload candidates.

[0231] 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, we will describe how to calculate the amount of resources required when processing a high arithmetic strength loop statement on an FPGA. 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.

[0232] Therefore, in this embodiment, the PLD processing pattern creation unit 215 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.

[0233] Returning to the flow of FIG. 12, in step S204, the PLD processing pattern creation unit 215 measures the number of loops in the loop statement of the application using a profiling tool such as gcov or gprof. In step S205, the PLD processing pattern creation unit 215 narrows down the loop statements to those with high arithmetic strength and a large number of loop counts.

[0234] In step S206, the PLD processing pattern creation unit 215 creates an OpenCL for offloading each of the narrowed down loop statements to the FPGA.

[0235] Here, we will provide additional information on converting loop statements to OpenCL (creating 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.

[0236] Next, several highly resource-efficient loop statements are selected, and offload patterns are created using these to measure performance. FPGA speedup can be achieved by concentrating FPGA resources on a single process, or by distributing FPGA resources across multiple processes. A certain number of patterns are created for the selected single loop statements, and then precompiled as a preliminary step to running them on the actual FPGA.

[0237] In step S207, the PLD processing pattern creation unit 215 precompiles the created OpenCL and calculates the amount of resources to be used ("first resource amount calculation").

[0238] In step S208, the PLD processing pattern creation unit 215 narrows down loop statements with high resource efficiency.

[0239] In step S209, the executable file creation unit 117 compiles OpenCL that offloads the narrowed-down loop statements.

[0240] In step S210, the performance measurement unit 116 measures the performance and power consumption of the compiled program ("first performance and power consumption measurement"). Since several candidate loop statements remain, the performance measurement unit 116 actually measures the performance and power consumption using them. When offloading processing to the FPGA, power consumption is also taken into consideration, so in addition to measuring performance, power consumption is measured (for details, see the subroutine in FIG. 13).

[0241] In step S211, the PLD processing pattern creation unit 215 creates a list of loop statements whose performance has been measured and whose performance has been improved compared to the CPU.

[0242] In step S212, the PLD processing pattern creation unit 215 creates an OpenCL for offloading by combining loop statements of the list. In step S213, the PLD processing pattern creation unit 215 calculates the amount of resources to be used by precompiling with the combined offload OpenCL ("second resource amount calculation"). Note that it is also possible to use the sum of the resource amounts in the precompilation before the first measurement without precompiling. In this way, the number of times of precompilation can be reduced.

[0243] In step S214, the executable file creation unit 117 compiles the combined offload OpenCL.

[0244] In step S215, the performance measurement unit 116 measures the performance of the compiled program ("second performance and power consumption measurement"). The performance measurement unit 116 compiles and measures the selected single-loop statement, and for the single-loop statements that have been speeded up, creates combination patterns and performs a second performance and power consumption measurement (for details, see the subroutine in FIG. 13).

[0245] In step S216, the production environment deployment unit 118 selects the pattern with the highest performance from the first and second measurements, and ends the process of this flow. From the multiple measured patterns, a pattern with a short time and low power consumption is selected as a solution.

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

[0247] 13 is a flowchart showing the performance / power consumption measurement process of the performance measuring unit 116. This flow is called and executed by the subroutine call of step S211 or step S215 in FIG.

[0248] In step S301, the power consumption measuring unit 116b measures the processing time and power consumption required during FPGA offloading.

[0249] In step S302, the evaluation value setting unit 116c sets an evaluation value based on the measured processing time and power consumption.

[0250] In step S303, the performance measuring unit 116 measures the performance and power consumption of a pattern with a high evaluation value, which is evaluated so that the higher the evaluation value of an individual, the higher the fitness, and then returns to step S211 or step S215 in FIG.

[0251] [Example of off-road pattern creation] FIG. 14 is a diagram showing an image of a search performed by the PLD processing pattern generating unit 215. As shown in FIG. The control unit (automatic offload function unit) 21 (see FIG. 11) analyzes the application code 130 (see FIG. 2) used by the user, and checks whether or not for statements can be executed in parallel from the code patterns 241 of the application code 130, as shown in FIG. 14. As shown by the symbol t in FIG. 14, if four for statements are found from the code pattern 241, one digit is assigned to each for statement, in this case four digits of 1 or 0 are assigned to the four for statements. Here, 1 is assigned if FPGA processing is performed, and 0 is assigned if FPGA processing is not performed (i.e. processing by the CPU).

[0252] [Flow from C code to search for the final OpenCL solution] The procedures A - F in FIG. 15 are diagrams explaining the flow from C code to search for the final OpenCL solution. The application code analysis unit 112 (see FIG. 11) performs syntax analysis on the "C code" shown in procedure A of FIG. 15 (<syntax analysis>: see reference u in FIG. 15), and the PLD processing specification unit 213 (see FIG. 11) identifies the "loop statements, variable information" shown in procedure B of FIG. 15 (see FIG. 14).

[0253] The arithmetic intensity calculation unit 214 (see FIG. 11) performs arithmetic intensity analysis using an arithmetic intensity analysis tool on the identified "loop statements, variable information". The PLD processing pattern creation unit 215 narrows down the loop statements with high arithmetic intensity as offload candidates. Further, the PLD processing pattern creation unit 215 performs profiling analysis using a profiling tool (<intensity analysis>: see reference v in FIG. 15) to further narrow down the loop statements with high arithmetic intensity and high loop counts.

[0254] Then, the PLD processing pattern creation unit 215 creates (OpenCLizes) OpenCL for offloading each narrowed-down loop statement to the FPGA. Furthermore, at the time of OpenCLization, high-speed techniques such as unfolding are introduced together with code splitting (described later).

[0255] <Specific example of "high arithmetic intensity, OpenCLization" (Part 1): Procedure C> For example, when four for loops (assignment of 4-digit 1 or 0) are found from the code pattern 241 of the application code 130 (see FIG. 14), three are narrowed down (selected) by arithmetic intensity analysis. That is, as shown by reference w in FIG. 15, the offload patterns "1000", "0010", and "0001" of three for loops are narrowed down from the four for loops.

[0256] <Example of "unfolding" executed together with code splitting during OpenCLization> When transferring data from FPGA to CPU, the loop statement written in the CPU program is [k=0; k<10; k++] { } In this example, we specify \pragma unroll at the top of the loop statement. \pragma unroll for(k=0; k<10; k++) } It is written as follows.

[0257] If you specify unroll using syntax that matches Intel or Xilinx (registered trademark) tools, such as \pragma unroll, the above example can be expanded to i=0, i=1, i=2 and executed in a pipeline. This will use 10 times the amount of resources, but may be faster. In addition, the number of loops to be expanded with unroll can be specified to be expanded to 5 instead of the total number of loops, in which case each loop will be expanded 2 times to 5. This completes the explanation of the "expansion" example.

[0258] Next, the PLD processing pattern creation unit 215 further narrows down the loop statements with high arithmetic strength narrowed down as offload candidates by using the resource amount. That is, the PLD processing pattern creation unit 215 calculates the resource amount, and the PLD processing pattern creation unit 215 analyzes the resource efficiency (=arithmetic strength / resource amount during FPGA processing, or arithmetic strength×number of loops / resource amount during FPGA processing) from the offload candidates of the loop statements with high arithmetic strength to extract loop statements with high resource efficiency.

[0259] At the symbol x in FIG. 15, the PLD processing pattern creation unit 215 compiles (<precompiles>) OpenCL for offloading the narrowed-down loop statements.

[0260] <"High arithmetic strength, OpenCL" concrete example (part 2)> As shown by symbol y in Figure 15, from the four offload patterns "1000", "0100", "0010", and "0001" narrowed down by the arithmetic intensity analysis, the above resource efficiency analysis narrows down the offload patterns to three: "1000", "0010", and "0001". The above is a description of "high arithmetic strength, OpenCL" shown in step C of FIG.

[0261] For the "loop statement with high resource efficiency" shown in step D of FIG. 15, the performance measurement unit 116 measures the performance of the compiled program ("first performance measurement"). Then, the PLD processing pattern creation unit 215 creates a list of loop statements whose performance has been measured and whose performance has been improved compared to the CPU. Thereafter, similarly, the resource amount is calculated, offload OpenCL compilation is performed, and the performance of the compiled program is measured.

[0262] <"High arithmetic strength, OpenCL" concrete example (part 3)> As shown by symbol y in Fig. 15, a first measurement is performed for three offload patterns "1000", "0010", and "0001". If the performance of "1000" and "0010" is found to be high among the three measurements, a second measurement is performed for the combination of "1000" and "0010".

[0263] At reference symbol z in FIG. 15, the executable file creation unit 117 compiles (<main compilation>) OpenCL for offloading the narrowed-down loop statements.

[0264] The "measurement of combination patterns" shown in step E of FIG. 15 refers to measuring a candidate loop statement alone and then measuring a verification pattern in combination with that statement.

[0265] <"High arithmetic strength, OpenCL" concrete example (part 4)> As shown by symbol aa in FIG. 15, a second measurement is made for "1010", which is a combination of "1000" and "0010". Two measurements are made, and as a result, "00010", which has the highest speed between the first and second measurements, is selected. In such a case, "0010" becomes the final solution. Here, there are cases where combination patterns cannot be measured due to resource limitations. In this case, it is possible to skip the combination and simply select the one with the highest speed from the individual results.

[0266] In the case of reference symbol bb in FIG. 15, the performance measuring unit 116 selects (<selects>) "0010" which has the best maximum speed and power consumption between the first and second measurements.

[0267] As a result of the above, "00010" (see symbol cc in FIG. 15) is selected as the "OpenCL final solution" shown in step F in FIG.

[0268] <Deployment> The final OpenCL solution, the PLD processing pattern with the highest processing performance, will be redeployed into a production environment and provided to users.

[0269] [Implementation example] An implementation example will be explained. FPGAs such as Intel PAC with Intel Arria10 GX FPGA can be used. FPGA processing can use the Intel Acceleration Stack (Intel FPGA SDK for OpenCL, Quartus Prime Version), etc. The Intel FPGA SDK for OpenCL is a high-level synthesis tool (HLS) that interprets standard OpenCL as well as Intel-specific #pragma directives. In the implementation example, the OpenCL code describing the kernel to be processed by the FPGA and the host program to be processed by the CPU is interpreted, information such as resource amounts is output, and wiring work for the FPGA is performed so that it can run on the FPGA. Even a small program of about 100 lines takes a long time, about three hours, to run on an actual FPGA. However, if the resource amounts are exceeded, an error will be displayed early on. Also, if the OpenCL code cannot be processed by the FPGA, an error will be displayed after several hours.

[0270] In the implementation example, when a request to use a C / C++ application is received, the system first analyzes the code of the C / C++ application to find for statements and understand the program structure, such as variable data used within the for statements. For syntax analysis, LLVM / Clang syntax analysis libraries can be used.

[0271] In the implementation example, next, to obtain a prediction of whether each loop statement can be effectively offloaded to FPGA, an arithmetic strength analysis tool is executed to obtain an index of arithmetic strength determined by the number of calculations, the number of accesses, etc. The ROSE framework, etc., can be used for the arithmetic strength analysis. Only the loop statements with the highest arithmetic strength are targeted. Next, we use a profiling tool such as gcov to obtain the loop count for each loop. We narrow down the candidates to the top a loop statements based on arithmetic strength x loop count.

[0272] In the implementation example, next, OpenCL code to be offloaded to FPGA is generated for each loop statement with high arithmetic strength. The OpenCL code is divided into the corresponding loop statement as an FPGA kernel and the rest as a CPU host program. When generating the FPGA kernel code, a fixed number b of loop statement unrolling may be performed as a speed-up technique. Although loop statement unrolling increases the amount of resources, it is effective in increasing speed. Therefore, the number of unrolling is limited to a fixed number b so that the amount of resources does not become excessive.

[0273] In the implementation example, next, the Intel FPGA SDK for OpenCL is used to precompile a number of OpenCL codes, and the amount of resources used, such as flip flops and look up tables, is calculated. The amount of resources used is displayed as a percentage of the total resource amount. Here, the resource efficiency of each loop statement is calculated from the arithmetic strength and resource amount, or the arithmetic strength, loop count, and resource amount. For example, a loop statement with arithmetic strength of 10 and resource amount of 0.5 has a resource efficiency of 10 / 0.5=20, while a loop statement with arithmetic strength of 3 and resource amount of 0.3 has a resource efficiency of 3 / 0.3=10, with the former being higher. Alternatively, the resource efficiency may be calculated by multiplying the number of loop counts. For each loop statement, c statements with the highest resource efficiency are selected.

[0274] In the implementation example, next, c loop statements are candidates and a pattern to be measured is created. For example, if the first and third loop statements are highly resource efficient, an OpenCL pattern is created to offload number 1 and another to offload number 3, then compiled and performance is measured. If speed can be improved by offloading multiple single-loop statements (for example, if both number 1 and number 3 are fast), an OpenCL pattern for that combination is created, compiled and performance is measured (for example, a pattern that offloads both number 1 and number 3).

[0275] When creating combinations of single loops, the amount of resources used is also a combination. For this reason, if it does not fit within the upper limit, that combination pattern will not be created. After creating d patterns, including combinations, performance is measured on a server equipped with an FPGA in the verification environment. To measure performance, sample processing specified by the application to be accelerated is performed. For example, if it is a Fourier transform application, performance is measured by benchmarking the conversion process on sample data. In the implementation example, finally, the fastest pattern among the multiple measurement patterns is selected as the solution.

[0276] [evaluation] Explain the evaluation. In the second embodiment, [FPGA automatic offloading of loop statements], similar to the first embodiment, [GPU automatic offloading of loop statements], a method is added to the existing implementation tool to determine the evaluation value of the measurement pattern so that the lower the power consumption, the higher the evaluation value is, and offloading is performed to confirm that low power consumption can be achieved.

[0277] <Evaluation target> In the second embodiment of [FPGA automatic offloading of loop statements], the evaluation target is MRI-Q for MRI (Magnetic Resonance Imaging) image processing. MRI-Q calculates the matrix Q that represents the scanner configuration used in 3D MRI reconstruction algorithms in non-Cartesian space. MRI-Q is written in C language, and during performance measurement, it executes 3D MRI image processing and measures the processing time with large (maximum) 64 x 64 x 64 size data. CPU processing is done using C language, and FPGA processing is done based on OpenCL.

[0278] <Evaluation method> The code of the target application is input, and loop offloading recognized by Clang etc. is attempted for the destination GPU or FPGA to determine the offloading pattern. During this process, the processing time and power consumption are measured. For the final offloading pattern, the change in power consumption over time is obtained, and the power reduction compared to processing entirely on the CPU is confirmed. In the second embodiment of [FPGA automatic offloading of loop statements], GA is not performed, and arithmetic strength and the like are used to narrow down the measurement patterns until they are reduced to four patterns. Offloadable loop statements: MRI-Q 16 Pattern compatibility: The evaluation value shown in formula (1), that is, (processing time) -1 / 2 ×(Electricity usage) -1 / 2 As shown in equation (1), the lower the processing time and power consumption, the higher the evaluation value and the higher the degree of compatibility.

[0279] <Evaluation environment> In the second embodiment, [FPGA automatic offload of loop statements], Intel PAC with Intel Arria10 GX FPGA (registered trademark) is used. The power consumption is measured by using ipmitool (registered trademark) of IPMI (Intelligent Platform Management Interface) of Dell (registered trademark) server.

[0280] <Results and Discussion> Figure 16 shows the power consumption (Watt) and processing time when MRI-Q is offloaded to the FPGA. The symbol dd in FIG. 16 shows a comparison of the power consumption Watts during each processing time for "all CPU processing" on the left side of FIG. 16 and "CPU and FPGA processing" on the right side of FIG. 16. In MRI-Q, the processing time for "CPU and GPU processing" on the right side of Fig. 16 is reduced from 14 seconds to 2 seconds compared to "all CPU processing" on the left side of Fig. 16, and the power consumption watts is also reduced from a maximum of about 122.2 W for "all CPU processing" to a maximum of about 112.0 W for "CPU and FPGA processing". As a result, the watt sec for "CPU and FPGA processing" is 223 Watt sec, about 1 / 8 of the 1694 Watt sec for "all CPU processing".

[0281] We also confirmed that power consumption was reduced for multiple applications. In the second embodiment, [FPGA automatic offload of loop statements], not only was the power consumption (Watt) reduced, but the synergistic effect of shortening the execution time also resulted in a significant reduction in power consumption. FPGAs are generally said to be power efficient, and we were able to confirm in our experiments that FPGAs consume less power. Therefore, when offloading in a mixed environment results in similar performance, it may be possible to select FPGAs.

[0282] As described above, in the second embodiment of [FPGA automatic offloading of loop statements], the method of including power consumption in the suitability level achieves automatic speedup and low power consumption by evaluating power consumption. In particular, when measuring FPGA automatic offloading in a verification environment, power consumption is acquired in addition to processing time, and patterns with short processing times and low power consumption are considered to have a high suitability level, and low power consumption is incorporated into the automatic code conversion. As described in the evaluation of FIG. 16, low power consumption was confirmed through the automatic offloading of existing applications, and the effectiveness of the method was confirmed.

[0283] [Automatic offloading in mixed environments] Among the various migration destinations such as GPU, FPGA, and many-core CPU, we will explain the technology to select a high-performance migration destination for offloading. An offload server 1 (see FIG. 1) and an offload server 1A (see FIG. 11) are combined (hereinafter, for convenience of explanation, they will be referred to as the offload servers 1 and 1A). The offload server 1, 1A offloads a specific process of an application to at least one of the GPU, the main core CPU, and the PLD.

[0284] The offload server 1, 1A includes a parallel processing pattern creation unit 214 (see FIG. 11) that creates a parallel processing pattern that excludes loop statements for a GPU or loop statements for a many-core CPU that generate a compilation error from being offloaded, and that specifies whether or not to perform parallel processing for loop statements for a GPU or loop statements for a many-core CPU that do not generate a compilation error, and a performance measurement unit 116 (see FIGS. 1 and 11) that compiles an application of a parallel processing pattern or a PLD processing pattern in a mixed environment of a GPU, a many-core CPU, and a PLD, places it on an accelerator verification device, and executes each performance measurement process when offloaded to the GPU, the many-core CPU, and the PLD.

[0285] The offload server 1, 1A also includes an evaluation value setting unit 116c (see Figures 1 and 11) that sets an evaluation value including processing time and power consumption based on the processing time and power consumption required when offloading the GPU, menu core CPU, and PLD measured by the performance measurement unit 116, and the evaluation value becomes higher as the processing time and power consumption becomes lower; and an executable file creation unit 117 (see Figures 1 and 11) that selects the GPU, menu core CPU, and PLD with the best processing time and power consumption based on the measurement results of the processing time and power consumption of the GPU, menu core CPU, and PLD, selects the parallel processing pattern or PLD processing pattern with the highest evaluation value from a plurality of parallel processing patterns or PLD processing patterns for the selected one, and compiles the parallel processing pattern or PLD processing pattern with the highest evaluation value to create an executable file.

[0286] The verification will be performed in the order of loop offload for many-core CPU, GPU, and FPGA, searching for a pattern that will provide high performance. With automatic offloading, it is expected that the pattern search will be performed as cheaply and quickly as possible. Therefore, FPGA, which requires a long verification time, will be performed last, and if a pattern that satisfies the user requirements has been found in the previous stage, FPGA verification will not be performed.

[0287] There is not a big difference between GPUs and many-core CPUs in terms of price or verification time, but the difference between a many-core CPU and a normal CPU is smaller than that of a GPU, which has a different device itself and has a separate memory space. Therefore, the verification order will be many-core CPUs first, and if a pattern that fully satisfies user requirements is found with the many-core CPU, GPU verification will not be performed. As described above, we have examined three migration destinations, GPU, FPGA, and many-core CPU, and automatically selected the fastest migration destination.

[0288] As described in the above embodiments, when automatically selecting a high-speed destination, not only short processing time but also low power consumption destinations are selected as candidates for automatic selection through actual measurements in a verification environment. For example, evaluation value=(processing time)-1 / 2 ×(Electricity usage) -1 / 2 The evaluation formula can be set so that the shorter the processing time and the lower the power consumption, the higher the score.

[0289] A typical example of the costs for a data center might be that the initial costs for hardware and development costs etc. are 1 / 3 of the total costs, operational costs for electricity and maintenance etc. are 1 / 3, and other costs for service orders etc. are 1 / 3. In this case, for example, if processing time is reduced to 1 / 5 and the number of hardware units, including CPUs and GPUs, is halved, the initial costs will also be reduced. Reducing power consumption by half will also lead to reduced operational costs. However, operational costs include many factors other than electricity, so halving power consumption does not necessarily halve operational costs. In addition, hardware prices vary from operator to operator, with volume discounts depending on the number of GPUs and FPGA servers introduced. For this reason, the evaluation formula needs to be set differently for each operator.

[0290] In this way, the appropriate offload destination is automatically selected taking into account not only the processing time but also the power consumption. Generally speaking, FPGAs are said to be more power efficient than CPUs and GPUs, so if actual measurements show that the processing time reduction after offloading is about the same, it is likely that the power-efficient FPGA will be selected as the offload destination.

[0291] [Hardware configuration] The offload server according to the first and second embodiments is realized by a computer 900, which is a physical device having a configuration as shown in FIG. 17 is a hardware configuration diagram showing an example of a computer that realizes the functions of the offload servers 1 and 1 A. The computer 900 has a CPU (Central Processing Unit) 901, a ROM (Read Only Memory) 902, a RAM 903, an HDD (Hard Disk Drive) 904, an input / output I / F (Interface) 905, a communication I / F 906, and a media I / F 907.

[0292] The CPU 901 operates based on a program stored in the ROM 902 or the HDD 904, and controls each processing unit of the offload servers 1, 1A shown in Figures 1 and 11. The ROM 902 stores a boot program executed by the CPU 901 when the computer 900 is started, programs related to the hardware of the computer 900, and the like.

[0293] The CPU 901 controls an input device 910 such as a mouse or a keyboard, and an output device 911 such as a display, via an input / output I / F 905. The CPU 901 acquires data from the input device 910 via the input / output I / F 905, and outputs generated data to the output device 911.

[0294] The HDD 904 stores programs executed by the CPU 901 and data used by the programs, etc. The communication I / F 906 receives data from other devices via a communication network (e.g., NW (Network) 920) and outputs the data to the CPU 901, and also transmits data generated by the CPU 901 to other devices via the communication network.

[0295] The media I / F 907 reads a program or data stored in the recording medium 912 and outputs it to the CPU 901 via the RAM 903. The CPU 901 loads a program related to a target process from the recording medium 912 onto the RAM 903 via the media I / F 907, and executes the loaded program. The recording medium 912 is an optical recording medium such as a DVD (Digital Versatile Disc) or a PD (Phase change rewritable Disc), a magneto-optical recording medium such as an MO (Magneto Optical disk), a magnetic recording medium, a conductive memory tape medium, a semiconductor memory, or the like.

[0296] For example, when the computer 900 functions as the offload server 1, 1A according to the first and second embodiments, the CPU 901 of the computer 900 executes a program loaded onto the RAM 903 to realize the function of the offload server 1, 1A. Furthermore, the HDD 904 stores data in the RAM 903. The CPU 901 reads and executes a program related to a target process from a recording medium 912. Additionally, the CPU 901 may read a program related to a target process from another device via a communication network (NW 920).

[0297] [effect] As described above, the offload server 1 according to the first embodiment includes an application code analysis unit 112 that analyzes the source code of an application, a data transfer specification unit 113 that specifies, among variables that need to be transferred between the CPU and GPU based on the results of the code analysis, that variables that are not mutually referenced or updated by the CPU processing and the GPU processing and that only return the results of GPU processing to the CPU, to be collectively transferred before and after the start and end of GPU processing, a parallel processing specification unit 114 that specifies loop statements of the application, specifies a parallel processing specification statement in the GPU for each specified loop statement, and compiles the specified loop statements, and specifies that loop statements that produce a compilation error are not to be offloaded, and specifies that loop statements that do not produce a compilation error are to be transferred in parallel a performance measurement unit 116 that compiles an application of the parallel processing pattern, places it on the accelerator verification device, and executes processing for measuring performance when offloaded to the accelerator; an evaluation value setting unit 116c that sets an evaluation value including the processing time and power consumption required at the time of offloading measured by the performance measurement unit 116, the evaluation value being higher as the processing time and power consumption are lower; and an execution file creation unit 117 that selects the parallel processing pattern with the highest evaluation value from the plurality of parallel processing patterns based on the measurement results of the processing time and power consumption, and compiles the parallel processing pattern with the highest evaluation value to create an execution file.

[0298] By doing this, instead of transferring instructions to the GPU (such as data copy) that are distributed throughout the program to the GPU individually, variables that can be transferred in bulk are collected together and transferred / instructed in a lump, reducing transfers between the CPU and GPU and further speeding up offloading.In addition, by evaluating not only the processing time during automatic offloading but also the power consumption, it is possible to reduce power consumption (low power consumption) while improving performance.

[0299] The offload server 1A according to the second embodiment includes an application code analysis unit 112 that analyzes the source code of an application, a PLD processing specification unit 213 that identifies loop statements in the application and creates and compiles a plurality of offload processing patterns that specify pipeline processing and parallel processing in the PLD in OpenCL for each identified loop statement, an arithmetic strength calculation unit 214 that calculates the arithmetic strength of the loop statement of the application, a PLD processing pattern creation unit 215 that narrows down loop statements whose arithmetic strength is higher than a predetermined threshold based on the arithmetic strength calculated by the arithmetic strength calculation unit 214 as offload candidates and creates a PLD processing pattern, and the performance measurement unit 116 that compiles an application of the selected PLD processing pattern, places it on the accelerator verification device 14, and executes processing for measuring performance when offloaded to the PLD; an evaluation value setting unit 116c that sets an evaluation value including the processing time and power consumption based on the processing time and power consumption required during offloading measured by the performance measurement unit 116, the evaluation value being higher as the processing time and power consumption are lower; and an executable file creation unit 117 that selects the PLD processing pattern with the highest evaluation value from the multiple PLD processing patterns based on the measurement results of the processing time and power consumption, and compiles the PLD processing pattern with the highest evaluation value to create an executable file.

[0300] In this way, the number of times performance measurements are performed can be reduced by narrowing down the patterns to be actually measured, then placing them in the verification environment, compiling them, and measuring performance on an actual PLD (e.g., FPGA). This allows automatic offloading of application loop statements to the PLD at high speed. In addition, by evaluating not only the processing time during automatic offloading but also the power consumption, it is possible to reduce power consumption (low power consumption) while improving performance.

[0301] The offload server 1, 1A offloads a specific process of an application to at least one of a GPU, a menu core CPU, and a PLD. The offload server 1, 1A includes an application code analysis unit 112 that analyzes the source code of the application, and a CPU (Central Processing Unit) that performs the offloading based on the result of the code analysis. Among the variables that require transfer between the application program (a multi-core unit) and the GPU or menu core CPU, for which the CPU processing or menu core CPU processing and the GPU processing are not mutually referenced or updated, and the result of the GPU processing or menu core CPU processing is simply returned to the CPU, a data transfer specification unit 113 specifies data transfer in a batch before and after the start and end of the GPU processing or menu core CPU processing, a parallel processing specification unit 114 specifies a loop statement for the GPU or a loop statement for the many-core CPU of the application, and specifies a parallel processing specification statement in the GPU for each specified loop statement, and compiles the specified loop statement, a PLD processing specification unit 213 specifies a loop statement for the PLD of the application, and creates pipeline processing and parallel processing in the PLD according to a plurality of offload processing patterns specified by OpenCL for each specified loop statement for the PLD, and compiles the specified loop statement, an arithmetic strength calculation unit 214 calculates the arithmetic strength of the loop statement for the PLD of the application, and an arithmetic a PLD processing pattern creation unit 215 which, based on the arithmetic strength calculated by the strength calculation unit 214, narrows down loop statements having arithmetic strengths higher than a predetermined threshold as offload candidates and creates a PLD processing pattern; a parallel processing pattern creation unit 214 which excludes loop statements for a GPU or loop statements for a many-core CPU which generate a compilation error from being offloaded, and creates a parallel processing pattern which specifies whether or not to perform parallel processing for loop statements for a GPU or loop statements for a many-core CPU which do not generate a compilation error; a performance measurement unit 116 which, in a mixed environment of a GPU, a menu core CPU, and a PLD, compiles an application of the parallel processing pattern or the PLD processing pattern, places it on an accelerator verification device, and executes each performance measurement process when offloaded to the GPU, the menu core CPU, and the PLD; and, based on the processing time and power consumption required when offloading the GPU, the menu core CPU, and the PLD measured by the performance measurement unit 116,The evaluation value setting unit 116c sets an evaluation value including processing time and power consumption, the lower the processing time and power consumption, the higher the evaluation value becomes; and an execution file creation unit 117 selects one of the GPU, the CPU core, and the PLD with the best processing time and power consumption based on the measurement results of the processing time and power consumption of the GPU, the CPU core, and the PLD, selects the parallel processing pattern or the PLD processing pattern with the highest evaluation value from a plurality of parallel processing patterns or PLD processing patterns for the selected one, and compiles the parallel processing pattern or the PLD processing pattern with the highest evaluation value to create an execution file.

[0302] By doing this, when there is a mixture of GPU, FPGA, and many-core CPU as migration destinations, it is possible to verify the three migration destinations of GPU, FPGA, and many-core CPU, and automatically select and offload the migration destination that offers superior performance and low power consumption.

[0303] [Other Effects] In the offload server 1 according to the first embodiment, the parallel processing designation unit 114 determines the number of loop statements that do not generate a compilation error as the gene length based on a genetic algorithm, the parallel processing pattern creation unit 115 maps the accelerator processing availability to the gene pattern by setting either 1 or 0 when GPU processing is performed and the other 0 or 1 when GPU processing is not performed, and prepares a designated number of gene patterns in which each gene value is randomly created as 1 or 0, the performance measurement unit 116 compiles application code that specifies a parallel processing designation statement for the GPU according to each individual, places it in the accelerator verification device 14, and executes processing for performance measurement in the accelerator verification device, the executable file creation unit 117 performs performance measurement for each individual, evaluates the individuals so that the shorter the processing time, the higher the fitness, and selects from each individual an individual with a fitness higher than a predetermined value as an individual with high performance, performs crossover and mutation processing on the selected individuals to create individuals of the next generation, and selects the parallel processing pattern with the highest performance as a solution after the processing for the designated number of generations is completed.

[0304] By doing this, first check for loop statements that can be parallelized, then use GA to repeatedly perform performance verification trials in a verification environment for parallelizable repetitive statements to search for an appropriate region. After narrowing down to parallelizable loop statements (e.g., for statements), parallel processing patterns that can be accelerated are retained and recombined in the form of gene parts, making it possible to efficiently search for patterns that can be accelerated from the vast number of possible parallel processing patterns.

[0305] In the offload server 1A according to the second embodiment, the PLD processing pattern creation unit 215 measures the number of loops in the loop statements of an application, and narrows down the loop statements having an arithmetic strength higher than a predetermined threshold and a loop count greater than a predetermined number as candidates for offloading.

[0306] In this way, by narrowing down loop statements with high arithmetic strength and a large number of loop counts, the loop statements can be further narrowed down, and automatic offloading of the loop statements of the application can be performed at higher speed.

[0307] In the offload server 1A according to the second embodiment, the PLD processing pattern creation unit 215 creates an OpenCL for offloading each of the narrowed-down loop statements to a PLD, precompiles the created OpenCL to calculate the amount of resources required for PLD processing, and further narrows down the offload candidates based on the calculated resource amounts.

[0308] In this way, the arithmetic strength, loop count, and resource amount of loop statements are analyzed, and loop statements with high resource efficiency are narrowed down as offload candidates, thereby making it possible to narrow down loop statements further while preventing excessive consumption of PLD (e.g., FPGA) resources, and to perform automatic offloading of loop statements of an application more quickly. Also, since the calculation of the resource amount for PLD processing takes only a few minutes up to the intermediate state stage such as HDL, the amount of resources to be used can be known in a short time even before compilation is completed.

[0309] The present invention provides an offload program for causing a computer to function as the offload server.

[0310] In this way, each function of the offload server 1 can be realized using a general computer.

[0311] Furthermore, among the processes described in each of the above embodiments, all or part of the processes described as being performed automatically can be performed manually, or all or part of the processes described as being performed manually can be performed automatically by a known method. In addition, the information including the processing procedures, control procedures, specific names, various data, and parameters shown in the above documents and drawings can be changed arbitrarily unless otherwise specified. In addition, each component of each device shown in the figure is a functional concept, and does not necessarily have to be physically configured as shown in the figure. In other words, the specific form of distribution and integration of each device is not limited to that shown in the figure, and all or part of them can be functionally or physically distributed and integrated in any unit according to various loads, usage conditions, etc.

[0312] In addition, the above-mentioned configurations, functions, processing units, processing means, etc. may be realized in part or in whole by hardware, for example, by designing them as integrated circuits. In addition, the above-mentioned configurations, functions, etc. may be realized by software for a processor to interpret and execute a program that realizes each function. Information on the programs, tables, files, etc. that realize each function can be stored in a memory, a recording device such as a hard disk or SSD (Solid State Drive), or a recording medium such as an IC (Integrated Circuit) card, SD (Secure Digital) card, or optical disk.

[0313] In addition, in this embodiment, a genetic algorithm (GA) method is used to find a solution to a combinatorial optimization problem within a limited optimization period, but any optimization method may be used, such as local search, dynamic programming, or a combination of these.

[0314] In addition, in this embodiment, an OpenACC compiler for C / C++ is used, but any compiler that can offload GPU processing may be used. For example, Java lambda (registered trademark) GPU processing or IBM Java 9 SDK (registered trademark) may be used. Note that the parallel processing specification statement depends on these development environments. For example, in Java (registered trademark), it has been possible to write parallel processing in lambda format since Java 8. IBM (registered trademark) provides a JIT compiler that offloads parallel processing descriptions in lambda format to the GPU. In Java, similar offloading is possible by using these and tuning the GA to determine whether or not to convert loop processing to lambda format.

[0315] In addition, in this embodiment, a for statement is exemplified as a repetitive statement (loop statement), but other statements such as a while statement and a do-while statement are also included. However, a for statement that specifies the continuation condition of the loop, etc., is more suitable. [Explanation of symbols]

[0316] 1,1A Offload Server 11,21 Control section 12 Input / output section 13 Storage section 14 Verification machine (accelerator verification device) 111 Application code designation section 112 Application Code Analysis Unit 113 Data transfer specification section 114 Parallel processing specification section 114a, 213a Off-road range extraction section 114b,213b Intermediate language file output section 115 Parallel Processing Pattern Creation Unit 116 Performance measurement section 116a Binary file placement section 116b Power usage measurement section (performance measurement section) 116c Evaluation value setting unit 117 Executable File Creation Department 118 Production Environment Deployment Department 119 Performance Measurement Test Extraction Execution Department 120 User Provision Department 130 Application Code 131 Test Case DB 132 Intermediate language files 151 Various Devices 152 CPU-GPU devices 153 Device with CPU-FPGA 154 Devices with CPU 215 PLD Processing Pattern Creation Department

Claims

1. An offload server that offloads a specific process of an application to a GPU (Graphics Processing Unit), an application code analysis unit that analyzes a source code of an application; a data transfer specification unit that specifies that, among variables that need to be transferred between a CPU (Central Processing Unit) and the GPU based on the results of the code analysis, those variables that are not mutually referenced or updated by the CPU processing and the GPU processing and that only return the results of the GPU processing to the CPU are to be transferred in a batch before the start and after the end of the GPU processing; a parallel processing specification unit that specifies loop statements of the application, and specifies a parallel processing specification statement for the GPU for each of the specified loop statements to compile the loop statements; a parallel processing pattern creation unit that creates a parallel processing pattern that excludes a loop statement that generates a compilation error from being subject to offloading and specifies whether or not a loop statement that does not generate a compilation error should be processed in parallel; a performance measurement unit that compiles the application of the parallel processing pattern, places it in an accelerator verification device, and executes a process for measuring performance when the application is offloaded to the GPU; an evaluation value setting unit that sets an evaluation value including a processing time and an amount of power consumption required during offloading based on the processing time and the amount of power consumption measured by the performance measurement unit, the evaluation value being higher as the processing time and the amount of power consumption are lower; an execution file creation unit that selects a parallel processing pattern with a highest evaluation value from the plurality of parallel processing patterns based on the measurement results of the processing time and the power consumption, and compiles the parallel processing pattern with the highest evaluation value to create an execution file; An offload server comprising:

2. An offload server that offloads a specific process of an application to a programmable logic device (PLD), an application code analysis unit that analyzes a source code of an application; a PLD processing specification unit that specifies loop statements of the application, and creates and compiles pipeline processing and parallel processing in the PLD for each of the specified loop statements according to a plurality of offload processing patterns specified in OpenCL (Open Computing Language) (trademark); an arithmetic strength calculation unit that calculates an arithmetic strength of a loop statement of the application; a PLD processing pattern creation unit that narrows down loop statements having arithmetic strengths higher than a predetermined threshold as offload candidates based on the arithmetic strengths calculated by the arithmetic strength calculation unit, and creates a PLD processing pattern; a performance measurement unit that compiles the application of the created PLD processing pattern, arranges the application in an accelerator verification device, and executes a process for performance measurement when the application is offloaded to the PLD; an evaluation value setting unit that sets an evaluation value including a processing time and an amount of power consumption required during offloading based on the processing time and the amount of power consumption measured by the performance measurement unit, the evaluation value being higher as the processing time and the amount of power consumption are lower; an execution file creation unit that selects a PLD processing pattern with a highest evaluation value from the plurality of PLD processing patterns based on the measurement results of the processing time and the power consumption, and compiles the PLD processing pattern with the highest evaluation value to create an execution file; An offload server comprising:

3. An offload server that offloads a specific process of an application to at least one of a graphics processing unit (GPU), a menu core CPU, and a programmable logic device (PLD), an application code analysis unit that analyzes a source code of an application; a data transfer specification unit that specifies that among variables that need to be transferred between a CPU (Central Processing Unit) and the GPU or the menu core CPU based on the results of the code analysis, variables that are not mutually referenced or updated between the CPU processing or the menu core CPU processing and the GPU processing, and that only return the results of the GPU processing or the menu core CPU processing to the CPU, should be transferred in a batch before and after the start and end of the GPU processing or the menu core CPU processing; a parallel processing specification unit that specifies a loop statement for a GPU or a loop statement for a many-core CPU of the application, and specifies a parallel processing specification statement for the GPU for each of the specified loop statements and compiles the loop statements; a PLD processing specification unit that specifies loop statements for a PLD of the application, and creates and compiles pipeline processing and parallel processing in the PLD for each of the specified loop statements for the PLD according to a plurality of offload processing patterns specified in OpenCL (Open Computing Language) (trademark); an arithmetic strength calculation unit that calculates arithmetic strength of a PLD-oriented loop statement of the application; a PLD processing pattern creation unit that narrows down loop statements having arithmetic strengths higher than a predetermined threshold as offload candidates based on the arithmetic strengths calculated by the arithmetic strength calculation unit, and creates a PLD processing pattern; a parallel processing pattern creation unit that creates a parallel processing pattern that excludes a loop statement for a GPU or a loop statement for a many-core CPU that generates a compilation error from being subject to offloading, and specifies whether or not to perform parallel processing on a loop statement for a GPU or a loop statement for a many-core CPU that does not generate a compilation error; a performance measurement unit that compiles the application of the parallel processing pattern or the PLD processing pattern in a mixed environment of the GPU, the menu core CPU, and the PLD, places the application in an accelerator verification device, and executes each performance measurement process when offloaded to the GPU, the menu core CPU, and the PLD; an evaluation value setting unit that sets an evaluation value including a processing time and an amount of power consumption based on the processing time and the amount of power consumption required for offloading the GPU, the menu core CPU, and the PLD measured by the performance measurement unit, the evaluation value being higher as the processing time and the amount of power consumption are lower; an executable file creation unit which selects one of the GPU, the menu core CPU, and the PLD having the best processing time and power consumption based on the measurement results of the processing times and power consumption of the GPU, the menu core CPU, and the PLD, selects a parallel processing pattern or a PLD processing pattern having a highest evaluation value from a plurality of the parallel processing patterns or PLD processing patterns for the selected one, and compiles the parallel processing pattern or the PLD processing pattern having the highest evaluation value to create an executable file; An offload server comprising:

4. An offload control method for an offload server that offloads a specific process of an application to a GPU (Graphics Processing Unit), comprising: The offload server includes: analyzing the source code of the application; a step of designating, based on the result of the code analysis, that among the variables that need to be transferred between a CPU (Central Processing Unit) and the GPU, those variables that are not mutually referenced or updated by the CPU processing and the GPU processing, and that only return the results of the GPU processing to the CPU, to be transferred in a lump before the start and after the end of the GPU processing; identifying loop statements of the application, and for each of the identified loop statements, specifying a parallel processing specification statement for the GPU and compiling the loop statements; creating a parallel processing pattern for specifying that a loop statement that generates a compilation error is not to be offloaded and that a loop statement that does not generate a compilation error is to be processed in parallel or not; A step of compiling the application of the parallel processing pattern, disposing it on an accelerator verification device, and executing a process for measuring performance when offloaded to the GPU; A step of setting an evaluation value including the processing time and the power consumption required during offloading based on the measured processing time and the power consumption, the evaluation value being higher as the processing time and the power consumption are lower; selecting a parallel processing pattern with a highest evaluation value from the plurality of parallel processing patterns based on the measurement results of the processing time and the power consumption, and compiling the parallel processing pattern with the highest evaluation value to create an executable file. An off-load control method comprising:

5. An offload control method for an offload server that offloads a specific process of an application to a programmable logic device (PLD), comprising: The offload server includes: analyzing the source code of the application; Identifying loop statements of the application, and for each of the identified loop statements, creating pipeline processing, parallel processing, and expansion processing in the PLD according to a plurality of offload processing patterns specified in OpenCL (Open Computing Language) (trademark) and compiling the same; calculating an arithmetic strength of a loop statement of the application; based on the calculated arithmetic strength, narrowing down loop statements having arithmetic strengths higher than a predetermined threshold as offload candidates, and creating a PLD processing pattern; a step of compiling the application of the created PLD processing pattern, disposing the application in an accelerator verification device, and executing a process for measuring performance when offloaded to the PLD; A step of setting an evaluation value including the processing time and the power consumption required during offloading based on the measured processing time and the power consumption, the evaluation value being higher as the processing time and the power consumption are lower; Based on the measurement results of the processing time and the power consumption, a PLD processing pattern with the highest evaluation value is selected from the plurality of PLD processing patterns, and the PLD processing pattern with the highest evaluation value is compiled to generate an executable file. An off-load control method comprising:

6. An offload control method for an offload server that offloads a specific process of an application to at least one of a graphics processing unit (GPU), a menu core CPU, and a programmable logic device (PLD), comprising: The offload server includes: analyzing the source code of the application; a step of designating, based on the result of the code analysis, that among the variables that need to be transferred between a CPU (Central Processing Unit) and the GPU or the menu core CPU, those variables that are not mutually referenced or updated between the CPU processing or the menu core CPU processing and the GPU processing, and that only return the results of the GPU processing or the menu core CPU processing to the CPU, to transfer the data in a lump before the start and after the end of the GPU processing or the menu core CPU processing; Identifying a loop statement for a GPU or a loop statement for a many-core CPU of the application, and compiling each of the identified loop statements by specifying a parallel processing specification statement for the GPU; Identifying loop statements for a PLD of the application, and for each of the identified loop statements for the PLD, creating pipeline processing and parallel processing in the PLD according to a plurality of offload processing patterns specified in OpenCL (Open Computing Language) (trademark) and compiling the same; calculating an arithmetic strength of a PLD-oriented loop statement of the application; a PLD processing pattern creation step of narrowing down loop statements having arithmetic strengths higher than a predetermined threshold as offload candidates based on the calculated arithmetic strengths and creating a PLD processing pattern; a parallel processing pattern creating step of creating a parallel processing pattern that excludes a loop statement for a GPU or a loop statement for a many-core CPU that generates a compilation error from being subject to offloading, and specifies whether or not to perform parallel processing on a loop statement for a GPU or a loop statement for a many-core CPU that does not generate a compilation error; In a mixed environment of the GPU, the menu core CPU, and the PLD, compiling the application of the parallel processing pattern or the PLD processing pattern, disposing the application on an accelerator verification device, and executing each performance measurement process when offloaded to the GPU, the menu core CPU, and the PLD; A step of setting an evaluation value including the processing time and the power consumption based on the measured processing time and the power consumption required for offloading the GPU, the menu core CPU, and the PLD, the evaluation value being higher as the processing time and the power consumption are lower; based on the measurement results of the processing times and power consumption of the GPU, the menu core CPU, and the PLD, selecting one of the GPU, the menu core CPU, and the PLD with the best processing time and power consumption, selecting a parallel processing pattern or PLD processing pattern with the highest evaluation value from a plurality of the parallel processing patterns or PLD processing patterns for the selected one, and compiling the parallel processing pattern or PLD processing pattern with the highest evaluation value to create an executable file. An off-load control method comprising:

7. An offload program for causing a computer to function as the offload server according to any one of claims 1 to 3.

Images