JP7521597B2 - OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM - Google Patents
OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM Download PDFInfo
- Publication number
- JP7521597B2 JP7521597B2 JP2022560579A JP2022560579A JP7521597B2 JP 7521597 B2 JP7521597 B2 JP 7521597B2 JP 2022560579 A JP2022560579 A JP 2022560579A JP 2022560579 A JP2022560579 A JP 2022560579A JP 7521597 B2 JP7521597 B2 JP 7521597B2
- Authority
- JP
- Japan
- Prior art keywords
- processing
- performance
- unit
- gpu
- offload
- Prior art date
- Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
- Active
Links
- 238000000034 method Methods 0.000 title claims description 220
- PWPJGUXAGUPAHP-UHFFFAOYSA-N lufenuron Chemical compound C1=C(Cl)C(OC(F)(F)C(C(F)(F)F)F)=CC(Cl)=C1NC(=O)NC(=O)C1=C(F)C=CC=C1F PWPJGUXAGUPAHP-UHFFFAOYSA-N 0.000 title 1
- 238000012545 processing Methods 0.000 claims description 706
- 238000005259 measurement Methods 0.000 claims description 192
- 230000008569 process Effects 0.000 claims description 172
- 238000004458 analytical method Methods 0.000 claims description 110
- 108090000623 proteins and genes Proteins 0.000 claims description 105
- 238000012546 transfer Methods 0.000 claims description 87
- 238000001514 detection method Methods 0.000 claims description 81
- 238000012795 verification Methods 0.000 claims description 68
- 238000012360 testing method Methods 0.000 claims description 40
- 230000035772 mutation Effects 0.000 claims description 39
- 239000000284 extract Substances 0.000 claims description 23
- 238000003860 storage Methods 0.000 claims description 9
- 239000011159 matrix material Substances 0.000 claims description 8
- 230000014509 gene expression Effects 0.000 claims description 7
- 238000003491 array Methods 0.000 claims description 2
- HPTJABJPZMULFH-UHFFFAOYSA-N 12-[(Cyclohexylcarbamoyl)amino]dodecanoic acid Chemical compound OC(=O)CCCCCCCCCCCNC(=O)NC1CCCCC1 HPTJABJPZMULFH-UHFFFAOYSA-N 0.000 claims 9
- 230000006870 function Effects 0.000 description 247
- 238000004519 manufacturing process Methods 0.000 description 36
- 238000011056 performance test Methods 0.000 description 23
- 238000010586 diagram Methods 0.000 description 22
- 238000000605 extraction Methods 0.000 description 22
- 239000008186 active pharmaceutical agent Substances 0.000 description 17
- 238000005516 engineering process Methods 0.000 description 15
- 230000003252 repetitive effect Effects 0.000 description 13
- 230000002068 genetic effect Effects 0.000 description 12
- 238000004549 pulsed laser deposition Methods 0.000 description 12
- 238000004364 calculation method Methods 0.000 description 8
- 238000002360 preparation method Methods 0.000 description 8
- 238000004891 communication Methods 0.000 description 7
- 230000003044 adaptive effect Effects 0.000 description 5
- 238000012790 confirmation Methods 0.000 description 5
- 238000011156 evaluation Methods 0.000 description 5
- 230000015572 biosynthetic process Effects 0.000 description 4
- 230000000694 effects Effects 0.000 description 4
- 238000003500 gene array Methods 0.000 description 4
- 238000010191 image analysis Methods 0.000 description 4
- 238000003786 synthesis reaction Methods 0.000 description 4
- 230000004888 barrier function Effects 0.000 description 3
- 230000008859 change Effects 0.000 description 3
- 238000013507 mapping Methods 0.000 description 3
- 238000005457 optimization Methods 0.000 description 3
- 238000003672 processing method Methods 0.000 description 3
- 239000004065 semiconductor Substances 0.000 description 3
- 238000005266 casting Methods 0.000 description 2
- 238000011161 development Methods 0.000 description 2
- 230000010354 integration Effects 0.000 description 2
- 230000004048 modification Effects 0.000 description 2
- 230000003287 optical effect Effects 0.000 description 2
- 238000010187 selection method Methods 0.000 description 2
- BUGBHKTXTAQXES-UHFFFAOYSA-N Selenium Chemical compound [Se] BUGBHKTXTAQXES-UHFFFAOYSA-N 0.000 description 1
- 230000001133 acceleration Effects 0.000 description 1
- 230000009471 action Effects 0.000 description 1
- 238000013459 approach Methods 0.000 description 1
- 230000008901 benefit Effects 0.000 description 1
- 230000000052 comparative effect Effects 0.000 description 1
- 238000013135 deep learning Methods 0.000 description 1
- 230000001419 dependent effect Effects 0.000 description 1
- 238000013461 design Methods 0.000 description 1
- 238000009826 distribution Methods 0.000 description 1
- 230000010429 evolutionary process Effects 0.000 description 1
- 238000003780 insertion Methods 0.000 description 1
- 230000037431 insertion Effects 0.000 description 1
- 238000010801 machine learning Methods 0.000 description 1
- 238000012986 modification Methods 0.000 description 1
- 238000012544 monitoring process Methods 0.000 description 1
- 238000004321 preservation Methods 0.000 description 1
- 229910052711 selenium Inorganic materials 0.000 description 1
- 239000011669 selenium Substances 0.000 description 1
Images
Classifications
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F8/00—Arrangements for software engineering
- G06F8/40—Transformation of program code
- G06F8/41—Compilation
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06N—COMPUTING ARRANGEMENTS BASED ON SPECIFIC COMPUTATIONAL MODELS
- G06N3/00—Computing arrangements based on biological models
- G06N3/12—Computing arrangements based on biological models using genetic models
Landscapes
- Engineering & Computer Science (AREA)
- Theoretical Computer Science (AREA)
- Physics & Mathematics (AREA)
- General Engineering & Computer Science (AREA)
- Biophysics (AREA)
- Software Systems (AREA)
- Health & Medical Sciences (AREA)
- Life Sciences & Earth Sciences (AREA)
- General Physics & Mathematics (AREA)
- Artificial Intelligence (AREA)
- Data Mining & Analysis (AREA)
- Evolutionary Biology (AREA)
- Genetics & Genomics (AREA)
- Bioinformatics & Cheminformatics (AREA)
- Biomedical Technology (AREA)
- Computational Linguistics (AREA)
- Bioinformatics & Computational Biology (AREA)
- Evolutionary Computation (AREA)
- General Health & Medical Sciences (AREA)
- Molecular Biology (AREA)
- Computing Systems (AREA)
- Mathematical Physics (AREA)
- Devices For Executing Special Programs (AREA)
Description
本発明は、機能処理をGPU(Graphics Processing Unit)等のアクセラレータに自動オフロードするオフロードサーバ、オフロード制御方法およびオフロードプログラムに関する。 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).
近年、CPUの半導体集積度が1.5年で2倍になるというムーアの法則が減速するのではないかと言われている。そのような状況から、少コアのCPUだけでなく、FPGA(Field Programmable Gate Array)やGPU(Graphics Processing Unit)等のデバイスの活用が増えている。例えば、Microsoft(登録商標)社はFPGAを使ってBingの検索効率を高めるといった取り組みをしており、Amazon(登録商標)社は、FPGA, GPU等をクラウドのインスタンスとして提供している。In recent years, it has been said that Moore's Law, which states that the semiconductor integration density of CPUs will double every 1.5 years, may be slowing down. In light of this situation, there has been an increase in the use of devices such as FPGAs (Field Programmable Gate Arrays) and GPUs (Graphics Processing Units) in addition to CPUs with fewer cores. For example, Microsoft (registered trademark) is working to improve the search efficiency of Bing using FPGAs, and Amazon (registered trademark) is providing FPGAs, GPUs, etc. as cloud instances.
少コアのCPU以外のデバイスをシステムで適切に活用するためには、デバイス特性を意識した設定やプログラム作成が必要であり、OpenMP(Open Multi-Processing)、OpenCL(Open Computing Language)、CUDA(Compute Unified Device Architecture)といった知識が必要になるため、大半のプログラマにとっては、スキルの壁が高い。 In order to properly utilize devices other than low-core CPUs in a system, it is necessary to configure and write programs that take into account the characteristics of the device, and knowledge of OpenMP (Open Multi-Processing), OpenCL (Open Computing Language), and CUDA (Compute Unified Device Architecture) is required, which creates a high skill barrier for most programmers.
少コアのCPU以外のGPUやFPGA、メニーコアCPU等のデバイスを活用するシステムは今後ますます増えていくと予想されるが、それらを最大限活用するには、技術的壁が高い。そこで、そのような壁を取り払い、少コアのCPU以外のデバイスを十分利用できるようにするため、プログラマが処理ロジックを記述したソフトウェアを、配置先の環境(FPGA、GPU、メニーコアCPU等)にあわせて、適応的に変換、設定し、環境に適合した動作をさせるような、プラットフォームが求められている。 It is expected that the number of systems that utilize devices other than few-core CPUs, such as GPUs, FPGAs, and many-core CPUs, will continue to increase in the future, but there are high technical barriers to making the most of them. Therefore, in order to remove such barriers and make full use of devices other than few-core CPUs, there is a demand for a platform that can adaptively convert and configure software in which a programmer writes processing logic to suit the environment in which it is deployed (FPGA, GPU, many-core CPU, etc.), allowing it to operate in accordance with the environment.
非特許文献1には、一度記述したコードを、配置先の環境に存在するGPUやFPGA、メニーコアCPU等を利用できるように、変換、リソース設定等を自動で行い、アプリケーションを高性能に動作させることを目的とした、環境適応ソフトウェアが記載されている。Non-patent
非特許文献2、3、4には、環境適応ソフトウェアの要素として、アプリケーションコードのループ文および機能ブロックを、FPGA、GPUに自動オフロードする方式が記載されている。Non-patent
GPUの並列計算パワーを画像処理でないものにも使うGPGPU(General Purpose GPU)を行うための環境としてCUDAが普及している。CUDAは、GPGPU向けのNVIDIA(登録商標)社の環境である。また、FPGA、メニーコアCPU、GPU等のヘテロなデバイスを同じように扱うための仕様としてOpenCLがあり、その開発環境も整いつつある。CUDA、OpenCLは、C言語の拡張を行いプログラムを行う形であり、プログラムの難度は高い(FPGA等のカーネルとCPUのホストとの間のメモリデータのコピーや解放の記述を明示的に行う等)。 CUDA has become popular as an environment for running GPGPUs (General Purpose GPUs), which use the parallel computing power of GPUs for purposes other than image processing. CUDA is an environment for GPGPUs developed by NVIDIA (registered trademark). There is also OpenCL, a specification for treating heterogeneous devices such as FPGAs, many-core CPUs, and GPUs in the same way, and its development environment is also in the process of being established. CUDA and OpenCL are programmed by extending the C language, and are difficult to program (such as explicitly writing the copying and release of memory data between the kernel of an FPGA or the like and the host CPU).
また、CUDAやOpenCLに比べて、より簡易にヘテロなデバイスを利用するための技術として、OpenACCやOpenMP等、コンパイラとしてPGIコンパイラやgcc(登録商標)等がある。このコンパイラは、指示行ベースで、並列処理等を行う箇所を指定して、指示行に従って、GPU、メニーコアCPU等に向けて実行ファイルを作成する。 In addition, compared to CUDA and OpenCL, there are technologies such as OpenACC and OpenMP that allow easier use of heterogeneous devices, and compilers such as the PGI compiler and gcc (registered trademark). These compilers specify the locations where parallel processing will be performed on a directive line basis, and create executable files for GPUs, many-core CPUs, etc. according to the directive lines.
上記、CUDA、OpenCL、OpenACC、OpenMP等の技術仕様を用いることで、FPGAやGPU、メニーコアCPUへオフロードすることは可能になっている。しかしながら、デバイス処理自体は行えるようになっても、高速化することには課題がある。例えば、マルチコアCPU向けに自動並列化機能を持つコンパイラとして、Intelコンパイラ(登録商標)等がある。これらは、自動並列化時に、コードの中のループ文中で並列処理可能な部分を抽出して、並列化している。しかし、メモリ処理等の影響で単に並列化可能ループ文を並列化しても性能がでないことも多い。FPGAやGPU等で高速化する際には、OpenCLやCUDAの技術者がチューニングを繰り返したり、OpenACCコンパイラ等を用いて適切な並列処理範囲を探索し試行することがされている。 By using the above technical specifications such as CUDA, OpenCL, OpenACC, and OpenMP, it is possible to offload to FPGAs, GPUs, and many-core CPUs. However, even if device processing itself can be performed, there are challenges in accelerating it. For example, Intel Compiler (registered trademark) is a compiler with an automatic parallelization function for multi-core CPUs. During automatic parallelization, these extract parts of the loop statements in the code that can be parallelized and parallelize them. However, due to the influence of memory processing, etc., simply parallelizing parallelizable loop statements often does not achieve the desired performance. When accelerating FPGAs, GPUs, etc., OpenCL and CUDA engineers repeatedly tune the system, or use OpenACC compilers, etc. to search for and try out the appropriate range of parallel processing.
このため、技術スキルが乏しいプログラマが、FPGAやGPU、メニーコアCPUを活用してアプリケーションを高速化することは難しいし、自動並列化技術等を使う場合も並列処理箇所探索の試行錯誤等の稼働が必要だった。現状、ヘテロなデバイスに対するオフロードは手動での取組みが主流である。 For this reason, it is difficult for programmers with little technical skill to use FPGAs, GPUs, or many-core CPUs to speed up applications, and even when using automatic parallelization technology, it is necessary to go through a process of trial and error to find parallel processing locations. Currently, offloading to heterogeneous devices is mainly done manually.
非特許文献1~4に記載の技術は、C言語プログラムのGPUやFPGAへのオフロードであり、Python(登録商標)、Java(登録商標)等の多様な移行元言語は想定されていない。
C言語だけでなく、Python、Javaと移行元言語が多様となった場合でも、アプリケーションプログラムを自動オフロードすることが要請されている。
The techniques described in Non-Patent
There is a demand for automatic offloading of application programs, even when the source languages are diverse, including not only C but also Python and Java.
このような点に鑑みて本発明がなされたのであり、移行元言語に合わせて、処理を検討したり実装する必要をなくし、移行元言語が多様となった場合でも、アプリケーションプログラムを自動でオフロードすることを課題とする。 This invention was made in consideration of these points, and its objective is to eliminate the need to consider and implement processing to match the source language, and to automatically offload application programs even when the source languages are diverse.
前記した課題を解決するため、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、前記アプリケーションプログラムは、Pythonアプリケーションプログラムであり、前記Pythonアプリケーションプログラムのソースコードを、Pythonを解析する構文解析ツールを用いて分析するアプリケーションコード分析部と、前記Pythonアプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、前記Pythonアプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、CUDAでの指示を追加したPythonコードをpyCUDAでインタプリットする際、CUDA文法を用いたGPU処理を指定してインタプリットする並列処理指定部と、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、前記並列処理パターンの前記Pythonアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するPythonアプリコードを含む最高性能の並列処理パターンを解とし、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、を備えることを特徴とするオフロードサーバとした。 In order to solve the above-mentioned problems, an offload server offloads a specific process of an application program to an accelerator , the application program being a Python application program, the offload server comprising: an application code analysis unit that analyzes source code of the Python application program using a syntax analysis tool that analyzes Python; a data transfer specification unit that analyzes reference relationships of variables used in loop statements of the Python application program, and for data that may be transferred outside the loop, performs data transfer specification using an explicit specification line that explicitly specifies data transfer outside the loop; a parallel processing specification unit that identifies loop statements of the Python application program, and when interpreting Python code with CUDA instructions added for each identified loop statement with pyCUDA, specifies GPU processing using CUDA grammar for interpretation; and a loop statement that generates a compilation error is excluded from offloading, and The offload server is characterized by comprising: 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 in which no pile error occurs; a performance measurement unit that compiles the Python application program of the parallel processing pattern, places it on an accelerator verification device, and executes processing for measuring performance when offloaded to the accelerator; and an executable file creation unit that selects a plurality of parallel processing patterns with high processing performance from the plurality of parallel processing patterns based on a performance measurement result, crosses the parallel processing patterns with high processing performance, creates a plurality of other parallel processing patterns by mutation processing, and performs new performance measurement. After a specified number of performance measurements, the parallel processing pattern with the highest processing performance is selected from the plurality of parallel processing patterns based on the performance measurement result, and after GA processing for a specified number of generations is completed, the parallel processing pattern with the highest processing performance that includes a Python application code that corresponds to the gene sequence with the highest performance is set as a solution, and the parallel processing pattern with the highest processing performance is compiled to create an executable file .
本発明によれば、移行元言語に合わせて、処理を検討したり実装する必要をなくし、移行元言語が多様となった場合でも、アプリケーションプログラムを自動でオフロードすることができる。 According to the present invention, there is no need to consider or implement processing depending on the source language, and application programs can be automatically offloaded even when the source languages are diverse.
次に、本発明を実施するための形態における、オフロードサーバ等について説明する。
以下、明細書の説明において、移行先環境としては、GPU、FPGA、メニーコアCPUの3つを想定した例について説明する。本発明は、プログラマブルロジックデバイス全般に適用可能である。
Next, an offload server and the like in an embodiment of the present invention will be described.
In the following description of the specification, three examples are assumed as destination environments: a GPU, an FPGA, and a many-core CPU. The present invention is applicable to programmable logic devices in general.
(多様移行元言語対応の基本的な考え方)
・移行元言語
本実施形態で対象とする多様な移行元言語としては、C言語、Python、Javaの3つとする。これら3つの言語は、毎月TIOBE(登録商標)が発表するプログラム言語の人気ランキングの上位3つであり、プログラマ人口が多い。また、C言語はコンパイル型、Pythonはインタプリタ型、Javaはその中間的方式と、方式上の多様性も3つでカバーされている。そのため、これら3つで共通的に利用できる方式であれば、より多くの言語への対応も容易と考える。
(Basic principles for dealing with diverse source languages)
Source Language The various source languages targeted in this embodiment are C, Python, and Java. These three languages are the top three in the popularity ranking of programming languages announced monthly by TIOBE (registered trademark), and have a large number of programmers. In addition, the C language is a compiled type, Python is an interpreted type, and Java is an intermediate type, so the three languages cover a diversity of formats. Therefore, if a format can be commonly used by these three languages, it is believed that it will be easy to support more languages.
本実施形態では、移行先環境が単なるCPUでない場合で、多様な移行元言語プログラムを、自動で高速にオフロードするために、検証環境の実機で性能測定し、進化計算手法等の手法と組み合わせて、徐々に高速なオフロードパターンを見つけるアプローチをとる。理由として、性能に関しては、コード構造だけでなく、処理するハードウェアのスペック、コンパイラやインタプリタ、データサイズ、ループ回数等の処理内容によって大きく変わるため、静的に予測する事が困難であり、動的な測定が必要だからである。実際に、市中には、ループ文を見つけコンパイル段階で並列化する自動並列化コンパイラがあるが、並列化可能ループ文の並列化だけでは性能を測定してみると低速になる場合も多いため、性能測定は必要である。In this embodiment, in order to automatically offload various source language programs at high speed when the destination environment is not simply a CPU, an approach is taken in which performance is measured on the actual machine in the verification environment, and a gradually faster offloading pattern is found by combining it with techniques such as evolutionary computing. The reason for this is that performance varies greatly depending not only on the code structure but also on the processing contents such as the specifications of the processing hardware, the compiler or interpreter, the data size, and the number of loops, making it difficult to predict statically and requiring dynamic measurement. In fact, there are automatic parallelizing compilers on the market that find loop statements and parallelize them at the compilation stage, but performance measurement is necessary because performance measurement often results in slow speeds when only parallelizing parallelizable loop statements.
・オフロードする対象
また、オフロードする対象については、アプリケーションプログラムのループ文および機能ブロックとするアプローチをとる。ループ文については、処理時間がかかるプログラムの処理の大半はループで費やされているという現状から、ループ文がオフロードのターゲットとして考えられる。一方、機能ブロックについては、特定処理を高速化する際に、処理内容や処理ハードウェアに適したアルゴリズムを用いることが多いため、個々のループ文の並列処理等に比べ、大きく高速化できる場合がある。行列積算やフーリエ変換等の頻繁に使われる機能ブロック単位で、GPU等の処理デバイスに応じたアルゴリズムで実装された処理(CUDAライブラリ等)に置換することで高速化する。
Targets for offloading The offloading target is the loop statements and function blocks of application programs. As for loop statements, most of the processing time of programs is currently spent in loops, so loop statements are considered as targets for offloading. On the other hand, when speeding up specific processing, algorithms suitable for the processing content and processing hardware are often used, so there are cases where the processing speed can be significantly increased compared to parallel processing of individual loop statements. The processing speed can be increased by replacing frequently used functional blocks such as matrix multiplication and Fourier transform with processing (CUDA library, etc.) implemented with algorithms suitable for processing devices such as GPUs.
・移行先環境
移行先環境としては、GPU、FPGA、メニーコアCPUの3つを想定し、これらが混在した環境でのC言語プログラムのオフロードも開示する。本発明の解決課題は、移行元言語が多様となった場合のアプリケーションの自動オフロードであるため、評価する移行先環境は限定されない。移行先環境は、一例としてGPUとし、FPGAやメニーコアCPUについては、GPUで共通的方式を確認できれば、その拡張で実現できる。
Destination environment Three destination environments are assumed: GPU, FPGA, and many-core CPU, and offloading of C language programs in a mixed environment is also disclosed. The problem to be solved by the present invention is automatic offloading of applications when the source language becomes diverse, so the destination environment to be evaluated is not limited. As an example of the destination environment, GPU is used, and FPGA and many-core CPU can be realized by extending it if a common method can be confirmed in GPU.
・共通的なGPUオフロード手法
共通的なGPUオフロード手法は、「ループ文のGPU自動オフロード(以下、ループ文オフロードという)」と「機能ブロックの自動オフロード(以下、機能ブロックオフロードという)」とに分けられ、それぞれ手法が異なる。
以下の説明において、第1の実施形態で「ループ文オフロード」を記載し、第2の実施形態で「機能ブロックの自動オフロード」を記載する。そして、第1の実施形態(「ループ文のGPU自動オフロード」)と第2の実施形態(「機能ブロックオフロード」)のそれぞれにおいて、構成と、共通処理とC言語とPythonとJavaとを説明する。目次で示すと下記である。
Common GPU offloading methods Common GPU offloading methods are divided into "GPU automatic offloading of loop statements (hereinafter referred to as loop statement offloading)" and "automatic offloading of function blocks (hereinafter referred to as function block offloading)", and each method is different.
In the following description, "loop statement offloading" will be described in the first embodiment, and "automatic offloading of function blocks" will be described in the second embodiment. Then, the configuration, common processing, C language, Python, and Java will be described for each of the first embodiment ("GPU automatic offloading of loop statements") and the second embodiment ("function block offloading"). The table of contents is as follows.
(目次)
・第1の実施形態(「ループ文オフロード」)の構成(図1)
共通(移行元言語において共通)処理(図2-図4)
共通フローチャート(図5A,図5B)
C言語の場合のフローチャート(図6A,図6B)
Pythonの場合の説明図(図7A,図7B,図8A,図8B)
Pythonの場合のフローチャート(図9A,図9B)
Javaの場合の説明図(図10A,図10B)
Javaの場合のフローチャート(図11A,図11B)
(table of contents)
Configuration of the first embodiment ("loop statement offload") (FIG. 1)
Common (common in source language) processing (Figure 2-Figure 4)
Common Flowchart (FIGS. 5A and 5B)
Flowchart for C language (Fig. 6A and Fig. 6B)
Diagram of Python (Fig. 7A, Fig. 7B, Fig. 8A, Fig. 8B)
Flowchart for Python (Figure 9A and Figure 9B)
Diagram of Java (Fig. 10A, Fig. 10B)
Flowchart for Java (Fig. 11A and Fig. 11B)
・第2の実施形態(「機能ブロックオフロード」)の構成(図12,図13)
共通フローチャート(図14,図15)
C言語の場合のフローチャート(図16,図17)
Pythonの場合のフローチャート(図18,図19)
Javaの場合のフローチャート(図20,図21)
Configuration of the second embodiment ("Function block offload") (FIGS. 12 and 13)
Common Flowchart (Fig. 14 and Fig. 15)
Flowchart for C language (Fig. 16, Fig. 17)
Flowchart for Python (Figure 18, Figure 19)
Flowchart for Java (Fig. 20, Fig. 21)
(第1の実施形態)
第1の実施形態は、ループ文オフロードについて記載する。
以下、第1の実施形態に係るオフロードサーバ1が、環境適応ソフトウェアシステムにおけるユーザ向けサービス利用のバックグラウンドで実行するオフロード処理を行う際の構成例について説明する。
サービスを提供する際は、初日は試し利用等の形でユーザにサービス提供し、そのバックグラウンドで画像分析等のオフロード処理を行い、翌日以降は画像分析をFPGAにオフロードしてリーズナブルな価格で見守りサービスを提供できるようにすることを想定する。
(First embodiment)
The first embodiment describes loop statement offloading.
An example of the configuration of the
When providing the service, it is expected that the service will be provided to users on a trial basis on the first day, with offloaded processing such as image analysis being performed in the background, and from the next day onwards, image analysis will be offloaded to the FPGA, making it possible to provide a monitoring service at a reasonable price.
図1は、本発明の第1の実施形態に係るオフロードサーバ1の構成例を示す機能ブロック図である。
オフロードサーバ1は、アプリケーションの特定処理をアクセラレータに自動的にオフロードする装置である。
図1に示すように、オフロードサーバ1は、制御部11と、入出力部12と、記憶部13と、検証用マシン14(Verification machine)(アクセラレータ検証用装置)と、を含んで構成される。
FIG. 1 is a functional block diagram showing an example of the configuration of an
The
As shown in FIG. 1, the
入出力部12は、クラウドレイヤ、ネットワークレイヤおよびデバイスレイヤに属する各デバイス等との間で情報の送受信を行うための通信インタフェースと、タッチパネルやキーボード等の入力装置や、モニタ等の出力装置との間で情報の送受信を行うための入出力インタフェースとから構成される。The input/
記憶部13は、ハードディスクやフラッシュメモリ、RAM(Random Access Memory)等により構成される。
この記憶部13には、テストケースDB(Test case database)131が記憶されるとともに、制御部11の各機能を実行させるためのプログラム(オフロードプログラム)や、制御部11の処理に必要な情報(例えば、中間言語ファイル(Intermediate file)132)が一時的に記憶される。
The
This
テストケースDB131は、検証対象ソフトに対応した試験項目のデータを格納する。試験項目のデータは、例えばMySQL等のデータベースシステムの場合、TPC-C等のトランザクション試験のデータである。 Test case DB131 stores data on test items corresponding to the software to be verified. For example, in the case of a database system such as MySQL, the test item data is data on transaction tests such as TPC-C.
制御部11は、オフロードサーバ1全体の制御を司る自動オフロード機能部(Automatic Offloading function)である。制御部11は、例えば、記憶部13に格納されたアプリケーションプログラム(オフロードプログラム)を不図示のCPU(Central Processing Unit)が、RAMに展開し実行することにより実現される。The
アプリケーションプログラムは、C言語、Python、およびJavaより選択される少なくとも一つを含む。 The application program includes at least one selected from C, Python, and Java.
制御部11は、アプリケーションコード指定部(Specify application code)111と、アプリケーションコード分析部(Analyze application code)112と、データ転送指定部113と、並列処理指定部114と、並列処理パターン作成部115と、性能測定部116と、実行ファイル作成部117と、本番環境配置部(Deploy final binary files to production environment)118と、性能測定テスト抽出実行部(Extract performance test cases and run automatically)119と、ユーザ提供部(Provide price and performance to a user to judge)120と、を備える。The
<アプリケーションコード指定部111>
アプリケーションコード指定部111は、入力されたアプリケーションコードの指定を行う。具体的には、アプリケーションコード指定部111は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部112に渡す。
<Application code specification unit 111>
The application code designation unit 111 designates the input application code. Specifically, the application code designation unit 111 passes the application code described in the received file to the application code analysis unit 112.
<アプリケーションコード分析部112>
アプリケーションコード分析部112は、処理機能のソースコードを分析し、ループ文やFFTライブラリ呼び出し等の構造を把握する。
<Application Code Analysis Unit 112>
The application code analysis unit 112 analyzes the source code of the processing function and understands the structure of loop statements, FFT library calls, and the like.
<データ転送指定部113>
データ転送指定部113は、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行う。
<Data
The data
データ転送指定部113は、CPUからGPUへのデータ転送を明示的に指定する明示的指定行と、GPUからCPUへのデータ転送を明示的に指定する明示的指定行と、同じ変数に関してCPUからGPUへの転送とGPUからCPUへの転送とが重なる場合、データコピーの往復をまとめて明示的に指定する明示的指定行と、を用いたデータ転送指定を行う。The data
データ転送指定部113は、CPUプログラム側で定義した変数とGPUプログラム側で参照する変数が重なる場合、CPUからGPUへのデータ転送の指示を行い、データ転送を指定する位置を、GPU処理するループ文かそれより上位のループ文で、該当変数の設定、定義を含まない最上位のループとする。また、データ転送指定部113は、GPUプログラム側で設定した変数とCPUプログラム側で参照する変数とが重なる場合、GPUからCPUへのデータ転送の指示を行い、データ転送を指定する位置を、GPU処理するループ文か、それより上位のループ文で、該当変数の参照、設定、定義を含まない最上位のループとする。When a variable defined on the CPU program side and a variable referenced on the GPU program side overlap, the data
<並列処理指定部114>
並列処理指定部114は、アプリケーションプログラムのループ文(繰り返し文)を特定し、各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする。
並列処理指定部114は、オフロード範囲抽出部(Extract offloadable area)114aと、中間言語ファイル出力部(Output intermediate file)114bと、を備える。
<Parallel
The parallel
The parallel
オフロード範囲抽出部114aは、ループ文やFFT等、GPU・FPGAにオフロード可能な処理を特定し、オフロード処理に応じた中間言語を抽出する。The offload range extraction unit 114a identifies processes that can be offloaded to a GPU or FPGA, such as loop statements and FFTs, and extracts an intermediate language corresponding to the offloaded process.
中間言語ファイル出力部114bは、抽出した中間言語ファイル132を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。The intermediate language
<並列処理パターン作成部115>
並列処理パターン作成部115は、コンパイルエラーが出るループ文(繰り返し文)に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。
<Parallel processing
The parallel processing
<性能測定部116>
性能測定部116は、並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン14に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する。
性能測定部116は、バイナリファイル配置部(Deploy binary files)116aを備える。バイナリファイル配置部116aは、GPU・FPGAを備えた検証用マシン14に、中間言語から導かれる実行ファイルをデプロイ(配置)する。
<
The
The
性能測定部116は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、オフロード範囲抽出部114aに戻す。この場合、オフロード範囲抽出部114aは、別の並列処理パターン抽出を行い、中間言語ファイル出力部114bは、抽出された中間言語をもとに、性能測定を試行する(後記図2の符号a参照)。The
<実行ファイル作成部117>
実行ファイル作成部117は、所定回数繰り返された、性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する。
<Executable
The executable
<本番環境配置部118>
本番環境配置部118は、作成した実行ファイルを、ユーザ向けの本番環境に配置する(「最終バイナリファイルの本番環境への配置」)。本番環境配置部118は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。
<Production
The production
<性能測定テスト抽出実行部119>
性能測定テスト抽出実行部119は、実行ファイル配置後、テストケースDB131から性能試験項目を抽出し、性能試験を実行する(「最終バイナリファイルの本番環境への配置」)。
性能測定テスト抽出実行部119は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースDB131から抽出し、抽出した性能試験を自動実行する。
<Performance Measurement Test
After placing the executable file, the performance measurement test
After arranging the executable file, the performance measurement test
<ユーザ提供部120>
ユーザ提供部120は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する(「価格・性能等の情報のユーザへの提供」)。テストケースDB131には、性能試験項目が格納されている。ユーザ提供部120は、テストケースDB131に格納された試験項目に対応した性能試験の実施結果に基づいて、価格、性能等のデータを、上記性能試験結果と共にユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。ここで、本番環境への一括デプロイには、非特許文献(Y. Yamato, M. Muroi, K. Tanaka and M. Uchimura, “Development of Template Management Technology for Easy Deployment of Virtual Resources on OpenStack,” Journal of Cloud Computing, Springer, 2014, 3:7, DOI: 10.1186/s13677-014-0007-3, 12 pages, June 2014.)の技術を、また、性能自動試験には、非特許文献(Y. Yamato, “Automatic verification technology of software patches for user virtual environments on IaaS cloud,” Journal of Cloud Computing, Springer, 2015, 4:4, DOI: 10.1186/s13677-015-0028-6, 14 pages, Feb. 2015.)の技術を用いればよい。
<
The
[遺伝的アルゴリズムの適用]
オフロードサーバ1は、オフロードの最適化にGAを用いることができる。GAを用いた場合のオフロードサーバ1の構成は下記の通りである。
すなわち、並列処理指定部114は、遺伝的アルゴリズムに基づき、コンパイルエラーが出ないループ文(繰り返し文)の数を遺伝子長とする。並列処理パターン作成部115は、アクセラレータ処理をする場合を1または0のいずれか一方、しない場合を他方の0または1として、アクセラレータ処理可否を遺伝子パターンにマッピングする。
[Application of genetic algorithm]
The
That is, the parallel
並列処理パターン作成部115は、遺伝子の各値を1か0にランダムに作成した指定個体数の遺伝子パターンを準備し、性能測定部116は、各個体に応じて、アクセラレータにおける並列処理指定文を指定したアプリケーションコードをコンパイルして、検証用マシン14に配置する。性能測定部116は、検証用マシン14において性能測定用処理を実行する。The parallel processing
ここで、性能測定部116は、途中世代で、以前と同じ並列処理パターンの遺伝子が生じた場合は、当該並列処理パターンに該当するアプリケーションコードのコンパイル、および、性能測定はせずに、性能測定値としては同じ値を使う。
また、性能測定部116は、コンパイルエラーが生じるアプリケーションコード、および、性能測定が所定時間で終了しないアプリケーションコードについては、タイムアウトの扱いとして、性能測定値を所定の時間(長時間)に設定する。
Here, if a gene with the same parallel processing pattern as before occurs in an intermediate generation, the
Furthermore, for application code that causes a compilation error and application code for which performance measurement does not end within a predetermined time, the
実行ファイル作成部117は、全個体に対して、性能測定を行い、処理時間の短い個体ほど適合度が高くなるように評価する。実行ファイル作成部117は、全個体から、適合度が所定値(例えば、全個数の上位n%、または全個数の上位m個 n,mは自然数)より高いものを性能の高い個体として選択し、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。実行ファイル作成部117は、指定世代数の処理終了後、最高性能の並列処理パターンを解として選択する。The executable
以下、上述のように構成されたオフロードサーバ1の自動オフロード動作について説明する。
[自動オフロード動作]
図2は、オフロードサーバ1のGAを用いた自動オフロード処理を示す図である。
図2に示すように、オフロードサーバ1は、環境適応ソフトウェアの要素技術に適用される。オフロードサーバ1は、制御部(自動オフロード機能部)11と、テストケースDB131と、中間言語ファイル132と、検証用マシン14と、を有している。
オフロードサーバ1は、ユーザが利用するアプリケーションコード(Application code)125を取得する。
The automatic offload operation of the
[Automatic offload operation]
FIG. 2 is a diagram showing an automatic offload process using the GA of the
2, the
The
ユーザは、例えば、各種デバイス(Device151、CPU-GPUを有する装置152、CPU-FPGAを有する装置153、CPUを有する装置154)の利用を契約した人である。
オフロードサーバ1は、機能処理をCPU-GPUを有する装置152、CPU-FPGAを有する装置153のアクセラレータに自動オフロードする。
The user is, for example, a person who has signed a contract to use various devices (
The
以下、図2のステップ番号を参照して各部の動作を説明する。
<ステップS11:Specify application code>
ステップS11において、アプリケーションコード指定部111(図1参照)は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部112に渡す。
The operation of each unit will be described below with reference to the step numbers in FIG.
<Step S11: Specify application code>
In
<ステップS12:Analyze application code>
ステップS12において、アプリケーションコード分析部112(図1参照)は、処理機能のソースコードを分析し、ループ文やFFTライブラリ呼び出し等の構造を把握する。
<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 structures of loop statements, FFT library calls, and the like.
<ステップS13:Extract offloadable area>
ステップS13において、並列処理指定部114(図1参照)は、アプリケーションのループ文(繰り返し文)を特定し、各繰り返し文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする。具体的には、オフロード範囲抽出部114a(図1参照)は、ループ文やFFT等、GPU・FPGAにオフロード可能な処理を特定し、オフロード処理に応じた中間言語を抽出する。
<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 a parallel processing specification statement in the accelerator. Specifically, the offload range extraction unit 114a (see FIG. 1) identifies processes that can be offloaded to the GPU/FPGA, such as loop statements and FFT, and extracts an intermediate language corresponding to the offloaded process.
<ステップS14:Output intermediate file>
ステップS14において、中間言語ファイル出力部114b(図1参照)は、中間言語ファイル132を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。
<Step S14: Output intermediate file>
In step S14, the intermediate language
<ステップS15:Compile error>
ステップS15において、並列処理パターン作成部115(図1参照)は、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。
<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 a compilation error from being offloaded, and specifies whether or not to perform parallel processing on repetitive statements that do not produce a compilation error.
<ステップS21:Deploy binary files>
ステップS21において、バイナリファイル配置部116a(図1参照)は、GPU・FPGAを備えた検証用マシン14に、中間言語から導かれる実行ファイルをデプロイする。
<Step S21: Deploy binary files>
In step S21, the binary
<ステップS22:Measure performances>
ステップS22において、性能測定部116(図1参照)は、配置したファイルを実行し、オフロードした際の性能を測定する。
オフロードする領域をより適切にするため、この性能測定結果は、オフロード範囲抽出部114aに戻され、オフロード範囲抽出部114aが、別パターンの抽出を行う。そして、中間言語ファイル出力部114bは、抽出された中間言語をもとに、性能測定を試行する(図2の符号a参照)。
<Step S22: Measure performance>
In step S22, the performance measurement unit 116 (see FIG. 1) executes the arranged file and measures the performance when offloaded.
In order to more appropriately determine the area to be offloaded, the performance measurement result is returned to the offload range extraction unit 114a, which extracts another pattern. The intermediate language
図2の符号aに示すように、制御部11は、上記ステップS12乃至ステップS22を繰り返し実行する。制御部11の自動オフロード機能をまとめると、下記である。すなわち、並列処理指定部114は、アプリケーションプログラムのループ文(繰り返し文)を特定し、各繰返し文に対して、GPUでの並列処理指定文を指定して、コンパイルする。そして、並列処理パターン作成部115は、コンパイルエラーが出るループ文を、オフロード対象外とし、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。そして、バイナリファイル配置部116aは、該当並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン14に配置し、性能測定部116が、検証用マシン14で性能測定用処理を実行する。実行ファイル作成部117は、所定回数繰り返された、性能測定結果をもとに、複数の並列処理パターンから最高処理性能のパターンを選択し、選択パターンをコンパイルして実行ファイルを作成する。As shown by the symbol a in FIG. 2, the
<ステップS23:Deploy final binary files to production environment>
ステップS23において、本番環境配置部118は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。
<Step S23: Deploy final binary files to production environment>
In step S23, the production
<ステップS24:Extract performance test cases and run automatically>
ステップS24において、性能測定テスト抽出実行部119は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースDB131から抽出し、抽出した性能試験を自動実行する。
<Step S24: Extract performance test cases and run automatically>
In step S24, after arranging the executable file, the performance measurement test
<ステップS25:Provide price and performance to a user to judge>
ステップS25において、ユーザ提供部120は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。
<Step S25: Provide price and performance to a user to judge>
In step S25, the
上記ステップS11~ステップS25は、例えばユーザのサービス利用のバックグラウンドで行われ、例えば、仮利用の初日の間に行う等を想定している。また、コスト低減のためにバックグラウンドで行う処理は、機能配置最適化とGPU・FPGAオフロードのみを対象としてもよい。 The above steps S11 to S25 are assumed to be performed, for example, in the background while the user is using the service, for example, during the first day of trial use. In addition, the processing performed in the background to reduce costs may only target function placement optimization and GPU/FPGA offloading.
上記したように、オフロードサーバ1の制御部(自動オフロード機能部)11は、環境適応ソフトウェアの要素技術に適用した場合、機能処理のオフロードのため、ユーザが利用するアプリケーションプログラムのソースコードから、オフロードする領域を抽出して中間言語を出力する(ステップS11~ステップS15)。制御部11は、中間言語から導かれる実行ファイルを、検証用マシン14に配置実行し、オフロード効果を検証する(ステップS21~ステップS22)。検証を繰り返し、適切なオフロード領域を定めたのち、制御部11は、実際にユーザに提供する本番環境に、実行ファイルをデプロイし、サービスとして提供する(ステップS23~ステップS25)。As described above, when applied to the elemental technology of the environment-adaptive software, the control unit (automatic offload function unit) 11 of the
[GAを用いたGPU自動オフロード]
GPU自動オフロードは、GPUに対して、図2のステップS12~ステップS22を繰り返し、最終的にステップS23でデプロイするオフロードコードを得るための処理である。
[Automatic GPU offloading using GA]
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 deployed in step S23.
GPUは、一般的にレイテンシーは保証しないが、並列処理によりスループットを高めることに向いたデバイスである。環境適応ソフトウェアの暗号化処理や、カメラ映像分析のための画像処理、大量センサデータ分析のための機械学習処理等が代表的であり、それらは、繰り返し処理が多い。そこで、アプリケーションの繰り返し文をGPUに自動でオフロードすることでの高速化を狙う。 GPUs generally do not guarantee latency, but are devices suited to increasing throughput through parallel processing. Typical applications include encryption processing in environmentally adaptive software, image processing for camera image analysis, and machine learning processing for analyzing large amounts of sensor data, which involve a lot of repetitive processing. Therefore, the aim is to increase speed by automatically offloading repetitive statements in applications to the GPU.
しかし、従来技術で記載の通り、高速化には適切な並列処理が必要である。特に、GPUを使う場合は、CPUとGPU間のメモリ転送のため、データサイズやループ回数が多くないと性能が出ないことが多い。また、メモリデータ転送のタイミング等により、並列高速化できる個々のループ文(繰り返し文)の組み合わせが、最速とならない場合等がある。例えば、10個のfor文(繰り返し文)で、1番、5番、10番の3つがCPUに比べて高速化できる場合に、1番、5番、10番の3つの組み合わせが最速になるとは限らない等である。However, as described in the prior art, appropriate parallel processing is necessary to increase speed. In particular, when using a GPU, due to memory transfer between the CPU and GPU, performance is often not achieved unless the data size and number of loops are large. Also, depending on the timing of memory data transfer, the combination of individual loop statements (repeated statements) that can be accelerated in parallel may not be the fastest. For example, if there are 10 "for" statements (repeated statements), and
適切な並列領域指定のため、PGIコンパイラを用いて、for文の並列可否を試行錯誤して最適化する試みがある。しかし、試行錯誤には多くの稼働がかかり、サービスとして提供する際に、ユーザの利用開始が遅くなり、コストも上がってしまう問題がある。 In order to specify appropriate parallel regions, there have been attempts to optimize the parallelism of for statements through trial and error using the PGI compiler. However, trial and error requires a lot of work, and when offered as a service, this can delay users' start-up and increase costs.
そこで、本実施形態では、並列化を想定していない汎用プログラムから、自動で適切なオフロード領域を抽出する。このため、最初に並列可能for文のチェックを行い、次に並列可能for文群に対してGAを用いて検証環境で性能検証試行を反復し適切な領域を探索すること、を実現する。並列可能for文に絞った上で、遺伝子の部分の形で、高速化可能な並列処理パターンを保持し組み換えていくことで、取り得る膨大な並列処理パターンから、効率的に高速化可能なパターンを探索できる。 In this embodiment, therefore, suitable offload areas are automatically extracted from general-purpose programs that do 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 group of parallelizable for statements to search for suitable areas. By narrowing down to parallelizable for statements and then retaining and recombining parallel processing patterns that can be accelerated in the form of genetic parts, it is possible to efficiently search for patterns that can be accelerated from the enormous number of possible parallel processing patterns.
[Simple GAによる制御部(自動オフロード機能部)11の探索イメージ]
図3は、Simple GAによる制御部(自動オフロード機能部)11の処理の探索イメージとfor文の遺伝子配列マッピングを示す図である。
GAは、生物の進化過程を模倣した組合せ最適化手法の一つである。GAのフローチャートは、初期化→評価→選択→交叉→突然変異→終了判定となっている。
本実施形態では、GAの中で、処理を単純にしたSimple GAを用いる。Simple GAは、遺伝子は1、0のみとし、ルーレット選択、一点交叉、突然変異は1箇所の遺伝子の値を逆にする等、単純化されたGAである。
[Search image of the control unit (automatic offload function unit) 11 using Simple GA]
FIG. 3 is a diagram showing a search image of the process of the control unit (automatic offload function unit) 11 by the Simple GA and a 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, a simple GA with simplified processing is used among GAs. The simple GA is a simplified GA in which genes are only 1 and 0, and roulette selection, one-point crossover, and mutation reverse the value of one gene.
<初期化>
初期化では、アプリケーションコードの全for文の並列可否をチェック後、並列可能for文を遺伝子配列にマッピングする。GPU処理する場合は1、GPU処理しない場合は0とする。遺伝子は、指定の個体数Mを準備し、1つのfor文にランダムに1、0の割り当てを行う。
具体的には、制御部(自動オフロード機能部)11(図1参照)は、ユーザが利用するアプリケーションコード(Application code)130(図2参照)を取得し、図3に示すように、アプリケーションコード130のコードパターン(Code patterns)141からfor文の並列可否をチェックする。図3に示すように、コードパターン141から3つのfor文が見つかった場合(図3の符号b参照)、各for文に対して1桁、ここでは3つのfor文に対し3桁の1または0を割り当てる。例えば、CPUで処理する場合0、GPUに出す場合1とする。ただし、この段階では1または0をランダムに割り当てる。
遺伝子長に該当するコードが3桁であり、3桁の遺伝子長のコードは23パターン、例えば100、101、…となる。なお、図3では、コードパターン141中の丸印(○印)をコードのイメージとして示している。
<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, it is set to 1, and if not, it 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 parallelized from the
The code corresponding to the gene length is three digits, and the three-digit gene length code has 2-3 patterns, for example, 100, 101, .... In Fig. 3, the circles (○ marks) in the
<評価>
評価では、デプロイとパフォーマンスの測定(Deploy & performance measurement)を行う(図3の符号c参照)。すなわち、性能測定部116(図1参照)は、遺伝子に該当するコードをコンパイルして検証用マシン14にデプロイして実行する。性能測定部116は、ベンチマーク性能測定を行う。性能が良いパターン(並列処理パターン)の遺伝子の適合度を高くする。
<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 gene, deploys it on the
<選択>
選択では、適合度に基づいて、高性能コードパターンを選択(Select high performance code patterns)する(図3の符号d参照)。性能測定部116(図1参照)は、適合度に基づいて、高適合度の遺伝子を、指定の個体数選択する。本実施形態では、適合度に応じたルーレット選択および最高適合度遺伝子のエリート選択を行う。
図3では、選択されたコードパターン(Select code patterns)142の中の丸印(○印)が、3つに減ったことを探索イメージとして示している。
<Select>
In the selection, high performance 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
<交叉>
交叉では、一定の交叉率Pcで、選択された個体間で一部の遺伝子をある一点で交換し、子の個体を作成する。
ルーレット選択された、あるパターン(並列処理パターン)と他のパターンとの遺伝子を交叉させる。一点交叉の位置は任意であり、例えば上記3桁のコードのうち2桁目で交叉させる。
<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 certain pattern (parallel processing pattern) selected by roulette wheel are crossed with those of another pattern. The position of the one-point crossover is arbitrary, and for example, the crossover is performed at the second digit of the above three-digit code.
<突然変異>
局所解を避けるため、突然変異を導入する。なお、演算量を削減するために突然変異を行わない態様でもよい。突然変異では、一定の突然変異率Pmで、個体の遺伝子の各値を0から1または1から0に変更する。
<Mutation>
In order to avoid local solutions, mutation is introduced. However, in order to reduce the amount of calculation, mutation may not be performed. In mutation, each value of an individual's genes is changed from 0 to 1 or from 1 to 0 at a certain mutation rate Pm.
<終了判定>
図3に示すように、交叉と突然変異後の次世代コードパターンの生成(Generate next generation code patterns after crossover & mutation)を行う(図3の符号e参照)。
終了判定では、指定の世代数T回、繰り返しを行った後に処理を終了し、最高適合度の遺伝子を解とする。
例えば、性能測定して、速い3つ101、010、001を選ぶ。この3つをGAにより、次の世代は、組み換えをして、例えば新しいパターン(並列処理パターン)101(一例)を作っていく。このとき、組み換えをしたパターンに、勝手に0を1にするなどの突然変異を入れる。上記を繰り返して、一番早いパターンを見付ける。指定世代(例えば、20世代)などを決めて、最終世代で残ったパターンを、最後の解とする。
<End Judgment>
As shown in FIG. 3, next generation code patterns after crossover & mutation are generated (see symbol e in FIG. 3).
In the termination determination, 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, 101, 010, and 001, are selected. These three are then recombined in the next generation using GA to create a new pattern (parallel processing pattern) 101 (one example), for example. At this time, a mutation is automatically introduced into 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.
<デプロイ(配置)>
最高適合度の遺伝子に該当する、最高処理性能の並列処理パターンで、本番環境に改めてデプロイして、ユーザに提供する。
<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.
GPUにオフロードできないfor文(ループ文;繰り返し文)が相当数存在する場合について説明する。例えば、for文が200個あっても、GPUにオフロードできるものは30個くらいである。ここでは、エラーになるものを除外し、この30個について、GAを行う。 We will explain the case where there are a considerable 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 run GA on these 30 statements.
例えば、C/C++コードに対するGPUの処理を行う仕様であるOpenACCには、ディレクティブ #pragma acc kernelsで指定して、GPU向けバイトコードを抽出し、実行によりGPUオフロードを可能とするコンパイラがある。Python, Javaの場合は、CUDAやJava Lambda式などでGPU処理を指定すればよい。For example, OpenACC, a specification for GPU processing of C/C++ code, has a compiler that allows you to specify the directive #pragma acc kernels to extract bytecode for the GPU and execute it to enable GPU offloading. In the case of Python and Java, you can specify GPU processing using CUDA or Java Lambda expressions.
また、C/C++を使った場合、C/C++のコードを分析し、for文を見付ける。for文を見付けると、OpenACCで並列処理の文法である #pragma acc kernelsを使ってfor文に対して書き込む。詳細には、何も入っていない #pragma acc kernels に、一つ一つfor文を入れてコンパイルして、エラーであれば、そのfor文はそもそも、GPU処理できないので、除外する。このようにして、残るfor文を見付ける。そして、エラーが出ないものを、長さ(遺伝子長)とする。エラーのないfor文が5つであれば、遺伝子長は5であり、エラーのないfor文が10であれば、遺伝子長は10である。なお、並列処理できないものは、前の処理を次の処理に使うようなデータに依存がある場合である。
以上が準備段階である。次にGA処理を行う。
Also, when C/C++ is used, the C/C++ code is analyzed to find for statements. When a for statement is found, it is written to the for statement using #pragma acc kernels, which is a parallel processing syntax in OpenACC. In detail, for statements are inserted one by one into an empty #pragma acc kernels and compiled. If an error occurs, the for statement is excluded since it cannot be processed by the GPU in the first place. In this way, the remaining for statements are found. The one that does not produce an error is then taken as the length (gene length). If there are five for statements without errors, the gene length is 5, and if there are 10 for statements without errors, the gene length is 10. Note that parallel processing is not possible when there is a data dependency such that the previous processing is used for the next processing.
The above is the preparation stage. Next, the GA process is carried out.
for文の数に対応する遺伝子長を有するコードパターンが得られている。始めはランダムに並列処理パターン100、010、001、…を割り当てる。そして、GA処理を行い、コンパイルする。その時に、オフロードできるfor文であるにもかかわらず、エラーがでることがある。for文が階層になっている(どちらか指定すればGPU処理できる)場合である。この場合は、エラーとなったfor文は、残してもよい。具体的には、処理時間が多くなった形にして、タイムアウトさせる方法がある。
A code pattern with a gene length corresponding to the number of for statements is obtained. First,
検証用マシン14でデプロイして、ベンチマーク、例えば画像処理であればその画像処理でベンチマークする、その処理時間が短い程、適応度が高いと評価する。例えば、処理時間の逆数、処理時間10秒かかるものは1、100秒かかるものは0.1、1秒のものは10とする。
適応度が高いものを選択して、例えば10個のなかから、3~5個を選択して、それを組み替えて新しいコードパターンを作る。作成途中で、前と同じものができる場合がある。この場合、同じベンチマークを行う必要はないので、前と同じデータを使う。本実施形態では、コードパターンと、その処理時間は記憶部13に保存しておく。
Deploy it on the
Those with high adaptability are selected, for example, 3 to 5 out of 10 are selected, and a new code pattern is created by rearranging them. During the creation process, it is possible that the same code pattern as before is created. In this case, since it is not necessary 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
以上で、Simple GAによる制御部(自動オフロード機能部)11の探索イメージについて説明した。次に、データ転送の一括処理手法について述べる。Above, we have explained the search image of the control unit (automatic offload function unit) 11 using Simple GA. Next, we will explain the batch processing method for data transfer.
[データ転送の一括処理手法]
上述したように、遺伝的アルゴリズムを用いることで、GPU処理で効果のある並列処理部を自動チューニングしている。しかしながら、CPU-GPUメモリ間のデータ転送によっては高性能化できないアプリケーションもあった。このため、スキルが無いユーザがGPUを使ってアプリケーションを高性能化することは難しいし、自動並列化技術等を使う場合も並列処理可否の試行錯誤が必要であり、高速化できない場合があった。
[Batch processing method for data transfer]
As mentioned above, the genetic algorithm is used to automatically tune the parallel processing units that are effective in GPU processing. However, there were some applications for which the performance could not be improved by data transfer between the CPU and GPU memory. For this reason, it is difficult for unskilled users to improve the performance of applications using GPUs, and even when using automatic parallelization technology, trial and error is required to determine whether parallel processing is possible, and there were cases where speed could not be improved.
そこで、本実施形態では、より多くのアプリケーションを、自動でGPUを用いて高性能化することを狙うとともに、GPUへのデータ転送回数を低減できる技術を提供する。 Therefore, in this embodiment, we aim to automatically use the GPU to improve the performance of more applications, while also providing technology that can reduce the number of data transfers to the GPU.
<基本的な考え方>
OpenACC等の仕様では、GPUでの並列処理を指定する指示行に加えて、CPUからGPUへのデータ転送やその逆を明示的に指定する指示行(以下、「明示的指示行」という)が定義されている。OpenACC等の明示的指示行は、CPUからGPUへのデータ転送のディレクティブ(directive:行頭に特殊な記号を記述した指示・指定コマンド)である「#pragma acc data copyin」、GPUからCPUへのデータ転送のディレクティブである「#pragma acc data copyout」、CPUからGPUへ再びCPUへのデータ転送のディレクティブである「#pragma acc data copy」等である。
<Basic Concept>
In specifications such as OpenACC, in addition to directives that specify parallel processing in the GPU, directives that explicitly specify data transfer from the CPU to the GPU and vice versa (hereinafter referred to as "explicit directives") are defined. Explicit directives such as OpenACC include "#pragma acc data copyin", which is a directive for data transfer from the CPU to the GPU (directive: an instruction/designation command with a special symbol written at the beginning of a line), "#pragma acc data copyout", which is a directive for data transfer from the GPU to the CPU, and "#pragma acc data copy", which is a directive for data transfer from the CPU to the GPU and back to the CPU.
本実施形態は、非効率なデータ転送を低減するため、明示的指示行を用いたデータ転送指定を、GAでの並列処理の抽出と合わせて行う。
本実施形態では、GAで生成された各個体について、ループ文の中で利用される変数データの参照関係を分析し、ループ毎に毎回データ転送するのではなくループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する。
In this embodiment, in order to reduce inefficient data transfer, data transfer is specified using an explicit directive line in combination with extraction of parallel processing in GA.
In this embodiment, for each individual generated by GA, the reference relationship of variable data used in the loop statement is analyzed, and for data that may be transferred outside the loop rather than every time a loop is performed, data transfer outside the loop is explicitly specified.
<具体例>
以下、具体的に処理を説明する。
データ転送の種類は、CPUからGPUへのデータ転送、および、GPUからCPUへのデータ転送がある。
<Example>
The process will be specifically described below.
The types of data transfer include data transfer from the CPU to the GPU and data transfer from the GPU to the CPU.
図4は、具体例の自動オフロード機能部が処理するアプリケーションプログラムのソースコードのループ文(繰り返し文)を示す図であり、CPUプログラム側で定義した変数とGPUプログラム側で参照する変数が重なる場合の例である。
具体例の自動オフロード機能部は、図1の制御部(自動オフロード機能部)11からデータ転送指定部113を取り去る、またはデータ転送指定部113を実行しない場合の例である。
FIG. 4 is a diagram showing a loop statement (repetitive statement) in the source code of an application program processed by a specific example of an automatic offload function unit, and is an example of a case where a variable defined on the CPU program side and a variable referenced on the GPU program side overlap.
The specific example of the automatic offload function unit is an example in which the data
具体例のCPUからGPUへのデータ転送を例に採る。
図4は、具体例のCPUからGPUへデータ転送する場合のループ文において、CPUプログラム側で定義した変数とGPUプログラム側で参照する変数が重なる場合の例である。なお、以下の記載および図4中のループ文の文頭の<1>~<4> は、説明の便宜上で付したものである(他図およびその説明においても同様)。
図4に示す具体例のループ文は、CPUプログラム側で記述され、
<1> ループ〔 for|do|while 〕 {
}
の中に、
<2> ループ〔 for|do|while 〕 {
}
があり、さらにその中に、
<3> ループ〔 for|do|while 〕 {
}
があり、さらにその中に、
<4> ループ〔 for〕{
}
がある。
Take data transfer from a CPU to a GPU as an example.
4 shows an example of a loop statement for transferring data from a CPU to a GPU, in which variables defined in the CPU program overlap with variables referenced in the GPU program. Note that the numbers <1> to <4> at the beginning of the loop statements in the following description and in FIG. 4 are added for the sake of convenience (the same applies to other figures and their explanations).
The loop statement of the specific example shown in FIG. 4 is written on the CPU program side,
<1> Loop [for | do | while] {
}
Among them,
<2> Loop [for | do | while] {
}
And within that,
<3> Loop [for | do | while] {
}
And within that,
<4> Loop〔for〕{
}
There is.
また、<1> ループ〔 for|do|while 〕 {
}で、変数aが設定され、<4> ループ〔 for|do|while 〕 {
}で、変数aが参照される。
Also, <1> Loop [for | do | while] {
}, the variable a is set, and the <4> loop [for | do | while] {
The variable a is referenced in }.
さらに、<3> ループ〔 for|do|while 〕 {
}で、PGIコンパイラによるfor文等の並列処理可能処理部を、OpenACCのディレクティブ #pragma acc kernels(並列処理指定文)で指定している(詳細後記)。
Furthermore, <3> Loop [for | do | while] {
}, parallel processing parts such as for statements by the PGI compiler are specified by the OpenACC directive #pragma acc kernels (parallel processing specification statement) (details will be described later).
図4に示す比較例のループ文では、図4の符号fに示すタイミングで毎回CPUからGPUにデータ転送する。このため、GPUへのデータ転送回数を低減することが求められる。
なお、GPUからCPUへのデータ転送も同様であり説明を省略する。
In the loop statement of the comparative example shown in Fig. 4, data is transferred from the CPU to the GPU every time at the timing indicated by the symbol f in Fig. 4. For this reason, it is desired to reduce the number of times data is transferred to the GPU.
Data transfer from the GPU to the CPU is similar, and a description thereof will be omitted.
以上述べたように、本実施形態では、できるだけ上位のループでデータ転送を一括して行うように、データ転送を明示的に指示することで、ループ毎に毎回データを転送する非効率な転送を避けることができる。As described above, in this embodiment, by explicitly instructing data transfer to be performed in bulk in as high-level a loop as possible, it is possible to avoid the inefficient transfer of data for each loop.
[GPUオフロード処理]
上述したデータ転送の一括処理手法により、オフロードに適切なループ文を抽出し、非効率なデータ転送を避けることができる。
ただし、上記データ転送の一括処理手法を用いても、GPUオフロードに向いていないプログラムも存在する。効果的なGPUオフロードには、オフロードする処理のループ回数が多いことが必要である。
[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-mentioned 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 process to be offloaded is large.
そこで、本実施形態では、本格的なオフロード処理探索の前段階として、プロファイリングツールを用いて、ループ回数を調査する。プロファイリングツールを用いると、各行の実行回数を調査できるため、例えば、5000万回以上のループを持つプログラムをオフロード処理探索の対象とする等、事前に振り分けることができる。以下、具体的に説明する(前記図3で述べた内容と一部重複する)。 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. Using a profiling tool, the number of executions of each line can be investigated, so that programs with loops of 50 million or more can be targeted for offload processing search in advance. A detailed explanation is given below (some of the content overlaps with that described in Figure 3 above).
本実施形態では、まず、制御部(自動オフロード機能部)11(図1参照)が、アプリケーションプログラムを分析し、for,do,while等のループ文を把握する。次に、サンプル処理を実行し、プロファイリングツールを用いて、各ループ文のループ回数を調査し、一定の値以上のループがあるか否かで、探索を本格的に行うか否かの判定を行う。In this embodiment, first, the control unit (automatic offload function unit) 11 (see FIG. 1) analyzes the application program and identifies loop statements such as for, do, and while. Next, a sample process is executed, and a profiling tool is used to check the number of loops in each loop statement, and a decision is made as to whether or not to perform a full search based on whether or not there are loops that exceed a certain value.
探索を本格的に行うと決まった場合は、GAの処理に入る(前記図3参照)。初期化ステップでは、アプリケーションコードの全ループ文の並列可否をチェックした後、並列可能ループ文をGPU処理する場合は1、しない場合は0として遺伝子配列にマッピングする。遺伝子は、指定の個体数が準備されるが、遺伝子の各値にはランダムに1,0の割り当てをする。 When it is decided to conduct a full-scale search, the GA process begins (see Figure 3 above). 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 are to be processed by the GPU, and 0 if they are not. A specified number of genes are prepared, and each gene value is randomly assigned a value of 1 or 0.
ここで、遺伝子に該当するコードでは、GPU処理すると指定されたループ文内の変数データ参照関係から、データ転送の明示的指示(OpenACCで指定するならば#pragma acc data copyin/copyout/copy)を追加する。Here, in the code corresponding to the gene, an explicit instruction for data transfer (#pragma acc data copyin/copyout/copy if specified with OpenACC) is added based on the variable data reference relationship within the loop statement specified for GPU processing.
評価ステップでは、遺伝子に該当するコードをコンパイルして検証用マシンにデプロイして実行し、ベンチマーク性能測定を行う。性能が良いパターンの遺伝子の適合度を高くする。遺伝子に該当するコードは、上述のように、並列処理指示行(例えば、図4の符号f参照)が挿入されている。In the evaluation step, the code corresponding to the genes is compiled, deployed and executed on a verification machine, and benchmark performance is measured. The fitness of genes with good performance patterns is increased. As described above, parallel processing instruction lines (for example, see symbol f in Figure 4) are inserted into the code corresponding to the genes.
選択ステップでは、適合度に基づいて、高適合度の遺伝子を、指定の個体数選択する。本実施形態では、適合度に応じたルーレット選択および最高適合度遺伝子のエリート選択を行う。交叉ステップでは、一定の交叉率Pcで、選択された個体間で一部の遺伝子をある一点で交換し、子の個体を作成する。突然変異ステップでは、一定の突然変異率Pmで、個体の遺伝子の各値を0から1または1から0に変更する。 In the selection step, a specified number of highly fit genes are selected based on fitness. In this embodiment, roulette selection and elite selection of the most fit genes are performed according to fitness. In the crossover step, some genes are exchanged at a certain point between the selected individuals at a constant crossover rate Pc to create offspring individuals. In the mutation step, the value of each gene of an individual is changed from 0 to 1 or from 1 to 0 at a constant mutation rate Pm.
突然変異ステップまで終わり、次の世代の遺伝子が指定個体数作成されると、初期化ステップと同様に、データ転送の明示的指示を追加し、評価、選択、交叉、突然変異ステップを繰り返す。 Once the mutation step is completed and the specified number of genes for the next generation have been created, explicit instructions for data transfer are added, as in the initialization step, and the evaluation, selection, crossover, and mutation steps are repeated.
最後に、終了判定ステップでは、指定の世代数、繰り返しを行った後に処理を終了し、最高適合度の遺伝子を解とする。最高適合度の遺伝子に該当する、最高性能のコードパターンで、本番環境に改めてデプロイして、ユーザに提供する。 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 best-performing code pattern that corresponds to the gene with the highest fitness is then redeployed to the production environment and provided to the user.
以下、オフロードサーバ1の実装を説明する。本実装は、本実施形態の有効性を確認するためのものである。
[実装]
C/C++アプリケーションを汎用のPGIコンパイラを用いて自動オフロードする実装を説明する。
本実装では、GPU自動オフロードの有効性確認が目的であるため、対象アプリケーションはC/C++言語のアプリケーションとし、GPU処理自体は、従来のPGIコンパイラを説明に用いる。
The following describes the implementation of the
[implementation]
This paper describes an implementation of automatic offloading of C/C++ applications using a general-purpose PGI compiler.
In this implementation, since the purpose is to confirm the effectiveness of GPU automatic offloading, the target application is an application written in C/C++ language, and the GPU processing itself is explained using a conventional PGI compiler.
C/C++言語は、OSS(Open Source Software)およびproprietaryソフトウェアの開発で、上位の人気を誇り、数多くのアプリケーションがC/C++言語で開発されている。一般ユーザが用いるアプリケーションプログラムのオフロードを確認するため、暗号処理や画像処理等のOSSの汎用アプリケーションを利用する。 The C/C++ language is one of the most popular languages for developing OSS (Open Source Software) and proprietary software, and many applications have been developed in C/C++. To verify the offloading of application programs used by general users, we use general-purpose OSS applications such as encryption and image processing.
GPU処理は、PGIコンパイラにより行う。PGIコンパイラは、OpenACCを解釈するC/C++/Fortran向けコンパイラである。本実施形態では、for文等の並列可能処理部を、OpenACCのディレクティブ #pragma acc kernels(並列処理指定文)で指定する。これにより、GPU向けバイトコードを抽出し、その実行によりGPUオフロードを可能としている。さらに、for文内のデータ同士に依存性があり並列処理できない処理やネストのfor文の異なる複数の階層を指定されている場合等の際に、エラーを出す。合わせて、#pragma acc data copyin/copyout/copy 等のディレクティブにより、明示的なデータ転送の指示が可能とする。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 the GPU 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 that makes parallel processing impossible, or when multiple levels of nested for statements with different levels are specified. Additionally, explicit data transfer instructions are possible using directives such as #pragma acc data copyin/copyout/copy.
上記 #pragma acc kernels(並列処理指定文)での指定に合わせて、OpenACCのcopyin 節の #pragma acc data copyout(a[…])の、上述した位置への挿入により、明示的なデータ転送の指示を行う。 In accordance with the specification in the above #pragma acc kernels (parallel processing specification statement), an explicit data transfer is instructed by inserting #pragma acc data copyout(a[…]) in the OpenACC copyin clause at the position mentioned above.
<実装の動作概要>
実装の動作概要を説明する。
高速化するC/C++アプリケーションプログラムとそれを性能測定するベンチマークツールを準備する。
<Implementation Overview>
The outline of the implementation will be explained.
Prepare the C/C++ application program to be accelerated and a benchmark tool to measure its performance.
実装では、C/C++アプリケーションプログラムの利用依頼があると、まず、C/C++アプリケーションのコードを解析して、for文を発見するとともに、for文内で使われる変数データ等の、プログラム構造を把握する。構文解析には、LLVM/Clangの構文解析ライブラリ(libClangのpython binding)等を使用する。 In the implementation, when a request is made to use a C/C++ application program, the system first analyzes the code of the C/C++ application to find for statements and understand the program structure, such as the variable data used within the for statements. For syntax analysis, the LLVM/Clang syntax analysis library (libClang python binding) etc. is used.
実装では、最初に、そのアプリケーションがGPUオフロード効果があるかの見込みを得るため、ベンチマークを実行し、上記構文解析で把握したfor文のループ回数を把握する。ループ回数把握には、GNUカバレッジのgcov等を用いる。プロファイリングツールとしては、「GNUプロファイラ(gprof)」、「GNUカバレッジ(gcov)」が知られている。双方とも各行の実行回数を調査できるため、どちらを用いてもよい。実行回数は、例えば、1000万回以上のループ回数を持つアプリケーションプログラムのみ対象とするようにできるが、この値は変更可能である。 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 application programs with a loop count of 10 million or more, for example, but this value can be changed.
CPU向け汎用アプリケーションプログラムは、並列化を想定して実装されているわけではない。そのため、まず、GPU処理自体が不可なfor文は排除する必要がある。そこで、各for文一つずつに対して、並列処理の#pragma acc kernels ディレクティブ挿入を試行し、コンパイル時にエラーが出るかの判定を行う。コンパイルエラーに関しては、幾つかの種類がある。for文の中で外部ルーチンが呼ばれている場合、ネストfor文で異なる階層が重複指定されている場合、break等でfor文を途中で抜ける処理がある場合、for文のデータにデータ依存性がある場合等がある。アプリケーションプログラムによって、コンパイル時エラーの種類は多彩であり、これ以外の場合もあるが、コンパイルエラーは処理対象外とし、#pragmaディレクティブは挿入しない。 General-purpose application programs 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, the #pragma acc kernels directive for parallel processing is inserted into each for statement one by one to 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 hierarchies are specified in a nested for statement, when there is a process that exits the for statement midway with break, when there is data dependency in the data in the for statement, etc. Compilation errors vary depending on the application program, and there may be other errors as well, but compilation errors are not processed and the #pragma directive is not inserted.
コンパイルエラーは自動対処が難しく、また対処しても効果が出ないことも多い。外部ルーチンコールの場合は、#pragma acc routineにより回避できる場合があるが、多くの外部コールはライブラリであり、それを含めてGPU処理してもそのコールがネックとなり性能が出ない。for文一つずつを試行するため、ネストのエラーに関しては、コンパイルエラーは生じない。また、break等で途中で抜ける場合は、並列処理にはループ回数を固定化する必要があり、プログラム改造が必要となる。データ依存が有る場合はそもそも並列処理自体ができない。 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 does not improve. Since each for statement is tried, no compilation errors will 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 will be necessary. If there is data dependency, parallel processing itself is not possible in the first place.
ここで、並列処理してもエラーが出ないループ文の数がaの場合、aが遺伝子長となる。遺伝子の1は並列処理ディレクティブ有、0は無に対応させ、長さaの遺伝子に、アプリケーションコードをマッピングする。 Here, if the number of loop statements that can be processed in parallel without causing an error is a, then a becomes the gene length. A gene of 1 corresponds to the presence of a parallel processing directive, and a corresponds to the absence of a parallel processing directive, and application code is mapped to a gene of length a.
次に、初期値として、指定個体数の遺伝子配列を準備する。遺伝子の各値は、図3で説明したように、0と1をランダムに割当てて作成する。準備された遺伝子配列に応じて、遺伝子の値が1の場合は並列処理を指定するディレクティブ #pragma acc kernels をC/C++コードに挿入する。この段階で、ある遺伝子に該当するコードの中で、GPUで処理させる部分が決まる。上記Clangで解析した、for文内の変数データの参照関係をもとに、上述したルールに基づいて、CPUからGPUへのデータ転送、その逆の場合のディレクティブ指定を行う。
具体的には、CPUからGPUへのデータ転送が必要な変数は、 #pragma acc data copyinで指定し(図示省略)、GPUからCPUへのデータ転送が必要な変数は、 #pragma acc data copyoutで指定する(図示省略)。同じ変数に関して、copyinとcopyoutが重なる場合は、#pragma acc data copyで纏め、記述をシンプルにする。
Next, a gene array for the specified number of individuals is prepared as an initial value. As explained in Figure 3, each gene value is created by randomly assigning 0 and 1. According to the prepared gene array, if the gene value is 1, the directive #pragma acc kernels is inserted into the C/C++ code to specify parallel processing. At this stage, the part of the code corresponding to a certain gene that is to be processed by the GPU is decided. Based on the reference relationship of the variable data in the for statement analyzed by Clang above, directives are specified for data transfer from the CPU to the GPU and vice versa, according to the rules mentioned above.
Specifically, variables that require data transfer from the CPU to the GPU are specified with #pragma acc data copyin (not shown), and variables that require data transfer from the GPU to the CPU are specified with #pragma acc data copyout (not shown). If copyin and copyout overlap for the same variable, they are combined with #pragma acc data copy to simplify the description.
並列処理およびデータ転送のディレクティブを挿入されたC/C++コードを、GPUを備えたマシン上のPGIコンパイラでコンパイルを行う。コンパイルした実行ファイルをデプロイし、ベンチマークツールで性能を測定する。 The C/C++ code with parallel processing and data transfer directives inserted is compiled with a PGI compiler on a machine with a GPU. The compiled executable is deployed and its performance is measured with a benchmark tool.
全個体数に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のGA処理を行い、次世代の個体群を作成する。 After measuring the benchmark performance for all individuals, the fitness of each gene sequence is set according to the benchmark processing time. Individuals to be kept 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.
次世代の個体に対して、ディレクティブ挿入、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。ここで、GA処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイル、性能測定をせず、以前と同じ測定値を用いる。 Directive insertion, compilation, performance measurement, fitness setting, selection, crossover, and mutation processes are performed on the next generation of individuals. If a gene with the same pattern as before is generated during the GA process, compilation and performance measurement are not performed on that individual, and the same measurement value as before is used.
指定世代数のGA処理終了後、最高性能の遺伝子配列に該当する、ディレクティブ付きC/C++コードを解とする。 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 determined as the solution.
この中で、個体数、世代数、交叉率、突然変異率、適合度設定、選択方法は、GAのパラメータであり、別途指定する。提案技術は、上記処理を自動化することで、従来、専門技術者の時間とスキルが必要だった、GPUオフロードの自動化を可能にする。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 specialized engineer.
[《ループ文オフロード:共通》フローチャート]
図5A-Bは、《ループ文オフロード:共通》フローチャートであり、図5Aと図5Bは、結合子で繋がれる。
[Flowchart for "Loop Statement Offloading: Common"]
5A and 5B are flowcharts for <<Loop Statement Offload: Common>>, and FIGS. 5A and 5B are connected with a connector.
<コード解析>
ステップS101で、アプリケーションコード分析部112(図1参照)は、アプリケーションプログラムのコード解析を行う。
<Code Analysis>
In step S101, the application code analysis unit 112 (see FIG. 1) performs a code analysis of the application program.
<ループ文特定>
ステップS102で、並列処理指定部114(図1参照)は、アプリケーションプログラムのループ文、参照関係を特定する。
<Loop statement identification>
In step S102, the parallel processing specification unit 114 (see FIG. 1) identifies loop statements and reference relationships in the application program.
<ループ文ループ回数>
ステップS103で、並列処理指定部114は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。
<Loop statement loop count>
In step S103, the parallel
<ループ文の並列処理可能性>
ステップS104で、並列処理指定部114は、各ループ文の並列処理可能性をチェックする。
<Parallel processing of loop statements>
In step S104, the parallel
<ループ文の繰り返し>
制御部(自動オフロード機能部)11は、ステップS105のループ始端とステップS108のループ終端間で、ステップS106-S107の処理についてループ文の数だけ繰り返す。
ステップS106で、並列処理指定部114は、各ループ文に対して、言語に応じた手法でGPU処理を指定してコンパイルまたはインタプリットする。
ステップS107で、並列処理指定部114は、エラー時は、該当for文からは、GPU処理指定を削除する。
ステップS109で、並列処理指定部114は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。
<Repeat loop statement>
The control unit (automatic offload function unit) 11 repeats the processing 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
In step S107, if an error occurs, the parallel
In step S109, the parallel
<指定個体数パターン準備>
次に、初期値として、並列処理指定部114は、指定個体数の遺伝子配列を準備する。ここでは、0と1をランダムに割当てて作成する。
ステップS110で、並列処理指定部114は、アプリケーションプログラムのコードを、遺伝子にマッピングする。0と1がランダムに割当てられた遺伝子配列を遺伝子にマッピングすることで、指定個体数パターンを準備する。
準備された遺伝子配列に応じて、遺伝子の値が1の場合は並列処理を指定するディレクティブをアプリケーションプログラムのコードに挿入する(例えば図3の#pragmaディレクティブ参照)。
<Preparation of designated population patterns>
Next, the parallel
In step S110, the parallel
According to the prepared gene sequence, a directive that specifies parallel processing when the gene value is 1 is inserted into the code of the application program (see, for example, the #pragma directive in FIG. 3).
制御部(自動オフロード機能部)11は、ステップS111のループ始端とステップS120のループ終端間で、ステップS112-S119の処理について指定世代数繰り返す。
また、上記指定世代数繰り返しにおいて、さらにステップS112のループ始端とステップS117のループ終端間で、ステップS113-S116の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。
The control unit (automatic offload function unit) 11 repeats the processing of steps S112-S119 for a designated number of generations between the loop start point of step S111 and the loop end point of step S120.
In addition, in the above-mentioned repetition of the specified number of generations, the processing of steps S113-S116 is further repeated a specified number of individuals between the loop start point of step S112 and the loop end point of step S117. 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.
<データ転送指定>
ステップS113で、データ転送指定部113は、変数参照関係から、言語に応じた手法でデータ転送を指定する。
<Data transfer specification>
In step S113, the data
<コンパイル>
ステップS114で、並列処理パターン作成部115(図1参照)は、遺伝子パターンに応じてGPU処理基盤でコンパイルまたはインタプリットする。すなわち、並列処理パターン作成部115は、作成したアプリケーションプログラムのコードを、GPUを備えた検証用マシン14上のPGIコンパイラでコンパイルまたはインタプリットを行う。
ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。
<Compile>
In step S114, the parallel processing pattern creation unit 115 (see FIG. 1) compiles or interprets the code of the created application program on the GPU processing platform in accordance with the gene pattern. That is, the parallel processing
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.
ステップS115で、性能測定部116(図1参照)は、CPU-GPU搭載の検証用マシン14に、実行ファイルをデプロイする。
ステップS116で、性能測定部116は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。
In step S115, the performance measurement unit 116 (see FIG. 1) deploys the executable file to the
In step S116, the
ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、GA処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
ステップS118で、実行ファイル作成部117(図1参照)は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。
Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same values are used. In other words, if a gene with the same pattern as before is generated during the GA process, compilation or performance measurement is not performed for that individual, and the same measured value as before is used.
In step S118, the executable file creation unit 117 (see FIG. 1) evaluates the individuals with shorter processing times so that they have a higher fitness, and selects the individuals with the highest performance.
ステップS119で、実行ファイル作成部117は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のGA処理を行い、次世代の個体群を作成する。
In step S119, the executable
That is, after benchmark performance is measured 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 selected individuals are subjected to GA processing including crossover, mutation, and direct copy to create the next generation population.
ステップS121で、実行ファイル作成部117は、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するC/C++コード(最高性能の並列処理パターン)を解とする。In step S121, after the GA processing for the specified number of generations is completed, the executable
<GAのパラメータ>
上記、個体数、世代数、交叉率、突然変異率、適合度設定、選択方法は、GAのパラメータである。GAのパラメータは、例えば、以下のように設定してもよい。
実行するSimple GAの、パラメータ、条件は例えば以下のようにできる。
遺伝子長:並列可能ループ文数
個体数M:遺伝子長以下
世代数T:遺伝子長以下
適合度:(処理時間)-1/2
この設定により、ベンチマーク処理時間が短い程、高適合度になる。また、適合度を、(処理時間)-1/2とすることで、処理時間が短い特定の個体の適合度が高くなり過ぎて、探索範囲が狭くなるのを防ぐことができる。また、性能測定が一定時間で終わらない場合は、タイムアウトさせ、処理時間1000秒等の時間(長時間)であるとして、適合度を計算する。このタイムアウト時間は、性能測定特性に応じて変更させればよい。
選択:ルーレット選択
ただし、世代での最高適合度遺伝子は交叉も突然変異もせず次世代に保存するエリート保存も合わせて行う。
交叉率Pc:0.9
突然変異率Pm:0.05
<GA parameters>
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 loop statements that can be parallelized Number of individuals M: Less than or equal to gene length Number of generations T: Less than or equal to gene length Fitness: (processing time) -1/2
With this setting, the shorter the benchmark processing time, the higher the fitness. Also, by setting the fitness to (processing time) -1/2 , it is possible to prevent the fitness of a specific individual with a short processing time from becoming too high, 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 crossover or mutation.
Crossover rate Pc: 0.9
Mutation rate Pm: 0.05
本実施形態では、gcov,gprof等を用いて、ループが多く実行時間がかかっているアプリケーションを事前に特定して、オフロード試行をする。これにより、効率的に高速化できるアプリケーションを見つけることができる。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 accelerated efficiently.
より短時間でオフロード部分を探索するためには、複数の検証用マシンで個体数分並列で性能測定することが考えられる。アプリケーションプログラムに応じて、タイムアウト時間を調整することも短時間化に繋がる。例えば、オフロード処理がCPUでの実行時間の2倍かかる場合はタイムアウトとする等である。また、個体数、世代数が多い方が、高性能な解を発見できる可能性が高まる。しかし、各パラメータを最大にする場合、個体数×世代数だけコンパイル、および性能ベンチマークを行う必要がある。このため、本番サービス利用開始までの時間がかかる。本実施形態では、GAとしては少ない個体数、世代数で行っているが、交叉率Pcを0.9と高い値にして広範囲を探索することで、ある程度の性能の解を早く発見するようにしている。 In order to search for the offloaded part in a shorter time, it is possible to measure the performance in parallel on multiple verification machines for the number of individuals. Adjusting the timeout time according to the application program can also shorten the time. For example, if the offload processing 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 chance 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.
以上、《ループ文オフロード:共通》について説明した、次に、《ループ文オフロード:C言語》について説明する。 We have explained "Loop Statement Offload: Common" above. Next, we will explain "Loop Statement Offload: C Language".
[《ループ文オフロード:C言語》]
ループ文オフロード:C言語について、基本的フローは、上記《ループ文オフロード:共通》と同様であり、言語非依存にできる。C言語に依存・非依存の処理について詳細に述べる。
ループ文オフロードの、コードの分析では、C言語を解析するClang等の構文解析ツールを用いて構文解析する。ループと変数の把握については、構文解析ツールの結果を管理する際は、言語に非依存に抽象的に管理できる。ループのGPU処理有無の遺伝子化についても、言語に非依存である。遺伝子情報のコード化では、遺伝子情報に合わせてGPUで実行するためのコードを作成するため、C言語の拡張文法であるOpenACCでGPU処理を指定したり、変数転送を指定したりする。
[<<Loop Statement Offload: C Language>>]
Loop statement offloading: Regarding C language, the basic flow is the same as the above <<Loop statement offloading: common>>, and it can be made language independent. We will now explain in detail the processing that is dependent and independent of the C language.
When analyzing code for offloading loop statements, syntax analysis is performed using a syntax analysis tool such as Clang, which analyzes the C language. When managing the results of the syntax analysis tool to understand loops and variables, they can be managed abstractly and independent of the language. Genetization of whether or not to use GPU processing for loops is also language independent. When encoding genetic information, GPU processing and variable transfer are specified using OpenACC, an extended grammar for the C language, to create code to be executed on the GPU according to the genetic information.
コンパイルは、OpenACCコードをPGIコンパイラ等でコンパイルする。性能測定は、言語に合わせて、Jenkins、Selenium等の自動測定ツールも用いて行う。次世代の遺伝子作成は、性能測定結果に合わせて適合度を設定し交叉等の処理を行うが、言語に非依存である。反復実行と最終解の決定も、言語に非依存である。
以上のように、ループ文オフロードでは、処理に関しては、ループと変数の管理とGAの遺伝子処理については言語に非依存に適用できる。
Compilation is performed by compiling OpenACC code with a PGI compiler or similar. Performance measurement is performed using automated measurement tools such as Jenkins and Selenium depending on the language. Creation of next-generation genes involves setting the fitness level according to the performance measurement results and performing processes such as crossover, but is language-independent. Iterative execution and determination of the final solution are also language-independent.
As described above, in loop statement offloading, the management of loops and variables and the genetic processing of GA can be applied independently of the language.
図6A-Bは、《ループ文オフロード:C言語》のフローチャートであり、図6Aと図6Bは、結合子で繋がれる。
C/C++向けOpenACCコンパイラを用いて以下の処理を行う。
6A and 6B are flowcharts of <<Loop Statement Offload: C Language>>, and FIGS. 6A and 6B are connected with a connector.
The following process is performed using the OpenACC compiler for C/C++.
<コード解析>
ステップS201で、アプリケーションコード分析部112(図1参照)は、C/C++アプリケーションプログラムのコード解析を行う。
<Code Analysis>
In step S201, the application code analysis unit 112 (see FIG. 1) performs code analysis of a C/C++ application program.
<ループ文特定>
ステップS202で、並列処理指定部114(図1参照)は、C/C++アプリケーションプログラムのループ文、参照関係を特定する。
<Loop statement identification>
In step S202, the parallel processing specification unit 114 (see FIG. 1) identifies loop statements and reference relationships in the C/C++ application program.
<ループ文ループ回数>
ステップS203で、並列処理指定部114は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。
<Loop statement loop count>
In step S203, the parallel
<ループ文の並列処理可能性>
ステップS204で、並列処理指定部114は、各ループ文の並列処理可能性をチェックする。
<Parallel processing of loop statements>
In step S204, the parallel
<ループ文の繰り返し>
制御部(自動オフロード機能部)11は、ステップS205のループ始端とステップS208のループ終端間で、ステップS206-S207の処理についてループ文の数だけ繰り返す。
ステップS206で、並列処理指定部114は、各ループ文に対して、OpenACC文法を用いて、#pragma acc kernelsでGPU処理を指定してコンパイルする。
ステップS207で、並列処理指定部114は、エラー時は、該当for文からは、#pragma acc kernelsを削除する。
ステップS209で、並列処理指定部114は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。
<Repeat loop statement>
The control unit (automatic offload function unit) 11 repeats the processing of steps S206-S207 between the loop start point of step S205 and the loop end point of step S208 for the number of loop statements.
In step S206, the parallel
In step S207, if an error occurs, the parallel
In step S209, the parallel
<指定個体数パターン準備>
次に、初期値として、並列処理指定部114は、指定個体数の遺伝子配列を準備する。ここでは、0と1をランダムに割当てて作成する。
ステップS210で、並列処理指定部114は、C/C++アプリコードを、遺伝子にマッピングする。0と1がランダムに割当てられた指定個体数の遺伝子配列を遺伝子にマッピングするすることで、指定個体数パターンを準備する。
準備された遺伝子配列に応じて、遺伝子の値が1の場合は並列処理を指定するディレクティブをC/C++アプリコードに挿入する(例えば図3の#pragmaディレクティブ参照)。
<Preparation of designated population patterns>
Next, the parallel
In step S210, the parallel
According to the prepared gene sequence, a directive that specifies parallel processing when the gene value is 1 is inserted into the C/C++ application code (see, for example, the #pragma directive in Figure 3).
制御部(自動オフロード機能部)11は、ステップS211のループ始端とステップS220のループ終端間で、ステップS212-S219の処理について指定世代数繰り返す。
また、上記指定世代数繰り返しにおいて、さらにステップS212のループ始端とステップS217のループ終端間で、ステップS213-S216の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。
The control unit (automatic offload function unit) 11 repeats the processing of steps S212-S219 for a designated number of generations between the loop start point of step S211 and the loop end point of step S220.
In addition, in the above-mentioned repetition of the specified number of generations, the processing of steps S213-S216 is further repeated a specified number of individuals between the loop start point of step S212 and the loop end point of step S217. In other words, within the repetition of the specified number of generations, the repetition of the specified number of individuals is processed in a nested state.
<データ転送指定>
ステップS213で、データ転送指定部113は、変数参照関係から、明示的指示行(#pragma acc data copy/in/out)を用いたデータ転送指定を行う。明示的指示行(#pragma acc data copy/in/out)を用いたデータ転送指定については、図4により説明した。
<Data transfer specification>
In step S213, the data
<コンパイル>
ステップS214で、並列処理パターン作成部115(図1参照)は、遺伝子パターンに応じてディレクティブ指定したC/C++コードをPGIコンパイラでコンパイルする。すなわち、並列処理パターン作成部115は、作成したC/C++アプリコードを、GPUを備えた検証用マシン14上のPGIコンパイラでコンパイルする。
ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。
<Compile>
In step S214, 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
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.
ステップS215で、性能測定部116(図1参照)は、CPU-GPU搭載の検証用マシン14に、実行ファイルをデプロイする。
ステップS216で、性能測定部116は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。
In step S215, the performance measurement unit 116 (see FIG. 1) deploys the executable file to the
In step S216, the
ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、GA処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
ステップS218で、実行ファイル作成部117(図1参照)は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。
Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same values are used. In other words, if a gene with the same pattern as before is generated during the GA process, compilation or performance measurement is not performed for that individual, and the same measured value as before is used.
In step S218, the executable file creation unit 117 (see FIG. 1) evaluates the individuals with shorter processing times so that they have a higher fitness, and selects the individuals with the highest performance.
ステップS219で、実行ファイル作成部117は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のGA処理を行い、次世代の個体群を作成する。
In step S219, the executable
That is, after benchmark performance is measured 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 selected individuals are subjected to GA processing including crossover, mutation, and direct copy to create the next generation population.
ステップS221で、実行ファイル作成部117は、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するC/C++コード(最高性能の並列処理パターン)を解とする。In step S221, after the GA processing for the specified number of generations is completed, the executable
以上、《ループ文オフロード:C言語》について説明した、次に、《ループ文オフロード:Python》について説明する。 We have explained "Loop statement offloading: C language" above. Next, we will explain "Loop statement offloading: Python".
[《ループ文オフロード:Python》]
ループ文オフロード:Pythonは、PythonコードをpyCUDAでインタプリットする方法(図7A-B参照)と、OpenACCを解釈するインタプリタpyACCを用いる方法(図8A-B、図9A-B参照)とがある。以下、順に説明する。
[Loop statement offloading: Python]
Loop statement offload: Python has a method of interpreting Python code with pyCUDA (see Figures 7A-B) and a method of using the interpreter pyACC that interprets OpenACC (see Figures 8A-B and 9A-B). The following describes each method in order.
<PythonコードをpyCUDAでインタプリットする方法>
ループ文オフロードの、コードの分析では、Pythonを解析するast等の構文解析ツールを用いて構文解析する。ループと変数の把握については、構文解析ツールの結果を管理する際は、言語に非依存で、抽象的に管理できる。
ループのGPU処理有無の遺伝子化についても、言語に非依存である。遺伝子情報のコード化では、遺伝子情報に合わせてGPUで実行するためのコードを作成するため、CUDA文法でGPU処理を指定したり、変数転送を指定したりする。
<How to interpret Python code with pyCUDA>
When offloading loop statements, code analysis is performed using a syntax analysis tool such as ast, which analyzes Python. When managing the results of the syntax analysis tool, loops and variables can be understood abstractly and independently of the language.
The coding of the loop with or without GPU processing is also language independent. In coding the genetic information, the GPU processing and variable transfer are specified in the CUDA grammar in order to create code for execution on the GPU according to the genetic information.
インタプリタは、CUDAでの指示を追加したPythonコードをpyCUDAでインタプリットする。性能測定は、言語に合わせて、Jenkins等の自動測定ツールも用いて行う。次世代の遺伝子作成は、性能測定結果に合わせて適合度を設定し交叉等の処理を行うが、言語に非依存である。反復実行と最終解の決定も、言語に非依存である。ここで、pyCUDAでなくpyACCというOpenACCを解釈するインタプリタを用いてもよい。その場合は、C言語と同様にOpenACC文法でループ文のGPU処理を指定すればよい(後記)。 The interpreter uses pyCUDA to interpret Python code with CUDA instructions added. Performance measurements are performed using an automated measurement tool such as Jenkins depending on the language. Creation of the next generation of genes involves setting the fitness level according to the performance measurement results and performing processes such as crossover, but is language independent. Iterative execution and determination of the final solution are also language independent. Here, an interpreter that interprets OpenACC, called pyACC, can be used instead of pyCUDA. In that case, GPU processing of the loop statement can be specified using OpenACC grammar, just like with C language (see below).
Python では、GPU処理指定したコードは実装によるが、CupyというオープンソースでNVIDIA GPUを利用するライブラリを用いた実装を説明する。
動作としては、PythonのコードでGPU処理を指定されるループ文はCupyライブラリを介して、NVIDIAのCUDAコマンドが実行され、NVIDIAのGPUで処理がされる。
In Python, the code that specifies GPU processing depends on the implementation, but we will explain an implementation using Cupy, an open source library that utilizes NVIDIA GPUs.
In operation, loop statements that specify GPU processing in Python code are executed by NVIDIA's CUDA command via the Cupy library, and processing is carried out on the NVIDIA GPU.
本実施形態では、C言語同様、Pythonのループ文に対して、GPU処理可否を遺伝的アルゴリズムにより選択し、適切なオフロードパターンを見つける。 In this embodiment, similar to the C language, for Python loop statements, a genetic algorithm is used to select whether GPU processing is possible and an appropriate offload pattern is found.
以下、Cupyを用いた例を記載する。
図7A-Bは、PythonコードをpyCUDAでインタプリットする方法を説明する図であり、図7Aは、変換元例を示し、図7Bは、変換後例(3階層のfor文の一番上をGPU処理指定する場合)を示す。
図7Aに示すように、Pythonのfor文は行列演算として指定される。Cupyは、CUDAコマンドを呼び、CUDAがGPUを実行する。
Below is an example using Cupy.
7A and 7B are diagrams explaining how to interpret Python code with pyCUDA. FIG. 7A shows an example of the original code, and FIG. 7B shows an example of the converted code (when the topmost part of a three-level for statement is specified for GPU processing).
As shown in Fig. 7A, a Python for statement is specified as a matrix operation. Cupy calls a CUDA command, and CUDA executes the GPU.
Cupyを用いる場合、C言語のOpenACCの\pragmaのようにオフロードするfor文を記載する形ではない。
CupyからCUDAを介してGPU処理する際は、GPUでの並列演算は行列演算である。GPU処理する箇所は、図7Bに示すように、
〔1~多重 for文の内側にある『配列[添字]』を右辺・左辺に持つ演算式〕は、
"[添字1][添字2]…" の部分を"[範囲開始1:範囲終了1,範囲開始2:範囲終了2, …]" 表現に書き換える。
《添字》が《範囲》に置き換わることによって、式全体が行列演算の式になる。
When using Cupy, it is not necessary to write a for statement to offload like the \pragma of OpenACC in C.
When processing on the GPU from Cupy via CUDA, the parallel calculation on the GPU is a matrix calculation. The part to be processed by the GPU is as shown in FIG. 7B.
[1~An arithmetic expression with an "array [subscript]" on the right and left sides inside a multiple for statement] is
Change the part "[index 1][index 2]…" to "[range start 1:range end 1, range start 2:range end 2, …]".
By replacing the "subscript" with a "range", the entire formula becomes a matrix operation formula.
以上、<PythonコードをpyCUDAでインタプリットする方法>について説明した。次に、<pyACCを用いる方法>について説明する。Above we have explained how to interpret Python code with pyCUDA. Next we will explain how to use pyACC.
<pyACCを用いる方法>
ループ文オフロード:Pythonは、上記pyCUDAでなくpyACCというOpenACCを解釈するインタプリタを用いてもよい。その場合は、C言語と同様にOpenACC文法でループ文のGPU処理を指定する。以下、pyACCを用いる方法について説明する。
<Method using pyACC>
Loop statement offload: Python may use an interpreter that interprets OpenACC called pyACC instead of the above pyCUDA. In that case, GPU processing of loop statements is specified using OpenACC syntax, just like C language. The method of using pyACC is explained below.
図8A-Bは、pyACC利用時のコードパターンを示す図であり、図8Aは、pyACC利用時のfor文を示し、図8Bは、図8Aのfor文から作成されるコードパターンを示す。
図8Bに示すコードパターンは、図3のコードパターンに置き換えて用いられる。
8A and 8B are diagrams showing code patterns when pyACC is used, where FIG. 8A shows a for statement when pyACC is used, and FIG. 8B shows a code pattern created from the for statement in FIG. 8A.
The code pattern shown in FIG. 8B is used to replace the code pattern in FIG.
図9A-Bは、《ループ文オフロード:Python》のフローチャートであり、図9Aと図9Bは、結合子で繋がれる。また、図8Bに示すコードパターンを、図3のコードパターンに置き換えて用いる。
C/C++向けOpenACCコンパイラを用いて以下の処理を行う。
9A and 9B are flowcharts of "Loop Statement Offload: Python", and Fig. 9A and Fig. 9B are connected with a connector. Also, the code pattern shown in Fig. 8B is used by replacing the code pattern in Fig. 3.
The following process is performed using the OpenACC compiler for C/C++.
<コード解析>
ステップS301で、アプリケーションコード分析部112(図1参照)は、Pythonアプリケーションプログラムのコード解析を行う。
<Code Analysis>
In step S301, the application code analysis unit 112 (see FIG. 1) performs code analysis of the Python application program.
<ループ文特定>
ステップS302で、並列処理指定部114(図1参照)は、Pythonアプリケーションプログラムのループ文、参照関係を特定する。
<Loop statement identification>
In step S302, the parallel processing specification unit 114 (see FIG. 1) identifies loop statements and reference relationships in the Python application program.
<ループ文ループ回数>
ステップS303で、並列処理指定部114は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。
<Loop statement loop count>
In step S303, the parallel
<ループ文の並列処理可能性>
ステップS304で、並列処理指定部114は、各ループ文の並列処理可能性をチェックする。
<Parallel processing of loop statements>
In step S304, the parallel
<ループ文の繰り返し>
制御部(自動オフロード機能部)11は、ステップS305のループ始端とステップS108のループ終端間で、ステップS306-S307の処理についてループ文の数だけ繰り返す。
ステップS306で、並列処理指定部114は、各ループ文に対して、GPU処理基盤に応じた手法でGPU処理を指定してインタプリットする。例えば、pyACC利用時はOpenACCの\pragmaacckernels、Cupy利用時は対象ループの計算を行列計算に変換して指定、pyCUDA直接利用時はCUDA文法等を用いる、が挙げられる。
<Repeat loop statement>
The control unit (automatic offload function unit) 11 repeats the processing of steps S306-S307 between the loop start point of step S305 and the loop end point of step S108 for the number of loop statements.
In step S306, the parallel
ステップS307で、並列処理指定部114は、エラー時は、該当for文からは、GPU処理を削除する。
ステップS309で、並列処理指定部114は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。
In step S307, if an error occurs, the parallel
In step S309, the parallel
<指定個体数パターン準備>
次に、初期値として、並列処理指定部114は、指定個体数の遺伝子配列を準備する。ここでは、0と1をランダムに割当てて作成する。
ステップS310で、並列処理指定部114は、Pythonアプリコードを、遺伝子にマッピングし、指定個体数パターン準備を行う。
準備された遺伝子配列に応じて、遺伝子の値が1の場合は並列処理を指定するディレクティブをPythonアプリコードに挿入する。
<Preparation of designated population patterns>
Next, the parallel
In step S310, the parallel
According to the prepared gene sequence, a directive specifying parallel processing when the gene value is 1 is inserted into the Python application code.
制御部(自動オフロード機能部)11は、ステップS311のループ始端とステップS320のループ終端間で、ステップS312-S319の処理について指定世代数繰り返す。
また、上記指定世代数繰り返しにおいて、さらにステップS312のループ始端とステップS317のループ終端間で、ステップS313-S316の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。
The control unit (automatic offload function unit) 11 repeats the processing of steps S312-S319 for a designated number of generations between the loop start point of step S311 and the loop end point of step S320.
In addition, in the above-mentioned repetition of the specified number of generations, the processing of steps S313-S316 is further repeated a specified number of individuals between the loop start point of step S312 and the loop end point of step S317. In other words, within the repetition of the specified number of generations, the repetition of the specified number of individuals is processed in a nested state.
<データ転送指定>
ステップS313で、データ転送指定部113は、変数参照関係から、GPU処理基盤に応じた手法でデータ転送を指定する。
<Data transfer specification>
In step S313, the data
<コンパイル>
ステップS314で、並列処理パターン作成部115(図1参照)は、遺伝子パターンに応じてディレクティブ指定したPythonアプリコードをGPU処理基盤でインタプリットする。
ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。
<Compile>
In step S314, the parallel processing pattern creation unit 115 (see FIG. 1) interprets the Python application code, which has been directive-specified according to the gene pattern, on the GPU processing platform.
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.
ステップS315で、性能測定部116(図1参照)は、CPU-GPU搭載の検証用マシン14に、実行ファイルをデプロイする。
ステップS316で、性能測定部116は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。
In step S315, the performance measurement unit 116 (see FIG. 1) deploys the executable file to the
In step S316, the
ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、GA処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
ステップS318で、実行ファイル作成部117(図1参照)は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。
Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same values are used. In other words, if a gene with the same pattern as before is generated during the GA process, compilation or performance measurement is not performed for that individual, and the same measured value as before is used.
In step S318, the executable file creation unit 117 (see FIG. 1) evaluates the individuals with shorter processing times so that they have a higher fitness, and selects the individuals with the highest performance.
ステップS319で、実行ファイル作成部117は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のGA処理を行い、次世代の個体群を作成する。
In step S319, the executable
That is, after benchmark performance is measured 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 selected individuals are subjected to GA processing including crossover, mutation, and direct copy to create the next generation population.
ステップS321で、実行ファイル作成部117は、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するPythonアプリコード(最高性能の並列処理パターン)を解とする。
以上、《ループ文オフロード:Python》について説明した、次に、《ループ文オフロード:Java》について説明する。
In step S321, the executable
Above we have explained "Loop Statement Offloading: Python". Next we will explain "Loop Statement Offloading: Java".
[《ループ文オフロード:Java》]
ループ文オフロード:Java では、ループ文オフロードの、コードの分析でJavaを解析するJavaParser等の構文解析ツールを用いて構文解析する。ループと変数の把握については、構文解析ツールの結果を管理する際は、言語に非依存に抽象的に管理できる。ループのGPU処理有無の遺伝子化についても、言語に非依存である。遺伝子情報のコード化では、遺伝子情報に合わせてGPUで実行するためのコードを作成するため、Javaのラムダ記述でGPU処理を指定する、あるいは変数転送を指定する。
[Loop statement offloading: Java]
Loop statement offloading: In Java, loop statement offloading is performed using a syntax analysis tool such as JavaParser, which analyzes Java code. When managing the results of the syntax analysis tool, understanding of loops and variables can be managed abstractly and independent of the language. Genetization of loops to be processed by GPU or not is also language independent. When encoding genetic information, GPU processing or variable transfer is specified in Java lambda notation to create code to be executed on the GPU according to the genetic information.
実行環境は、Javaのラムダ記述での並列化をGPUに対して行うことができるIBM JDK(登録商標)を用いる。IBM JDKはJavaのラムダ記述に従って並列処理をGPUに対して実行する仮想マシンである。
性能測定は、言語に合わせて、Jenkins(登録商標)等の自動測定ツールも用いて行う。次世代の遺伝子作成は、性能測定結果に合わせて適合度を設定し交叉等の処理を行うが、言語に非依存である。反復実行と最終解の決定も、言語に非依存である。
The execution environment uses IBM JDK (registered trademark), which can perform parallelization on the GPU using Java lambda descriptions. IBM JDK is a virtual machine that executes parallel processing on the GPU according to Java lambda descriptions.
Performance measurement is performed using an automatic measurement tool such as Jenkins (registered trademark) according to the language. The next generation of genes is created by setting the fitness according to the performance measurement results and performing processes such as crossover, but is independent of the language. Iterative execution and the determination of the final solution are also independent of the language.
図10A-Bは、IBM JDK 利用時のコードパターンを示す図であり、図10Aは、IBM JDK 利用時のfor文を示し、図10Bは、図10Aのfor文から作成されるコードパターンを示す。
図10Bに示すコードパターンは、図3のコードパターンに置き換えて用いられる。
10A and 10B are diagrams showing code patterns when the IBM JDK is used, where FIG. 10A shows a for statement when the IBM JDK is used, and FIG. 10B shows a code pattern created from the for statement in FIG. 10A.
The code pattern shown in FIG. 10B is used to replace the code pattern in FIG.
図11A-Bは、《ループ文オフロード:Java》のフローチャートであり、図11Aと図11Bは、結合子で繋がれる。また、図10Bに示すコードパターンを、図3のコードパターンに置き換えて用いる。 Figures 11A-B are flowcharts for "Loop Statement Offload: Java", and Figures 11A and 11B are connected with a connector. The code pattern shown in Figure 10B is used by replacing the code pattern in Figure 3.
<コード解析>
ステップS401で、アプリケーションコード分析部112(図1参照)は、Javaアプリケーションプログラムのコード解析を行う。
<Code Analysis>
In step S401, the application code analysis unit 112 (see FIG. 1) performs code analysis of the Java application program.
<ループ文特定>
ステップS402で、並列処理指定部114(図1参照)は、Javaアプリケーションプログラムのループ文、参照関係を特定する。
<Loop statement identification>
In step S402, the parallel processing specification unit 114 (see FIG. 1) identifies loop statements and reference relationships in the Java application program.
<ループ文ループ回数>
ステップS403で、並列処理指定部114は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。
<Loop statement loop count>
In step S403, the parallel
<ループ文の並列処理可能性>
ステップS404で、並列処理指定部114は、各ループ文の並列処理可能性をチェックする。
<Parallel processing of loop statements>
In step S404, the parallel
<ループ文の繰り返し>
制御部(自動オフロード機能部)11は、ステップS405のループ始端とステップS408のループ終端間で、ステップS406-S407の処理についてループ文の数だけ繰り返す。
ステップS406で、並列処理指定部114は、各ループ文に対して、Javaのlambda式を用いて、java.util.Stream.IntStream.range(0,n).parallel()forEach(i -> {});でGPU処理を指定してコンパイルする。
<Repeat loop statement>
The control unit (automatic offload function unit) 11 repeats the processing of steps S406-S407 between the loop start point of step S405 and the loop end point of step S408 for the number of loop statements.
In step S406, the parallel
ステップS407で、並列処理指定部114は、エラー時は、該当for文からは、java.util.Stream.IntStream.range(0,n).parallel()forEach(i -> {});を削除する。
ステップS409で、並列処理指定部114は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。
In step S407, when an error occurs, the parallel
In step S409, the parallel
<指定個体数パターン準備>
次に、初期値として、並列処理指定部114は、指定個体数の遺伝子配列を準備する。ここでは、0と1をランダムに割当てて作成する。
ステップS410で、並列処理指定部114は、Javaアプリコードを、遺伝子にマッピングする。0と1がランダムに割当てられた指定個体数の遺伝子配列を遺伝子にマッピングするすることで、指定個体数パターンを準備する。
準備された遺伝子配列に応じて、遺伝子の値が1の場合は並列処理を指定するディレクティブをJavaアプリコードに挿入する。
<Preparation of designated population patterns>
Next, the parallel
In step S410, the parallel
According to the prepared gene sequence, a directive specifying parallel processing when the gene value is 1 is inserted into the Java application code.
制御部(自動オフロード機能部)11は、ステップS411のループ始端とステップS420のループ終端間で、ステップS412-S419の処理について指定世代数繰り返す。
また、上記指定世代数繰り返しにおいて、さらにステップS412のループ始端とステップS417のループ終端間で、ステップS413-S416の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。
The control unit (automatic offload function unit) 11 repeats the processing of steps S412-S419 for a designated number of generations between the loop start point of step S411 and the loop end point of step S420.
In addition, in the above-mentioned repetition of the specified number of generations, the processing of steps S413-S416 is further repeated a specified number of individuals between the loop start point of step S412 and the loop end point of step S417. In other words, within the repetition of the specified number of generations, the repetition of the specified number of individuals is processed in a nested state.
<データ転送指定>
ステップS413で、データ転送指定部113は、変数参照関係から、Javaの記述でデータ転送を指定する。
<Data transfer specification>
In step S413, the data
<コンパイル>
ステップS414で、並列処理パターン作成部115(図1参照)は、遺伝子パターンに応じてディレクティブ指定したJavaアプリコードをIBM JDKでビルドする。
ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。
<Compile>
In step S414, the parallel processing pattern creation unit 115 (see FIG. 1) builds the Java application code with directives specified according to the gene pattern using IBM JDK.
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.
ステップS415で、性能測定部116(図1参照)は、CPU-GPU搭載の検証用マシン14に、実行ファイルをデプロイする。
ステップS416で、性能測定部116は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。
In step S415, the performance measurement unit 116 (see FIG. 1) deploys the executable file to the
In step S416, the
ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、GA処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
ステップS418で、実行ファイル作成部117(図1参照)は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。
Here, in the intermediate generation, genes with the same pattern as before are not measured, and the same values are used. In other words, if a gene with the same pattern as before is generated during the GA process, compilation or performance measurement is not performed for that individual, and the same measured value as before is used.
In step S418, the executable file creation unit 117 (see FIG. 1) evaluates the individuals with shorter processing times so that they have a higher fitness, and selects the individuals with the highest performance.
ステップS419で、実行ファイル作成部117は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のGA処理を行い、次世代の個体群を作成する。
In step S419, the executable
That is, after benchmark performance is measured 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 selected individuals are subjected to GA processing including crossover, mutation, and direct copy to create the next generation population.
ステップS421で、実行ファイル作成部117は、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するJavaアプリコード(最高性能の並列処理パターン)を解とする。
以上、第1の実施形態(「ループ文オフロード」)について説明した。
In step S421, the executable
The first embodiment ("loop statement offload") has been described above.
(第2の実施形態)
第2の実施形態は、機能ブロックオフロードについて記載する。
図12~図13を参照して、機能ブロックオフロードの全体構成および動作を説明し、以下、機能ブロックオフロード:共通(図14,図15)、機能ブロックオフロード:C言語(図16,図17)、機能ブロックオフロード:Python(図18,図19)、機能ブロックオフロード:Java(図20,図21)を順に説明する。
Second Embodiment
The second embodiment describes function block offloading.
The overall configuration and operation of function block offload will be described with reference to Figures 12 and 13, and then the following will be described in order: function block offload: common (Figures 14 and 15), function block offload: C language (Figures 16 and 17), function block offload: Python (Figures 18 and 19), and function block offload: Java (Figures 20 and 21).
図13は、本発明の第2の実施形態に係るオフロードサーバ200の構成例を示す機能ブロック図である。図2と同一構成部分には、同一符号を付して重複箇所の説明を省略する。
オフロードサーバ200は、アプリケーションプログラムの特定処理をアクセラレータに自動的にオフロードする装置である。
13 is a functional block diagram showing a configuration example of an
The
アプリケーションプログラムは、C言語、Python、およびJavaより選択される少なくとも一つを含む。 The application program includes at least one selected from C, Python, and Java.
図12に示すように、オフロードサーバ200は、制御部210と、入出力部12と、記憶部130と、検証用マシン14 (アクセラレータ検証用装置)と、を含んで構成される。As shown in FIG. 12, the
入出力部12は、各機器等との間で情報の送受信を行うための通信インタフェースと、タッチパネルやキーボード等の入力装置や、モニタ等の出力装置との間で情報の送受信を行うための入出力インタフェースとから構成される。The input/
記憶部130は、ハードディスクやフラッシュメモリ、RAM(Random Access Memory)等により構成され、制御部210の各機能を実行させるためのプログラム(オフロードプログラム)や、制御部210の処理に必要な情報(例えば、中間言語ファイル(Intermediate file)132)が一時的に記憶される。The
記憶部13は、コードパターンDB(Code pattern database)230(後記)、テストケースDB(Test case database)131を備える。
The
テストケースDB131には、性能試験項目が格納される。テストケースDB131は、高速化するアプリケーションの性能を測定するような試験を行うための情報が格納される。例えば、画像分析処理の深層学習アプリケーションであれば、サンプルの画像とそれを実行する試験項目である。 Test case DB131 stores performance test items. Test case DB131 stores information for conducting tests to measure the performance of applications that are being accelerated. For example, in the case of a deep learning application for image analysis processing, the test items are sample images and the execution of the images.
検証用マシン14は、環境適応ソフトウェアの検証用環境として、CPU(Central Processing Unit)、GPU、FPGAを備える。The
<コードパターンDB230>
・GPUライブラリ、IPコアの記憶
コードパターンDB230は、GPUやFPGA等にオフロード可能なライブラリおよびIPコア(後記)を記憶する。すなわち、コードパターンDB230は、後記<処理B-1>のために、特定のライブラリ、機能ブロックを高速化するGPU用ライブラリ(GPUライブラリ)やFPGA用IPコア(IPコア)とそれに関連する情報を保持する。例えば、コードパターンDB230は、FFT等算術計算等のライブラリリスト(外部ライブラリリスト)を保持する。
<Code Pattern DB230>
Storage of GPU libraries and IP cores The
・CUDAライブラリの記憶
コードパターンDB230は、GPUライブラリとして、例えばCUDAライブラリと当該CUDAライブラリを利用するためのライブラリ利用手順とを記憶する。すなわち、後記<処理C-1>において、置換するライブラリやIPコアをGPUやFPGAに実装し、ホスト側(CPU)プログラムと繋ぐ場合、ライブラリ利用手順も含めて登録しておき、その手順に従って利用する。例えば、CUDAライブラリでは、C言語コードからCUDAライブラリを利用する手順がライブラリとともに公開されているため、コードパターンDB230にライブラリ利用手順も含めて登録しておく。
Storage of CUDA library The
・クラス、構造体の記憶
コードパターンDB230は、ホストで計算する場合に記述が同様になる処理のクラスまたは構造体を記憶する。すなわち、後記<処理B-2>において、登録されていないライブラリ呼び出し以外の機能処理を検出するため、構文解析にてソースコードの定義記述からクラス、構造体等を検出する。コードパターンDB230は、後記<処理B-2>のために、ホストで計算する場合に記述が同様になる処理のクラスまたは構造体を登録しておく。なお、クラスまたは構造体の機能処理に対して、高速化するライブラリやIPコアがあることは、類似性検出ツール(後記)で検出する。
Storage of classes and structures The
・OpenCLコードの記憶
コードパターンDB230は、IPコア関連の情報としてOpenCLコードを記憶する。コードパターンDB230に、OpenCLコードを記憶しておくことで、OpenCLコードから、OpenCLインタフェースを用いたCPUとFPGAの接続および、FPGAへのIPコア実装が、XilinxやIntel等のFPGAベンダの高位合成ツール(後記)を介して行うことができる。
Storage of OpenCL code The
<制御部210>
制御部210は、オフロードサーバ200全体の制御を司る自動オフロード機能部(Automatic Offloading function)であり、記憶部130に格納されたプログラム(オフロードプログラム)を不図示のCPUが、RAMに展開し実行することにより実現される。
<
The
特に、制御部210は、CPU向けの既存プログラムコードの中にFPGAやGPUへオフロードすることで処理を高速化できる機能ブロックを検出し、検出した機能ブロックをGPU向けライブラリやFPGA向けIPコア等に置き換えることで高速化をする機能ブロックのオフロード処理を行う。In particular, the
制御部210は、アプリケーションコード指定部(Specify application code)111と、アプリケーションコード分析部(Analyze application code)112と、置換機能検出部213と、置換処理部214と、オフロードパターン作成部215と、性能測定部116と、実行ファイル作成部117と、本番環境配置部(Deploy final binary files to production environment)118と、性能測定テスト抽出実行部(Extract performance test cases and run automatically)119と、ユーザ提供部(Provide price and performance to a user to judge)120と、を備える。The
<アプリケーションコード指定部111>
アプリケーションコード指定部111は、入力されたアプリケーションコードの指定を行う。具体的には、アプリケーションコード指定部111は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部112に渡す。
<Application code specification unit 111>
The application code designation unit 111 designates the input application code. Specifically, the application code designation unit 111 passes the application code described in the received file to the application code analysis unit 112.
<アプリケーションコード分析部112>
アプリケーションコード分析部112は、後記<処理A-1>において、アプリケーションプログラムのソースコードを分析して、当該ソースコードに含まれる外部ライブラリの呼び出しを検出する。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
<Application Code Analysis Unit 112>
In <Process A-1> described later, the application code analysis unit 112 analyzes the source code of the application program and detects calls to external libraries contained in the source code. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code to analyze library calls and functional processing contained in the code, along with loop statement structures, etc.
上述したコード分析は、オフロードするデバイスを想定した分析が必要になるため、一般化は難しい。ただし、ループ文や変数の参照関係等のコードの構造を把握したり、機能ブロックとしてFFT処理を行う機能ブロックであることや、FFT処理を行うライブラリを呼び出している等を把握することは可能である。機能ブロックの判断は、オフロードサーバが自動判断することは難しい。これもDeckard等の類似性検出ツールを用いて類似度判定等で把握することは可能である。ここで、Clangは、C/C++向けツールであるが、解析する言語に合わせたツールを選ぶ必要がある。 The code analysis described above is difficult to generalize, as it requires analysis that takes into account the device to be offloaded. However, it is possible to grasp the code structure, such as loop statements and variable reference relationships, and to determine that a functional block performs FFT processing and that it calls a library that performs FFT processing. It is difficult for the offload server to automatically determine the functional block. This can also be grasped by using a similarity detection tool such as Deckard to determine the similarity. Here, Clang is a tool for C/C++, but it is necessary to choose a tool that suits the language to be analyzed.
また、アプリケーションコード分析部112は、後記<処理A-2>において、ソースコードからクラスまたは構造体のコードを検出する。 In addition, in <Process A-2> described below, the application code analysis unit 112 detects class or structure code from the source code.
<置換機能検出部213>
置換機能検出部213は、後記<処理B-1>において、検出された呼び出しをキーにして、コードパターンDB230からGPUライブラリおよびIPコアを取得する。具体的には、置換機能検出部213は、検出したライブラリ呼び出しに対して、ライブラリ名をキーとして、コードパターンDB230と照合することで、GPU、FPGAにオフロードできるオフロード可能処理を抽出する。
<Replacement function detection unit 213>
In <Process B-1> described later, the replacement function detection unit 213 uses the detected call as a key to acquire a GPU library and an IP core from the
ここで、コードパターンDB230は、GPUライブラリとして、例えばCUDAライブラリと当該CUDAライブラリを利用するためのライブラリ利用手順とを記憶している。そして、置換機能検出部213は、ライブラリ利用手順をもとに、コードパターンDB230からCUDAライブラリを取得する。Here, the
置換機能検出部213は、後記<処理B-2>において、検出されたクラスまたは構造体(後記)の定義記述コードをキーにして、コードパターンDB230からGPUライブラリおよびIPコアを取得する。具体的には、置換機能検出部213は、コピーコードやコピー後変更した定義記述コードを検出する類似性検出ツールを用いて、置換元コードに含まれるクラスや構造体に対して、コードパターンDB230から類似のクラスまたは構造体に紐づいて管理されているGPU、FPGAにオフロードできるGPUライブラリおよびIPコアを抽出する。In <Process B-2> described below, the replacement function detection unit 213 uses the definition description code of the detected class or structure (described below) as a key to obtain a GPU library and an IP core from the
<置換処理部214>
置換処理部214は、後記<処理C-1>において、アプリケーションプログラムのソースコードの置換元の処理記述を、置換機能検出部213が取得した置換先のライブラリおよびIPコアの処理記述に置換する。具体的には、置換処理部214は、抽出したオフロード可能処理を、GPU向けのライブラリやFPGA向けのIPコア等に置換する。
また、置換処理部214は、置換したライブラリおよびIPコアの処理記述を、オフロード対象の機能ブロックとして、GPUやFPGA等にオフロードする。具体的には、置換処理部214は、GPU向けのライブラリやFPGA向けのIPコア等に置換した機能ブロックを、CPUプログラムとのインタフェースを作成することでオフロードする。置換処理部214は、CUDA,OpenCL等の中間言語ファイル132を出力する。
<
In <Processing C-1> described later, the
The
置換処理部214は、後記<処理C-2>において、アプリケーションプログラムのソースコードの置換元の処理記述を、取得したライブラリおよびIPコアの処理記述に置換するとともに、置換元と置換先で引数、戻り値の数または型が異なる場合に、その確認を通知する。
In the process C-2 described below, the
置換処理部214は、《機能ブロックオフロード:C言語》では、CUDAのライブラリ呼び出しを、PGIコンパイラに指定する。
置換処理部214は、《機能ブロックオフロード:Python》では、CUDAのライブラリ呼び出しを、pyCudaで指定する。
置換処理部214は、《機能ブロックオフロード:Java》では、CUDAのライブラリ呼び出しを、Jcudaで指定する。
In the case of <<Function block offload: C language>>, the
In <<Function block offload: Python>>, the
In <<Function Block Offload: Java>>, the
<オフロードパターン作成部215>
オフロードパターン作成部215は、1以上のオフロードするパターンを作成する。具体的には、ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出する。
<Offload
The offload
ここで、コードパターンDB230は、IPコア関連の情報としてOpenCLコードを記憶している。オフロードパターン作成部215は、FPGA等のPLDにオフロードする場合は、OpenCLコードをもとにOpenCLインタフェースを用いてホストとPLDとを接続するとともに、OpenCLコードをもとにPLDへのIPコアの実装を行う。Here, the
OpenCLのAPIに沿う、カーネルプログラムとホストプログラムのインタフェース記述について述べる。なお、下記説明は、後記[処理C](ホスト側とのインタフェースの整合)の<処理C-1>の具体例に対応する。This section describes the interface description between the kernel program and the host program in accordance with the OpenCL API. Note that the following explanation corresponds to a specific example of <Process C-1> in [Process C] (matching the interface with the host side) described below.
OpenCLのC言語の文法に沿って作成したカーネルは、OpenCLのC言語のランタイムAPIを利用して、作成するホスト(例えば、CPU)側のプログラムによりデバイス(例えば、FPGA)で実行される。カーネル関数hello()をホスト側から呼び出す部分は、OpenCLランタイムAPIの一つであるclEnqueueTask()を呼び出すことである。
ホストコードで記述するOpenCLの初期化、実行、終了の基本フローは、下記ステップ1~13である。このステップ1~13のうち、ステップ1~10がカーネル関数hello()をホスト側から呼び出すまでの手続(準備)であり、ステップ11でカーネルの実行となる。
A kernel created according to the OpenCL C language syntax is executed on a device (e.g., FPGA) by a program created on the host (e.g., CPU) side using the OpenCL C language runtime API. The part where the kernel function hello() is called from the host side is to call clEnqueueTask(), which is one of the OpenCL runtime APIs.
The basic flow of OpenCL initialization, execution, and termination described in the host code consists of the following
1.プラットフォーム特定
OpenCLランタイムAPIで定義されているプラットフォーム特定機能を提供する関数clGetPlatformIDs()を用いて、OpenCLが動作するプラットフォームを特定する。
1. Platform Specific
The platform on which OpenCL runs is identified using a function clGetPlatformIDs( ) that provides platform identification functionality defined in the OpenCL runtime API.
2.デバイス特定
OpenCLランタイムAPIで定義されているデバイス特定機能を提供する関数clGetDeviceIDs()を用いて、プラットフォームで使用するGPU等のデバイスを特定する。
2. Device Identification
A device such as a GPU to be used on the platform is identified using a function clGetDeviceIDs( ) that provides a device identification function defined in the OpenCL runtime API.
3.コンテキスト作成
OpenCLランタイムAPIで定義されているコンテキスト作成機能を提供する関数clCreateContext()を用いて、OpenCLを動作させる実行環境となるOpenCLコンテキストを作成する。
3. Creating a context
An OpenCL context that serves as an execution environment for running OpenCL is created using a function clCreateContext( ) that provides a context creation function defined in the OpenCL runtime API.
4.コマンドキュー作成
OpenCLランタイムAPIで定義されているコマンドキュー作成機能を提供する関数clCreateCommandQueue()を用いて、デバイスを制御する準備であるコマンドキューを作成する。OpenCLでは、コマンドキューを通して、ホストからデバイスに対する働きかけ(カーネル実行コマンドやホスト-デバイス間のメモリコピーコマンドの発行)を実行する。
4. Creating a command queue
A command queue is created in preparation for controlling the device using the function clCreateCommandQueue(), which provides a command queue creation function 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.
5.メモリオブジェクト作成
OpenCLランタイムAPIで定義されているデバイス上にメモリを確保する機能を提供する関数clCreateBuffer()を用いて、ホスト側からメモリオブジェクトを参照できるようにするメモリオブジェクトを作成する。
5. Creating a memory object
A memory object that allows the host side to reference the memory object is created using a function clCreateBuffer( ) that provides a function for allocating memory on a device defined in the OpenCL runtime API.
6.カーネルファイル読み込み
デバイスで実行するカーネルは、その実行自体をホスト側のプログラムで制御する。このため、ホストプログラムは、まずカーネルプログラムを読み込む必要がある。カーネルプログラムには、OpenCLコンパイラで作成したバイナリデータや、OpenCL C言語で記述されたソースコードがある。このカーネルファイルを読み込む(記述省略)。なお、カーネルファイル読み込みでは、OpenCLランタイムAPIは使用しない。
6. Loading a 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. The kernel program includes binary data created by the OpenCL compiler and 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.
7.プログラムオブジェクト作成
OpenCLでは、カーネルプログラムをプログラムオブジェクトとして認識する。この手続きがプログラムオブジェクト作成である。
OpenCLランタイムAPIで定義されているプログラムオブジェクト作成機能を提供する関数clCreateProgramWithSource()を用いて、ホスト側からメモリオブジェクトを参照できるようにするプログラムオブジェクトを作成する。カーネルプログラムのコンパイル済みバイナリ列から作成する場合は、clCreateProgramWithBinary()を使用する。
7. Creating a program object
In OpenCL, a kernel program is recognized as a program object. This procedure is called program object creation.
Create a program object that allows the host to reference the memory object by using the function clCreateProgramWithSource(), which provides the program object creation function defined in the OpenCL runtime API. When creating a program object from a compiled binary sequence of a kernel program, use clCreateProgramWithBinary().
8.ビルド
ソースコードとして登録したプログラムオブジェクトを OpenCL Cコンパイラ・リンカを使いビルドする。
OpenCLランタイムAPIで定義されているOpenCL Cコンパイラ・リンカによるビルドを実行する関数clBuildProgram()を用いて、プログラムオブジェクトをビルドする。なお、clCreateProgramWithBinary()でコンパイル済みのバイナリ列からプログラムオブジェクトを生成した場合、このコンパイル手続は不要である。
8. Build Build the program object registered as source code using the OpenCL C compiler and linker.
A program object is built using the function clBuildProgram(), which executes a build by the OpenCL C compiler and linker defined in the OpenCL runtime API. Note that if a program object is generated from a compiled binary string using clCreateProgramWithBinary(), this compilation procedure is not necessary.
9.カーネルオブジェクト作成
OpenCLランタイムAPIで定義されているカーネルオブジェクト作成機能を提供する関数clCreateKernel()を用いて、カーネルオブジェクトを作成する。1つのカーネルオブジェクトは、1つのカーネル関数に対応するので、カーネルオブジェクト作成時には、カーネル関数の名前(hello)を指定する。また、複数のカーネル関数を1つのプログラムオブジェクトとして記述した場合、1つのカーネルオブジェクトは、1つのカーネル関数に1対1で対応するので、clCreateKernel()を複数回呼び出す。
9. Creating a kernel object
A kernel object is created using the function clCreateKernel(), which provides a kernel object creation function defined in the OpenCL runtime API. Since one kernel object corresponds to one kernel function, the name of the kernel function (hello) is specified when creating the kernel object. Furthermore, if multiple kernel functions are written as one program object, one kernel object corresponds one-to-one to one kernel function, so clCreateKernel() is called multiple times.
10.カーネル引数設定
OpenCLランタイムAPIで定義されているカーネルへ引数を与える(カーネル関数が持つ引数へ値を渡す)機能を提供する関数clSetKernel()を用いて、カーネル引数を設定する。
以上、上記ステップ1~10で準備が整い、ホスト側からデバイスでカーネルを実行するステップ11に入る。
10. Kernel argument settings
The kernel arguments are set using the function clSetKernel(), which provides a function for giving arguments to the kernel defined in the OpenCL runtime API (passing values to arguments held by the kernel function).
As described above, the preparation is completed in
11.カーネル実行
カーネル実行(コマンドキューへ投入)は、デバイスに対する働きかけとなるので、コマンドキューへのキューイング関数となる。
OpenCLランタイムAPIで定義されているカーネル実行機能を提供する関数clEnqueueTask()を用いて、カーネルhelloをデバイスで実行するコマンドをキューイングする。カーネルhelloを実行するコマンドがキューイングされた後、デバイス上の実行可能な演算ユニットで実行されることになる。
11. Kernel Execution Kernel execution (submission to the command queue) is an action on the device, so it is a queuing function for the command queue.
A command to execute the kernel hello on the device is queued using a function clEnqueueTask() that provides a kernel execution function defined in the OpenCL runtime API. After the command to execute the kernel hello is queued, it will be executed on an executable computing unit on the device.
12.メモリオブジェクトからの読み込み
OpenCLランタイムAPIで定義されているデバイス側のメモリからホスト側のメモリへデータをコピーする機能を提供する関数clEnqueueReadBuffer()を用いて、デバイス側のメモリ領域からホスト側のメモリ領域にデータをコピーする。また、ホスト側からデバイス側のメモリへデータをコピーする機能を提供する関数clEnqueueWrightBuffer()を用いて、ホスト側のメモリ領域からデバイス側のメモリ領域にデータをコピーする。なお、これらの関数は、デバイスに対する働きかけとなるので、一度コマンドキューへコピーコマンドがキューイングされてからデータコピーが始まることになる。
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 a function 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 a function for copying data from the host memory to the device memory. Note that these functions act on the device, so the copy command is queued in the command queue once before data copying begins.
13.オブジェクト解放
最後に、ここまでに作成してきた各種オブジェクトを解放する。
以上、OpenCL C言語に沿って作成されたカーネルの、デバイス実行について説明した。
13. Releasing objects Finally, release the various objects that you have created up to this point.
The above describes device execution of a kernel written in accordance with the OpenCL C language.
<性能測定部116>
性能測定部116は、作成された処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン14に配置し、GPUやFPGA等にオフロードした際の性能測定用処理を実行する。
性能測定部116は、バイナリファイル配置部(Deploy binary files)116aを備える。バイナリファイル配置部116aは、GPUやFPGAを備えた検証用マシン14に、中間言語から導かれるバイナリファイルをデプロイ(配置)する。
<
The
The
性能測定部116は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、バイナリファイル配置部116aに戻す。この場合、性能測定部116は、抽出された別の処理パターンを用いて、抽出された中間言語をもとに、性能測定を試行する(後記図13の符号g参照)。The
性能測定の具体例について述べる。
オフロードパターン作成部215は、GPUやFPGAにオフロード可能な機能ブロックをオフロードする処理パターンを作成し、作成された処理パターンの中間言語を、実行ファイル作成部117がコンパイルする。性能測定部116は、コンパイルされたプログラムの性能を測定する(「1回目の性能測定」)。
A specific example of performance measurement will be described.
The offload
そして、オフロードパターン作成部215は、性能測定された中でCPUに比べ高性能化された処理パターンをリスト化する。オフロードパターン作成部215は、リストの処理パターンを組み合わせてオフロードする新たな処理パターンを作成する。オフロードパターン作成部215は、組み合わせたオフロード処理パターンと中間言語を作成し、中間言語を、実行ファイル作成部117がコンパイルする。
性能測定部116は、コンパイルされたプログラムの性能を測定する(「2回目の性能測定」)。
The offload
The
<実行ファイル作成部117>
実行ファイル作成部117は、オフロードする処理パターンの中間言語をコンパイルして実行ファイルを作成する。一定数繰り返された、性能測定結果をもとに、1以上の処理パターンから最高処理性能の処理パターンを選択し、最高処理性能の処理パターンをコンパイルして最終実行ファイルを作成する。
<Executable
The executable
<本番環境配置部118>
本番環境配置部118は、作成した実行ファイルを、ユーザ向けの本番環境に配置する(「最終バイナリファイルの本番環境への配置」)。本番環境配置部118は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。
<Production
The production
<性能測定テスト抽出実行部119>
性能測定テスト抽出実行部119は、実行ファイル配置後、テストケースDB131から性能試験項目を抽出し、性能試験を実行する。
性能測定テスト抽出実行部119は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースDB131から抽出し、抽出した性能試験を自動実行する。
<Performance Measurement Test
After arranging the executable file, the performance measurement test
After arranging the executable file, the performance measurement test
<ユーザ提供部120>
ユーザ提供部120は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する(「価格・性能等の情報のユーザへの提供」)。テストケースDB131には、アプリケーションの性能を測定する試験を自動で行うためのデータが格納されている。ユーザ提供部120は、テストケースDB131の試験データを実行した結果と、システムに用いられるリソース(仮想マシンや、FPGAインスタンス、GPUインスタンス等)の各単価から決まるシステム全体の価格をユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。
<
The
以下、上述のように構成されたオフロードサーバ200の機能ブロックオフロード処理について説明する。
Below, we will explain the functional block offload processing of the
上記、機能ブロックのオフロードの処理の概要と考慮点について説明する。
FPGAに関しては、ハードウェア回路設計に多大な時間がかかることもあり、一度設計した機能を、IPコア(Intellectual Property Core)という形で再利用可能にすることが多い。IPコアとは、FPGA、IC、LSIなどの半導体を構成するための部分的な回路情報であり、特に機能単位でまとめられている。IPコアは、暗号化/復号化処理、FFT(Fast Fourier Transform)等の算術演算、画像処理、音声処理等が代表的な機能例である。IPコアは、ライセンス料を支払うものが多いが、一部はフリーで提供されているものもある。
The outline of the above-mentioned function block offloading process and points to consider will be described below.
With regard to FPGAs, since hardware circuit design can take a lot of time, functions that have already been designed are often made reusable in the form of IP cores (Intellectual Property Cores). IP cores are partial circuit information for configuring semiconductors such as FPGAs, ICs, and LSIs, and are organized by function. Typical examples of IP core functions include encryption/decryption processing, arithmetic operations such as FFT (Fast Fourier Transform), image processing, and audio processing. Many IP cores require a license fee, but some are provided free of charge.
第2の実施形態では、FPGAに関しては、IPコアを自動オフロードに利用する。また、GPUに関しては、IPコアという言い方ではないものの、FFT、線形代数演算等が代表的な機能例であり、CUDAを用いて実装されたcuFFTやcuBLAS等がGPU向けライブラリとしてフリーで提供されている。本第2の実施形態では、GPUに関してこれらのライブラリを活用する。In the second embodiment, for FPGAs, IP cores are used for automatic offloading. For GPUs, although they are not called IP cores, FFT and linear algebra operations are typical examples of functions, and cuFFT and cuBLAS implemented using CUDA are provided free of charge as libraries for GPUs. In this second embodiment, these libraries are utilized for GPUs.
本第2の実施形態では、CPU向けに作られた既存プログラムコードの中で、FFT処理等、GPU、FPGAにオフロードすることで高速化できるような機能ブロックが含まれる場合に、GPU向けライブラリやFPGA向けIPコア等に置き換えることでの高速化を図る。In this second embodiment, when existing program code created for a CPU contains functional blocks such as FFT processing that can be accelerated by offloading them to a GPU or FPGA, the acceleration is achieved by replacing them with GPU libraries or FPGA IP cores, etc.
[機能ブロックのオフロード処理概要]
第2の実施形態のオフロードサーバ200は、環境適応ソフトウェアの要素技術としてユーザアプリケーションロジックのGPU、FPGA自動オフロードに適用した例である。
図13は、オフロードサーバ200の機能ブロックのオフロード処理を示す図である。
図13に示すように、オフロードサーバ200は、環境適応ソフトウェアの要素技術に適用される。オフロードサーバ200は、制御部(自動オフロード機能部)11と、コードパターンDB230、テストケースDB131と、中間言語ファイル132と、検証用マシン14と、を有している。
オフロードサーバ200は、ユーザが利用するアプリケーションコード(Application code)130を取得する。
[Outline of offload processing of functional blocks]
The
FIG. 13 is a diagram showing the offload process of the functional blocks of the
13, the
The
ユーザは、例えば、各種デバイス(Device)151、CPU-GPUを有する装置152、CPU-FPGAを有する装置153、CPUを有する装置154を利用する。オフロードサーバ200は、機能処理をCPU-GPUを有する装置152、CPU-FPGAを有する装置153のアクセラレータに自動オフロードする。
A user uses, for example,
以下、図13のステップ番号を参照して各部の動作を説明する。
<ステップS31:Specify application code>
ステップS31において、アプリケーションコード指定部111(図12参照)は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部112に渡す。
The operation of each unit will be described below with reference to the step numbers in FIG.
<Step S31: Specify application code>
In step S 31 , the application code designation unit 111 (see FIG. 12 ) passes the application code described in the received file to the application code analysis unit 112 .
<ステップS32:Analyze application code>(コード分析)
ステップS32において、アプリケーションコード分析部112(図12参照)は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
<Step S32: Analyze application code>
In step S32, the application code analysis unit 112 (see FIG. 12) uses a syntax analysis tool such as Clang to analyze the source code to analyze the library calls and functional processes contained in the code, as well as loop statement structures and the like.
<ステップS33:Extract offloadable area>(オフロード可能処理抽出)
ステップS33において、置換機能検出部213(図12参照)は、把握したライブラリ呼び出しや機能処理について、コードパターンDB230と照合することで、GPU、FPGAにオフロードできるオフロード可能処理を抽出する。
<Step S33: Extract offloadable area> (Extract offloadable processing)
In step S33, the replacement function detection unit 213 (see FIG. 12) compares the identified library calls and functional processes with the
<ステップS34:Output intermediate file>(オフロード用中間ファイル出力)
ステップS34において、置換処理部214(図12参照)は、抽出したオフロード可能処理を、GPU向けのライブラリやFPGA向けのIPコア等に置換する。置換処理部214は、GPU向けのライブラリやFPGA向けのIPコア等に置換した機能ブロックを、CPUプログラムとのインタフェースを作成することでオフロードする。置換処理部214は、CUDA,OpenCL等の中間言語ファイル132を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。
ここで、オフロード可能処理が直ちに高速化につながるか、またコスト効果が十分であるかは分からないので、オフロードパターン作成部215は、後述する検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出する。
<Step S34: Output intermediate file> (Output intermediate file for offloading)
In step S34, the replacement processing unit 214 (see FIG. 12) replaces the extracted offloadable processes with libraries for GPUs, IP cores for FPGAs, etc. The
Here, since it is not known whether offloadable processing will immediately lead to faster processing or whether the cost-effectiveness will be sufficient, the offload
<ステップS21:Deploy binary files>(デプロイ、性能測定試行)
ステップS21において、バイナリファイル配置部116a(図12参照)は、GPU、FPGAを備えた検証用マシン14に、中間言語から導かれる実行ファイルをデプロイする。バイナリファイル配置部116aは、配置したファイルを起動し、想定するテストケースを実行して、オフロードした際の性能を測定する。
<Step S21: Deploy binary files> (Deployment, performance measurement trial)
In step S21, the binary
<ステップS22:Measure performances>
ステップS22において、性能測定部116(図12参照)は、配置したファイルを実行し、オフロードした際の性能を測定する。
<Step S22: Measure performance>
In step S22, the performance measurement unit 116 (see FIG. 12) executes the arranged file and measures the performance when offloaded.
図13の符号gに示すように、制御部210は、上記ステップS12乃至ステップS22を繰り返し実行する。制御部210の自動オフロード機能をまとめると、下記である。すなわち、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。置換機能検出部213は、検出したライブラリ呼び出しや機能処理について、コードパターンDB230と照合することで、GPU、FPGAにオフロードできるオフロード可能処理を抽出する。置換処理部214は、抽出したオフロード可能処理を、GPU向けのライブラリやFPGA向けのIPコア等に置換する。そして、オフロードパターン作成部215は、GPU向けのライブラリやFPGA向けのIPコア等に置換した機能ブロックを、CPUプログラムとのインタフェースを作成することでオフロードする。As shown by the symbol g in FIG. 13, the
<ステップS23:Deploy final binary files to production environment>
ステップS23において、本番環境配置部118は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。
<Step S23: Deploy final binary files to production environment>
In step S23, the production
<ステップS24:Extract performance test cases and run automatically>
ステップS24において、性能測定テスト抽出実行部119は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースDB131から抽出し、抽出した性能試験を自動実行する。
<Step S24: Extract performance test cases and run automatically>
In step S24, after arranging the executable file, the performance measurement test
<ステップS25:Provide price and performance to a user to judge>
ステップS25において、ユーザ提供部120は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。
<Step S25: Provide price and performance to a user to judge>
In step S25, the
上記ステップS11~ステップS25は、ユーザのサービス利用のバックグラウンドで行われ、例えば、仮利用の初日の間に行う等を想定している。The above steps S11 to S25 are carried out in the background while the user is using the service, and are assumed to be carried out, for example, during the first day of trial use.
上記したように、オフロードサーバ200の制御部(自動オフロード機能部)210は、環境適応ソフトウェアの要素技術に適用した場合、機能処理のオフロードのため、ユーザが利用するアプリケーションプログラムのソースコードから、オフロードする領域を抽出して中間言語を出力する(ステップS31~ステップS34)。制御部210は、中間言語から導かれる実行ファイルを、検証用マシン14に配置実行し、オフロード効果を検証する(ステップS21~ステップS22)。検証を繰り返し、適切なオフロード領域を定めたのち、制御部210は、実際にユーザに提供する本番環境に、実行ファイルをデプロイし、サービスとして提供する(ステップS23~ステップS25)。As described above, when applied to the elemental technology of the environment-adaptive software, the control unit (automatic offload function unit) 210 of the
一般に、性能に関しては、最大性能になる設定を一回で自動発見するのは難しい。このため、オフロードパターンを、性能測定を検証環境で何度か繰り返すことにより試行し、高速化できるパターンを見つけることが本発明の特徴である。 In general, it is difficult to automatically discover settings that maximize performance in one go. For this reason, a feature of the present invention is that it tries out offload patterns by repeatedly measuring performance in a verification environment several times to find a pattern that can increase speed.
[機能ブロックのオフロード処理詳細]
機能ブロックのオフロードについては、機能ブロックの検出(以下、「処理A」という)、その機能ブロックがオフロード用の既存ライブラリ/IPコア等があるかを検出(以下、「処理B」という)、機能ブロックをライブラリ/IPコア等と置換した際にホスト側とのインタフェースの整合(以下、「処理C」という)、の3つ要素を考慮する必要がある。上記3つ要素の考慮点に従い、機能ブロックのオフロード処理について詳細に述べる。
[Function block offload processing details]
For offloading of a functional block, three elements must be considered: detection of the functional block (hereinafter referred to as "Process A"), detection of whether the functional block has an existing library/IP core, etc. for offloading (hereinafter referred to as "Process B"), and consistency of the interface with the host side when replacing the functional block with the library/IP core, etc. (hereinafter referred to as "Process C"). The offloading process of the functional block will be described in detail according to the considerations of the above three elements.
[処理A](機能ブロックの検出)
「処理A」(機能ブロックの検出)は、ライブラリの関数呼び出しを行い、ライブラリの関数呼び出しを機能ブロックとする<処理A-1>と、登録されていないライブラリの関数呼び出しである場合、クラス、構造体等を検出して機能ブロックとする<処理A-2>と、に分けられる。すなわち、<処理A-1>は、既存のライブラリの関数呼び出しを検出して機能ブロックとするものであり、<処理A-2>は、<処理A-1>において機能ブロックを検出しない場合に、クラスまたは構造体から機能ブロックを抽出するものである。
[Process A] (detection of functional blocks)
"Processing A" (detection of function blocks) is divided into <Processing A-1>, which calls a function in a library and makes the function call in the library a function block, and <Processing A-2>, which detects a class, structure, etc. and makes the function block if the function call is in an unregistered library. That is, <Processing A-1> detects a function call in an existing library and makes it a function block, and <Processing A-2> extracts a function block from a class or structure if a function block is not detected in <Processing A-1>.
<処理A-1>
アプリケーションコード分析部112は、構文解析を用いて、ソースコードから外部のライブラリの関数呼び出しを行っていることを検知する。詳細には、下記の通りである。コードパターンDB230は、FFT等算術計算等のライブラリリストを保持している。アプリケーションコード分析部112は、ソースコードを構文解析し、コードパターンDB230が保持しているライブラリリストと照合して、外部のライブラリの関数呼び出しを行っていることを検知する。
<Process A-1>
The application code analysis unit 112 uses syntax analysis to detect whether a function of an external library is being called from source code. The details are as follows. The
<処理A-2>
アプリケーションコード分析部112は、登録されていないライブラリ呼び出し以外の機能処理を機能ブロックとして検出するため、構文解析を用いて、ソースコードの定義記述からクラスまたは構造体の機能処理を検出する。アプリケーションコード分析部112は、例えば、C言語のstructを使って定義されるいくつかの変数をひとまとまりにした型である構造体(structure)や、インスタンス化したオブジェクトの型が値型である構造体に対して参照型であるクラス(class)を検出する。また、アプリケーションコード分析部112は、例えばJava(登録商標)において構造体に代替使用されるクラスを検出する。
<Process A-2>
In order to detect functional processes other than unregistered library calls as functional blocks, the application code analysis unit 112 detects functional processes of classes or structures from the definition description of the source code using syntax analysis. For example, the application code analysis unit 112 detects structures that are types that group together several variables defined using the struct of the C language, and classes in which the type of an instantiated object is a reference type for a structure that is a value type. The application code analysis unit 112 also detects classes that are used as substitutes for structures in Java (registered trademark), for example.
[処理B](オフロード可能機能の検出)
[処理B](オフロード可能機能の検出)は、<処理A-1>を受け、コードパターンDB230を参照して置換可能GPUライブラリ、IPコアを取得する<処理B-1>と、<処理A-2>を受け、アプリコードの置換元の処理記述を、置換先のGPUライブラリ、IPコア処理記述に置換する<処理B-2>と、に分けられる。すなわち、<処理B-1>は、ライブラリ名をキーに、コードパターンDB230から置換可能GPUライブラリ、IPコアを取得するものである。<処理B-2>は、クラス、構造体等のコードをキーに、置換可能GPUライブラリ・IPコアを検出し、アプリコードの置換元の処理記述を、置換先のGPUライブラリ、IPコア処理記述に置換するものである。
[Process B] (Detection of offloadable functions)
[Processing B] (detection of offloadable functions) is divided into <Processing B-1>, which receives <Processing A-1> and refers to the
処理Bの前提として、コードパターンDB230には、特定のライブラリ、機能ブロックを高速化するGPU用ライブラリやFPGA用IPコアとそれに関連する情報が保持されている。また、コードパターンDB230には、置換元のライブラリ、機能ブロックについては、機能名とともにコードや実行ファイルが登録されている。As a premise for process B, the
<処理B-1>
置換機能検出部213は、<処理A-1>でアプリケーションコード分析部112が検出したライブラリ呼び出しに対して、ライブラリ名をキーに、コードパターンDB230を検索し、コードパターンDB230から、置換可能GPUライブラリ(高速化できるGPU用ライブラリ)やFPGA用IPコアを取得する。
<Process B-1>
The replacement function detection unit 213 searches the
<処理B-1>の例を記載する。
置換機能検出部213は、例えば、置換元の処理が2D FFTの処理(非特許文献4等にコードがある)であった場合は、その外部ライブラリ名をキーに、2D FFTを処理するFPGA処理として、OpenCLコードを検出する(ホストプログラム、カーネルプログラム)等)。なお、OpenCLコードは、コードパターンDB230に記憶されている。
An example of <Process B-1> is described below.
For example, if the processing to be replaced is 2D FFT processing (code is available in
置換機能検出部213は、例えば、置換元の処理が2D FFTの処理であった場合は、GPUライブラリとして検出されたcuFFTの中の関数呼び出しに置換する。なお、GPUライブラリは、コードパターンDB230に記憶されている。
For example, if the original process is a 2D FFT process, the replacement function detection unit 213 replaces it with a function call in cuFFT detected as a GPU library. The GPU library is stored in the
<処理B-2>
置換機能検出部213は、<処理A-2>でアプリケーションコード分析部112が検出したクラス、構造体等のコードをキーに、コードパターンDB230を検索し、コードパターンDB230から、類似性検出ツールを用いて置換可能GPUライブラリ(高速化できるGPU用ライブラリ)やFPGA用IPコアを取得する。類似性検出ツールとは、Deckard等、コピーコードやコピー後変更したコードの検出を対象とするツールである。置換機能検出部213が、類似性検出ツールを用いることで、行列計算のコード等、CPUで計算する場合は記述が同様になる処理や、他者のコードをコピーして変更した処理等を一部検出できる。なお、類似性検出ツールは、新規に独立に作成したようなクラス等については検出が困難となるため対象外である。
<Process B-2>
The replacement function detection unit 213 searches the
<処理B-2>の例を記載する。
置換機能検出部213は、置換元CPUコードに検知されたクラスや構造体に対して、Deckard等の類似性検知ツールを用いて、コードパターンDB230に登録された類似クラスや構造体を検索する。例えば、置換元の処理(非特許文献4等にコードがある)が2D FFTのクラスであった場合は、その類似クラスとしてコードパターンDB230に登録されたクラスが2D FFTのクラスが検出される。コードパターンDB230には、2D FFTをオフロード可能なIPコアやGPUライブラリが登録されている。そのため、<処理B-1>と同様に、2D FFTに対して、OpenCLコード(ホストプログラム、カーネルプログラム等)やGPUライブラリを検出する。
An example of <Process B-2> is described below.
The replacement function detection unit 213 searches for similar classes and structures registered in the
[処理C](ホスト側とのインタフェースの整合)
[処理C](ホスト側とのインタフェースの整合)は、<処理C-1>と、<処理C-2>とを有する。<処理C-1>は、<処理B-1>を受け、アプリコードの置換元の処理記述を、置換先のGPUライブラリ、IPコア処理記述に置換するとともに、GPUライブラリ、IPコア呼び出しのためのインタフェース処理を記述する。<処理C-2>は、<処理B-2>を受け、アプリコードの置換元の処理記述を、置換先のGPUライブラリ、IPコア処理記述に置換するとともに、GPUライブラリ、IPコア呼び出しのためのインタフェース処理を記述する。ここで、上記GPUライブラリ、IPコア呼び出しのためのインタフェース処理の記述が、「ホスト側とのインタフェースの整合」に対応する。
[Process C] (Interface matching with the host side)
[Processing C] (matching the interface with the host side) includes <Processing C-1> and <Processing C-2>. <Processing C-1> receives <Processing B-1>, and replaces the processing description of the application code to be replaced with the GPU library and IP core processing description of the replacement destination, and also describes the interface processing for calling the GPU library and IP core. <Processing C-2> receives <Processing B-2>, and replaces the processing description of the application code to be replaced with the GPU library and IP core processing description of the replacement destination, and also describes the interface processing for calling the GPU library and IP core. Here, the description of the interface processing for calling the GPU library and IP core corresponds to "matching the interface with the host side".
<処理C-1>
置換処理部214は、アプリコードの置換元の処理記述を、置換先のGPUライブラリ、IPコア処理記述に置換する。そして、置換処理部214は、GPUライブラリ、IPコア呼び出しのためのインタフェース処理を記述し(OpenCL API等)、作成したパターンをコンパイルする。
<Process C-1>
The
<処理C-1>について、より詳細に説明する。
置換機能検出部213は、<処理A-1>で検出したライブラリ呼び出しに対して、<処理B-1>で該当するライブラリやIPコアを検索している。このため、置換処理部214は、置換するライブラリやIPコアをGPUやFPGAに実装し、ホスト側(CPU)プログラムと繋ぐインタフェース処理を行う。
<Process C-1> will now be described in more detail.
The replacement function detection unit 213 searches for a corresponding library or IP core in <Process B-1> for the library call detected in <Process A-1>. For this reason, the
ここで、GPU用ライブラリの場合は、CUDA等のライブラリを想定しており、C言語コードからCUDAライブラリを利用する手法がライブラリとともに公開されている。そこで、コードパターンDB230に、ライブラリ利用手法も含めて登録しておき、置換処理部214は、コードパターンDB230に登録されたライブラリ利用手法に従って、アプリコードの置換元の処理記述を、置換先のGPUライブラリに置換するとともに、GPUライブラリで利用する関数の呼び出し等の所定記述を行う。
In the case of a GPU library, a library such as CUDA is assumed, and a method for using the CUDA library from C language code is published together with the library. Therefore, the library usage method is also registered in the
FPGA用IPコアの場合は、HDL(Hardware Description Language)等が想定される。この場合、IPコア関連の情報としてOpenCLコードもコードパターンDB230に保持されている。置換処理部214は、FPGAとのインタフェース処理を、高位合成ツール(例えば、Xilinx Vivado, Intel HLS Compiler等)を介して行うことができる。置換処理部214は、例えば、OpenCLコードから、OpenCLインタフェースを用いたCPUとFPGAの接続を、高位合成ツールを介して行う。同様に、置換処理部214は、FPGAへのIPコア実装を、XilinxやIntel等のFPGAベンダの高位合成ツールを介して行う。In the case of an IP core for FPGA, HDL (Hardware Description Language) and the like are assumed. In this case, OpenCL code is also stored in the
<処理C-2>
置換処理部214は、アプリコードの置換元の処理記述を、置換先のGPUライブラリ、IPコア処理記述に置換する。そして、置換処理部214は、置換元と置換先で引数や戻り値の数や型が異なる場合に、ユーザに確認し、GPUライブラリ、IPコア呼び出しのためのインタフェース処理を記述(OpenCL API等)するとともに、作成したパターンをコンパイルする。すなわち、<処理C-2>では、置換処理部214は、<処理A-2>で検出したクラス、構造体等に対して、<処理B-2>で高速化できるライブラリやIPコアを検索している。このため、置換処理部214は、<処理C-2>では該当するライブラリやIPコアをGPUやFPGAに実装する。
<Process C-2>
The
<処理C-2>について、より詳細に説明する。
<処理C-1>では、特定のライブラリ呼び出しに対して高速化するライブラリやIPコアであるため、インタフェース部分の生成等は必要になるものの、GPU、FPGAとホスト側プログラムの想定する引数、戻り値の数や型は合っていた。しかし、<処理B-2>は、類似性等で判断しているため、引数や戻り値の数や型等の基本的な部分が合っている保証はない。ライブラリやIPコアは、既存ノウハウであり、引数、戻り値の数や型が合っていない場合であっても、変更が頻繁にできるものではない。このため、オフロードを依頼するユーザに対して、元のコードの引数や戻り値の数や型について、ライブラリやIPコアに合わせて変更するか否かを確認する。そして、確認了承後にオフロード性能試験を試行する。
<Process C-2> will now be described in more detail.
In <Process C-1>, because it is a library or IP core that speeds up a specific library call, it is necessary to generate an interface, but the number and type of arguments and return values expected by the GPU, FPGA, and the host program match. However, in <Process B-2>, since it is judged based on similarity, there is no guarantee that basic parts such as the number and type of arguments and return values match. Libraries and IP cores are existing know-how, and even if the number and type of arguments and return values do not match, they cannot be changed frequently. For this reason, the user who requests offloading is asked whether the number and type of arguments and return values of the original code will be changed to match the library or IP core. After confirmation and approval, an offload performance test is attempted.
型の違いについて、floatとdouble等キャストすればよいだけであれば、処理パターン作成時にキャストする処理を追加し、特にユーザ確認せずに性能測定試行に入ってもよい。また、引数や戻り値で、元のプログラムとライブラリやIPコアで数が異なる場合、例えば、CPUプログラムで引数1,2が必須で引数3がオプションであり、ライブラリやIPコアで引数1,2が必須の場合等は、オプション引数3は省略しても問題はない。このような場合は、ユーザに確認せず、処理パターン作成時にオプション引数は自動で無しとして扱うなどしてもよい。なお、引数や戻り値の数や型が完全に合っている場合は、<処理C-1>と同様の処理でよい。
Regarding type differences, if casting between float and double etc. is all that is required, a casting process can be added when creating the processing pattern, and performance measurement trials can begin without user confirmation. Also, if the number of arguments or return values differs between the original program and the library or IP core, for example,
[《機能ブロックオフロード:共通》フローチャート]
次に、図14および図15を参照してオフロードサーバ200の《機能ブロックオフロード:共通》の動作概要を説明する。
[Function Block Offload: Common Flowchart]
Next, an outline of the operation of the <<Function block offload: common>> of the
・<処理A-1>と<処理B-1>と<処理C-1>のフローチャート
図14は、オフロードサーバ200の制御部(自動オフロード機能部)210が、《機能ブロックオフロード:共通》のオフロード処理において<処理A-1>と<処理B-1>と<処理C-1>とを実行する場合のフローチャートである。
ステップS501でアプリケーションコード分析部112(図12参照)は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
Flowchart of <Processing A-1>, <Processing B-1>, and <Processing C-1> FIG. 14 is a flowchart when the control unit (automatic offload function unit) 210 of the
In step S501, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application program to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code, including the loop statement structure and the like, as well as library calls and functional processing included in the code.
ステップS502で置換機能検出部213(図12参照)は、アプリケーションプログラムの外部ライブラリ呼び出しを検出する。In step S502, the replacement function detection unit 213 (see Figure 12) detects an external library call of the application program.
ステップS503で置換機能検出部213は、コードパターンDB230から、ライブラリ名をキーに、置換可能GPUライブラリを取得する。具体的には、置換機能検出部213は、把握した外部ライブラリ呼び出しについて、コードパターンDB230と照合することで、検出した置換可能GPUライブラリ・IPコアを、GPU、FPGAにオフロードできるオフロード可能な機能ブロックとして取得する。In step S503, the replacement function detection unit 213 obtains a replaceable GPU library from the
ステップS504で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリの処理記述に置換する。In step S504, the
ステップS505で置換処理部214は、置換したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。
In step S505, the
ステップS506で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
ステップS507で実行ファイル作成部117は、作成したパターンをコンパイルまたはインタプリットする。
In step S506, the
In step S507, the executable
ステップS508で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。
ステップS509で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。
In step S508, the
In step S509, the executable
ステップS510で実行ファイル作成部117は、作成した組合せパターンをコンパイルまたはインタプリットする。
ステップS511で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。
In step S510, the executable
In step S511, the
ステップS512で本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S512, the production
・<処理A-2>と<処理B-2>と<処理C-2>のフローチャート
図15は、オフロードサーバ200の制御部(自動オフロード機能部)210が、機能ブロックのオフロード処理において<処理A-2>と<処理B-2>と<処理C-2>とを実行する場合のフローチャートである。なお、<処理A-2>からの処理は、<処理A-1>からの処理と並行して行えばよい。
ステップS601でアプリケーションコード分析部112(図12参照)は、アプリケーションのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
15 is a flowchart of the case where the control unit (automatic offload function unit) 210 of the
In step S601, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code to analyze the library calls and functional processing included in the code, together with the loop statement structure, etc.
ステップS602で置換機能検出部213(図12参照)は、ソースコードからクラスまたは構造体の定義記述コードを検出する。In step S602, the replacement function detection unit 213 (see Figure 12) detects the definition description code of a class or structure from the source code.
ステップS603で置換機能検出部213は、コードパターンDB230から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能GPUライブラリを取得する。In step S603, the replacement function detection unit 213 uses a similarity detection tool to obtain a replaceable GPU library from the
ステップS604で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリ処理記述に置換する。In step S604, the
ステップS605で置換処理部214は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。
In step S605, the
ステップS606で置換機能検出部213は、置換または確認したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。 In step S606, the replacement function detection unit 213 offloads the processing description of the replaced or confirmed GPU library to the GPU as a functional block to be offloaded.
ステップS607で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S607, the
ステップS608で実行ファイル作成部117は、作成したパターンをコンパイルまたはインタプリットする。
ステップS609で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。
In step S608, the executable
In step S609, the
ステップS610で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。In step S610, the executable
ステップS611で実行ファイル作成部117は、作成した組合せパターンをコンパイルまたはインタプリットする。
ステップS612で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。
In step S611, the executable
In step S612, the
ステップS613で本番環境配置部118は、本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S613, the production
[《機能ブロックオフロード:C言語》]
機能ブロックオフロード:C言語の、コードの分析では、C言語を解析するClang(登録商標)等の構文解析ツールを用いて構文解析する。機能ブロックの把握については、構文解析ツールの結果を用いて、次処理のマッチング探索に用いるため、言語に非依存の機能ブロックとして管理する。
[Function block offload: C language]
Function block offload: In the analysis of C language code, a syntax analysis tool such as Clang (registered trademark) that analyzes C language is used to analyze the syntax. The results of the syntax analysis tool are used to grasp the function blocks, and they are managed as language-independent function blocks to be used for matching search in the next process.
オフロード可能機能ブロックの探索では、ライブラリ等の名前一致でのマッチングと、Deckard(登録商標)等のC言語機能ブロックの類似性検出ツールを用いた類似性検知による、探索が行われる。オフロード機能ブロックへの置換は、CUDAライブラリ呼び出し等、その言語からのオフロード機能利用に合わせた処理に置換する必要がある。 When searching for offloadable function blocks, the search is performed by matching the names of libraries, etc., and by detecting similarities using a similarity detection tool for C language function blocks such as Deckard (registered trademark). Replacement with an offload function block requires replacement with processing that matches the use of offload functions from that language, such as calling a CUDA library.
コンパイルは、CUDAライブラリ呼び出し等のC言語コードをPGIコンパイラ等でコンパイルする。性能測定は、言語に合わせて、Jenkins(登録商標)等の自動測定ツールも用いて行う。オフロード可能機能ブロックが複数の際は反復実行され、最高性能のパターンが最終解として決定される。 Compilation involves compiling C language code, such as CUDA library calls, with a PGI compiler or similar. Performance measurement is performed using an automated measurement tool, such as Jenkins (registered trademark), depending on the language. When there are multiple offloadable function blocks, they are executed repeatedly, and the pattern with the best performance is determined as the final solution.
このように、機能ブロック文オフロードでは、処理に関しては、機能ブロックの管理と機能ブロックの名前一致でのマッチングについては言語に非依存に適用できる。 In this way, with function block statement offloading, in terms of processing, function block management and name matching of function blocks can be applied independently of the language.
[《機能ブロックオフロード:C言語》フローチャート]
次に、図16および図17を参照してオフロードサーバ200の《機能ブロックオフロード:C言語》の動作概要を説明する。
[Function block offload: C language flowchart]
Next, an outline of the operation of the
・<処理A-1>と<処理B-1>と<処理C-1>のフローチャート
図16は、オフロードサーバ200の制御部(自動オフロード機能部)210が、《機能ブロックオフロード:C言語》のオフロード処理において<処理A-1>と<処理B-1>と<処理C-1>とを実行する場合のフローチャートである。
ステップS701でアプリケーションコード分析部112(図12参照)は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
Flowchart of <Processing A-1>, <Processing B-1>, and <Processing C-1> FIG. 16 is a flowchart when the control unit (automatic offload function unit) 210 of the
In step S701, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application program to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code, including the loop statement structure and the like, as well as library calls and functional processing included in the code.
ステップS702で置換機能検出部213(図12参照)は、アプリケーションの外部ライブラリ呼び出しを検出する。In step S702, the replacement function detection unit 213 (see Figure 12) detects an external library call of the application.
ステップS703で置換機能検出部213は、コードパターンDB230から、ライブラリ名をキーに、置換可能GPUライブラリを取得する。具体的には、置換機能検出部213は、把握した外部ライブラリ呼び出しについて、コードパターンDB230と照合することで、検出した置換可能GPUライブラリを、GPU、FPGAにオフロードできるオフロード可能な機能ブロックとして取得する。In step S703, the replacement function detection unit 213 obtains a replaceable GPU library from the
ステップS704で置換処理部214は、アプリケーションコードの置換元の処理記述を、置換先のGPUライブラリの処理記述に置換する。In step S704, the
ステップS705で置換処理部214は、置換したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。
In step S705, the
ステップS706で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S706, the
ステップS707で置換処理部214は、CUDAのライブラリ呼び出しを、PGIコンパイラに指定する。
In step S707, the
ステップS708で実行ファイル作成部117は、作成したパターンをPGIコンパイラでコンパイルする。
In step S708, the executable
ステップS709で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。
ステップS710で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。
In step S709, the
In step S710, the executable
ステップS711で実行ファイル作成部117は、作成した組合せパターンをPGIコンパイラでコンパイルする。
In step S711, the executable
ステップS712で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。In step S712, the
ステップS713で本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S713, the production
・<処理A-2>と<処理B-2>と<処理C-2>のフローチャート
図17は、オフロードサーバ200の制御部(自動オフロード機能部)210が、機能ブロックのオフロード処理において<処理A-2>と<処理B-2>と<処理C-2>とを実行する場合のフローチャートである。なお、<処理A-2>からの処理は、<処理A-1>からの処理と並行して行えばよい。
ステップS801でアプリケーションコード分析部112(図12参照)は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
17 is a flowchart of the case where the control unit (automatic offload function unit) 210 of the
In step S801, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application program to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code to analyze the library calls and functional processing included in the code, together with the loop statement structure and the like.
ステップS802で置換機能検出部213(図12参照)は、ソースコードからクラスまたは構造体の定義記述コードを検出する。In step S802, the replacement function detection unit 213 (see Figure 12) detects the definition description code of a class or structure from the source code.
ステップS803で置換機能検出部213は、コードパターンDB230から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能GPUライブラリを取得する。In step S803, the replacement function detection unit 213 uses a similarity detection tool to obtain a replaceable GPU library from the
ステップS804で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリの処理記述に置換する。In step S804, the
ステップS805で置換処理部214は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。
In step S805, the
ステップS806で置換機能検出部213は、置換または確認したGPUライブラリの置換元の処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。In step S806, the replacement function detection unit 213 offloads the processing description of the replaced or confirmed GPU library to the GPU as a functional block to be offloaded.
ステップS807で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S807, the
ステップS808で置換処理部214は、CUDAのライブラリ呼び出しを、PGIコンパイラに指定する。
In step S808, the
ステップS809で実行ファイル作成部117は、作成したパターンをPGIコンパイラでコンパイルする。
In step S809, the executable
ステップS810で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。In step S810, the
ステップS811で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。In step S811, the executable
ステップS812で実行ファイル作成部117は、作成した組合せパターンをPGIコンパイラでコンパイルする。
In step S812, the executable
ステップS813で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。In step S813, the
ステップS814で本番環境配置部118は、本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S814, the production
以上、《機能ブロックオフロード:C言語》について説明した、次に、《機能ブロックオフロード:Python》について説明する。 We have explained "Function Block Offload: C Language" above. Next, we will explain "Function Block Offload: Python".
[《機能ブロックオフロード:Python》]
機能ブロックオフロード:Pythonは、機能ブロックオフロードの、コードの分析では、Pythonを解析するast(登録商標)等の構文解析ツールを用いて構文解析する。機能ブロックの把握については、構文解析ツールの結果を用いて、次処理のマッチング探索に用いるため、言語に非依存の機能ブロックとして管理する。
[Function Block Offload: Python]
Function block offload: In the analysis of the code of function block offload, Python is parsed using a syntax analysis tool such as ast (registered trademark) that analyzes Python. The results of the syntax analysis tool are used to grasp the function blocks, and they are managed as language-independent function blocks to be used for matching search in the next process.
オフロード可能機能ブロックの探索では、ライブラリ等の名前一致でのマッチングと、CloneDigger(登録商標)等のPython機能ブロックの類似性検出ツールを用いた類似性検知による、探索が行われる。 When searching for offloadable function blocks, the search is performed by matching the names of libraries, etc., and by similarity detection using a Python function block similarity detection tool such as CloneDigger (registered trademark).
オフロード機能ブロックへの置換は、GPU処理のpyCUDAでの呼び出し等、その言語からのオフロード機能利用に合わせた処理に置換する必要がある。インタプリタは、CUDAに合わせたPythonコードをpyCUDAでインタプリットする。 When replacing with an offload function block, it is necessary to replace it with a process that matches the use of the offload function from the language, such as calling GPU processing with pyCUDA. The interpreter interprets Python code that is compatible with CUDA using pyCUDA.
性能測定は、言語に合わせて、Jenkins(登録商標)等の自動測定ツールも用いて行う。オフロード可能機能ブロックが複数の際は反復実行され、最高性能のパターンが最終解として決定される。 Performance measurements are also performed using automated measurement tools such as Jenkins (registered trademark) depending on the language. When there are multiple offloadable function blocks, they are executed repeatedly and the pattern with the best performance is determined as the final solution.
[《機能ブロックオフロード:Python》フローチャート]
次に、図18および図19を参照してオフロードサーバ200の《機能ブロック:Python》の動作概要を説明する。
[Function Block Offloading: Python Flowchart]
Next, an outline of the operation of the <<functional block: Python>> of the
・<処理A-1>と<処理B-1>と<処理C-1>のフローチャート
図18は、オフロードサーバ200の制御部(自動オフロード機能部)210が、《機能ブロック:Python》のオフロード処理において<処理A-1>と<処理B-1>と<処理C-1>とを実行する場合のフローチャートである。
ステップS901でアプリケーションコード分析部112(図12参照)は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
Flowchart of <Processing A-1>, <Processing B-1>, and <Processing C-1> Figure 18 is a flowchart when the control unit (automatic offload function unit) 210 of the
In step S901, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application program to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code, including the loop statement structure and the like, as well as library calls and functional processing included in the code.
ステップS902で置換機能検出部213(図12参照)は、アプリケーションの外部ライブラリ呼び出しを検出する。In step S902, the replacement function detection unit 213 (see Figure 12) detects an external library call of the application.
ステップS903で置換機能検出部213は、コードパターンDB230から、ライブラリ名をキーに、置換可能GPUライブラリを取得する。具体的には、置換機能検出部213は、把握した外部ライブラリ呼び出しについて、コードパターンDB230と照合することで、検出した置換可能GPUライブラリを、GPU、FPGAにオフロードできるオフロード可能な機能ブロックとして取得する。In step S903, the replacement function detection unit 213 obtains a replaceable GPU library from the
ステップS904で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリの処理記述に置換する。In step S904, the
ステップS905で置換処理部214は、置換したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。
In step S905, the
ステップS906で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S906, the
ステップS907で置換処理部214は、CUDAのライブラリ呼び出しを、pyCudaで指定する。
In step S907, the
ステップS908で実行ファイル作成部117は、作成したパターンをpyCudaでインタプリットする。
In step S908, the executable
ステップS909で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。
ステップS910で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。
In step S909, the
In step S910, the executable
ステップS911で実行ファイル作成部117は、作成した組合せパターンpyCudaでインタプリットする。
ステップS912で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。
In step S911, the executable
In step S912, the
ステップS913で本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S913, the production
・<処理A-2>と<処理B-2>と<処理C-2>のフローチャート
図19は、オフロードサーバ200の制御部(自動オフロード機能部)210が、機能ブロックのオフロード処理において<処理A-2>と<処理B-2>と<処理C-2>とを実行する場合のフローチャートである。なお、<処理A-2>からの処理は、<処理A-1>からの処理と並行して行えばよい。
ステップS1001でアプリケーションコード分析部112(図12参照)は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
19 is a flowchart of the case where the control unit (automatic offload function unit) 210 of the
In step S1001, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application program to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code, including the loop statement structure and the like, as well as library calls and functional processing included in the code.
ステップS1002で置換機能検出部213(図12参照)は、ソースコードからクラスまたは構造体の定義記述コードを検出する。In step S1002, the replacement function detection unit 213 (see Figure 12) detects definition description code of a class or structure from the source code.
ステップS1003で置換機能検出部213は、コードパターンDB230から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能GPUライブラリを取得する。In step S1003, the replacement function detection unit 213 uses a similarity detection tool to obtain a replaceable GPU library from the
ステップS1004で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリ処理記述に置換する。In step S1004, the
ステップS1005で置換処理部214は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。
In step S1005, the
ステップS1006で置換機能検出部213は、置換または確認したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。 In step S1006, the replacement function detection unit 213 offloads the processing description of the replaced or confirmed GPU library to the GPU as a functional block to be offloaded.
ステップS1007で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S1007, the
ステップS1008で置換処理部214は、CUDAのライブラリ呼び出しを、pyCudaで指定する。
In step S1008, the
ステップS1009で実行ファイル作成部117は、作成したパターンをpyCudaでインタプリットする。
In step S1009, the executable
ステップS1010で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。In step S1010, the
ステップS1011で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。In step S1011, the executable
ステップS1012で実行ファイル作成部117は、作成した組合せパターンをpyCudaでインタプリットする。
ステップS1013で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。
In step S1012, the executable
In step S1013, the
ステップS1014で本番環境配置部118は、本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S1014, the production
以上、《機能ブロックオフロード:Python》について説明した、次に、《機能ブロックオフロード:Java》について説明する。 We have explained "Function Block Offload: Python" above. Next, we will explain "Function Block Offload: Java".
[《機能ブロックオフロード:Java》]
機能ブロックオフロード:Javaの、コードの分析では、Javaを解析するJavaParser(登録商標)等の構文解析ツールを用いて構文解析する。機能ブロックの把握については、構文解析ツールの結果を用いて、次処理のマッチング探索に用いるため、言語に非依存の機能ブロックとして管理する。
[Function Block Offload: Java]
Function block offload: In the analysis of Java code, a syntax analysis tool such as JavaParser (registered trademark) is used to analyze the syntax of Java. The results of the syntax analysis tool are used to grasp the function blocks, and they are managed as language-independent function blocks to be used in the matching search for the next process.
オフロード可能機能ブロックの探索では、ライブラリ等の名前一致でのマッチングと、Deckard(登録商標)等のJava機能ブロックの類似性検出ツールを用いた類似性検知による、探索が行われる。オフロード機能ブロックへの置換は、GPU処理のCUDAライブラリの呼び出し等、その言語からのオフロード機能利用に合わせた処理に置換する必要がある。 When searching for offloadable function blocks, searches are performed by matching the names of libraries, etc., and by detecting similarities using similarity detection tools for Java function blocks such as Deckard (registered trademark). Replacement with an offload function block requires replacement with processing that matches the use of offload functions from that language, such as calling the CUDA library for GPU processing.
実行環境は、Javaのラムダ記述での処理をGPUに対して行うことができるIBM JDK(登録商標)を用いる。性能測定は、言語に合わせて、Jenkins(登録商標)等の自動測定ツールも用いて行う。オフロード可能機能ブロックが複数の際は反復実行され、最高性能のパターンが最終解として決定される。 The execution environment uses IBM JDK (registered trademark), which allows processing using Java lambda statements on the GPU. Performance measurements are also performed using automated measurement tools such as Jenkins (registered trademark) depending on the language. When there are multiple offloadable function blocks, they are executed repeatedly, and the pattern with the best performance is determined as the final solution.
[《機能ブロックオフロード:Java》フローチャート]
次に、図20および図21を参照してオフロードサーバ200の《機能ブロック:Java》の動作概要を説明する。
["Function Block Offload: Java" Flowchart]
Next, an outline of the operation of the <<functional block: Java>> of the
・<処理A-1>と<処理B-1>と<処理C-1>のフローチャート
図20は、オフロードサーバ200の制御部(自動オフロード機能部)210が、《機能ブロック:Java》のオフロード処理において<処理A-1>と<処理B-1>と<処理C-1>とを実行する場合のフローチャートである。
ステップS1101でアプリケーションコード分析部112(図12参照)は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
Flowchart of <Processing A-1>, <Processing B-1>, and <Processing C-1> FIG. 20 is a flowchart when the control unit (automatic offload function unit) 210 of the
In step S1101, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application program to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code to analyze the library calls and functional processing included in the code, together with the loop statement structure and the like.
ステップS1102で置換機能検出部213(図12参照)は、アプリケーションプログラムの外部ライブラリ呼び出しを検出する。 In step S1102, the replacement function detection unit 213 (see Figure 12) detects an external library call of the application program.
ステップS1103で置換機能検出部213は、コードパターンDB230から、ライブラリ名をキーに、置換可能GPUライブラリを取得する。具体的には、置換機能検出部213は、把握した外部ライブラリ呼び出しについて、コードパターンDB230と照合することで、検出した置換可能GPUライブラリを、GPU、FPGAにオフロードできるオフロード可能な機能ブロックとして取得する。In step S1103, the replacement function detection unit 213 obtains a replaceable GPU library from the
ステップS1104で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリ、IPコアの処理記述に置換する。In step S1104, the
ステップS1105で置換処理部214は、置換したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。
In step S1105, the
ステップS1106で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S1106, the
ステップS1107で置換処理部214は、CUDAのライブラリ呼び出しを、Jcudaで指定する。
In step S1107, the
ステップS1108で実行ファイル作成部117は、作成したパターンをJcudaでビルドする。
In step S1108, the executable
ステップS1109で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。
ステップS1110で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。
In step S1109, the
In step S1110, the executable
ステップS1111で実行ファイル作成部117は、作成した組合せパターンをJcudaでビルドする。
ステップS1112で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。
In step S1111, the executable
In step S1112, the
ステップS1113で本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S1113, the production
・<処理A-2>と<処理B-2>と<処理C-2>のフローチャート
図21は、オフロードサーバ200の制御部(自動オフロード機能部)210が、機能ブロックのオフロード処理において<処理A-2>と<処理B-2>と<処理C-2>とを実行する場合のフローチャートである。なお、<処理A-2>からの処理は、<処理A-1>からの処理と並行して行えばよい。
ステップS1201でアプリケーションコード分析部112(図12参照)は、アプリケーションのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部112は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。
21 is a flowchart of the case where the control unit (automatic offload function unit) 210 of the
In step S1201, the application code analysis unit 112 (see FIG. 12) analyzes the source code of the application to be offloaded. Specifically, the application code analysis unit 112 uses a syntax analysis tool such as Clang to analyze the source code to analyze the library calls and functional processing included in the code, together with the loop statement structure, etc.
ステップS1202で置換機能検出部213(図12参照)は、ソースコードからクラスまたは構造体の定義記述コードを検出する。In step S1202, the replacement function detection unit 213 (see Figure 12) detects definition description code of a class or structure from the source code.
ステップS1203で置換機能検出部213は、コードパターンDB230から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能GPUライブラリを取得する。In step S1203, the replacement function detection unit 213 uses a similarity detection tool to obtain a replaceable GPU library from the
ステップS1204で置換処理部214は、アプリケーションのソースコードの置換元の処理記述を、置換先のGPUライブラリ処理記述に置換する。In step S1204, the
ステップS1205で置換処理部214は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。
In step S1205, the
ステップS1206で置換機能検出部213は、置換または確認したGPUライブラリの処理記述を、オフロード対象の機能ブロックとして、GPUにオフロードする。 In step S1206, the replacement function detection unit 213 offloads the processing description of the replaced or confirmed GPU library to the GPU as a functional block to be offloaded.
ステップS1207で置換処理部214は、GPUライブラリ呼び出しのためのインタフェース処理を記述する。
In step S1207, the
ステップS1108で置換処理部214は、CUDAのライブラリ呼び出しを、Jcudaで指定する。
In step S1108, the
ステップS1209で実行ファイル作成部117は、作成したパターンをJcudaでビルドする。
In step S1209, the executable
ステップS1210で性能測定部116は、作成したパターンを検証環境で性能測定する(「1回目の性能測定」)。In step S1210, the
ステップS1211で実行ファイル作成部117は、1回目測定時に高速化できたパターンについて組合せパターンを作成する。In step S1211, the executable
ステップS1212で実行ファイル作成部117は、作成した組合せパターンをJcudaでビルドする。
In step S1212, the executable
ステップS1213で性能測定部116は、作成した組合せパターンを検証環境で性能測定する(「2回目の性能測定」)。In step S1213, the
ステップS1214で本番環境配置部118は、本番環境配置部118は、1回目と2回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。
In step S1214, the production
[実装例]
第1の実施形態(「ループ文オフロード」)および第2の実施形態(「機能ブロックオフロード」)の実装例を説明する。
[Implementation example]
An implementation example of the first embodiment ("loop statement offload") and the second embodiment ("function block offload") will be described.
<利用ツール>
対象アプリケーションプログラムはC/C++言語、Python、Javaのアプリケーションとする。
GPU処理は、C/C++言語はPGIコンパイラ19.10を用いる。PGIコンパイラは、OpenACCを解釈するC/C++向けコンパイラである。PGIコンパイラは、cuFFT等のCUDAライブラリの呼び出しも処理が可能である。
<Tools>
The target application programs are C/C++, Python, and Java applications.
For GPU processing, PGI Compiler 19.10 is used for C/C++ language. The PGI Compiler is a C/C++ compiler that interprets OpenACC. The PGI Compiler can also process calls to CUDA libraries such as cuFFT.
Pythonは、pyCUDA 2019.1.2を用いる。pyCUDAは、PythonからGPUに処理実行するためのインタプリタである。あるいは、Pythonには、PyACCを用いる。PyACCは、PythonからOpenACCを解釈実行するためのインタプリタである。 For Python, pyCUDA 2019.1.2 is used. pyCUDA is an interpreter for executing processing on the GPU from Python. Alternatively, for Python, PyACC is used. PyACC is an interpreter for interpreting and executing OpenACC from Python.
Javaは、IBM JDK(登録商標)を用いる。IBM JDKはJavaのラムダ記述に従って並列処理をGPUに対して実行する仮想マシンである。 For Java, IBM JDK (registered trademark) is used. IBM JDK is a virtual machine that executes parallel processing on the GPU according to Java lambda notation.
<構文解析>
C/C++言語の構文解析には、LLVM/Clang 6.0の構文解析ライブラリ(libClang(登録商標)のpython binding) を用いる。Pythonの構文解析には、astを用いる。Javaの構文解析には、Java Parser(登録商標)を用いる。
<Syntax analysis>
For C/C++ language syntax analysis, we use the LLVM/Clang 6.0 syntax analysis library (libClang (registered trademark) python binding). For Python syntax analysis, we use ast. For Java syntax analysis, we use Java Parser (registered trademark).
<類似性検出ツール>
類似性検出ツールには、C/C++言語、Javaには、Deckard v2.0(登録商標)を用いる。Deckardは、機能ブロックのオフロードの適用領域拡大のため、ライブラリ呼び出し以外にも、コードコピーし変更した機能等のオフロードを実現するため、照合対象となる部分コードと、DBに登録されたコードの類似性を判定する。Pythonには、CloneDigger(登録商標)を用いる。
<Similarity detection tool>
The similarity detection tool is Deckard v2.0 (registered trademark) for C/C++ and Java. Deckard determines the similarity between the partial code to be compared and the code registered in the DB in order to expand the scope of application of function block offloading, not only for library calls but also for offloading functions that have been copied and modified. CloneDigger (registered trademark) is used for Python.
<コードパターンDB>
照合に用いるコードパターンDB230(図12参照)は、MySQL8を用いる。呼び出しているライブラリ名をキーに、高速化できるライブラリ等を検索するためのレコードを保持する。ライブラリには、それに紐づく名前やコードや実行ファイルが保持される。実行ファイルはその利用手法等も登録されている。コードパターンDB230には、ライブラリ等を類似性検出技術で検出するための、比較用コードとの対応関係も保持される。
<Code Pattern DB>
The code pattern DB 230 (see FIG. 12) used for matching uses MySQL 8. It holds records for searching for libraries etc. that can increase speed using the called library name as a key. The library holds the name, code and executable file associated with it. For the executable file, the method of use etc. are also registered. The
<実装動作>
実装の動作概要について述べる。
実装は、アプリケーションの利用依頼があると、構文解析ライブラリを用いてコード解析を行う。次に、機能ブロックオフロード、ループ文オフロードの順に試行を行う。これは、ループ文と機能ブロックに関しては、アルゴリズム含めて処理内容に合わせてオフロードする機能ブロックオフロードの方が高速化できるからである。機能ブロックオフロードが可能であった場合は、後半のループ文オフロードはオフロード可能であった機能ブロック部分を抜いたコードに対して試行する。
性能測定の結果、最高性能のパターンを解とする。
<Implementation Operation>
An overview of the implementation will be given below.
When a request to use an application is received, the implementation uses a syntax analysis library to analyze the code. Next, function block offloading and loop statement offloading are attempted in that order. This is because, for loop statements and function blocks, function block offloading, which offloads according to the processing content including the algorithm, can achieve higher speeds. If function block offloading is possible, the latter half of the loop statement offloading is attempted on the code excluding the function block parts that were offloadable.
As a result of the performance measurement, the pattern with the best performance is determined as the solution.
以上、ライブラリ呼び出しの場合について記載した。
置換機能検出部213(図12参照)が、類似性検出ツールを用いて類似性検出を行う場合について説明する。類似性検出を行う場合には、上記置換記述と並行して処理がされる。すなわち、置換機能検出部213が、類似性検出を行う場合、実装例では、<処理B-2>でDeckardを用いて、検出されたクラス、構造体等の部分コードとコードパターンDB230に登録された比較用コードとの類似性検出を行う。そして、置換機能検出部213は、閾値超えの機能ブロックと該当するGPU用ライブラリやFPGA用IPコアを検出する。置換機能検出部213は、<処理B-1>の場合と同様に、実行ファイルやOpenCLを取得する。実装例では、次にC-1の場合と同様に実行用ファイルを作成するが、特に置換元のコードと置換するライブラリやIPコアの引数や戻り値、型等のインタフェースが異なる場合は、オフロードを依頼したユーザに対して、置換先ライブラリやIPコアに合わせて、インタフェースを変更してよいか確認し、確認後に実行用ファイルを作成する。
The above describes the case of library calls.
A case where the replacement function detection unit 213 (see FIG. 12) performs similarity detection using a similarity detection tool will be described. When similarity detection is performed, processing is performed in parallel with the above replacement description. That is, when the replacement function detection unit 213 performs similarity detection, in the implementation example, Deckard is used in <Process B-2> to detect similarity between partial code such as a detected class or structure and a comparison code registered in the
この時点で、検証環境のGPUやFPGAで性能測定できる実行用ファイルが作成される。機能ブロックオフロードについては、置換する機能ブロックが一つの場合は、その一つをオフロードするかしないかだけである。複数ある場合は、一つずつオフロードする/しないを検証パターンとして作成し、性能を測定し高速な解を検出する。これは、高速化可能とされていても実測してみないとその条件で高速になるかわからないためである。例えば、5つオフロード可能な機能ブロックがあり、1回目測定の結果、2番と4番のオフロードが高速化できた場合は、2番と4番両方をオフロードするパターンで2回目測定を行い、2番と4番単独でオフロードする場合より高速となっている場合は、解として選択する。
At this point, an executable file is created that can measure performance on the GPU or FPGA in the verification environment. Regarding function block offloading, if there is only one function block to be replaced, it is only a matter of whether or not to offload that one block. If there are multiple blocks, verification patterns are created for offloading/not offloading each block one by one, and performance is measured to find a fast solution. This is because even if it is said that speed can be increased, it is not known whether it will be faster under those conditions without actually measuring it. For example, if there are five function blocks that can be offloaded, and the first measurement shows that offloading
[ハードウェア構成]
第1および第2の実施形態に係るオフロードサーバは、例えば図22に示すような構成の物理装置であるコンピュータ900によって実現される。
図22は、オフロードサーバ1,200の機能を実現するコンピュータの一例を示すハードウェア構成図である。コンピュータ900は、CPU(Central Processing Unit)901、ROM(Read Only Memory)902、RAM903、HDD(Hard Disk Drive)904、入出力I/F(Interface)905、通信I/F906およびメディアI/F907を有する。
[Hardware configuration]
The offload server according to the first and second embodiments is realized by a
22 is a hardware configuration diagram showing an example of a computer that realizes the functions of the
CPU901は、ROM902またはHDD904に記憶されたプログラムに基づき作動し、図1、図12に示すオフロードサーバ1,200の各処理部による制御を行う。ROM902は、コンピュータ900の起動時にCPU901により実行されるブートプログラムや、コンピュータ900のハードウェアに係るプログラム等を記憶する。The
CPU901は、入出力I/F905を介して、マウスやキーボード等の入力装置910、および、ディスプレイ等の出力装置911を制御する。CPU901は、入出力I/F905を介して、入力装置910からデータを取得するともに、生成したデータを出力装置911へ出力する。The
HDD904は、CPU901により実行されるプログラムおよび当該プログラムによって使用されるデータ等を記憶する。通信I/F906は、通信網(例えば、NW(Network)920)を介して他の装置からデータを受信してCPU901へ出力し、また、CPU901が生成したデータを、通信網を介して他の装置へ送信する。The
メディアI/F907は、記録媒体912に格納されたプログラムまたはデータを読み取り、RAM903を介してCPU901へ出力する。CPU901は、目的の処理に係るプログラムを、メディアI/F907を介して記録媒体912からRAM903上にロードし、ロードしたプログラムを実行する。記録媒体912は、DVD(Digital Versatile Disc)、PD(Phase change rewritable Disk)等の光学記録媒体、MO(Magneto Optical disk)等の光磁気記録媒体、磁気記録媒体、導体メモリテープ媒体又は半導体メモリ等である。The media I/
例えば、コンピュータ900が第1および第2の実施形態に係るオフロードサーバ1,200として機能する場合、コンピュータ900のCPU901は、RAM903上にロードされたプログラムを実行することによりオフロードサーバ1,200の機能を実現する。また、HDD904には、RAM903内のデータが記憶される。CPU901は、目的の処理に係るプログラムを記録媒体912から読み取って実行する。この他、CPU901は、他の装置から通信網(NW920)を介して目的の処理に係るプログラムを読み込んでもよい。For example, when the
[効果]
以下、本発明に係るオフロードサーバ等の効果について説明する。
第1の実施形態に係るオフロードサーバ1(図1参照)は、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、アプリケーションプログラムは、C言語、Python、およびJavaより選択される少なくとも一つであり、オフロードサーバ1は、アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部112と、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部113と、アプリケーションプログラムのループ文を特定し、特定した各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部114と、コンパイルエラーが出る繰り返し文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部115と、を備える。また、並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン14に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部116と、性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、を備える。
[effect]
The effects of the offload server and the like according to the present invention will be described below.
An offload server 1 (see FIG. 1 ) according to the first embodiment is an offload server that offloads specific processing of an application program to an accelerator, and the application program is at least one selected from C language, Python, and Java. The
このようにすることにより、オフロードサーバ1は、移行元言語がC言語、Python、Javaを含む多様な言語の場合でも共通的な方式でGPUに自動オフロードすることができる。これにより、移行元言語に合わせて、処理を検討したり実装する必要がなくなり、コストダウンを図ることができる。
In this way, the
さらに、オフロードサーバ1は、移行元言語がC言語、Python、またはJavaのいずれの場合であっても、CPU-GPU間のデータ転送回数を低減しつつ、アプリケーションプログラムの特定処理をアクセラレータに自動でオフロードすることで、全体の処理能力を向上させることができる。これにより、CUDA等のスキルが無いユーザでもGPUを使い高性能処理ができる。また、従来GPUでの高性能化が検討されていない汎用的なCPU向けアプリケーションを高性能化できる。また、高性能計算用サーバでない汎用的マシンのGPUにオフロードすることができる。
Furthermore, the
第2の実施形態に係るオフロードサーバ200(図12参照)は、アプリケーションプログラムの特定処理をGPUまたはPLDにオフロードするオフロードサーバであって、アプリケーションプログラムは、C言語、Python、およびJavaより選択される少なくとも一つであり、オフロードサーバ200は、GPUまたはPLDにオフロード可能なライブラリおよびIPコアを記憶するコードパターンDB230と、アプリケーションプログラムのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するアプリケーションコード分析部112と、検出された外部ライブラリ呼び出しをキーにして、コードパターンDB230からライブラリおよびIPコアを取得する置換機能検出部213と、アプリケーションプログラムのソースコードの置換元の処理記述を、置換機能検出部213が取得した置換先のライブラリおよびIPコアの置換先の処理記述として置換するとともに、置換したライブラリおよびIPコアの処理記述を、オフロード対象の機能ブロックとして、GPUまたはPLDにオフロードする置換処理部214と、ホストプログラムとのインタフェースを作成するオフロードパターン作成部215と、作成されたGPUまたはPLD処理パターンの前記アプリケーションをコンパイルして、実行ファイルを作成する実行ファイル作成部117と、作成された実行ファイルをアクセラレータ検証用装置に配置し、GPUまたはPLDにオフロードした際の性能測定用処理を実行する性能測定部116と、を備え、実行ファイル作成部117は、性能測定用処理による性能測定結果をもとに、複数のGPUまたはPLD処理パターンから最高処理性能のGPUまたはPLD処理パターンを選択し、最高処理性能のGPUまたはPLD処理パターンをコンパイルして、最終実行ファイルを作成する。The
このようにすることにより、オフロードサーバ200は、移行元言語がC言語、Python、Javaを含む多様な言語の場合でも共通的な方式でGPUに自動オフロードすることができる。これにより、移行元言語に合わせて、処理を検討したり実装する必要をなり、コストダウンを図ることができる。
In this way, the
さらに、オフロードサーバ200は、移行元言語がC言語、Python、またはJavaのいずれの場合であっても、アプリケーションコードの置換元の処理記述を、置換先のライブラリおよびIPコア処理記述に置換して、オフロード可能な機能ブロックとして、GPUやPLD(FPGA等)にオフロードする。すなわち、個々のループ文でなく、行列積算やフーリエ変換等のより大きな単位で、FPGAやGPU等ハードウェア向けのアルゴリズム含めて実装された機能ブロックをオフロードする。これにより、GPUやPLD(FPGA等)への自動オフロードにおいて、機能ブロックの単位でオフロードすることで、オフロード処理の高速化を図ることができる。その結果、GPU、FPGA、IoTデバイス等環境が多様になる中で、アプリケーションを環境に合わせて適応させることが可能になり、高性能にアプリケーションを動作させることができる。
Furthermore, the
第1および第2の実施形態に係るオフロードサーバ1,200において、アプリケーションプログラムが、C言語の場合、ループ文のGPU処理をOpenACC文法で指定し、機能ブロックのGPU処理をCUDAライブラリ呼び出すようにしてC言語コンパイラを用いてGPUオフロードすることを特徴とする。In the offload server 1,200 according to the first and second embodiments, when the application program is written in C language, the GPU processing of the loop statement is specified in OpenACC grammar, and the GPU processing of the functional block is offloaded by GPU offloading using a C language compiler by calling the CUDA library.
このようにすることにより、移行元言語がC言語の場合に共通的な方式で、GPUに自動オフロードすることができる。 By doing this, it is possible to automatically offload to the GPU using a common method when the source language is C.
第1および第2の実施形態に係るオフロードサーバ1,200において、アプリケーションプログラムが、Pythonの場合、ループ文のGPU処理をCUDA文法で指定し、機能ブロックのGPU処理をCUDAライブラリ呼び出すようにしてpyCUDAを用いてGPUオフロードすることを特徴とする。In the offload server 1,200 according to the first and second embodiments, when the application program is Python, the GPU processing of the loop statement is specified in CUDA grammar, and the GPU processing of the functional block is offloaded using pyCUDA by calling the CUDA library.
このようにすることにより、移行元言語がPythonの場合に共通的な方式で、GPUに自動オフロードすることができる。 By doing this, it is possible to automatically offload to the GPU using a common method when the source language is Python.
第1および第2の実施形態に係るオフロードサーバ1,200において、アプリケーションプログラムが、Pythonの場合、ループ文のGPU処理をOpenACC文法で指定し、機能ブロックのGPU処理をCUDAライブラリ呼び出すようにしてpyACCを用いてGPUオフロードすることを特徴とする。In the offload server 1,200 according to the first and second embodiments, when the application program is Python, the GPU processing of the loop statement is specified in OpenACC grammar, and the GPU processing of the functional block is offloaded using pyACC by calling the CUDA library.
このようにすることにより、移行元言語がPythonの場合にC言語の場合と同様に、共通的な方式でGPUに自動オフロードすることができる。 By doing this, when the source language is Python, it is possible to automatically offload to the GPU using a common method, just as in the case of C.
第1および第2の実施形態に係るオフロードサーバ1,200において、アプリケーションプログラムが、Javaの場合、ループ文のGPU処理をJavaのラムダ文法で指定し、機能ブロックのGPU処理をCUDAライブラリ呼び出すようにしてJava仮想マシンを用いてGPUオフロードすることを特徴とする。In the offload server 1,200 according to the first and second embodiments, when the application program is Java, the GPU processing of the loop statement is specified using Java lambda grammar, and the GPU processing of the functional block is called by the CUDA library, thereby performing GPU offloading using a Java virtual machine.
このようにすることにより、移行元言語がJavaの場合に共通的な方式でGPUに自動オフロードすることができる。 By doing this, when the source language is Java, it is possible to automatically offload to the GPU in a common manner.
本発明は、コンピュータを、上記オフロードサーバとして機能させるためのオフロードプログラムとした。 The present invention is an offload program for causing a computer to function as the above-mentioned offload server.
このようにすることにより、一般的なコンピュータを用いて、上記オフロードサーバ200の各機能を実現させることができる。
By doing this, it is possible to realize each function of the
[変形例]
第1の実施形態に係るオフロードサーバ1と第2の実施形態に係るオフロードサーバ200とを組み合わせて、データ転送指定部113(図1参照)は、アプリケーションコード分析部112が行うコード解析結果をもとに、機能ブロックオフロード、ループ文オフロードの順に試行するようにデータ転送を行うととともに、機能ブロックオフロードが可能であった場合は、オフロード可能であった機能ブロック部分を抜いたコードに対して、ループ文オフロードを試行するデータ転送を行う構成としてもよい。
[Modification]
By combining the
この構成により、オフロードサーバは、まず、アプリケーションプログラムの利用依頼があると、構文解析ライブラリを用いてコード解析を行い、次に、機能ブロックオフロード、ループ文オフロードの順に試行を行う。機能ブロックオフロードが可能であった場合は、オフロード可能であった機能ブロック部分を抜いたコードに対して、ループ文オフロードを試行し、性能測定の結果、最高性能のパターンを解とする。これにより、ループ文と機能ブロックに関しては、アルゴリズム含めて処理内容に合わせてオフロードする機能ブロックオフロードの方が高速化できる。機能ブロックオフロード、ループ文オフロードの順に試行を行うことで、処理の高速化を図ることができ、全体の処理能力を向上させることができる。 With this configuration, when an offload server receives a request to use an application program, it first performs code analysis using the syntax analysis library, and then attempts to offload function blocks and then loop statements. If function block offloading is possible, it attempts to offload loop statements on the code excluding the function block portions that were offloadable, and the pattern with the highest performance as a result of performance measurements is determined to be the solution. As a result, with regard to loop statements and function blocks, function block offloading, which offloads according to the processing content including algorithms, can achieve faster processing. By attempting function block offloading and then loop statement offloading, it is possible to speed up processing and improve overall processing capacity.
また、上記各実施形態において説明した各処理のうち、自動的に行われるものとして説明した処理の全部又は一部を手作業で行うこともでき、あるいは、手作業で行われるものとして説明した処理の全部又は一部を公知の方法で自動的に行うこともできる。この他、上述文書中や図面中に示した処理手順、制御手順、具体的名称、各種のデータやパラメータを含む情報については、特記する場合を除いて任意に変更することができる。
また、図示した各装置の各構成要素は機能概念的なものであり、必ずしも物理的に図示の如く構成されていることを要しない。すなわち、各装置の分散・統合の具体的形態は図示のものに限られず、その全部又は一部を、各種の負荷や使用状況などに応じて、任意の単位で機能的又は物理的に分散・統合して構成することができる。
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 depending on various loads, usage conditions, etc.
1,200 オフロードサーバ
11,210 制御部
12 入出力部
13,130 記憶部
14 検証用マシン (アクセラレータ検証用装置)
15 商用環境
111 アプリケーションコード指定部
112 アプリケーションコード分析部
113 データ転送指定部
114 並列処理指定部
114a オフロード範囲抽出部
114b 中間言語ファイル出力部
115 並列処理パターン作成部
116 性能測定部
116a バイナリファイル配置部
117 実行ファイル作成部
118 本番環境配置部
119 性能測定テスト抽出実行部
120 ユーザ提供部
125 アプリケーションコード
131 テストケースDB
132 中間言語ファイル
151 各種デバイス
152 CPU-GPUを有する装置
153 CPU-FPGAを有する装置
154 CPUを有する装置
213 置換機能検出部
214 置換処理部
215 オフロードパターン作成部
230 コードパターンDB
1,200 Offload server 11,210
15 Commercial environment 111 Application code specification unit 112 Application
132
Claims (8)
前記アプリケーションプログラムは、Pythonアプリケーションプログラムであり、
前記Pythonアプリケーションプログラムのソースコードを、Pythonを解析する構文解析ツールを用いて分析するアプリケーションコード分析部と、
前記Pythonアプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、
前記Pythonアプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、CUDAでの指示を追加したPythonコードをpyCUDAでインタプリットする際、CUDA文法を用いたGPU処理を指定してインタプリットする並列処理指定部と、
コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、
前記並列処理パターンの前記Pythonアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、
性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するPythonアプリコードを含む最高性能の並列処理パターンを解とし、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、
を備えることを特徴とするオフロードサーバ。 An offload server that offloads a specific process of an application program to an accelerator,
the application program is a Python application program;
an application code analysis unit that analyzes a source code of the Python application program using a syntax analysis tool that analyzes Python;
a data transfer specification unit that analyzes reference relationships between variables used in a loop statement of the Python application program, and specifies data transfer using an explicit specification line that explicitly specifies data transfer outside the loop for data that may be transferred outside the loop;
a parallel processing specification unit that specifies a loop statement of the Python application program, and when interpreting Python code with CUDA instructions added for each of the specified loop statements using pyCUDA, specifies GPU processing using CUDA grammar and performs interpretation;
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 Python application program of the parallel processing pattern, places it in an accelerator verification device, and executes a process for performance measurement when offloaded to the accelerator;
an executable file creation unit which selects a plurality of parallel processing patterns with high processing performance from the plurality of parallel processing patterns based on a performance measurement result, crosses the parallel processing patterns with high processing performance, creates a plurality of other parallel processing patterns by mutation processing, and performs new performance measurement; after a designated number of performance measurements, selects a parallel processing pattern with the highest processing performance from the plurality of parallel processing patterns based on the performance measurement result; after GA processing for a designated number of generations is completed, sets the parallel processing pattern with the highest processing performance including a Python application code corresponding to the gene sequence with the highest performance as a solution; and compiles the parallel processing pattern with the highest processing performance to create an executable file;
An offload server comprising:
前記アプリケーションプログラムは、Pythonアプリケーションプログラムであり、
前記Pythonアプリケーションプログラムのソースコードを、Pythonを解析する構文解析ツールを用いて分析するアプリケーションコード分析部と、
前記Pythonアプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、Cupyがデータ転送指定を行うデータ転送指定部と、
PythonのコードでGPU処理を指定されるループ文について、前記Cupyがライブラリを介して、CUDAコマンドを呼び、CUDAがGPU処理する並列処理指定部と、
コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、
前記並列処理パターンの前記Pythonアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、
性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するPythonアプリコードを含む最高性能の並列処理パターンを解とし、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、を備え、
前記並列処理指定部は、CUDAを介してGPU処理する際、for文の配列を右辺と左辺に持つ演算式を、範囲開始と範囲終了の表現に書き換えて、式全体を行列演算の式に変換して指定する
ことを特徴とするオフロードサーバ。 An offload server that offloads a specific process of an application program to an accelerator including a GPU (Graphics Processing Unit),
the application program is a Python application program;
an application code analysis unit that analyzes a source code of the Python application program using a syntax analysis tool that analyzes Python;
A data transfer specification unit that analyzes the reference relationship of variables used in a loop statement of the Python application program and specifies data transfer for data that may be transferred outside the loop;
A parallel processing specification section in which, for a loop statement in which GPU processing is specified in Python code, the Cupy calls a CUDA command via a library and CUDA performs GPU processing;
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 Python application program of the parallel processing pattern, places it in an accelerator verification device, and executes a process for performance measurement when offloaded to the accelerator;
an executable file creation unit which selects a plurality of parallel processing patterns with high processing performance from the plurality of parallel processing patterns based on a performance measurement result, crosses the parallel processing patterns with high processing performance, creates a plurality of other parallel processing patterns by mutation processing, and performs new performance measurement, and after a designated number of performance measurements, selects a parallel processing pattern with the highest processing performance from the plurality of parallel processing patterns based on the performance measurement result, and after GA processing for a designated number of generations is completed, sets the parallel processing pattern with the highest processing performance including a Python application code corresponding to the gene sequence with the highest performance as a solution, and compiles the parallel processing pattern with the highest processing performance to create an executable file,
The parallel processing designation unit, when performing GPU processing via CUDA, rewrites an arithmetic expression having arrays of a for statement on the right and left sides into an expression of a range start and a range end, and converts the entire expression into a matrix operation expression and specifies it.
前記アプリケーションプログラムは、Pythonアプリケーションプログラムであり、
前記Pythonアプリケーションプログラムのソースコードを、Pythonを解析する構文解析ツールを用いて分析するアプリケーションコード分析部と、
前記Pythonアプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、
前記Pythonアプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、CUDAでの指示を追加したPythonコードをOpenACCを解釈するpyACCでインタプリットする際、前記OpenACCの\pragmaacckernelsを指定する並列処理指定部と、
コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、
前記並列処理パターンの前記Pythonアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、
性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するPythonアプリコードを含む最高性能の並列処理パターンを解とし、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、
を備えることを特徴とするオフロードサーバ。 An offload server that offloads a specific process of an application program to an accelerator,
the application program is a Python application program;
an application code analysis unit that analyzes a source code of the Python application program using a syntax analysis tool that analyzes Python;
a data transfer specification unit that analyzes reference relationships between variables used in a loop statement of the Python application program, and specifies data transfer using an explicit specification line that explicitly specifies data transfer outside the loop for data that may be transferred outside the loop;
A parallel processing specification unit that specifies \pragmaacckernels of OpenACC when identifying loop statements of the Python application program and interpreting Python code to which a CUDA instruction has been added for each of the identified loop statements using pyACC, which interprets OpenACC ;
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 Python application program of the parallel processing pattern, places it in an accelerator verification device, and executes a process for performance measurement when offloaded to the accelerator;
an executable file creation unit which selects a plurality of parallel processing patterns with high processing performance from the plurality of parallel processing patterns based on a performance measurement result, crosses the parallel processing patterns with high processing performance, creates a plurality of other parallel processing patterns by mutation processing, and performs new performance measurement; after a designated number of performance measurements, selects a parallel processing pattern with the highest processing performance from the plurality of parallel processing patterns based on the performance measurement result; after GA processing for a designated number of generations is completed, sets the parallel processing pattern with the highest processing performance including a Python application code corresponding to the gene sequence with the highest performance as a solution; and compiles the parallel processing pattern with the highest processing performance to create an executable file;
An offload server comprising:
前記アプリケーションプログラムは、Pythonアプリケーションプログラムであり、
前記GPUまたは前記PLDにオフロード可能なGPUライブラリおよびIPコアを記憶する記憶部と、
前記Pythonアプリケーションプログラムのソースコードを、Pythonを解析する構文解析ツールを用いて分析するアプリケーションコード分析部と、
検出された外部ライブラリ呼び出しをキーにして、前記記憶部から前記GPUライブラリおよび前記IPコアを取得する置換機能検出部と、
GPUのライブラリ呼び出しを、pyCUDAで指定し、前記Pythonアプリケーションプログラムのソースコードの置換元の処理記述を、前記置換機能検出部が取得した前記GPUライブラリおよび前記IPコアの置換先の処理記述として置換する際、CUDAでの指示を追加したPythonコードをpyCUDAでインタプリットして、GPU処理のpyCUDAでの呼び出しに合わせた処理に置換し、
置換した前記GPUライブラリおよび前記IPコアの処理記述を、オフロード対象の機能ブロックとして、前記GPUまたは前記PLDにオフロードする置換処理部と、
ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出するオフロードパターン作成部と、
作成されたGPUまたはPLD処理パターンの前記Pythonアプリケーションプログラムをコンパイルして、実行ファイルを作成する実行ファイル作成部と、
作成された前記実行ファイルをアクセラレータ検証用装置に配置し、前記GPUまたは前記PLDにオフロードした際の性能測定用処理を実行する性能測定部と、を備え、
前記実行ファイル作成部は、前記性能測定用処理による性能測定結果をもとに、複数の前記GPUまたはPLD処理パターンから最高処理性能の前記GPUまたはPLD処理パターンを選択し、最高処理性能の前記GPUまたはPLD処理パターンをコンパイルして、最終実行ファイルを作成する
ことを特徴とするオフロードサーバ。 An offload server that offloads a specific process of an application program to a graphics processing unit (GPU) or a programmable logic device (PLD),
the application program is a Python application program;
A storage unit that stores a GPU library and an IP core that can be offloaded to the GPU or the PLD;
an application code analysis unit that analyzes a source code of the Python application program using a syntax analysis tool that analyzes Python;
a replacement function detection unit that acquires the GPU library and the IP core from the storage unit using the detected external library call as a key;
When a GPU library call is specified by pyCUDA and a source process description of the source code of the Python application program is replaced with a destination process description of the GPU library and the IP core acquired by the replacement function detection unit, the Python code to which a CUDA instruction has been added is interpreted by pyCUDA and replaced with a process that matches the call of the GPU process by pyCUDA;
a replacement processing unit that offloads the replaced processing description of the GPU library and the IP core to the GPU or the PLD as a functional block to be offloaded;
an offload pattern creation unit that creates an interface with a host program, and extracts an offload pattern that will be faster by testing whether or not to offload through performance measurement in a verification environment;
an executable file creation unit that compiles the Python application program of the created GPU or PLD processing pattern to create an executable file;
a performance measurement unit that places the created executable file in an accelerator verification device and executes a process for measuring performance when the executable file is offloaded to the GPU or the PLD;
The offload server is characterized in that the executable file creation unit selects the GPU or PLD processing pattern with the highest processing performance from the multiple GPU or PLD processing patterns based on the performance measurement results from the performance measurement process, compiles the GPU or PLD processing pattern with the highest processing performance, and creates a final executable file.
ことを特徴とする請求項1または請求項4にオフロードサーバ。 5. The offload server according to claim 1 , wherein the performance measurement unit executes the process for measuring the performance when the offload is performed by using an automatic measurement tool including Jenkins.
ことを特徴とする請求項4にオフロードサーバ。 5. The offload server according to claim 4 , wherein said performance measurement unit executes the execution repeatedly when there are a plurality of offloadable function blocks, and determines the pattern with the highest performance as the final solution.
前記アプリケーションプログラムは、Pythonアプリケーションプログラムであり、
前記Pythonアプリケーションプログラムのソースコードを、Pythonを解析する構文解析ツールを用いて分析するステップと、
前記Pythonアプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うステップと、
前記Pythonアプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、CUDAでの指示を追加したPythonコードをpyCUDAでインタプリットする際、CUDA文法を用いたGPU処理を指定してインタプリットするステップと、
コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成するステップと、
前記並列処理パターンの前記Pythonアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行するステップと、
性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、指定世代数のGA処理終了後、最高性能の遺伝子配列に該当するPythonアプリコードを含む最高性能の並列処理パターンを解とし、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成するステップと、を実行する
ことを特徴とするオフロード制御方法。 An offload control method for an offload server that offloads a specific process of an application program to an accelerator, comprising:
the application program is a Python application program;
analyzing the source code of the Python application program with a parser that parses Python;
A step of analyzing reference relationships of variables used in a loop statement of the Python application program, and for data that may be transferred outside the loop, specifying data transfer using an explicit specification line that explicitly specifies data transfer outside the loop;
Identifying loop statements of the Python application program, and when interpreting Python code with CUDA instructions added for each of the identified loop statements using pyCUDA, interpreting the Python code by specifying GPU processing using CUDA grammar ;
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;
Compiling the Python application program of the parallel processing pattern, distributing the Python application program on an accelerator verification device, and executing a process for measuring performance when the Python application program is offloaded to the accelerator;
an offload control method comprising the steps of: selecting a plurality of parallel processing patterns with high processing performance from the plurality of parallel processing patterns based on performance measurement results; crossing the parallel processing patterns with high processing performance; creating another plurality of parallel processing patterns by mutation processing; and performing new performance measurements; after a designated number of performance measurements, selecting a parallel processing pattern with the highest processing performance from the plurality of parallel processing patterns based on the performance measurement results; after completion of GA processing for a designated number of generations, determining a parallel processing pattern with the highest processing performance that includes a Python application code that corresponds to a gene sequence with the highest performance as a solution; and compiling the parallel processing pattern with the highest processing performance to create an executable file.
Applications Claiming Priority (1)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
PCT/JP2020/041413 WO2022097245A1 (en) | 2020-11-05 | 2020-11-05 | Offload server, offload control method, and offload program |
Publications (2)
Publication Number | Publication Date |
---|---|
JPWO2022097245A1 JPWO2022097245A1 (en) | 2022-05-12 |
JP7521597B2 true JP7521597B2 (en) | 2024-07-24 |
Family
ID=81456976
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
JP2022560579A Active JP7521597B2 (en) | 2020-11-05 | 2020-11-05 | OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM |
Country Status (2)
Country | Link |
---|---|
JP (1) | JP7521597B2 (en) |
WO (1) | WO2022097245A1 (en) |
Families Citing this family (1)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
US20230221932A1 (en) * | 2022-01-12 | 2023-07-13 | Vmware, Inc. | Building a unified machine learning (ml)/ artificial intelligence (ai) acceleration framework across heterogeneous ai accelerators |
Citations (1)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
JP2020137017A (en) | 2019-02-22 | 2020-08-31 | 日本電信電話株式会社 | Optimum arrangement method for software of off-road server, and program |
-
2020
- 2020-11-05 WO PCT/JP2020/041413 patent/WO2022097245A1/en active Application Filing
- 2020-11-05 JP JP2022560579A patent/JP7521597B2/en active Active
Patent Citations (1)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
JP2020137017A (en) | 2019-02-22 | 2020-08-31 | 日本電信電話株式会社 | Optimum arrangement method for software of off-road server, and program |
Non-Patent Citations (5)
Title |
---|
ALEKSANDROV, Dmitry,Programming the GPU in Java,Java Magazine [online],2020年01月11日,[検索日 2023.12.27], インターネット: <URL: https://blogs.oracle.com/javamagazine/post/programming-the-gpu-in-java> |
inducer77,pycuda 2020.1,PyPI [online],2020年10月08日,[検索日 2023.12.27], インターネット: <URL: https://pypi.org/project/pycuda/2020.1/> |
Max Strange,pyACC,GitHub [online],2019年07月17日,[検索日 2023.12.27], インターネット: <URL: https://github.com/MaxStrange/pyACC/blob/master/README.md> |
samacoba,pythonで簡単にGPU計算ができるCupyを紹介,Qiita [online],2018年07月14日,1~9頁,[検索日 2024.03.22], インターネット:<URL: https://qiita.com/samacoba/items/d18e6cf09f544477aff4> |
山登 庸次,アプリケーション機能ブロックのGPU,FPGA自動オフロード手法の評価,電子情報通信学会技術研究報告 Vol.119 No.482,日本,一般社団法人電子情報通信学会,2020年03月09日,59~66頁 |
Also Published As
Publication number | Publication date |
---|---|
JPWO2022097245A1 (en) | 2022-05-12 |
WO2022097245A1 (en) | 2022-05-12 |
Similar Documents
Publication | Publication Date | Title |
---|---|---|
Abadi et al. | Tensorflow: Large-scale machine learning on heterogeneous distributed systems | |
JP7063289B2 (en) | Optimal software placement method and program for offload servers | |
CN103858099B (en) | The method and system applied for execution, the circuit with machine instruction | |
Pérez et al. | Simplifying programming and load balancing of data parallel applications on heterogeneous systems | |
Zhang et al. | Retiarii: A deep learning {Exploratory-Training} framework | |
JP5479942B2 (en) | Parallelization method, system, and program | |
CN112106023A (en) | Offload server and offload program | |
US12050894B2 (en) | Offload server, offload control method, and offload program | |
JP7521597B2 (en) | OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM | |
Yamato | Proposal and evaluation of GPU offloading parts reconfiguration during applications operations for environment adaptation | |
JP6992911B2 (en) | Offload server and offload program | |
JP7544142B2 (en) | OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM | |
Yamato | Study and evaluation of automatic offloading method in mixed offloading destination environment | |
JP7363930B2 (en) | Offload server, offload control method and offload program | |
JP7363931B2 (en) | Offload server, offload control method and offload program | |
JP7380823B2 (en) | Offload server, offload control method and offload program | |
WO2023002546A1 (en) | Offload server, offload control method, and offload program | |
JP7473003B2 (en) | OFFLOAD SERVER, OFFLOAD CONTROL METHOD, AND OFFLOAD PROGRAM | |
WO2024147197A1 (en) | Offload server, offload control method, and offload program | |
WO2023144926A1 (en) | Offload server, offload control method, and offload program | |
Fumero et al. | Using compiler snippets to exploit parallelism on heterogeneous hardware: a Java reduction case study | |
JP7184180B2 (en) | offload server and offload program | |
WO2023228369A1 (en) | Offload server, offload control method, and offload program | |
WO2024079886A1 (en) | Offload server, offload control method, and offload program | |
US20240192934A1 (en) | Framework for development and deployment of portable software over heterogenous compute systems |
Legal Events
Date | Code | Title | Description |
---|---|---|---|
A621 | Written request for application examination |
Free format text: JAPANESE INTERMEDIATE CODE: A621 Effective date: 20230302 |
|
A131 | Notification of reasons for refusal |
Free format text: JAPANESE INTERMEDIATE CODE: A131 Effective date: 20240109 |
|
A521 | Request for written amendment filed |
Free format text: JAPANESE INTERMEDIATE CODE: A523 Effective date: 20240306 |
|
A131 | Notification of reasons for refusal |
Free format text: JAPANESE INTERMEDIATE CODE: A131 Effective date: 20240402 |
|
A521 | Request for written amendment filed |
Free format text: JAPANESE INTERMEDIATE CODE: A523 Effective date: 20240529 |
|
TRDD | Decision of grant or rejection written | ||
A01 | Written decision to grant a patent or to grant a registration (utility model) |
Free format text: JAPANESE INTERMEDIATE CODE: A01 Effective date: 20240611 |
|
A61 | First payment of annual fees (during grant procedure) |
Free format text: JAPANESE INTERMEDIATE CODE: A61 Effective date: 20240624 |
|
R150 | Certificate of patent or registration of utility model |
Ref document number: 7521597 Country of ref document: JP Free format text: JAPANESE INTERMEDIATE CODE: R150 |