WO2022097245A1

WO2022097245A1 - オフロードサーバ、オフロード制御方法およびオフロードプログラム

Info

Publication number: WO2022097245A1
Application number: PCT/JP2020/041413
Authority: WO
Inventors: 庸次山登
Original assignee: 日本電信電話株式会社
Priority date: 2020-11-05
Filing date: 2020-11-05
Publication date: 2022-05-12
Also published as: JPWO2022097245A1

Abstract

オフロードサーバ（１）は、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、オフロードサーバ（１）は、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部（１１３）と、アプリケーションプログラムのループ文を特定し、特定した各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部（１１４）と、を備える。

Description

オフロードサーバ、オフロード制御方法およびオフロードプログラム

　本発明は、機能処理をＧＰＵ（Graphics Processing Unit）等のアクセラレータに自動オフロードするオフロードサーバ、オフロード制御方法およびオフロードプログラムに関する。

　近年、ＣＰＵの半導体集積度が1.5年で２倍になるというムーアの法則が減速するのではないかと言われている。そのような状況から、少コアのＣＰＵだけでなく、ＦＰＧＡ（Field Programmable Gate Array）やＧＰＵ（Graphics Processing Unit）等のデバイスの活用が増えている。例えば、Microsoft（登録商標）社はＦＰＧＡを使ってBingの検索効率を高めるといった取り組みをしており、Amazon（登録商標）社は、ＦＰＧＡ, ＧＰＵ等をクラウドのインスタンスとして提供している。

　少コアのＣＰＵ以外のデバイスをシステムで適切に活用するためには、デバイス特性を意識した設定やプログラム作成が必要であり、OpenMP（Open Multi-Processing）、OpenCL（Open Computing Language）、ＣＵＤＡ（Compute Unified Device Architecture）といった知識が必要になるため、大半のプログラマにとっては、スキルの壁が高い。

　少コアのＣＰＵ以外のＧＰＵやＦＰＧＡ、メニーコアＣＰＵ等のデバイスを活用するシステムは今後ますます増えていくと予想されるが、それらを最大限活用するには、技術的壁が高い。そこで、そのような壁を取り払い、少コアのＣＰＵ以外のデバイスを十分利用できるようにするため、プログラマが処理ロジックを記述したソフトウェアを、配置先の環境（ＦＰＧＡ、ＧＰＵ、メニーコアＣＰＵ等）にあわせて、適応的に変換、設定し、環境に適合した動作をさせるような、プラットフォームが求められている。

　非特許文献１には、一度記述したコードを、配置先の環境に存在するＧＰＵやＦＰＧＡ、メニーコアＣＰＵ等を利用できるように、変換、リソース設定等を自動で行い、アプリケーションを高性能に動作させることを目的とした、環境適応ソフトウェアが記載されている。

　非特許文献２、３、４には、環境適応ソフトウェアの要素として、アプリケーションコードのループ文および機能ブロックを、ＦＰＧＡ、ＧＰＵに自動オフロードする方式が記載されている。

　ＧＰＵの並列計算パワーを画像処理でないものにも使うＧＰＧＰＵ（General Purpose GPU）を行うための環境としてＣＵＤＡが普及している。ＣＵＤＡは、ＧＰＧＰＵ向けのNVIDIA（登録商標）社の環境である。また、ＦＰＧＡ、メニーコアＣＰＵ、ＧＰＵ等のヘテロなデバイスを同じように扱うための仕様としてOpenCLがあり、その開発環境も整いつつある。ＣＵＤＡ、OpenCLは、Ｃ言語の拡張を行いプログラムを行う形であり、プログラムの難度は高い（ＦＰＧＡ等のカーネルとＣＰＵのホストとの間のメモリデータのコピーや解放の記述を明示的に行う等）。

　また、ＣＵＤＡやOpenCLに比べて、より簡易にヘテロなデバイスを利用するための技術として、OpenACCやOpenMP等、コンパイラとしてＰＧＩコンパイラやgcc（登録商標）等がある。このコンパイラは、指示行ベースで、並列処理等を行う箇所を指定して、指示行に従って、ＧＰＵ、メニーコアＣＰＵ等に向けて実行ファイルを作成する。

　上記、ＣＵＤＡ、OpenCL、OpenACC、OpenMP等の技術仕様を用いることで、ＦＰＧＡやＧＰＵ、メニーコアＣＰＵへオフロードすることは可能になっている。しかしながら、デバイス処理自体は行えるようになっても、高速化することには課題がある。例えば、マルチコアＣＰＵ向けに自動並列化機能を持つコンパイラとして、Intelコンパイラ（登録商標）等がある。これらは、自動並列化時に、コードの中のループ文中で並列処理可能な部分を抽出して、並列化している。しかし、メモリ処理等の影響で単に並列化可能ループ文を並列化しても性能がでないことも多い。ＦＰＧＡやＧＰＵ等で高速化する際には、OpenCLやＣＵＤＡの技術者がチューニングを繰り返したり、OpenACCコンパイラ等を用いて適切な並列処理範囲を探索し試行することがされている。

　このため、技術スキルが乏しいプログラマが、ＦＰＧＡやＧＰＵ、メニーコアＣＰＵを活用してアプリケーションを高速化することは難しいし、自動並列化技術等を使う場合も並列処理箇所探索の試行錯誤等の稼働が必要だった。現状、ヘテロなデバイスに対するオフロードは手動での取組みが主流である。

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

　非特許文献１～４に記載の技術は、Ｃ言語プログラムのＧＰＵやＦＰＧＡへのオフロードであり、Python（登録商標）、Java（登録商標）等の多様な移行元言語は想定されていない。
　Ｃ言語だけでなく、Python、Javaと移行元言語が多様となった場合でも、アプリケーションプログラムを自動オフロードすることが要請されている。

　このような点に鑑みて本発明がなされたのであり、移行元言語に合わせて、処理を検討したり実装する必要をなくし、移行元言語が多様となった場合でも、アプリケーションプログラムを自動でオフロードすることを課題とする。

　前記した課題を解決するため、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、前記アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部と、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、を備えることを特徴とするオフロードサーバとした。

　本発明によれば、移行元言語に合わせて、処理を検討したり実装する必要をなくし、移行元言語が多様となった場合でも、アプリケーションプログラムを自動でオフロードすることができる。

本発明の第１の実施形態に係るオフロードサーバの構成例を示す機能ブロック図である。本発明の第１の実施形態に係るオフロードサーバのＧＡを用いた自動オフロード処理を示す図である。本発明の第１の実施形態に係るオフロードサーバのSimple GAによる制御部（自動オフロード機能部）の処理の探索イメージとfor文の遺伝子配列マッピングを示す図である。本発明の第１実施形態に係るオフロードサーバの自動オフロード機能部が処理するアプリケーションプログラムのソースコードのループ文（繰り返し文）を示す図である。本発明の第１の実施形態に係るオフロードサーバの《ループ文オフロード：共通》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバの《ループ文オフロード：共通》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバの《ループ文オフロード：Ｃ言語》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバの《ループ文オフロード：Ｃ言語》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバのPythonコードをpyCUDAでインタプリットする方法による《ループ文オフロード：Python》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバのPythonコードをpyCUDAでインタプリットする方法による《ループ文オフロード：Python》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバのpyACC利用時のfor文を示す図である。図８Ａのfor文から作成されるコードパターンを示す図である。本発明の第１の実施形態に係るオフロードサーバのpyACCを用いる方法による《ループ文オフロード：Python》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバのpyACCを用いる方法によるループ文オフロード：Python》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバのIBM JDK 利用時のfor文を示す図である。図８Ａのfor文から作成されるコードパターンを示す図である。本発明の第１の実施形態に係るオフロードサーバのpyACCを用いる方法による《ループ文オフロード：Java》のフローチャートである。本発明の第１の実施形態に係るオフロードサーバのpyACCを用いる方法による《ループ文オフロード：Java》のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの構成例を示す機能ブロック図である。本発明の第２の実施形態に係るオフロードサーバの機能ブロックのオフロード処理を示す図である。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、《機能ブロックオフロード：共通》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、《機能ブロックオフロード：Ｃ言語》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、《機能ブロック：Python》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、《機能ブロック：Java》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。本発明の第２の実施形態に係るオフロードサーバの制御部（自動オフロード機能部）が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。本発明の実施形態に係るオフロードサーバの機能を実現するコンピュータの一例を示すハードウェア構成図である。

　次に、本発明を実施するための形態における、オフロードサーバ等について説明する。
　以下、明細書の説明において、移行先環境としては、ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵの３つを想定した例について説明する。本発明は、プログラマブルロジックデバイス全般に適用可能である。

（多様移行元言語対応の基本的な考え方）
・移行元言語
　本実施形態で対象とする多様な移行元言語としては、Ｃ言語、Python、Javaの３つとする。これら３つの言語は、毎月TIOBE（登録商標）が発表するプログラム言語の人気ランキングの上位３つであり、プログラマ人口が多い。また、Ｃ言語はコンパイル型、Pythonはインタプリタ型、Javaはその中間的方式と、方式上の多様性も３つでカバーされている。そのため、これら３つで共通的に利用できる方式であれば、より多くの言語への対応も容易と考える。

　本実施形態では、移行先環境が単なるＣＰＵでない場合で、多様な移行元言語プログラムを、自動で高速にオフロードするために、検証環境の実機で性能測定し、進化計算手法等の手法と組み合わせて、徐々に高速なオフロードパターンを見つけるアプローチをとる。理由として、性能に関しては、コード構造だけでなく、処理するハードウェアのスペック、コンパイラやインタプリタ、データサイズ、ループ回数等の処理内容によって大きく変わるため、静的に予測する事が困難であり、動的な測定が必要だからである。実際に、市中には、ループ文を見つけコンパイル段階で並列化する自動並列化コンパイラがあるが、並列化可能ループ文の並列化だけでは性能を測定してみると低速になる場合も多いため、性能測定は必要である。

・オフロードする対象
　また、オフロードする対象については、アプリケーションプログラムのループ文および機能ブロックとするアプローチをとる。ループ文については、処理時間がかかるプログラムの処理の大半はループで費やされているという現状から、ループ文がオフロードのターゲットとして考えられる。一方、機能ブロックについては、特定処理を高速化する際に、処理内容や処理ハードウェアに適したアルゴリズムを用いることが多いため、個々のループ文の並列処理等に比べ、大きく高速化できる場合がある。行列積算やフーリエ変換等の頻繁に使われる機能ブロック単位で、ＧＰＵ等の処理デバイスに応じたアルゴリズムで実装された処理（ＣＵＤＡライブラリ等）に置換することで高速化する。

・移行先環境
　移行先環境としては、ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵの３つを想定し、これらが混在した環境でのＣ言語プログラムのオフロードも開示する。本発明の解決課題は、移行元言語が多様となった場合のアプリケーションの自動オフロードであるため、評価する移行先環境は限定されない。移行先環境は、一例としてＧＰＵとし、ＦＰＧＡやメニーコアＣＰＵについては、ＧＰＵで共通的方式を確認できれば、その拡張で実現できる。

・共通的なＧＰＵオフロード手法
　共通的なＧＰＵオフロード手法は、「ループ文のＧＰＵ自動オフロード（以下、ループ文オフロードという）」と「機能ブロックの自動オフロード（以下、機能ブロックオフロードという）」とに分けられ、それぞれ手法が異なる。
　以下の説明において、第１の実施形態で「ループ文オフロード」を記載し、第２の実施形態で「機能ブロックの自動オフロード」を記載する。そして、第１の実施形態（「ループ文のＧＰＵ自動オフロード」）と第２の実施形態（「機能ブロックオフロード」）のそれぞれにおいて、構成と、共通処理とＣ言語とPythonとJavaとを説明する。目次で示すと下記である。

（目次）
・第１の実施形態（「ループ文オフロード」）の構成（図１）
　共通（移行元言語において共通）処理（図２－図４）
　共通フローチャート（図５Ａ，図５Ｂ）
　Ｃ言語の場合のフローチャート（図６Ａ，図６Ｂ）
　Pythonの場合の説明図（図７Ａ，図７Ｂ，図８Ａ，図８Ｂ）
　Pythonの場合のフローチャート（図９Ａ，図９Ｂ）
　Javaの場合の説明図（図１０Ａ，図１０Ｂ）
　Javaの場合のフローチャート（図１１Ａ，図１１Ｂ）

・第２の実施形態（「機能ブロックオフロード」）の構成（図１２，図１３）
　共通フローチャート（図１４，図１５）
　Ｃ言語の場合のフローチャート（図１６，図１７）
　Pythonの場合のフローチャート（図１８，図１９）
　Javaの場合のフローチャート（図２０，図２１）

（第１の実施形態）
　第１の実施形態は、ループ文オフロードについて記載する。
　以下、第１の実施形態に係るオフロードサーバ１が、環境適応ソフトウェアシステムにおけるユーザ向けサービス利用のバックグラウンドで実行するオフロード処理を行う際の構成例について説明する。
　サービスを提供する際は、初日は試し利用等の形でユーザにサービス提供し、そのバックグラウンドで画像分析等のオフロード処理を行い、翌日以降は画像分析をＦＰＧＡにオフロードしてリーズナブルな価格で見守りサービスを提供できるようにすることを想定する。

　図１は、本発明の第１の実施形態に係るオフロードサーバ１の構成例を示す機能ブロック図である。
　オフロードサーバ１は、アプリケーションの特定処理をアクセラレータに自動的にオフロードする装置である。
　図１に示すように、オフロードサーバ１は、制御部１１と、入出力部１２と、記憶部１３と、検証用マシン１４（Verification machine）(アクセラレータ検証用装置)と、を含んで構成される。

　入出力部１２は、クラウドレイヤ、ネットワークレイヤおよびデバイスレイヤに属する各デバイス等との間で情報の送受信を行うための通信インタフェースと、タッチパネルやキーボード等の入力装置や、モニタ等の出力装置との間で情報の送受信を行うための入出力インタフェースとから構成される。

　記憶部１３は、ハードディスクやフラッシュメモリ、ＲＡＭ（Random Access Memory）等により構成される。
　この記憶部１３には、テストケースＤＢ（Test case database）１３１が記憶されるとともに、制御部１１の各機能を実行させるためのプログラム（オフロードプログラム）や、制御部１１の処理に必要な情報（例えば、中間言語ファイル(Intermediate file)１３２）が一時的に記憶される。

　テストケースＤＢ１３１は、検証対象ソフトに対応した試験項目のデータを格納する。試験項目のデータは、例えばMySQL等のデータベースシステムの場合、TPC-C等のトランザクション試験のデータである。

　制御部１１は、オフロードサーバ１全体の制御を司る自動オフロード機能部（Automatic Offloading function）である。制御部１１は、例えば、記憶部１３に格納されたアプリケーションプログラム（オフロードプログラム）を不図示のＣＰＵ（Central Processing Unit）が、ＲＡＭに展開し実行することにより実現される。

　アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つを含む。

　制御部１１は、アプリケーションコード指定部（Specify application code）１１１と、アプリケーションコード分析部（Analyze application code）１１２と、データ転送指定部１１３と、並列処理指定部１１４と、並列処理パターン作成部１１５と、性能測定部１１６と、実行ファイル作成部１１７と、本番環境配置部（Deploy final binary files to production environment）１１８と、性能測定テスト抽出実行部（Extract performance test cases and run automatically）１１９と、ユーザ提供部（Provide price and performance to a user to judge）１２０と、を備える。

<アプリケーションコード指定部１１１>
　アプリケーションコード指定部１１１は、入力されたアプリケーションコードの指定を行う。具体的には、アプリケーションコード指定部１１１は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部１１２に渡す。

<アプリケーションコード分析部１１２>
　アプリケーションコード分析部１１２は、処理機能のソースコードを分析し、ループ文やＦＦＴライブラリ呼び出し等の構造を把握する。

<データ転送指定部１１３>
　データ転送指定部１１３は、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行う。

　データ転送指定部１１３は、ＣＰＵからＧＰＵへのデータ転送を明示的に指定する明示的指定行と、ＧＰＵからＣＰＵへのデータ転送を明示的に指定する明示的指定行と、同じ変数に関してＣＰＵからＧＰＵへの転送とＧＰＵからＣＰＵへの転送とが重なる場合、データコピーの往復をまとめて明示的に指定する明示的指定行と、を用いたデータ転送指定を行う。

　データ転送指定部１１３は、ＣＰＵプログラム側で定義した変数とＧＰＵプログラム側で参照する変数が重なる場合、ＣＰＵからＧＰＵへのデータ転送の指示を行い、データ転送を指定する位置を、ＧＰＵ処理するループ文かそれより上位のループ文で、該当変数の設定、定義を含まない最上位のループとする。また、データ転送指定部１１３は、ＧＰＵプログラム側で設定した変数とＣＰＵプログラム側で参照する変数とが重なる場合、ＧＰＵからＣＰＵへのデータ転送の指示を行い、データ転送を指定する位置を、ＧＰＵ処理するループ文か、それより上位のループ文で、該当変数の参照、設定、定義を含まない最上位のループとする。

<並列処理指定部１１４>
　並列処理指定部１１４は、アプリケーションプログラムのループ文（繰り返し文）を特定し、各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする。
　並列処理指定部１１４は、オフロード範囲抽出部（Extract offloadable area）１１４ａと、中間言語ファイル出力部（Output intermediate file）１１４ｂと、を備える。

　オフロード範囲抽出部１１４ａは、ループ文やＦＦＴ等、ＧＰＵ・ＦＰＧＡにオフロード可能な処理を特定し、オフロード処理に応じた中間言語を抽出する。

　中間言語ファイル出力部１１４ｂは、抽出した中間言語ファイル１３２を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。

<並列処理パターン作成部１１５>
　並列処理パターン作成部１１５は、コンパイルエラーが出るループ文（繰り返し文）に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。

<性能測定部１１６>
　性能測定部１１６は、並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン１４に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する。
　性能測定部１１６は、バイナリファイル配置部（Deploy binary files）１１６ａを備える。バイナリファイル配置部１１６ａは、ＧＰＵ・ＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイ(配置)する。

　性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、オフロード範囲抽出部１１４ａに戻す。この場合、オフロード範囲抽出部１１４ａは、別の並列処理パターン抽出を行い、中間言語ファイル出力部１１４ｂは、抽出された中間言語をもとに、性能測定を試行する（後記図２の符号ａ参照）。

<実行ファイル作成部１１７>
　実行ファイル作成部１１７は、所定回数繰り返された、性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する。

<本番環境配置部１１８>
　本番環境配置部１１８は、作成した実行ファイルを、ユーザ向けの本番環境に配置する（「最終バイナリファイルの本番環境への配置」）。本番環境配置部１１８は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。

<性能測定テスト抽出実行部１１９>
　性能測定テスト抽出実行部１１９は、実行ファイル配置後、テストケースＤＢ１３１から性能試験項目を抽出し、性能試験を実行する（「最終バイナリファイルの本番環境への配置」）。
　性能測定テスト抽出実行部１１９は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースＤＢ１３１から抽出し、抽出した性能試験を自動実行する。

<ユーザ提供部１２０>
　ユーザ提供部１２０は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する（「価格・性能等の情報のユーザへの提供」）。テストケースＤＢ１３１には、性能試験項目が格納されている。ユーザ提供部１２０は、テストケースＤＢ１３１に格納された試験項目に対応した性能試験の実施結果に基づいて、価格、性能等のデータを、上記性能試験結果と共にユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。ここで、本番環境への一括デプロイには、非特許文献（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.）の技術を用いればよい。

［遺伝的アルゴリズムの適用］
　オフロードサーバ１は、オフロードの最適化にＧＡを用いることができる。ＧＡを用いた場合のオフロードサーバ１の構成は下記の通りである。
　すなわち、並列処理指定部１１４は、遺伝的アルゴリズムに基づき、コンパイルエラーが出ないループ文（繰り返し文）の数を遺伝子長とする。並列処理パターン作成部１１５は、アクセラレータ処理をする場合を１または０のいずれか一方、しない場合を他方の０または１として、アクセラレータ処理可否を遺伝子パターンにマッピングする。

　並列処理パターン作成部１１５は、遺伝子の各値を１か０にランダムに作成した指定個体数の遺伝子パターンを準備し、性能測定部１１６は、各個体に応じて、アクセラレータにおける並列処理指定文を指定したアプリケーションコードをコンパイルして、検証用マシン１４に配置する。性能測定部１１６は、検証用マシン１４において性能測定用処理を実行する。

　ここで、性能測定部１１６は、途中世代で、以前と同じ並列処理パターンの遺伝子が生じた場合は、当該並列処理パターンに該当するアプリケーションコードのコンパイル、および、性能測定はせずに、性能測定値としては同じ値を使う。
　また、性能測定部１１６は、コンパイルエラーが生じるアプリケーションコード、および、性能測定が所定時間で終了しないアプリケーションコードについては、タイムアウトの扱いとして、性能測定値を所定の時間（長時間）に設定する。

　実行ファイル作成部１１７は、全個体に対して、性能測定を行い、処理時間の短い個体ほど適合度が高くなるように評価する。実行ファイル作成部１１７は、全個体から、適合度が所定値（例えば、全個数の上位ｎ％、または全個数の上位ｍ個　ｎ，ｍは自然数）より高いものを性能の高い個体として選択し、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。実行ファイル作成部１１７は、指定世代数の処理終了後、最高性能の並列処理パターンを解として選択する。

　以下、上述のように構成されたオフロードサーバ１の自動オフロード動作について説明する。
［自動オフロード動作］
　図２は、オフロードサーバ１のＧＡを用いた自動オフロード処理を示す図である。
　図２に示すように、オフロードサーバ１は、環境適応ソフトウェアの要素技術に適用される。オフロードサーバ１は、制御部（自動オフロード機能部）１１と、テストケースＤＢ１３１と、中間言語ファイル１３２と、検証用マシン１４と、を有している。
　オフロードサーバ１は、ユーザが利用するアプリケーションコード（Application code）１２５を取得する。

　ユーザは、例えば、各種デバイス（Device１５１、ＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３、ＣＰＵを有する装置１５４）の利用を契約した人である。
　オフロードサーバ１は、機能処理をＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３のアクセラレータに自動オフロードする。

　以下、図２のステップ番号を参照して各部の動作を説明する。
<ステップＳ１１：Specify application code>
　ステップＳ１１において、アプリケーションコード指定部１１１（図１参照）は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部１１２に渡す。

<ステップＳ１２：Analyze application code>
　ステップＳ１２において、アプリケーションコード分析部１１２（図１参照）は、処理機能のソースコードを分析し、ループ文やＦＦＴライブラリ呼び出し等の構造を把握する。

<ステップＳ１３：Extract offloadable area>
　ステップＳ１３において、並列処理指定部１１４（図１参照）は、アプリケーションのループ文（繰り返し文）を特定し、各繰り返し文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする。具体的には、オフロード範囲抽出部１１４ａ（図１参照）は、ループ文やＦＦＴ等、ＧＰＵ・ＦＰＧＡにオフロード可能な処理を特定し、オフロード処理に応じた中間言語を抽出する。

<ステップＳ１４：Output intermediate file>
　ステップＳ１４において、中間言語ファイル出力部１１４ｂ（図１参照）は、中間言語ファイル１３２を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。

<ステップＳ１５：Compile error>
　ステップＳ１５において、並列処理パターン作成部１１５（図１参照）は、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。

<ステップＳ２１：Deploy binary files>
　ステップＳ２１において、バイナリファイル配置部１１６ａ（図１参照）は、ＧＰＵ・ＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイする。

<ステップＳ２２：Measure performances>
　ステップＳ２２において、性能測定部１１６（図１参照）は、配置したファイルを実行し、オフロードした際の性能を測定する。
　オフロードする領域をより適切にするため、この性能測定結果は、オフロード範囲抽出部１１４ａに戻され、オフロード範囲抽出部１１４ａが、別パターンの抽出を行う。そして、中間言語ファイル出力部１１４ｂは、抽出された中間言語をもとに、性能測定を試行する（図２の符号ａ参照）。

　図２の符号ａに示すように、制御部１１は、上記ステップＳ１２乃至ステップＳ２２を繰り返し実行する。制御部１１の自動オフロード機能をまとめると、下記である。すなわち、並列処理指定部１１４は、アプリケーションプログラムのループ文（繰り返し文）を特定し、各繰返し文に対して、ＧＰＵでの並列処理指定文を指定して、コンパイルする。そして、並列処理パターン作成部１１５は、コンパイルエラーが出るループ文を、オフロード対象外とし、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。そして、バイナリファイル配置部１１６ａは、該当並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン１４に配置し、性能測定部１１６が、検証用マシン１４で性能測定用処理を実行する。実行ファイル作成部１１７は、所定回数繰り返された、性能測定結果をもとに、複数の並列処理パターンから最高処理性能のパターンを選択し、選択パターンをコンパイルして実行ファイルを作成する。

<ステップＳ２３：Deploy final binary files to production environment>
　ステップＳ２３において、本番環境配置部１１８は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。

<ステップＳ２４：Extract performance test cases and run automatically>
　ステップＳ２４において、性能測定テスト抽出実行部１１９は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースＤＢ１３１から抽出し、抽出した性能試験を自動実行する。

<ステップＳ２５：Provide price and performance to a user to judge>
　ステップＳ２５において、ユーザ提供部１２０は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。

　上記ステップＳ１１～ステップＳ２５は、例えばユーザのサービス利用のバックグラウンドで行われ、例えば、仮利用の初日の間に行う等を想定している。また、コスト低減のためにバックグラウンドで行う処理は、機能配置最適化とＧＰＵ・ＦＰＧＡオフロードのみを対象としてもよい。

　上記したように、オフロードサーバ１の制御部（自動オフロード機能部）１１は、環境適応ソフトウェアの要素技術に適用した場合、機能処理のオフロードのため、ユーザが利用するアプリケーションプログラムのソースコードから、オフロードする領域を抽出して中間言語を出力する（ステップＳ１１～ステップＳ１５）。制御部１１は、中間言語から導かれる実行ファイルを、検証用マシン１４に配置実行し、オフロード効果を検証する（ステップＳ２１～ステップＳ２２）。検証を繰り返し、適切なオフロード領域を定めたのち、制御部１１は、実際にユーザに提供する本番環境に、実行ファイルをデプロイし、サービスとして提供する（ステップＳ２３～ステップＳ２５）。

［ＧＡを用いたＧＰＵ自動オフロード］
　ＧＰＵ自動オフロードは、ＧＰＵに対して、図２のステップＳ１２～ステップＳ２２を繰り返し、最終的にステップＳ２３でデプロイするオフロードコードを得るための処理である。

　ＧＰＵは、一般的にレイテンシーは保証しないが、並列処理によりスループットを高めることに向いたデバイスである。環境適応ソフトウェアの暗号化処理や、カメラ映像分析のための画像処理、大量センサデータ分析のための機械学習処理等が代表的であり、それらは、繰り返し処理が多い。そこで、アプリケーションの繰り返し文をＧＰＵに自動でオフロードすることでの高速化を狙う。

　しかし、従来技術で記載の通り、高速化には適切な並列処理が必要である。特に、ＧＰＵを使う場合は、ＣＰＵとＧＰＵ間のメモリ転送のため、データサイズやループ回数が多くないと性能が出ないことが多い。また、メモリデータ転送のタイミング等により、並列高速化できる個々のループ文（繰り返し文）の組み合わせが、最速とならない場合等がある。例えば、１０個のfor文（繰り返し文）で、１番、５番、１０番の３つがＣＰＵに比べて高速化できる場合に、１番、５番、１０番の３つの組み合わせが最速になるとは限らない等である。

　適切な並列領域指定のため、ＰＧＩコンパイラを用いて、for文の並列可否を試行錯誤して最適化する試みがある。しかし、試行錯誤には多くの稼働がかかり、サービスとして提供する際に、ユーザの利用開始が遅くなり、コストも上がってしまう問題がある。

　そこで、本実施形態では、並列化を想定していない汎用プログラムから、自動で適切なオフロード領域を抽出する。このため、最初に並列可能for文のチェックを行い、次に並列可能for文群に対してＧＡを用いて検証環境で性能検証試行を反復し適切な領域を探索すること、を実現する。並列可能for文に絞った上で、遺伝子の部分の形で、高速化可能な並列処理パターンを保持し組み換えていくことで、取り得る膨大な並列処理パターンから、効率的に高速化可能なパターンを探索できる。

［Simple GAによる制御部（自動オフロード機能部）１１の探索イメージ］
　図３は、Simple GAによる制御部（自動オフロード機能部）１１の処理の探索イメージとfor文の遺伝子配列マッピングを示す図である。
　ＧＡは、生物の進化過程を模倣した組合せ最適化手法の一つである。ＧＡのフローチャートは、初期化→評価→選択→交叉→突然変異→終了判定となっている。
　本実施形態では、ＧＡの中で、処理を単純にしたSimple GAを用いる。Simple GAは、遺伝子は１、０のみとし、ルーレット選択、一点交叉、突然変異は１箇所の遺伝子の値を逆にする等、単純化されたＧＡである。

<初期化>
　初期化では、アプリケーションコードの全for文の並列可否をチェック後、並列可能for文を遺伝子配列にマッピングする。ＧＰＵ処理する場合は１、ＧＰＵ処理しない場合は０とする。遺伝子は、指定の個体数Ｍを準備し、１つのfor文にランダムに１、０の割り当てを行う。
　具体的には、制御部（自動オフロード機能部）１１（図１参照）は、ユーザが利用するアプリケーションコード（Application code）１３０（図２参照）を取得し、図３に示すように、アプリケーションコード１３０のコードパターン（Code patterns）１４１からfor文の並列可否をチェックする。図３に示すように、コードパターン１４１から３つのfor文が見つかった場合（図３の符号ｂ参照）、各for文に対して１桁、ここでは３つのfor文に対し３桁の１または０を割り当てる。例えば、ＣＰＵで処理する場合０、ＧＰＵに出す場合１とする。ただし、この段階では１または０をランダムに割り当てる。
　遺伝子長に該当するコードが３桁であり、３桁の遺伝子長のコードは２^３パターン、例えば１００、１０１、…となる。なお、図３では、コードパターン１４１中の丸印（○印）をコードのイメージとして示している。

<評価>
　評価では、デプロイとパフォーマンスの測定（Deploy & performance measurement）を行う（図３の符号ｃ参照）。すなわち、性能測定部１１６（図１参照）は、遺伝子に該当するコードをコンパイルして検証用マシン１４にデプロイして実行する。性能測定部１１６は、ベンチマーク性能測定を行う。性能が良いパターン（並列処理パターン）の遺伝子の適合度を高くする。

<選択>
　選択では、適合度に基づいて、高性能コードパターンを選択（Select high performance code patterns）する（図３の符号ｄ参照）。性能測定部１１６（図１参照）は、適合度に基づいて、高適合度の遺伝子を、指定の個体数選択する。本実施形態では、適合度に応じたルーレット選択および最高適合度遺伝子のエリート選択を行う。
　図３では、選択されたコードパターン（Select code patterns）１４２の中の丸印（○印）が、３つに減ったことを探索イメージとして示している。

<交叉>
　交叉では、一定の交叉率Ｐｃで、選択された個体間で一部の遺伝子をある一点で交換し、子の個体を作成する。
　ルーレット選択された、あるパターン（並列処理パターン）と他のパターンとの遺伝子を交叉させる。一点交叉の位置は任意であり、例えば上記３桁のコードのうち２桁目で交叉させる。

<突然変異>
　局所解を避けるため、突然変異を導入する。なお、演算量を削減するために突然変異を行わない態様でもよい。突然変異では、一定の突然変異率Ｐｍで、個体の遺伝子の各値を０から１または１から０に変更する。

<終了判定>
　図３に示すように、交叉と突然変異後の次世代コードパターンの生成（Generate next generation code patterns after crossover & mutation）を行う（図３の符号ｅ参照）。
　終了判定では、指定の世代数Ｔ回、繰り返しを行った後に処理を終了し、最高適合度の遺伝子を解とする。
　例えば、性能測定して、速い３つ１０１、０１０、００１を選ぶ。この３つをＧＡにより、次の世代は、組み換えをして、例えば新しいパターン（並列処理パターン）１０１（一例）を作っていく。このとき、組み換えをしたパターンに、勝手に０を１にするなどの突然変異を入れる。上記を繰り返して、一番早いパターンを見付ける。指定世代（例えば、２０世代）などを決めて、最終世代で残ったパターンを、最後の解とする。

<デプロイ（配置）>
　最高適合度の遺伝子に該当する、最高処理性能の並列処理パターンで、本番環境に改めてデプロイして、ユーザに提供する。

　ＧＰＵにオフロードできないfor文（ループ文；繰り返し文）が相当数存在する場合について説明する。例えば、for文が２００個あっても、ＧＰＵにオフロードできるものは３０個くらいである。ここでは、エラーになるものを除外し、この３０個について、ＧＡを行う。

　例えば、C/C++コードに対するＧＰＵの処理を行う仕様であるOpenＡＣＣには、ディレクティブ #pragma acc kernelsで指定して、ＧＰＵ向けバイトコードを抽出し、実行によりＧＰＵオフロードを可能とするコンパイラがある。Python, Javaの場合は、ＣＵＤＡやJava Lambda式などでＧＰＵ処理を指定すればよい。

　また、C/C++を使った場合、C/C++のコードを分析し、for文を見付ける。for文を見付けると、OpenＡＣＣで並列処理の文法である #pragma acc kernelsを使ってfor文に対して書き込む。詳細には、何も入っていない #pragma acc kernels に、一つ一つfor文を入れてコンパイルして、エラーであれば、そのfor文はそもそも、ＧＰＵ処理できないので、除外する。このようにして、残るfor文を見付ける。そして、エラーが出ないものを、長さ（遺伝子長）とする。エラーのないfor文が５つであれば、遺伝子長は５であり、エラーのないfor文が１０であれば、遺伝子長は１０である。なお、並列処理できないものは、前の処理を次の処理に使うようなデータに依存がある場合である。
　以上が準備段階である。次にＧＡ処理を行う。

　for文の数に対応する遺伝子長を有するコードパターンが得られている。始めはランダムに並列処理パターン１００、０１０、００１、…を割り当てる。そして、ＧＡ処理を行い、コンパイルする。その時に、オフロードできるfor文であるにもかかわらず、エラーがでることがある。for文が階層になっている（どちらか指定すればＧＰＵ処理できる）場合である。この場合は、エラーとなったfor文は、残してもよい。具体的には、処理時間が多くなった形にして、タイムアウトさせる方法がある。

　検証用マシン１４でデプロイして、ベンチマーク、例えば画像処理であればその画像処理でベンチマークする、その処理時間が短い程、適応度が高いと評価する。例えば、処理時間の逆数、処理時間１０秒かかるものは１、１００秒かかるものは０．１、１秒のものは１０とする。
　適応度が高いものを選択して、例えば１０個のなかから、３～５個を選択して、それを組み替えて新しいコードパターンを作る。作成途中で、前と同じものができる場合がある。この場合、同じベンチマークを行う必要はないので、前と同じデータを使う。本実施形態では、コードパターンと、その処理時間は記憶部１３に保存しておく。

　以上で、Simple GAによる制御部（自動オフロード機能部）１１の探索イメージについて説明した。次に、データ転送の一括処理手法について述べる。

［データ転送の一括処理手法］
　上述したように、遺伝的アルゴリズムを用いることで、ＧＰＵ処理で効果のある並列処理部を自動チューニングしている。しかしながら、ＣＰＵ－ＧＰＵメモリ間のデータ転送によっては高性能化できないアプリケーションもあった。このため、スキルが無いユーザがＧＰＵを使ってアプリケーションを高性能化することは難しいし、自動並列化技術等を使う場合も並列処理可否の試行錯誤が必要であり、高速化できない場合があった。

　そこで、本実施形態では、より多くのアプリケーションを、自動でＧＰＵを用いて高性能化することを狙うとともに、ＧＰＵへのデータ転送回数を低減できる技術を提供する。

<基本的な考え方>
　OpenＡＣＣ等の仕様では、ＧＰＵでの並列処理を指定する指示行に加えて、ＣＰＵからＧＰＵへのデータ転送やその逆を明示的に指定する指示行（以下、「明示的指示行」という）が定義されている。OpenＡＣＣ等の明示的指示行は、ＣＰＵからＧＰＵへのデータ転送のディレクティブ（directive：行頭に特殊な記号を記述した指示・指定コマンド）である「#pragma acc data copyin」、ＧＰＵからＣＰＵへのデータ転送のディレクティブである「#pragma acc data copyout」、ＣＰＵからＧＰＵへ再びＣＰＵへのデータ転送のディレクティブである「#pragma acc data copy」等である。

　本実施形態は、非効率なデータ転送を低減するため、明示的指示行を用いたデータ転送指定を、ＧＡでの並列処理の抽出と合わせて行う。
　本実施形態では、ＧＡで生成された各個体について、ループ文の中で利用される変数データの参照関係を分析し、ループ毎に毎回データ転送するのではなくループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する。

<具体例>
　以下、具体的に処理を説明する。
　データ転送の種類は、ＣＰＵからＧＰＵへのデータ転送、および、ＧＰＵからＣＰＵへのデータ転送がある。

　図４は、具体例の自動オフロード機能部が処理するアプリケーションプログラムのソースコードのループ文（繰り返し文）を示す図であり、ＣＰＵプログラム側で定義した変数とＧＰＵプログラム側で参照する変数が重なる場合の例である。
　具体例の自動オフロード機能部は、図１の制御部（自動オフロード機能部）１１からデータ転送指定部１１３を取り去る、またはデータ転送指定部１１３を実行しない場合の例である。

　具体例のＣＰＵからＧＰＵへのデータ転送を例に採る。
　図４は、具体例のＣＰＵからＧＰＵへデータ転送する場合のループ文において、ＣＰＵプログラム側で定義した変数とＧＰＵプログラム側で参照する変数が重なる場合の例である。なお、以下の記載および図４中のループ文の文頭の<1>～<4> は、説明の便宜上で付したものである（他図およびその説明においても同様）。
　図４に示す具体例のループ文は、ＣＰＵプログラム側で記述され、
<1> ループ〔 for｜do｜while 〕 {
}
の中に、
　<2> ループ〔 for｜do｜while 〕 {
　}
があり、さらにその中に、
　　<3> ループ〔 for｜do｜while 〕 {
　　}
があり、さらにその中に、
　　　<4> ループ〔 for〕{
　　　}
がある。

　また、<1> ループ〔 for｜do｜while 〕 {
}で、変数ａが設定され、<4> ループ〔 for｜do｜while 〕 {
}で、変数ａが参照される。

　さらに、<3> ループ〔 for｜do｜while 〕 {
}で、ＰＧＩコンパイラによるfor文等の並列処理可能処理部を、OpenＡＣＣのディレクティブ #pragma acc kernels（並列処理指定文）で指定している（詳細後記）。

　図４に示す比較例のループ文では、図４の符号ｆに示すタイミングで毎回ＣＰＵからＧＰＵにデータ転送する。このため、ＧＰＵへのデータ転送回数を低減することが求められる。
　なお、ＧＰＵからＣＰＵへのデータ転送も同様であり説明を省略する。

　以上述べたように、本実施形態では、できるだけ上位のループでデータ転送を一括して行うように、データ転送を明示的に指示することで、ループ毎に毎回データを転送する非効率な転送を避けることができる。

［ＧＰＵオフロード処理］
　上述したデータ転送の一括処理手法により、オフロードに適切なループ文を抽出し、非効率なデータ転送を避けることができる。
　ただし、上記データ転送の一括処理手法を用いても、ＧＰＵオフロードに向いていないプログラムも存在する。効果的なＧＰＵオフロードには、オフロードする処理のループ回数が多いことが必要である。

　そこで、本実施形態では、本格的なオフロード処理探索の前段階として、プロファイリングツールを用いて、ループ回数を調査する。プロファイリングツールを用いると、各行の実行回数を調査できるため、例えば、５０００万回以上のループを持つプログラムをオフロード処理探索の対象とする等、事前に振り分けることができる。以下、具体的に説明する（前記図３で述べた内容と一部重複する）。

　本実施形態では、まず、制御部（自動オフロード機能部）１１（図１参照）が、アプリケーションプログラムを分析し、for，do，while等のループ文を把握する。次に、サンプル処理を実行し、プロファイリングツールを用いて、各ループ文のループ回数を調査し、一定の値以上のループがあるか否かで、探索を本格的に行うか否かの判定を行う。

　探索を本格的に行うと決まった場合は、ＧＡの処理に入る（前記図３参照）。初期化ステップでは、アプリケーションコードの全ループ文の並列可否をチェックした後、並列可能ループ文をＧＰＵ処理する場合は１、しない場合は０として遺伝子配列にマッピングする。遺伝子は、指定の個体数が準備されるが、遺伝子の各値にはランダムに１，０の割り当てをする。

　ここで、遺伝子に該当するコードでは、ＧＰＵ処理すると指定されたループ文内の変数データ参照関係から、データ転送の明示的指示（OpenACCで指定するならば#pragma acc data copyin/copyout/copy）を追加する。

　評価ステップでは、遺伝子に該当するコードをコンパイルして検証用マシンにデプロイして実行し、ベンチマーク性能測定を行う。性能が良いパターンの遺伝子の適合度を高くする。遺伝子に該当するコードは、上述のように、並列処理指示行（例えば、図４の符号ｆ参照）が挿入されている。

　選択ステップでは、適合度に基づいて、高適合度の遺伝子を、指定の個体数選択する。本実施形態では、適合度に応じたルーレット選択および最高適合度遺伝子のエリート選択を行う。交叉ステップでは、一定の交叉率Ｐｃで、選択された個体間で一部の遺伝子をある一点で交換し、子の個体を作成する。突然変異ステップでは、一定の突然変異率Ｐｍで、個体の遺伝子の各値を０から１または１から０に変更する。

　突然変異ステップまで終わり、次の世代の遺伝子が指定個体数作成されると、初期化ステップと同様に、データ転送の明示的指示を追加し、評価、選択、交叉、突然変異ステップを繰り返す。

　最後に、終了判定ステップでは、指定の世代数、繰り返しを行った後に処理を終了し、最高適合度の遺伝子を解とする。最高適合度の遺伝子に該当する、最高性能のコードパターンで、本番環境に改めてデプロイして、ユーザに提供する。

　以下、オフロードサーバ１の実装を説明する。本実装は、本実施形態の有効性を確認するためのものである。
［実装］
　C/C++アプリケーションを汎用のＰＧＩコンパイラを用いて自動オフロードする実装を説明する。
　本実装では、ＧＰＵ自動オフロードの有効性確認が目的であるため、対象アプリケーションはC/C++言語のアプリケーションとし、ＧＰＵ処理自体は、従来のＰＧＩコンパイラを説明に用いる。

　C/C++言語は、ＯＳＳ（Open Source Software）およびproprietaryソフトウェアの開発で、上位の人気を誇り、数多くのアプリケーションがC/C++言語で開発されている。一般ユーザが用いるアプリケーションプログラムのオフロードを確認するため、暗号処理や画像処理等のＯＳＳの汎用アプリケーションを利用する。

　ＧＰＵ処理は、ＰＧＩコンパイラにより行う。ＰＧＩコンパイラは、OpenＡＣＣを解釈するC/C++/Fortran向けコンパイラである。本実施形態では、for文等の並列可能処理部を、OpenＡＣＣのディレクティブ #pragma acc kernels（並列処理指定文）で指定する。これにより、ＧＰＵ向けバイトコードを抽出し、その実行によりＧＰＵオフロードを可能としている。さらに、for文内のデータ同士に依存性があり並列処理できない処理やネストのfor文の異なる複数の階層を指定されている場合等の際に、エラーを出す。合わせて、#pragma acc data copyin/copyout/copy 等のディレクティブにより、明示的なデータ転送の指示が可能とする。

　上記 #pragma acc kernels（並列処理指定文）での指定に合わせて、OpenＡＣＣのcopyin 節の #pragma acc data copyout(a[…])の、上述した位置への挿入により、明示的なデータ転送の指示を行う。

<実装の動作概要>
　実装の動作概要を説明する。
　高速化するC/C++アプリケーションプログラムとそれを性能測定するベンチマークツールを準備する。

　実装では、C/C++アプリケーションプログラムの利用依頼があると、まず、C/C++アプリケーションのコードを解析して、for文を発見するとともに、for文内で使われる変数データ等の、プログラム構造を把握する。構文解析には、LLVM/Clangの構文解析ライブラリ（libClangのpython binding）等を使用する。

　実装では、最初に、そのアプリケーションがＧＰＵオフロード効果があるかの見込みを得るため、ベンチマークを実行し、上記構文解析で把握したfor文のループ回数を把握する。ループ回数把握には、GNUカバレッジのgcov等を用いる。プロファイリングツールとしては、「GNUプロファイラ(gprof)」、「GNUカバレッジ(gcov)」が知られている。双方とも各行の実行回数を調査できるため、どちらを用いてもよい。実行回数は、例えば、１０００万回以上のループ回数を持つアプリケーションプログラムのみ対象とするようにできるが、この値は変更可能である。

　ＣＰＵ向け汎用アプリケーションプログラムは、並列化を想定して実装されているわけではない。そのため、まず、ＧＰＵ処理自体が不可なfor文は排除する必要がある。そこで、各for文一つずつに対して、並列処理の#pragma acc kernels ディレクティブ挿入を試行し、コンパイル時にエラーが出るかの判定を行う。コンパイルエラーに関しては、幾つかの種類がある。for文の中で外部ルーチンが呼ばれている場合、ネストfor文で異なる階層が重複指定されている場合、break等でfor文を途中で抜ける処理がある場合、for文のデータにデータ依存性がある場合等がある。アプリケーションプログラムによって、コンパイル時エラーの種類は多彩であり、これ以外の場合もあるが、コンパイルエラーは処理対象外とし、#pragmaディレクティブは挿入しない。

　コンパイルエラーは自動対処が難しく、また対処しても効果が出ないことも多い。外部ルーチンコールの場合は、#pragma acc routineにより回避できる場合があるが、多くの外部コールはライブラリであり、それを含めてＧＰＵ処理してもそのコールがネックとなり性能が出ない。for文一つずつを試行するため、ネストのエラーに関しては、コンパイルエラーは生じない。また、break等で途中で抜ける場合は、並列処理にはループ回数を固定化する必要があり、プログラム改造が必要となる。データ依存が有る場合はそもそも並列処理自体ができない。

　ここで、並列処理してもエラーが出ないループ文の数がａの場合、ａが遺伝子長となる。遺伝子の１は並列処理ディレクティブ有、０は無に対応させ、長さａの遺伝子に、アプリケーションコードをマッピングする。

　次に、初期値として、指定個体数の遺伝子配列を準備する。遺伝子の各値は、図３で説明したように、０と１をランダムに割当てて作成する。準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブ #pragma acc kernels をC/C++コードに挿入する。この段階で、ある遺伝子に該当するコードの中で、ＧＰＵで処理させる部分が決まる。上記Clangで解析した、for文内の変数データの参照関係をもとに、上述したルールに基づいて、ＣＰＵからＧＰＵへのデータ転送、その逆の場合のディレクティブ指定を行う。
　具体的には、ＣＰＵからＧＰＵへのデータ転送が必要な変数は、 #pragma acc data copyinで指定し（図示省略）、ＧＰＵからＣＰＵへのデータ転送が必要な変数は、 #pragma acc data copyoutで指定する（図示省略）。同じ変数に関して、copyinとcopyoutが重なる場合は、#pragma acc data copyで纏め、記述をシンプルにする。

　並列処理およびデータ転送のディレクティブを挿入されたC/C++コードを、ＧＰＵを備えたマシン上のＰＧＩコンパイラでコンパイルを行う。コンパイルした実行ファイルをデプロイし、ベンチマークツールで性能を測定する。

　全個体数に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。

　次世代の個体に対して、ディレクティブ挿入、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。ここで、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイル、性能測定をせず、以前と同じ測定値を用いる。

　指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当する、ディレクティブ付きC/C++コードを解とする。

　この中で、個体数、世代数、交叉率、突然変異率、適合度設定、選択方法は、ＧＡのパラメータであり、別途指定する。提案技術は、上記処理を自動化することで、従来、専門技術者の時間とスキルが必要だった、ＧＰＵオフロードの自動化を可能にする。

［《ループ文オフロード：共通》フローチャート］
　図５Ａ－Ｂは、《ループ文オフロード：共通》フローチャートであり、図５Ａと図５Ｂは、結合子で繋がれる。

<コード解析>
　ステップＳ１０１で、アプリケーションコード分析部１１２（図１参照）は、アプリケーションプログラムのコード解析を行う。

<ループ文特定>
　ステップＳ１０２で、並列処理指定部１１４（図１参照）は、アプリケーションプログラムのループ文、参照関係を特定する。

<ループ文ループ回数>
　ステップＳ１０３で、並列処理指定部１１４は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。

<ループ文の並列処理可能性>
　ステップＳ１０４で、並列処理指定部１１４は、各ループ文の並列処理可能性をチェックする。

<ループ文の繰り返し>
　制御部（自動オフロード機能部）１１は、ステップＳ１０５のループ始端とステップＳ１０８のループ終端間で、ステップＳ１０６－Ｓ１０７の処理についてループ文の数だけ繰り返す。
　ステップＳ１０６で、並列処理指定部１１４は、各ループ文に対して、言語に応じた手法でＧＰＵ処理を指定してコンパイルまたはインタプリットする。
　ステップＳ１０７で、並列処理指定部１１４は、エラー時は、該当for文からは、ＧＰＵ処理指定を削除する。
　ステップＳ１０９で、並列処理指定部１１４は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。

<指定個体数パターン準備>
　次に、初期値として、並列処理指定部１１４は、指定個体数の遺伝子配列を準備する。ここでは、０と１をランダムに割当てて作成する。
　ステップＳ１１０で、並列処理指定部１１４は、アプリケーションプログラムのコードを、遺伝子にマッピングする。０と１がランダムに割当てられた遺伝子配列を遺伝子にマッピングすることで、指定個体数パターンを準備する。
　準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをアプリケーションプログラムのコードに挿入する（例えば図３の#pragmaディレクティブ参照）。

　制御部（自動オフロード機能部）１１は、ステップＳ１１１のループ始端とステップＳ１２０のループ終端間で、ステップＳ１１２－Ｓ１１９の処理について指定世代数繰り返す。
　また、上記指定世代数繰り返しにおいて、さらにステップＳ１１２のループ始端とステップＳ１１７のループ終端間で、ステップＳ１１３－Ｓ１１６の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。

<データ転送指定>
　ステップＳ１１３で、データ転送指定部１１３は、変数参照関係から、言語に応じた手法でデータ転送を指定する。

<コンパイル>
　ステップＳ１１４で、並列処理パターン作成部１１５（図１参照）は、遺伝子パターンに応じてＧＰＵ処理基盤でコンパイルまたはインタプリットする。すなわち、並列処理パターン作成部１１５は、作成したアプリケーションプログラムのコードを、ＧＰＵを備えた検証用マシン１４上のＰＧＩコンパイラでコンパイルまたはインタプリットを行う。
　ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。

　ステップＳ１１５で、性能測定部１１６（図１参照）は、ＣＰＵ-ＧＰＵ搭載の検証用マシン１４に、実行ファイルをデプロイする。
　ステップＳ１１６で、性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。

　ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
　ステップＳ１１８で、実行ファイル作成部１１７（図１参照）は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。

　ステップＳ１１９で、実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
　すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。

　ステップＳ１２１で、実行ファイル作成部１１７は、指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当するC/C++コード（最高性能の並列処理パターン）を解とする。

<ＧＡのパラメータ>
　上記、個体数、世代数、交叉率、突然変異率、適合度設定、選択方法は、ＧＡのパラメータである。ＧＡのパラメータは、例えば、以下のように設定してもよい。
　実行するSimple GAの、パラメータ、条件は例えば以下のようにできる。
　遺伝子長：並列可能ループ文数
　個体数Ｍ：遺伝子長以下
　世代数Ｔ：遺伝子長以下
　適合度：(処理時間)^-1/2
　この設定により、ベンチマーク処理時間が短い程、高適合度になる。また、適合度を、(処理時間)^-1/2とすることで、処理時間が短い特定の個体の適合度が高くなり過ぎて、探索範囲が狭くなるのを防ぐことができる。また、性能測定が一定時間で終わらない場合は、タイムアウトさせ、処理時間１０００秒等の時間（長時間）であるとして、適合度を計算する。このタイムアウト時間は、性能測定特性に応じて変更させればよい。
　選択：ルーレット選択
　ただし、世代での最高適合度遺伝子は交叉も突然変異もせず次世代に保存するエリート保存も合わせて行う。
　交叉率Ｐｃ：０．９
　突然変異率Ｐｍ：０．０５

　本実施形態では、gcov，gprof等を用いて、ループが多く実行時間がかかっているアプリケーションを事前に特定して、オフロード試行をする。これにより、効率的に高速化できるアプリケーションを見つけることができる。

　より短時間でオフロード部分を探索するためには、複数の検証用マシンで個体数分並列で性能測定することが考えられる。アプリケーションプログラムに応じて、タイムアウト時間を調整することも短時間化に繋がる。例えば、オフロード処理がＣＰＵでの実行時間の２倍かかる場合はタイムアウトとする等である。また、個体数、世代数が多い方が、高性能な解を発見できる可能性が高まる。しかし、各パラメータを最大にする場合、個体数×世代数だけコンパイル、および性能ベンチマークを行う必要がある。このため、本番サービス利用開始までの時間がかかる。本実施形態では、ＧＡとしては少ない個体数、世代数で行っているが、交叉率Ｐｃを０．９と高い値にして広範囲を探索することで、ある程度の性能の解を早く発見するようにしている。

　以上、《ループ文オフロード：共通》について説明した、次に、《ループ文オフロード：Ｃ言語》について説明する。

［《ループ文オフロード：Ｃ言語》］
　ループ文オフロード：Ｃ言語について、基本的フローは、上記《ループ文オフロード：共通》と同様であり、言語非依存にできる。Ｃ言語に依存・非依存の処理について詳細に述べる。
　ループ文オフロードの、コードの分析では、Ｃ言語を解析するClang等の構文解析ツールを用いて構文解析する。ループと変数の把握については、構文解析ツールの結果を管理する際は、言語に非依存に抽象的に管理できる。ループのＧＰＵ処理有無の遺伝子化についても、言語に非依存である。遺伝子情報のコード化では、遺伝子情報に合わせてＧＰＵで実行するためのコードを作成するため、Ｃ言語の拡張文法であるOpenACCでＧＰＵ処理を指定したり、変数転送を指定したりする。

　コンパイルは、OpenACCコードをＰＧＩコンパイラ等でコンパイルする。性能測定は、言語に合わせて、Jenkins、Selenium等の自動測定ツールも用いて行う。次世代の遺伝子作成は、性能測定結果に合わせて適合度を設定し交叉等の処理を行うが、言語に非依存である。反復実行と最終解の決定も、言語に非依存である。
　以上のように、ループ文オフロードでは、処理に関しては、ループと変数の管理とＧＡの遺伝子処理については言語に非依存に適用できる。

　図６Ａ－Ｂは、《ループ文オフロード：Ｃ言語》のフローチャートであり、図６Ａと図６Ｂは、結合子で繋がれる。
　C/C++向けOpenＡＣＣコンパイラを用いて以下の処理を行う。

<コード解析>
　ステップＳ２０１で、アプリケーションコード分析部１１２（図１参照）は、C/C++アプリケーションプログラムのコード解析を行う。

<ループ文特定>
　ステップＳ２０２で、並列処理指定部１１４（図１参照）は、C/C++アプリケーションプログラムのループ文、参照関係を特定する。

<ループ文ループ回数>
　ステップＳ２０３で、並列処理指定部１１４は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。

<ループ文の並列処理可能性>
　ステップＳ２０４で、並列処理指定部１１４は、各ループ文の並列処理可能性をチェックする。

<ループ文の繰り返し>
　制御部（自動オフロード機能部）１１は、ステップＳ２０５のループ始端とステップＳ２０８のループ終端間で、ステップＳ２０６－Ｓ２０７の処理についてループ文の数だけ繰り返す。
　ステップＳ２０６で、並列処理指定部１１４は、各ループ文に対して、OpenＡＣＣ文法を用いて、#pragma acc kernelsでＧＰＵ処理を指定してコンパイルする。
　ステップＳ２０７で、並列処理指定部１１４は、エラー時は、該当for文からは、#pragma acc kernelsを削除する。
　ステップＳ２０９で、並列処理指定部１１４は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。

<指定個体数パターン準備>
　次に、初期値として、並列処理指定部１１４は、指定個体数の遺伝子配列を準備する。ここでは、０と１をランダムに割当てて作成する。
　ステップＳ２１０で、並列処理指定部１１４は、C/C++アプリコードを、遺伝子にマッピングする。０と１がランダムに割当てられた指定個体数の遺伝子配列を遺伝子にマッピングするすることで、指定個体数パターンを準備する。
　準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをC/C++アプリコードに挿入する（例えば図３の#pragmaディレクティブ参照）。

　制御部（自動オフロード機能部）１１は、ステップＳ２１１のループ始端とステップＳ２２０のループ終端間で、ステップＳ２１２－Ｓ２１９の処理について指定世代数繰り返す。
　また、上記指定世代数繰り返しにおいて、さらにステップＳ２１２のループ始端とステップＳ２１７のループ終端間で、ステップＳ２１３－Ｓ２１６の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。

<データ転送指定>
　ステップＳ２１３で、データ転送指定部１１３は、変数参照関係から、明示的指示行（#pragma acc data copy/in/out）を用いたデータ転送指定を行う。明示的指示行（#pragma acc data copy/in/out）を用いたデータ転送指定については、図４により説明した。

<コンパイル>
　ステップＳ２１４で、並列処理パターン作成部１１５（図１参照）は、遺伝子パターンに応じてディレクティブ指定したC/C++コードをＰＧＩコンパイラでコンパイルする。すなわち、並列処理パターン作成部１１５は、作成したC/C++アプリコードを、ＧＰＵを備えた検証用マシン１４上のＰＧＩコンパイラでコンパイルする。
　ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。

　ステップＳ２１５で、性能測定部１１６（図１参照）は、ＣＰＵ-ＧＰＵ搭載の検証用マシン１４に、実行ファイルをデプロイする。
　ステップＳ２１６で、性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。

　ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
　ステップＳ２１８で、実行ファイル作成部１１７（図１参照）は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。

　ステップＳ２１９で、実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
　すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。

　ステップＳ２２１で、実行ファイル作成部１１７は、指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当するC/C++コード（最高性能の並列処理パターン）を解とする。

　以上、《ループ文オフロード：Ｃ言語》について説明した、次に、《ループ文オフロード：Python》について説明する。

［《ループ文オフロード：Python》］
　ループ文オフロード：Pythonは、PythonコードをpyCUDAでインタプリットする方法（図７Ａ－Ｂ参照）と、OpenACCを解釈するインタプリタpyACCを用いる方法（図８Ａ－Ｂ、図９Ａ－Ｂ参照）とがある。以下、順に説明する。

<PythonコードをpyCUDAでインタプリットする方法>
　ループ文オフロードの、コードの分析では、Pythonを解析するast等の構文解析ツールを用いて構文解析する。ループと変数の把握については、構文解析ツールの結果を管理する際は、言語に非依存で、抽象的に管理できる。
　ループのＧＰＵ処理有無の遺伝子化についても、言語に非依存である。遺伝子情報のコード化では、遺伝子情報に合わせてＧＰＵで実行するためのコードを作成するため、ＣＵＤＡ文法でＧＰＵ処理を指定したり、変数転送を指定したりする。

　インタプリタは、ＣＵＤＡでの指示を追加したPythonコードをpyCUDAでインタプリットする。性能測定は、言語に合わせて、Jenkins等の自動測定ツールも用いて行う。次世代の遺伝子作成は、性能測定結果に合わせて適合度を設定し交叉等の処理を行うが、言語に非依存である。反復実行と最終解の決定も、言語に非依存である。ここで、pyCUDAでなくpyACCというOpenACCを解釈するインタプリタを用いてもよい。その場合は、Ｃ言語と同様にOpenACC文法でループ文のＧＰＵ処理を指定すればよい（後記）。

　Python では、ＧＰＵ処理指定したコードは実装によるが、CupyというオープンソースでNVIDIA GPUを利用するライブラリを用いた実装を説明する。
　動作としては、PythonのコードでＧＰＵ処理を指定されるループ文はCupyライブラリを介して、NVIDIAのＣＵＤＡコマンドが実行され、NVIDIAのＧＰＵで処理がされる。

　本実施形態では、Ｃ言語同様、Pythonのループ文に対して、ＧＰＵ処理可否を遺伝的アルゴリズムにより選択し、適切なオフロードパターンを見つける。

　以下、Cupyを用いた例を記載する。
　図７Ａ－Ｂは、PythonコードをpyCUDAでインタプリットする方法を説明する図であり、図７Ａは、変換元例を示し、図７Ｂは、変換後例（３階層のfor文の一番上をＧＰＵ処理指定する場合）を示す。
　図７Ａに示すように、Pythonのfor文は行列演算として指定される。Cupyは、ＣＵＤＡコマンドを呼び、ＣＵＤＡがＧＰＵを実行する。

　Cupyを用いる場合、Ｃ言語のOpenACCの\pragmaのようにオフロードするfor文を記載する形ではない。
　CupyからＣＵＤＡを介してＧＰＵ処理する際は、ＧＰＵでの並列演算は行列演算である。ＧＰＵ処理する箇所は、図７Ｂに示すように、
〔1～多重 for文の内側にある『配列[添字]』を右辺・左辺に持つ演算式〕は、
"[添字1][添字2]…" の部分を"[範囲開始1:範囲終了1,範囲開始2:範囲終了2, …]" 表現に書き換える。
《添字》が《範囲》に置き換わることによって、式全体が行列演算の式になる。

　以上、<PythonコードをpyCUDAでインタプリットする方法>について説明した。次に、<pyACCを用いる方法>について説明する。

<pyACCを用いる方法>
　ループ文オフロード：Pythonは、上記pyCUDAでなくpyACCというOpenACCを解釈するインタプリタを用いてもよい。その場合は、Ｃ言語と同様にOpenACC文法でループ文のＧＰＵ処理を指定する。以下、pyACCを用いる方法について説明する。

　図８Ａ－Ｂは、pyACC利用時のコードパターンを示す図であり、図８Ａは、pyACC利用時のfor文を示し、図８Ｂは、図８Ａのfor文から作成されるコードパターンを示す。
　図８Ｂに示すコードパターンは、図３のコードパターンに置き換えて用いられる。

　図９Ａ－Ｂは、《ループ文オフロード：Python》のフローチャートであり、図９Ａと図９Ｂは、結合子で繋がれる。また、図８Ｂに示すコードパターンを、図３のコードパターンに置き換えて用いる。
　C/C++向けOpenＡＣＣコンパイラを用いて以下の処理を行う。

<コード解析>
　ステップＳ３０１で、アプリケーションコード分析部１１２（図１参照）は、Pythonアプリケーションプログラムのコード解析を行う。

<ループ文特定>
　ステップＳ３０２で、並列処理指定部１１４（図１参照）は、Pythonアプリケーションプログラムのループ文、参照関係を特定する。

<ループ文ループ回数>
　ステップＳ３０３で、並列処理指定部１１４は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。

<ループ文の並列処理可能性>
　ステップＳ３０４で、並列処理指定部１１４は、各ループ文の並列処理可能性をチェックする。

<ループ文の繰り返し>
　制御部（自動オフロード機能部）１１は、ステップＳ３０５のループ始端とステップＳ１０８のループ終端間で、ステップＳ３０６－Ｓ３０７の処理についてループ文の数だけ繰り返す。
　ステップＳ３０６で、並列処理指定部１１４は、各ループ文に対して、ＧＰＵ処理基盤に応じた手法でＧＰＵ処理を指定してインタプリットする。例えば、pyACC利用時はOpenACCの\pragmaacckernels、Cupy利用時は対象ループの計算を行列計算に変換して指定、pyCUDA直接利用時はＣＵＤＡ文法等を用いる、が挙げられる。

　ステップＳ３０７で、並列処理指定部１１４は、エラー時は、該当for文からは、ＧＰＵ処理を削除する。
　ステップＳ３０９で、並列処理指定部１１４は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。

<指定個体数パターン準備>
　次に、初期値として、並列処理指定部１１４は、指定個体数の遺伝子配列を準備する。ここでは、０と１をランダムに割当てて作成する。
　ステップＳ３１０で、並列処理指定部１１４は、Pythonアプリコードを、遺伝子にマッピングし、指定個体数パターン準備を行う。
　準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをPythonアプリコードに挿入する。

　制御部（自動オフロード機能部）１１は、ステップＳ３１１のループ始端とステップＳ３２０のループ終端間で、ステップＳ３１２－Ｓ３１９の処理について指定世代数繰り返す。
　また、上記指定世代数繰り返しにおいて、さらにステップＳ３１２のループ始端とステップＳ３１７のループ終端間で、ステップＳ３１３－Ｓ３１６の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。

<データ転送指定>
　ステップＳ３１３で、データ転送指定部１１３は、変数参照関係から、ＧＰＵ処理基盤に応じた手法でデータ転送を指定する。

<コンパイル>
　ステップＳ３１４で、並列処理パターン作成部１１５（図１参照）は、遺伝子パターンに応じてディレクティブ指定したPythonアプリコードをＧＰＵ処理基盤でインタプリットする。
　ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。

　ステップＳ３１５で、性能測定部１１６（図１参照）は、ＣＰＵ－ＧＰＵ搭載の検証用マシン１４に、実行ファイルをデプロイする。
　ステップＳ３１６で、性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。

　ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
　ステップＳ３１８で、実行ファイル作成部１１７（図１参照）は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。

　ステップＳ３１９で、実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
　すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。

　ステップＳ３２１で、実行ファイル作成部１１７は、指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当するPythonアプリコード（最高性能の並列処理パターン）を解とする。
　以上、《ループ文オフロード：Python》について説明した、次に、《ループ文オフロード：Java》について説明する。

［《ループ文オフロード：Java》］
　ループ文オフロード：Java では、ループ文オフロードの、コードの分析でJavaを解析するJavaParser等の構文解析ツールを用いて構文解析する。ループと変数の把握については、構文解析ツールの結果を管理する際は、言語に非依存に抽象的に管理できる。ループのＧＰＵ処理有無の遺伝子化についても、言語に非依存である。遺伝子情報のコード化では、遺伝子情報に合わせてＧＰＵで実行するためのコードを作成するため、Javaのラムダ記述でＧＰＵ処理を指定する、あるいは変数転送を指定する。

　実行環境は、Javaのラムダ記述での並列化をＧＰＵに対して行うことができるIBM JDK（登録商標）を用いる。IBM JDKはJavaのラムダ記述に従って並列処理をＧＰＵに対して実行する仮想マシンである。
　性能測定は、言語に合わせて、Jenkins（登録商標）等の自動測定ツールも用いて行う。次世代の遺伝子作成は、性能測定結果に合わせて適合度を設定し交叉等の処理を行うが、言語に非依存である。反復実行と最終解の決定も、言語に非依存である。

　図１０Ａ－Ｂは、IBM JDK 利用時のコードパターンを示す図であり、図１０Ａは、IBM JDK 利用時のfor文を示し、図１０Ｂは、図１０Ａのfor文から作成されるコードパターンを示す。
　図１０Ｂに示すコードパターンは、図３のコードパターンに置き換えて用いられる。

　図１１Ａ－Ｂは、《ループ文オフロード：Java》のフローチャートであり、図１１Ａと図１１Ｂは、結合子で繋がれる。また、図１０Ｂに示すコードパターンを、図３のコードパターンに置き換えて用いる。

<コード解析>
　ステップＳ４０１で、アプリケーションコード分析部１１２（図１参照）は、Javaアプリケーションプログラムのコード解析を行う。

<ループ文特定>
　ステップＳ４０２で、並列処理指定部１１４（図１参照）は、Javaアプリケーションプログラムのループ文、参照関係を特定する。

<ループ文ループ回数>
　ステップＳ４０３で、並列処理指定部１１４は、ベンチマークツールを動作させ、ループ文ループ回数を把握し、閾値振分けする。

<ループ文の並列処理可能性>
　ステップＳ４０４で、並列処理指定部１１４は、各ループ文の並列処理可能性をチェックする。

<ループ文の繰り返し>
　制御部（自動オフロード機能部）１１は、ステップＳ４０５のループ始端とステップＳ４０８のループ終端間で、ステップＳ４０６－Ｓ４０７の処理についてループ文の数だけ繰り返す。
　ステップＳ４０６で、並列処理指定部１１４は、各ループ文に対して、Javaのlambda式を用いて、java.util.Stream.IntStream.range(0,n).parallel()forEach(i -> {});でＧＰＵ処理を指定してコンパイルする。

　ステップＳ４０７で、並列処理指定部１１４は、エラー時は、該当for文からは、java.util.Stream.IntStream.range(0,n).parallel()forEach(i -> {});を削除する。
　ステップＳ４０９で、並列処理指定部１１４は、コンパイルエラーが出ないfor文の数をカウントし、遺伝子長とする。

<指定個体数パターン準備>
　次に、初期値として、並列処理指定部１１４は、指定個体数の遺伝子配列を準備する。ここでは、０と１をランダムに割当てて作成する。
　ステップＳ４１０で、並列処理指定部１１４は、Javaアプリコードを、遺伝子にマッピングする。０と１がランダムに割当てられた指定個体数の遺伝子配列を遺伝子にマッピングするすることで、指定個体数パターンを準備する。
　準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをJavaアプリコードに挿入する。

　制御部（自動オフロード機能部）１１は、ステップＳ４１１のループ始端とステップＳ４２０のループ終端間で、ステップＳ４１２－Ｓ４１９の処理について指定世代数繰り返す。
　また、上記指定世代数繰り返しにおいて、さらにステップＳ４１２のループ始端とステップＳ４１７のループ終端間で、ステップＳ４１３－Ｓ４１６の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。

<データ転送指定>
　ステップＳ４１３で、データ転送指定部１１３は、変数参照関係から、Javaの記述でデータ転送を指定する。

<コンパイル>
　ステップＳ４１４で、並列処理パターン作成部１１５（図１参照）は、遺伝子パターンに応じてディレクティブ指定したJavaアプリコードをIBM JDKでビルドする。
　ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。

　ステップＳ４１５で、性能測定部１１６（図１参照）は、ＣＰＵ-ＧＰＵ搭載の検証用マシン１４に、実行ファイルをデプロイする。
　ステップＳ４１６で、性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。

　ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
　ステップＳ４１８で、実行ファイル作成部１１７（図１参照）は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。

　ステップＳ４１９で、実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
　すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。

　ステップＳ４２１で、実行ファイル作成部１１７は、指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当するJavaアプリコード（最高性能の並列処理パターン）を解とする。
　以上、第１の実施形態（「ループ文オフロード」）について説明した。

（第２の実施形態）
　第２の実施形態は、機能ブロックオフロードについて記載する。
　図１２～図１３を参照して、機能ブロックオフロードの全体構成および動作を説明し、以下、機能ブロックオフロード：共通（図１４，図１５）、機能ブロックオフロード：Ｃ言語（図１６，図１７）、機能ブロックオフロード：Python（図１８，図１９）、機能ブロックオフロード：Java（図２０，図２１）を順に説明する。

　図１３は、本発明の第２の実施形態に係るオフロードサーバ２００の構成例を示す機能ブロック図である。図２と同一構成部分には、同一符号を付して重複箇所の説明を省略する。
　オフロードサーバ２００は、アプリケーションプログラムの特定処理をアクセラレータに自動的にオフロードする装置である。

　図１２に示すように、オフロードサーバ２００は、制御部２１０と、入出力部１２と、記憶部１３０と、検証用マシン１４ (アクセラレータ検証用装置)と、を含んで構成される。

　入出力部１２は、各機器等との間で情報の送受信を行うための通信インタフェースと、タッチパネルやキーボード等の入力装置や、モニタ等の出力装置との間で情報の送受信を行うための入出力インタフェースとから構成される。

　記憶部１３０は、ハードディスクやフラッシュメモリ、ＲＡＭ（Random Access Memory）等により構成され、制御部２１０の各機能を実行させるためのプログラム（オフロードプログラム）や、制御部２１０の処理に必要な情報（例えば、中間言語ファイル(Intermediate file)１３２）が一時的に記憶される。

　記憶部１３は、コードパターンＤＢ（Code pattern database）２３０（後記）、テストケースＤＢ（Test case database）１３１を備える。

　テストケースＤＢ１３１には、性能試験項目が格納される。テストケースＤＢ１３１は、高速化するアプリケーションの性能を測定するような試験を行うための情報が格納される。例えば、画像分析処理の深層学習アプリケーションであれば、サンプルの画像とそれを実行する試験項目である。

　検証用マシン１４は、環境適応ソフトウェアの検証用環境として、ＣＰＵ（Central Processing Unit）、ＧＰＵ、ＦＰＧＡを備える。

<コードパターンＤＢ２３０>
・ＧＰＵライブラリ、ＩＰコアの記憶
　コードパターンＤＢ２３０は、ＧＰＵやＦＰＧＡ等にオフロード可能なライブラリおよびＩＰコア（後記）を記憶する。すなわち、コードパターンＤＢ２３０は、後記<処理Ｂ－１>のために、特定のライブラリ、機能ブロックを高速化するＧＰＵ用ライブラリ（ＧＰＵライブラリ）やＦＰＧＡ用ＩＰコア（ＩＰコア）とそれに関連する情報を保持する。例えば、コードパターンＤＢ２３０は、ＦＦＴ等算術計算等のライブラリリスト（外部ライブラリリスト）を保持する。

・ＣＵＤＡライブラリの記憶
　コードパターンＤＢ２３０は、ＧＰＵライブラリとして、例えばＣＵＤＡライブラリと当該ＣＵＤＡライブラリを利用するためのライブラリ利用手順とを記憶する。すなわち、後記<処理Ｃ－１>において、置換するライブラリやＩＰコアをＧＰＵやＦＰＧＡに実装し、ホスト側（ＣＰＵ）プログラムと繋ぐ場合、ライブラリ利用手順も含めて登録しておき、その手順に従って利用する。例えば、ＣＵＤＡライブラリでは、Ｃ言語コードからＣＵＤＡライブラリを利用する手順がライブラリとともに公開されているため、コードパターンＤＢ２３０にライブラリ利用手順も含めて登録しておく。

・クラス、構造体の記憶
　コードパターンＤＢ２３０は、ホストで計算する場合に記述が同様になる処理のクラスまたは構造体を記憶する。すなわち、後記<処理Ｂ－２>において、登録されていないライブラリ呼び出し以外の機能処理を検出するため、構文解析にてソースコードの定義記述からクラス、構造体等を検出する。コードパターンＤＢ２３０は、後記<処理Ｂ－２>のために、ホストで計算する場合に記述が同様になる処理のクラスまたは構造体を登録しておく。なお、クラスまたは構造体の機能処理に対して、高速化するライブラリやＩＰコアがあることは、類似性検出ツール（後記）で検出する。

・OpenCLコードの記憶
　コードパターンＤＢ２３０は、ＩＰコア関連の情報としてOpenCLコードを記憶する。コードパターンＤＢ２３０に、OpenCLコードを記憶しておくことで、OpenCLコードから、OpenCLインタフェースを用いたＣＰＵとＦＰＧＡの接続および、ＦＰＧＡへのＩＰコア実装が、XilinxやIntel等のＦＰＧＡベンダの高位合成ツール（後記）を介して行うことができる。

<制御部２１０>
　制御部２１０は、オフロードサーバ２００全体の制御を司る自動オフロード機能部（Automatic Offloading function）であり、記憶部１３０に格納されたプログラム（オフロードプログラム）を不図示のＣＰＵが、ＲＡＭに展開し実行することにより実現される。

　特に、制御部２１０は、ＣＰＵ向けの既存プログラムコードの中にＦＰＧＡやＧＰＵへオフロードすることで処理を高速化できる機能ブロックを検出し、検出した機能ブロックをＧＰＵ向けライブラリやＦＰＧＡ向けＩＰコア等に置き換えることで高速化をする機能ブロックのオフロード処理を行う。

　制御部２１０は、アプリケーションコード指定部（Specify application code）１１１と、アプリケーションコード分析部（Analyze application code）１１２と、置換機能検出部２１３と、置換処理部２１４と、オフロードパターン作成部２１５と、性能測定部１１６と、実行ファイル作成部１１７と、本番環境配置部（Deploy final binary files to production environment）１１８と、性能測定テスト抽出実行部（Extract performance test cases and run automatically）１１９と、ユーザ提供部（Provide price and performance to a user to judge）１２０と、を備える。

<アプリケーションコード分析部１１２>
　アプリケーションコード分析部１１２は、後記<処理Ａ－１>において、アプリケーションプログラムのソースコードを分析して、当該ソースコードに含まれる外部ライブラリの呼び出しを検出する。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　上述したコード分析は、オフロードするデバイスを想定した分析が必要になるため、一般化は難しい。ただし、ループ文や変数の参照関係等のコードの構造を把握したり、機能ブロックとしてＦＦＴ処理を行う機能ブロックであることや、ＦＦＴ処理を行うライブラリを呼び出している等を把握することは可能である。機能ブロックの判断は、オフロードサーバが自動判断することは難しい。これもDeckard等の類似性検出ツールを用いて類似度判定等で把握することは可能である。ここで、Clangは、C/C++向けツールであるが、解析する言語に合わせたツールを選ぶ必要がある。

　また、アプリケーションコード分析部１１２は、後記<処理Ａ－２>において、ソースコードからクラスまたは構造体のコードを検出する。

<置換機能検出部２１３>
　置換機能検出部２１３は、後記<処理Ｂ－１>において、検出された呼び出しをキーにして、コードパターンＤＢ２３０からＧＰＵライブラリおよびＩＰコアを取得する。具体的には、置換機能検出部２１３は、検出したライブラリ呼び出しに対して、ライブラリ名をキーとして、コードパターンＤＢ２３０と照合することで、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能処理を抽出する。

　ここで、コードパターンＤＢ２３０は、ＧＰＵライブラリとして、例えばＣＵＤＡライブラリと当該ＣＵＤＡライブラリを利用するためのライブラリ利用手順とを記憶している。そして、置換機能検出部２１３は、ライブラリ利用手順をもとに、コードパターンＤＢ２３０からＣＵＤＡライブラリを取得する。

　置換機能検出部２１３は、後記<処理Ｂ－２>において、検出されたクラスまたは構造体（後記）の定義記述コードをキーにして、コードパターンＤＢ２３０からＧＰＵライブラリおよびＩＰコアを取得する。具体的には、置換機能検出部２１３は、コピーコードやコピー後変更した定義記述コードを検出する類似性検出ツールを用いて、置換元コードに含まれるクラスや構造体に対して、コードパターンＤＢ２３０から類似のクラスまたは構造体に紐づいて管理されているＧＰＵ、ＦＰＧＡにオフロードできるＧＰＵライブラリおよびＩＰコアを抽出する。

<置換処理部２１４>
　置換処理部２１４は、後記<処理Ｃ－１>において、アプリケーションプログラムのソースコードの置換元の処理記述を、置換機能検出部２１３が取得した置換先のライブラリおよびＩＰコアの処理記述に置換する。具体的には、置換処理部２１４は、抽出したオフロード可能処理を、ＧＰＵ向けのライブラリやＦＰＧＡ向けのＩＰコア等に置換する。
　また、置換処理部２１４は、置換したライブラリおよびＩＰコアの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵやＦＰＧＡ等にオフロードする。具体的には、置換処理部２１４は、ＧＰＵ向けのライブラリやＦＰＧＡ向けのＩＰコア等に置換した機能ブロックを、ＣＰＵプログラムとのインタフェースを作成することでオフロードする。置換処理部２１４は、ＣＵＤＡ,OpenCL等の中間言語ファイル１３２を出力する。

　置換処理部２１４は、後記<処理Ｃ－２>において、アプリケーションプログラムのソースコードの置換元の処理記述を、取得したライブラリおよびＩＰコアの処理記述に置換するとともに、置換元と置換先で引数、戻り値の数または型が異なる場合に、その確認を通知する。

　置換処理部２１４は、《機能ブロックオフロード：Ｃ言語》では、ＣＵＤＡのライブラリ呼び出しを、ＰＧＩコンパイラに指定する。
　置換処理部２１４は、《機能ブロックオフロード：Python》では、ＣＵＤＡのライブラリ呼び出しを、pyCudaで指定する。
　置換処理部２１４は、《機能ブロックオフロード：Java》では、ＣＵＤＡのライブラリ呼び出しを、Jcudaで指定する。

<オフロードパターン作成部２１５>
　オフロードパターン作成部２１５は、１以上のオフロードするパターンを作成する。具体的には、ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出する。

　ここで、コードパターンＤＢ２３０は、ＩＰコア関連の情報としてOpenCLコードを記憶している。オフロードパターン作成部２１５は、ＦＰＧＡ等のＰＬＤにオフロードする場合は、OpenCLコードをもとにOpenCLインタフェースを用いてホストとＰＬＤとを接続するとともに、OpenCLコードをもとにＰＬＤへのＩＰコアの実装を行う。

　OpenCLのＡＰＩに沿う、カーネルプログラムとホストプログラムのインタフェース記述について述べる。なお、下記説明は、後記［処理Ｃ］（ホスト側とのインタフェースの整合）の<処理Ｃ－１>の具体例に対応する。

　OpenCLのＣ言語の文法に沿って作成したカーネルは、OpenCLのＣ言語のランタイムＡＰＩを利用して、作成するホスト（例えば、ＣＰＵ）側のプログラムによりデバイス（例えば、ＦＰＧＡ）で実行される。カーネル関数hello()をホスト側から呼び出す部分は、OpenCLランタイムＡＰＩの一つであるclEnqueueTask()を呼び出すことである。
　ホストコードで記述するOpenCLの初期化、実行、終了の基本フローは、下記ステップ１～１３である。このステップ１～１３のうち、ステップ１～１０がカーネル関数hello()をホスト側から呼び出すまでの手続（準備）であり、ステップ１１でカーネルの実行となる。

１．プラットフォーム特定
　OpenCLランタイムＡＰＩで定義されているプラットフォーム特定機能を提供する関数clGetPlatformIDs()を用いて、OpenCLが動作するプラットフォームを特定する。

２．デバイス特定
　OpenCLランタイムＡＰＩで定義されているデバイス特定機能を提供する関数clGetDeviceIDs()を用いて、プラットフォームで使用するＧＰＵ等のデバイスを特定する。

３．コンテキスト作成
　OpenCLランタイムＡＰＩで定義されているコンテキスト作成機能を提供する関数clCreateContext()を用いて、OpenCLを動作させる実行環境となるOpenCLコンテキストを作成する。

４．コマンドキュー作成
　OpenCLランタイムＡＰＩで定義されているコマンドキュー作成機能を提供する関数clCreateCommandQueue()を用いて、デバイスを制御する準備であるコマンドキューを作成する。OpenCLでは、コマンドキューを通して、ホストからデバイスに対する働きかけ（カーネル実行コマンドやホスト－デバイス間のメモリコピーコマンドの発行）を実行する。

５．メモリオブジェクト作成
　OpenCLランタイムＡＰＩで定義されているデバイス上にメモリを確保する機能を提供する関数clCreateBuffer()を用いて、ホスト側からメモリオブジェクトを参照できるようにするメモリオブジェクトを作成する。

６．カーネルファイル読み込み
　デバイスで実行するカーネルは、その実行自体をホスト側のプログラムで制御する。このため、ホストプログラムは、まずカーネルプログラムを読み込む必要がある。カーネルプログラムには、OpenCLコンパイラで作成したバイナリデータや、OpenCL　Ｃ言語で記述されたソースコードがある。このカーネルファイルを読み込む（記述省略）。なお、カーネルファイル読み込みでは、OpenCLランタイムＡＰＩは使用しない。

７．プログラムオブジェクト作成
　OpenCLでは、カーネルプログラムをプログラムオブジェクトとして認識する。この手続きがプログラムオブジェクト作成である。
　OpenCLランタイムＡＰＩで定義されているプログラムオブジェクト作成機能を提供する関数clCreateProgramWithSource()を用いて、ホスト側からメモリオブジェクトを参照できるようにするプログラムオブジェクトを作成する。カーネルプログラムのコンパイル済みバイナリ列から作成する場合は、clCreateProgramWithBinary()を使用する。

８．ビルド
　ソースコードとして登録したプログラムオブジェクトを　OpenCL　Ｃコンパイラ・リンカを使いビルドする。
　OpenCLランタイムＡＰＩで定義されているOpenCL　Ｃコンパイラ・リンカによるビルドを実行する関数clBuildProgram()を用いて、プログラムオブジェクトをビルドする。なお、clCreateProgramWithBinary()でコンパイル済みのバイナリ列からプログラムオブジェクトを生成した場合、このコンパイル手続は不要である。

９．カーネルオブジェクト作成
　OpenCLランタイムＡＰＩで定義されているカーネルオブジェクト作成機能を提供する関数clCreateKernel()を用いて、カーネルオブジェクトを作成する。１つのカーネルオブジェクトは、１つのカーネル関数に対応するので、カーネルオブジェクト作成時には、カーネル関数の名前(hello)を指定する。また、複数のカーネル関数を１つのプログラムオブジェクトとして記述した場合、１つのカーネルオブジェクトは、１つのカーネル関数に１対１で対応するので、clCreateKernel()を複数回呼び出す。

１０．カーネル引数設定
　OpenCLランタイムＡＰＩで定義されているカーネルへ引数を与える（カーネル関数が持つ引数へ値を渡す）機能を提供する関数clSetKernel()を用いて、カーネル引数を設定する。
　以上、上記ステップ１～１０で準備が整い、ホスト側からデバイスでカーネルを実行するステップ１１に入る。

１１．カーネル実行
　カーネル実行（コマンドキューへ投入）は、デバイスに対する働きかけとなるので、コマンドキューへのキューイング関数となる。
　OpenCLランタイムＡＰＩで定義されているカーネル実行機能を提供する関数clEnqueueTask()を用いて、カーネルhelloをデバイスで実行するコマンドをキューイングする。カーネルhelloを実行するコマンドがキューイングされた後、デバイス上の実行可能な演算ユニットで実行されることになる。

１２．メモリオブジェクトからの読み込み
　OpenCLランタイムＡＰＩで定義されているデバイス側のメモリからホスト側のメモリへデータをコピーする機能を提供する関数clEnqueueReadBuffer()を用いて、デバイス側のメモリ領域からホスト側のメモリ領域にデータをコピーする。また、ホスト側からデバイス側のメモリへデータをコピーする機能を提供する関数clEnqueueWrightBuffer()を用いて、ホスト側のメモリ領域からデバイス側のメモリ領域にデータをコピーする。なお、これらの関数は、デバイスに対する働きかけとなるので、一度コマンドキューへコピーコマンドがキューイングされてからデータコピーが始まることになる。

１３．オブジェクト解放
　最後に、ここまでに作成してきた各種オブジェクトを解放する。
　以上、OpenCL　Ｃ言語に沿って作成されたカーネルの、デバイス実行について説明した。

<性能測定部１１６>
　性能測定部１１６は、作成された処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン１４に配置し、ＧＰＵやＦＰＧＡ等にオフロードした際の性能測定用処理を実行する。
　性能測定部１１６は、バイナリファイル配置部（Deploy binary files）１１６ａを備える。バイナリファイル配置部１１６ａは、ＧＰＵやＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれるバイナリファイルをデプロイ(配置)する。

　性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、バイナリファイル配置部１１６ａに戻す。この場合、性能測定部１１６は、抽出された別の処理パターンを用いて、抽出された中間言語をもとに、性能測定を試行する（後記図１３の符号ｇ参照）。

　性能測定の具体例について述べる。
　オフロードパターン作成部２１５は、ＧＰＵやＦＰＧＡにオフロード可能な機能ブロックをオフロードする処理パターンを作成し、作成された処理パターンの中間言語を、実行ファイル作成部１１７がコンパイルする。性能測定部１１６は、コンパイルされたプログラムの性能を測定する（「１回目の性能測定」）。

　そして、オフロードパターン作成部２１５は、性能測定された中でＣＰＵに比べ高性能化された処理パターンをリスト化する。オフロードパターン作成部２１５は、リストの処理パターンを組み合わせてオフロードする新たな処理パターンを作成する。オフロードパターン作成部２１５は、組み合わせたオフロード処理パターンと中間言語を作成し、中間言語を、実行ファイル作成部１１７がコンパイルする。
　性能測定部１１６は、コンパイルされたプログラムの性能を測定する（「２回目の性能測定」）。

<実行ファイル作成部１１７>
　実行ファイル作成部１１７は、オフロードする処理パターンの中間言語をコンパイルして実行ファイルを作成する。一定数繰り返された、性能測定結果をもとに、１以上の処理パターンから最高処理性能の処理パターンを選択し、最高処理性能の処理パターンをコンパイルして最終実行ファイルを作成する。

<性能測定テスト抽出実行部１１９>
　性能測定テスト抽出実行部１１９は、実行ファイル配置後、テストケースＤＢ１３１から性能試験項目を抽出し、性能試験を実行する。
　性能測定テスト抽出実行部１１９は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースＤＢ１３１から抽出し、抽出した性能試験を自動実行する。

<ユーザ提供部１２０>
　ユーザ提供部１２０は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する（「価格・性能等の情報のユーザへの提供」）。テストケースＤＢ１３１には、アプリケーションの性能を測定する試験を自動で行うためのデータが格納されている。ユーザ提供部１２０は、テストケースＤＢ１３１の試験データを実行した結果と、システムに用いられるリソース（仮想マシンや、ＦＰＧＡインスタンス、ＧＰＵインスタンス等）の各単価から決まるシステム全体の価格をユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。

　以下、上述のように構成されたオフロードサーバ２００の機能ブロックオフロード処理について説明する。

　上記、機能ブロックのオフロードの処理の概要と考慮点について説明する。
　ＦＰＧＡに関しては、ハードウェア回路設計に多大な時間がかかることもあり、一度設計した機能を、ＩＰコア（Intellectual Property Core）という形で再利用可能にすることが多い。ＩＰコアとは、ＦＰＧＡ、ＩＣ、ＬＳＩなどの半導体を構成するための部分的な回路情報であり、特に機能単位でまとめられている。ＩＰコアは、暗号化／復号化処理、ＦＦＴ（Fast Fourier Transform）等の算術演算、画像処理、音声処理等が代表的な機能例である。ＩＰコアは、ライセンス料を支払うものが多いが、一部はフリーで提供されているものもある。

　第２の実施形態では、ＦＰＧＡに関しては、ＩＰコアを自動オフロードに利用する。また、ＧＰＵに関しては、ＩＰコアという言い方ではないものの、ＦＦＴ、線形代数演算等が代表的な機能例であり、ＣＵＤＡを用いて実装されたcuFFTやcuBLAS等がＧＰＵ向けライブラリとしてフリーで提供されている。本第２の実施形態では、ＧＰＵに関してこれらのライブラリを活用する。

　本第２の実施形態では、ＣＰＵ向けに作られた既存プログラムコードの中で、ＦＦＴ処理等、ＧＰＵ、ＦＰＧＡにオフロードすることで高速化できるような機能ブロックが含まれる場合に、ＧＰＵ向けライブラリやＦＰＧＡ向けＩＰコア等に置き換えることでの高速化を図る。

［機能ブロックのオフロード処理概要］
　第２の実施形態のオフロードサーバ２００は、環境適応ソフトウェアの要素技術としてユーザアプリケーションロジックのＧＰＵ、ＦＰＧＡ自動オフロードに適用した例である。
　図１３は、オフロードサーバ２００の機能ブロックのオフロード処理を示す図である。
　図１３に示すように、オフロードサーバ２００は、環境適応ソフトウェアの要素技術に適用される。オフロードサーバ２００は、制御部（自動オフロード機能部）１１と、コードパターンＤＢ２３０、テストケースＤＢ１３１と、中間言語ファイル１３２と、検証用マシン１４と、を有している。
　オフロードサーバ２００は、ユーザが利用するアプリケーションコード（Application code）１３０を取得する。

　ユーザは、例えば、各種デバイス（Device）１５１、ＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３、ＣＰＵを有する装置１５４を利用する。オフロードサーバ２００は、機能処理をＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３のアクセラレータに自動オフロードする。

　以下、図１３のステップ番号を参照して各部の動作を説明する。
<ステップＳ３１：Specify application code>
　ステップＳ３１において、アプリケーションコード指定部１１１（図１２参照）は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部１１２に渡す。

<ステップＳ３２：Analyze application code>（コード分析）
　ステップＳ３２において、アプリケーションコード分析部１１２（図１２参照）は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

<ステップＳ３３：Extract offloadable area>（オフロード可能処理抽出）
　ステップＳ３３において、置換機能検出部２１３（図１２参照）は、把握したライブラリ呼び出しや機能処理について、コードパターンＤＢ２３０と照合することで、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能処理を抽出する。

<ステップＳ３４：Output intermediate file>（オフロード用中間ファイル出力）
　ステップＳ３４において、置換処理部２１４（図１２参照）は、抽出したオフロード可能処理を、ＧＰＵ向けのライブラリやＦＰＧＡ向けのＩＰコア等に置換する。置換処理部２１４は、ＧＰＵ向けのライブラリやＦＰＧＡ向けのＩＰコア等に置換した機能ブロックを、ＣＰＵプログラムとのインタフェースを作成することでオフロードする。置換処理部２１４は、ＣＵＤＡ,OpenCL等の中間言語ファイル１３２を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。
　ここで、オフロード可能処理が直ちに高速化につながるか、またコスト効果が十分であるかは分からないので、オフロードパターン作成部２１５は、後述する検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出する。

<ステップＳ２１：Deploy binary files>（デプロイ、性能測定試行）
　ステップＳ２１において、バイナリファイル配置部１１６ａ（図１２参照）は、ＧＰＵ、ＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイする。バイナリファイル配置部１１６ａは、配置したファイルを起動し、想定するテストケースを実行して、オフロードした際の性能を測定する。

<ステップＳ２２：Measure performances>
　ステップＳ２２において、性能測定部１１６（図１２参照）は、配置したファイルを実行し、オフロードした際の性能を測定する。

　図１３の符号ｇに示すように、制御部２１０は、上記ステップＳ１２乃至ステップＳ２２を繰り返し実行する。制御部２１０の自動オフロード機能をまとめると、下記である。すなわち、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。置換機能検出部２１３は、検出したライブラリ呼び出しや機能処理について、コードパターンＤＢ２３０と照合することで、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能処理を抽出する。置換処理部２１４は、抽出したオフロード可能処理を、ＧＰＵ向けのライブラリやＦＰＧＡ向けのＩＰコア等に置換する。そして、オフロードパターン作成部２１５は、ＧＰＵ向けのライブラリやＦＰＧＡ向けのＩＰコア等に置換した機能ブロックを、ＣＰＵプログラムとのインタフェースを作成することでオフロードする。

　上記ステップＳ１１～ステップＳ２５は、ユーザのサービス利用のバックグラウンドで行われ、例えば、仮利用の初日の間に行う等を想定している。

　上記したように、オフロードサーバ２００の制御部（自動オフロード機能部）２１０は、環境適応ソフトウェアの要素技術に適用した場合、機能処理のオフロードのため、ユーザが利用するアプリケーションプログラムのソースコードから、オフロードする領域を抽出して中間言語を出力する（ステップＳ３１～ステップＳ３４）。制御部２１０は、中間言語から導かれる実行ファイルを、検証用マシン１４に配置実行し、オフロード効果を検証する（ステップＳ２１～ステップＳ２２）。検証を繰り返し、適切なオフロード領域を定めたのち、制御部２１０は、実際にユーザに提供する本番環境に、実行ファイルをデプロイし、サービスとして提供する（ステップＳ２３～ステップＳ２５）。

　一般に、性能に関しては、最大性能になる設定を一回で自動発見するのは難しい。このため、オフロードパターンを、性能測定を検証環境で何度か繰り返すことにより試行し、高速化できるパターンを見つけることが本発明の特徴である。

［機能ブロックのオフロード処理詳細］
　機能ブロックのオフロードについては、機能ブロックの検出（以下、「処理Ａ」という）、その機能ブロックがオフロード用の既存ライブラリ／ＩＰコア等があるかを検出（以下、「処理Ｂ」という）、機能ブロックをライブラリ／ＩＰコア等と置換した際にホスト側とのインタフェースの整合（以下、「処理Ｃ」という）、の３つ要素を考慮する必要がある。上記３つ要素の考慮点に従い、機能ブロックのオフロード処理について詳細に述べる。

［処理Ａ］（機能ブロックの検出）
　「処理Ａ」（機能ブロックの検出）は、ライブラリの関数呼び出しを行い、ライブラリの関数呼び出しを機能ブロックとする<処理Ａ－１>と、登録されていないライブラリの関数呼び出しである場合、クラス、構造体等を検出して機能ブロックとする<処理Ａ－２>と、に分けられる。すなわち、<処理Ａ－１>は、既存のライブラリの関数呼び出しを検出して機能ブロックとするものであり、<処理Ａ－２>は、<処理Ａ－１>において機能ブロックを検出しない場合に、クラスまたは構造体から機能ブロックを抽出するものである。

　<処理Ａ－１>
　アプリケーションコード分析部１１２は、構文解析を用いて、ソースコードから外部のライブラリの関数呼び出しを行っていることを検知する。詳細には、下記の通りである。コードパターンＤＢ２３０は、ＦＦＴ等算術計算等のライブラリリストを保持している。アプリケーションコード分析部１１２は、ソースコードを構文解析し、コードパターンＤＢ２３０が保持しているライブラリリストと照合して、外部のライブラリの関数呼び出しを行っていることを検知する。

　<処理Ａ－２>
　アプリケーションコード分析部１１２は、登録されていないライブラリ呼び出し以外の機能処理を機能ブロックとして検出するため、構文解析を用いて、ソースコードの定義記述からクラスまたは構造体の機能処理を検出する。アプリケーションコード分析部１１２は、例えば、Ｃ言語のstructを使って定義されるいくつかの変数をひとまとまりにした型である構造体（structure）や、インスタンス化したオブジェクトの型が値型である構造体に対して参照型であるクラス（class）を検出する。また、アプリケーションコード分析部１１２は、例えばJava（登録商標）において構造体に代替使用されるクラスを検出する。

［処理Ｂ］（オフロード可能機能の検出）
　［処理Ｂ］（オフロード可能機能の検出）は、<処理Ａ－１>を受け、コードパターンＤＢ２３０を参照して置換可能ＧＰＵライブラリ、ＩＰコアを取得する<処理Ｂ－１>と、<処理Ａ－２>を受け、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換する<処理Ｂ－２>と、に分けられる。すなわち、<処理Ｂ－１>は、ライブラリ名をキーに、コードパターンＤＢ２３０から置換可能ＧＰＵライブラリ、ＩＰコアを取得するものである。<処理Ｂ－２>は、クラス、構造体等のコードをキーに、置換可能ＧＰＵライブラリ・ＩＰコアを検出し、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換するものである。

　処理Ｂの前提として、コードパターンＤＢ２３０には、特定のライブラリ、機能ブロックを高速化するＧＰＵ用ライブラリやＦＰＧＡ用ＩＰコアとそれに関連する情報が保持されている。また、コードパターンＤＢ２３０には、置換元のライブラリ、機能ブロックについては、機能名とともにコードや実行ファイルが登録されている。

　<処理Ｂ－１>
　置換機能検出部２１３は、<処理Ａ－１>でアプリケーションコード分析部１１２が検出したライブラリ呼び出しに対して、ライブラリ名をキーに、コードパターンＤＢ２３０を検索し、コードパターンＤＢ２３０から、置換可能ＧＰＵライブラリ（高速化できるＧＰＵ用ライブラリ）やＦＰＧＡ用ＩＰコアを取得する。

　<処理Ｂ－１>の例を記載する。
　置換機能検出部２１３は、例えば、置換元の処理が2D FFTの処理（非特許文献４等にコードがある）であった場合は、その外部ライブラリ名をキーに、2D FFTを処理するＦＰＧＡ処理として、OpenCLコードを検出する（ホストプログラム、カーネルプログラム）等）。なお、OpenCLコードは、コードパターンＤＢ２３０に記憶されている。

　置換機能検出部２１３は、例えば、置換元の処理が2D FFTの処理であった場合は、ＧＰＵライブラリとして検出されたcuFFTの中の関数呼び出しに置換する。なお、ＧＰＵライブラリは、コードパターンＤＢ２３０に記憶されている。

　<処理Ｂ－２>
　置換機能検出部２１３は、<処理Ａ－２>でアプリケーションコード分析部１１２が検出したクラス、構造体等のコードをキーに、コードパターンＤＢ２３０を検索し、コードパターンＤＢ２３０から、類似性検出ツールを用いて置換可能ＧＰＵライブラリ（高速化できるＧＰＵ用ライブラリ）やＦＰＧＡ用ＩＰコアを取得する。類似性検出ツールとは、Deckard等、コピーコードやコピー後変更したコードの検出を対象とするツールである。置換機能検出部２１３が、類似性検出ツールを用いることで、行列計算のコード等、ＣＰＵで計算する場合は記述が同様になる処理や、他者のコードをコピーして変更した処理等を一部検出できる。なお、類似性検出ツールは、新規に独立に作成したようなクラス等については検出が困難となるため対象外である。

　<処理Ｂ－２>の例を記載する。
　置換機能検出部２１３は、置換元ＣＰＵコードに検知されたクラスや構造体に対して、Deckard等の類似性検知ツールを用いて、コードパターンＤＢ２３０に登録された類似クラスや構造体を検索する。例えば、置換元の処理（非特許文献４等にコードがある）が2D FFTのクラスであった場合は、その類似クラスとしてコードパターンＤＢ２３０に登録されたクラスが2D FFTのクラスが検出される。コードパターンＤＢ２３０には、2D FFTをオフロード可能なＩＰコアやＧＰＵライブラリが登録されている。そのため、<処理Ｂ－１>と同様に、2D FFTに対して、OpenCLコード（ホストプログラム、カーネルプログラム等）やＧＰＵライブラリを検出する。

［処理Ｃ］（ホスト側とのインタフェースの整合）
　［処理Ｃ］（ホスト側とのインタフェースの整合）は、<処理Ｃ－１>と、<処理Ｃ－２>とを有する。<処理Ｃ－１>は、<処理Ｂ－１>を受け、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換するとともに、ＧＰＵライブラリ、ＩＰコア呼び出しのためのインタフェース処理を記述する。<処理Ｃ－２>は、<処理Ｂ－２>を受け、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換するとともに、ＧＰＵライブラリ、ＩＰコア呼び出しのためのインタフェース処理を記述する。ここで、上記ＧＰＵライブラリ、ＩＰコア呼び出しのためのインタフェース処理の記述が、「ホスト側とのインタフェースの整合」に対応する。

　<処理Ｃ－１>
　置換処理部２１４は、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換する。そして、置換処理部２１４は、ＧＰＵライブラリ、ＩＰコア呼び出しのためのインタフェース処理を記述し（OpenCL API等）、作成したパターンをコンパイルする。

　<処理Ｃ－１>について、より詳細に説明する。
　置換機能検出部２１３は、<処理Ａ－１>で検出したライブラリ呼び出しに対して、<処理Ｂ－１>で該当するライブラリやＩＰコアを検索している。このため、置換処理部２１４は、置換するライブラリやＩＰコアをＧＰＵやＦＰＧＡに実装し、ホスト側（ＣＰＵ）プログラムと繋ぐインタフェース処理を行う。

　ここで、ＧＰＵ用ライブラリの場合は、ＣＵＤＡ等のライブラリを想定しており、Ｃ言語コードからＣＵＤＡライブラリを利用する手法がライブラリとともに公開されている。そこで、コードパターンＤＢ２３０に、ライブラリ利用手法も含めて登録しておき、置換処理部２１４は、コードパターンＤＢ２３０に登録されたライブラリ利用手法に従って、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリに置換するとともに、ＧＰＵライブラリで利用する関数の呼び出し等の所定記述を行う。

　ＦＰＧＡ用ＩＰコアの場合は、ＨＤＬ（Hardware Description Language）等が想定される。この場合、ＩＰコア関連の情報としてOpenCLコードもコードパターンＤＢ２３０に保持されている。置換処理部２１４は、ＦＰＧＡとのインタフェース処理を、高位合成ツール（例えば、Xilinx Vivado, Intel HLS Compiler等）を介して行うことができる。置換処理部２１４は、例えば、OpenCLコードから、OpenCLインタフェースを用いたＣＰＵとＦＰＧＡの接続を、高位合成ツールを介して行う。同様に、置換処理部２１４は、ＦＰＧＡへのＩＰコア実装を、XilinxやIntel等のＦＰＧＡベンダの高位合成ツールを介して行う。

　<処理Ｃ－２>
　置換処理部２１４は、アプリコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換する。そして、置換処理部２１４は、置換元と置換先で引数や戻り値の数や型が異なる場合に、ユーザに確認し、ＧＰＵライブラリ、ＩＰコア呼び出しのためのインタフェース処理を記述（OpenCL API等）するとともに、作成したパターンをコンパイルする。すなわち、<処理Ｃ－２>では、置換処理部２１４は、<処理Ａ－２>で検出したクラス、構造体等に対して、<処理Ｂ－２>で高速化できるライブラリやＩＰコアを検索している。このため、置換処理部２１４は、<処理Ｃ－２>では該当するライブラリやＩＰコアをＧＰＵやＦＰＧＡに実装する。

　<処理Ｃ－２>について、より詳細に説明する。
　<処理Ｃ－１>では、特定のライブラリ呼び出しに対して高速化するライブラリやＩＰコアであるため、インタフェース部分の生成等は必要になるものの、ＧＰＵ、ＦＰＧＡとホスト側プログラムの想定する引数、戻り値の数や型は合っていた。しかし、<処理Ｂ－２>は、類似性等で判断しているため、引数や戻り値の数や型等の基本的な部分が合っている保証はない。ライブラリやＩＰコアは、既存ノウハウであり、引数、戻り値の数や型が合っていない場合であっても、変更が頻繁にできるものではない。このため、オフロードを依頼するユーザに対して、元のコードの引数や戻り値の数や型について、ライブラリやＩＰコアに合わせて変更するか否かを確認する。そして、確認了承後にオフロード性能試験を試行する。

　型の違いについて、floatとdouble等キャストすればよいだけであれば、処理パターン作成時にキャストする処理を追加し、特にユーザ確認せずに性能測定試行に入ってもよい。また、引数や戻り値で、元のプログラムとライブラリやＩＰコアで数が異なる場合、例えば、ＣＰＵプログラムで引数１，２が必須で引数３がオプションであり、ライブラリやＩＰコアで引数１，２が必須の場合等は、オプション引数３は省略しても問題はない。このような場合は、ユーザに確認せず、処理パターン作成時にオプション引数は自動で無しとして扱うなどしてもよい。なお、引数や戻り値の数や型が完全に合っている場合は、<処理Ｃ－１>と同様の処理でよい。

［《機能ブロックオフロード：共通》フローチャート］
　次に、図１４および図１５を参照してオフロードサーバ２００の《機能ブロックオフロード：共通》の動作概要を説明する。

・<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>のフローチャート
　図１４は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、《機能ブロックオフロード：共通》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。
　ステップＳ５０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ５０２で置換機能検出部２１３（図１２参照）は、アプリケーションプログラムの外部ライブラリ呼び出しを検出する。

　ステップＳ５０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、ライブラリ名をキーに、置換可能ＧＰＵライブラリを取得する。具体的には、置換機能検出部２１３は、把握した外部ライブラリ呼び出しについて、コードパターンＤＢ２３０と照合することで、検出した置換可能ＧＰＵライブラリ・ＩＰコアを、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能な機能ブロックとして取得する。

　ステップＳ５０４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリの処理記述に置換する。

　ステップＳ５０５で置換処理部２１４は、置換したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ５０６で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。
　ステップＳ５０７で実行ファイル作成部１１７は、作成したパターンをコンパイルまたはインタプリットする。

　ステップＳ５０８で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。
　ステップＳ５０９で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ５１０で実行ファイル作成部１１７は、作成した組合せパターンをコンパイルまたはインタプリットする。
　ステップＳ５１１で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ５１２で本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

・<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>のフローチャート
　図１５は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。なお、<処理Ａ－２>からの処理は、<処理Ａ－１>からの処理と並行して行えばよい。
　ステップＳ６０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ６０２で置換機能検出部２１３（図１２参照）は、ソースコードからクラスまたは構造体の定義記述コードを検出する。

　ステップＳ６０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能ＧＰＵライブラリを取得する。

　ステップＳ６０４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリ処理記述に置換する。

　ステップＳ６０５で置換処理部２１４は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。

　ステップＳ６０６で置換機能検出部２１３は、置換または確認したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ６０７で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ６０８で実行ファイル作成部１１７は、作成したパターンをコンパイルまたはインタプリットする。
　ステップＳ６０９で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。

　ステップＳ６１０で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ６１１で実行ファイル作成部１１７は、作成した組合せパターンをコンパイルまたはインタプリットする。
　ステップＳ６１２で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ６１３で本番環境配置部１１８は、本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

［《機能ブロックオフロード：Ｃ言語》］
　機能ブロックオフロード：Ｃ言語の、コードの分析では、Ｃ言語を解析するClang（登録商標）等の構文解析ツールを用いて構文解析する。機能ブロックの把握については、構文解析ツールの結果を用いて、次処理のマッチング探索に用いるため、言語に非依存の機能ブロックとして管理する。

　オフロード可能機能ブロックの探索では、ライブラリ等の名前一致でのマッチングと、Deckard（登録商標）等のＣ言語機能ブロックの類似性検出ツールを用いた類似性検知による、探索が行われる。オフロード機能ブロックへの置換は、ＣＵＤＡライブラリ呼び出し等、その言語からのオフロード機能利用に合わせた処理に置換する必要がある。

　コンパイルは、ＣＵＤＡライブラリ呼び出し等のＣ言語コードをＰＧＩコンパイラ等でコンパイルする。性能測定は、言語に合わせて、Jenkins（登録商標）等の自動測定ツールも用いて行う。オフロード可能機能ブロックが複数の際は反復実行され、最高性能のパターンが最終解として決定される。

　このように、機能ブロック文オフロードでは、処理に関しては、機能ブロックの管理と機能ブロックの名前一致でのマッチングについては言語に非依存に適用できる。

［《機能ブロックオフロード：Ｃ言語》フローチャート］
　次に、図１６および図１７を参照してオフロードサーバ２００の《機能ブロックオフロード：Ｃ言語》の動作概要を説明する。

・<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>のフローチャート
　図１６は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、《機能ブロックオフロード：Ｃ言語》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。
　ステップＳ７０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ７０２で置換機能検出部２１３（図１２参照）は、アプリケーションの外部ライブラリ呼び出しを検出する。

　ステップＳ７０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、ライブラリ名をキーに、置換可能ＧＰＵライブラリを取得する。具体的には、置換機能検出部２１３は、把握した外部ライブラリ呼び出しについて、コードパターンＤＢ２３０と照合することで、検出した置換可能ＧＰＵライブラリを、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能な機能ブロックとして取得する。

　ステップＳ７０４で置換処理部２１４は、アプリケーションコードの置換元の処理記述を、置換先のＧＰＵライブラリの処理記述に置換する。

　ステップＳ７０５で置換処理部２１４は、置換したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ７０６で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ７０７で置換処理部２１４は、ＣＵＤＡのライブラリ呼び出しを、ＰＧＩコンパイラに指定する。

　ステップＳ７０８で実行ファイル作成部１１７は、作成したパターンをＰＧＩコンパイラでコンパイルする。

　ステップＳ７０９で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。
　ステップＳ７１０で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ７１１で実行ファイル作成部１１７は、作成した組合せパターンをＰＧＩコンパイラでコンパイルする。

　ステップＳ７１２で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ７１３で本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

・<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>のフローチャート
　図１７は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。なお、<処理Ａ－２>からの処理は、<処理Ａ－１>からの処理と並行して行えばよい。
　ステップＳ８０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ８０２で置換機能検出部２１３（図１２参照）は、ソースコードからクラスまたは構造体の定義記述コードを検出する。

　ステップＳ８０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能ＧＰＵライブラリを取得する。

　ステップＳ８０４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリの処理記述に置換する。

　ステップＳ８０５で置換処理部２１４は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。

　ステップＳ８０６で置換機能検出部２１３は、置換または確認したＧＰＵライブラリの置換元の処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ８０７で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ８０８で置換処理部２１４は、ＣＵＤＡのライブラリ呼び出しを、ＰＧＩコンパイラに指定する。

　ステップＳ８０９で実行ファイル作成部１１７は、作成したパターンをＰＧＩコンパイラでコンパイルする。

　ステップＳ８１０で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。

　ステップＳ８１１で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ８１２で実行ファイル作成部１１７は、作成した組合せパターンをＰＧＩコンパイラでコンパイルする。

　ステップＳ８１３で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ８１４で本番環境配置部１１８は、本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

　以上、《機能ブロックオフロード：Ｃ言語》について説明した、次に、《機能ブロックオフロード：Python》について説明する。

［《機能ブロックオフロード：Python》］
　機能ブロックオフロード：Pythonは、機能ブロックオフロードの、コードの分析では、Pythonを解析するast（登録商標）等の構文解析ツールを用いて構文解析する。機能ブロックの把握については、構文解析ツールの結果を用いて、次処理のマッチング探索に用いるため、言語に非依存の機能ブロックとして管理する。

　オフロード可能機能ブロックの探索では、ライブラリ等の名前一致でのマッチングと、CloneDigger（登録商標）等のPython機能ブロックの類似性検出ツールを用いた類似性検知による、探索が行われる。

　オフロード機能ブロックへの置換は、ＧＰＵ処理のpyCUDAでの呼び出し等、その言語からのオフロード機能利用に合わせた処理に置換する必要がある。インタプリタは、ＣＵＤＡに合わせたPythonコードをpyCUDAでインタプリットする。

　性能測定は、言語に合わせて、Jenkins（登録商標）等の自動測定ツールも用いて行う。オフロード可能機能ブロックが複数の際は反復実行され、最高性能のパターンが最終解として決定される。

［《機能ブロックオフロード：Python》フローチャート］
　次に、図１８および図１９を参照してオフロードサーバ２００の《機能ブロック：Python》の動作概要を説明する。

・<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>のフローチャート
　図１８は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、《機能ブロック：Python》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。
　ステップＳ９０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ９０２で置換機能検出部２１３（図１２参照）は、アプリケーションの外部ライブラリ呼び出しを検出する。

　ステップＳ９０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、ライブラリ名をキーに、置換可能ＧＰＵライブラリを取得する。具体的には、置換機能検出部２１３は、把握した外部ライブラリ呼び出しについて、コードパターンＤＢ２３０と照合することで、検出した置換可能ＧＰＵライブラリを、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能な機能ブロックとして取得する。

　ステップＳ９０４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリの処理記述に置換する。

　ステップＳ９０５で置換処理部２１４は、置換したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ９０６で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ９０７で置換処理部２１４は、ＣＵＤＡのライブラリ呼び出しを、pyCudaで指定する。

　ステップＳ９０８で実行ファイル作成部１１７は、作成したパターンをpyCudaでインタプリットする。

　ステップＳ９０９で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。
　ステップＳ９１０で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ９１１で実行ファイル作成部１１７は、作成した組合せパターンpyCudaでインタプリットする。
　ステップＳ９１２で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ９１３で本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

・<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>のフローチャート
　図１９は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。なお、<処理Ａ－２>からの処理は、<処理Ａ－１>からの処理と並行して行えばよい。
　ステップＳ１００１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ１００２で置換機能検出部２１３（図１２参照）は、ソースコードからクラスまたは構造体の定義記述コードを検出する。

　ステップＳ１００３で置換機能検出部２１３は、コードパターンＤＢ２３０から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能ＧＰＵライブラリを取得する。

　ステップＳ１００４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリ処理記述に置換する。

　ステップＳ１００５で置換処理部２１４は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。

　ステップＳ１００６で置換機能検出部２１３は、置換または確認したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ１００７で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ１００８で置換処理部２１４は、ＣＵＤＡのライブラリ呼び出しを、pyCudaで指定する。

　ステップＳ１００９で実行ファイル作成部１１７は、作成したパターンをpyCudaでインタプリットする。

　ステップＳ１０１０で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。

　ステップＳ１０１１で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ１０１２で実行ファイル作成部１１７は、作成した組合せパターンをpyCudaでインタプリットする。
　ステップＳ１０１３で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ１０１４で本番環境配置部１１８は、本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

　以上、《機能ブロックオフロード：Python》について説明した、次に、《機能ブロックオフロード：Java》について説明する。

［《機能ブロックオフロード：Java》］
　機能ブロックオフロード：Javaの、コードの分析では、Javaを解析するJavaParser（登録商標）等の構文解析ツールを用いて構文解析する。機能ブロックの把握については、構文解析ツールの結果を用いて、次処理のマッチング探索に用いるため、言語に非依存の機能ブロックとして管理する。

　オフロード可能機能ブロックの探索では、ライブラリ等の名前一致でのマッチングと、Deckard（登録商標）等のJava機能ブロックの類似性検出ツールを用いた類似性検知による、探索が行われる。オフロード機能ブロックへの置換は、ＧＰＵ処理のＣＵＤＡライブラリの呼び出し等、その言語からのオフロード機能利用に合わせた処理に置換する必要がある。

　実行環境は、Javaのラムダ記述での処理をＧＰＵに対して行うことができるIBM JDK（登録商標）を用いる。性能測定は、言語に合わせて、Jenkins（登録商標）等の自動測定ツールも用いて行う。オフロード可能機能ブロックが複数の際は反復実行され、最高性能のパターンが最終解として決定される。

［《機能ブロックオフロード：Java》フローチャート］
　次に、図２０および図２１を参照してオフロードサーバ２００の《機能ブロック：Java》の動作概要を説明する。

・<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>のフローチャート
　図２０は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、《機能ブロック：Java》のオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。
　ステップＳ１１０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションプログラムのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ１１０２で置換機能検出部２１３（図１２参照）は、アプリケーションプログラムの外部ライブラリ呼び出しを検出する。

　ステップＳ１１０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、ライブラリ名をキーに、置換可能ＧＰＵライブラリを取得する。具体的には、置換機能検出部２１３は、把握した外部ライブラリ呼び出しについて、コードパターンＤＢ２３０と照合することで、検出した置換可能ＧＰＵライブラリを、ＧＰＵ、ＦＰＧＡにオフロードできるオフロード可能な機能ブロックとして取得する。

　ステップＳ１１０４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコアの処理記述に置換する。

　ステップＳ１１０５で置換処理部２１４は、置換したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ１１０６で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ１１０７で置換処理部２１４は、ＣＵＤＡのライブラリ呼び出しを、Jcudaで指定する。

　ステップＳ１１０８で実行ファイル作成部１１７は、作成したパターンをJcudaでビルドする。

　ステップＳ１１０９で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。
　ステップＳ１１１０で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ１１１１で実行ファイル作成部１１７は、作成した組合せパターンをJcudaでビルドする。
　ステップＳ１１１２で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ１１１３で本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

・<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>のフローチャート
　図２１は、オフロードサーバ２００の制御部（自動オフロード機能部）２１０が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。なお、<処理Ａ－２>からの処理は、<処理Ａ－１>からの処理と並行して行えばよい。
　ステップＳ１２０１でアプリケーションコード分析部１１２（図１２参照）は、アプリケーションのオフロードしたいソースコードの分析を行う。具体的には、アプリケーションコード分析部１１２は、Clang等の構文解析ツールを用いて、ループ文構造等とともに、コードに含まれるライブラリ呼び出しや、機能処理を分析するソースコードの分析を行う。

　ステップＳ１２０２で置換機能検出部２１３（図１２参照）は、ソースコードからクラスまたは構造体の定義記述コードを検出する。

　ステップＳ１２０３で置換機能検出部２１３は、コードパターンＤＢ２３０から、類似性検出ツールを用いて、クラスまたは構造体の定義記述コードをキーにして、置換可能ＧＰＵライブラリを取得する。

　ステップＳ１２０４で置換処理部２１４は、アプリケーションのソースコードの置換元の処理記述を、置換先のＧＰＵライブラリ処理記述に置換する。

　ステップＳ１２０５で置換処理部２１４は、置換元と置換先で引数、戻り値の数や型が異なる場合に、ユーザに確認する。

　ステップＳ１２０６で置換機能検出部２１３は、置換または確認したＧＰＵライブラリの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵにオフロードする。

　ステップＳ１２０７で置換処理部２１４は、ＧＰＵライブラリ呼び出しのためのインタフェース処理を記述する。

　ステップＳ１１０８で置換処理部２１４は、ＣＵＤＡのライブラリ呼び出しを、Jcudaで指定する。

　ステップＳ１２０９で実行ファイル作成部１１７は、作成したパターンをJcudaでビルドする。

　ステップＳ１２１０で性能測定部１１６は、作成したパターンを検証環境で性能測定する（「１回目の性能測定」）。

　ステップＳ１２１１で実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組合せパターンを作成する。

　ステップＳ１２１２で実行ファイル作成部１１７は、作成した組合せパターンをJcudaでビルドする。

　ステップＳ１２１３で性能測定部１１６は、作成した組合せパターンを検証環境で性能測定する（「２回目の性能測定」）。

　ステップＳ１２１４で本番環境配置部１１８は、本番環境配置部１１８は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。

［実装例］
　第１の実施形態（「ループ文オフロード」）および第２の実施形態（「機能ブロックオフロード」）の実装例を説明する。

<利用ツール>
　対象アプリケーションプログラムはC/C++言語、Python、Javaのアプリケーションとする。
　ＧＰＵ処理は、C/C++言語はＰＧＩコンパイラ19.10を用いる。ＰＧＩコンパイラは、OpenACCを解釈するC/C++向けコンパイラである。ＰＧＩコンパイラは、cuFFT等のＣＵＤＡライブラリの呼び出しも処理が可能である。

　Pythonは、pyCUDA 2019.1.2を用いる。pyCUDAは、PythonからＧＰＵに処理実行するためのインタプリタである。あるいは、Pythonには、PyACCを用いる。PyACCは、PythonからOpenACCを解釈実行するためのインタプリタである。

　Javaは、IBM JDK（登録商標）を用いる。IBM JDKはJavaのラムダ記述に従って並列処理をＧＰＵに対して実行する仮想マシンである。

<構文解析>
　C/C++言語の構文解析には、LLVM/Clang 6.0の構文解析ライブラリ(libClang（登録商標）のpython binding) を用いる。Pythonの構文解析には、astを用いる。Javaの構文解析には、Java Parser（登録商標）を用いる。

<類似性検出ツール>
　類似性検出ツールには、C/C++言語、Javaには、Deckard v2.0（登録商標）を用いる。Deckardは、機能ブロックのオフロードの適用領域拡大のため、ライブラリ呼び出し以外にも、コードコピーし変更した機能等のオフロードを実現するため、照合対象となる部分コードと、ＤＢに登録されたコードの類似性を判定する。Pythonには、CloneDigger（登録商標）を用いる。

<コードパターンＤＢ>
　照合に用いるコードパターンＤＢ２３０（図１２参照）は、MySQL8を用いる。呼び出しているライブラリ名をキーに、高速化できるライブラリ等を検索するためのレコードを保持する。ライブラリには、それに紐づく名前やコードや実行ファイルが保持される。実行ファイルはその利用手法等も登録されている。コードパターンＤＢ２３０には、ライブラリ等を類似性検出技術で検出するための、比較用コードとの対応関係も保持される。

<実装動作>
　実装の動作概要について述べる。
　実装は、アプリケーションの利用依頼があると、構文解析ライブラリを用いてコード解析を行う。次に、機能ブロックオフロード、ループ文オフロードの順に試行を行う。これは、ループ文と機能ブロックに関しては、アルゴリズム含めて処理内容に合わせてオフロードする機能ブロックオフロードの方が高速化できるからである。機能ブロックオフロードが可能であった場合は、後半のループ文オフロードはオフロード可能であった機能ブロック部分を抜いたコードに対して試行する。
　性能測定の結果、最高性能のパターンを解とする。

　以上、ライブラリ呼び出しの場合について記載した。
　置換機能検出部２１３（図１２参照）が、類似性検出ツールを用いて類似性検出を行う場合について説明する。類似性検出を行う場合には、上記置換記述と並行して処理がされる。すなわち、置換機能検出部２１３が、類似性検出を行う場合、実装例では、<処理Ｂ－２>でDeckardを用いて、検出されたクラス、構造体等の部分コードとコードパターンＤＢ２３０に登録された比較用コードとの類似性検出を行う。そして、置換機能検出部２１３は、閾値超えの機能ブロックと該当するＧＰＵ用ライブラリやＦＰＧＡ用ＩＰコアを検出する。置換機能検出部２１３は、<処理Ｂ－１>の場合と同様に、実行ファイルやOpenCLを取得する。実装例では、次にＣ-１の場合と同様に実行用ファイルを作成するが、特に置換元のコードと置換するライブラリやＩＰコアの引数や戻り値、型等のインタフェースが異なる場合は、オフロードを依頼したユーザに対して、置換先ライブラリやＩＰコアに合わせて、インタフェースを変更してよいか確認し、確認後に実行用ファイルを作成する。

　この時点で、検証環境のＧＰＵやＦＰＧＡで性能測定できる実行用ファイルが作成される。機能ブロックオフロードについては、置換する機能ブロックが一つの場合は、その一つをオフロードするかしないかだけである。複数ある場合は、一つずつオフロードする／しないを検証パターンとして作成し、性能を測定し高速な解を検出する。これは、高速化可能とされていても実測してみないとその条件で高速になるかわからないためである。例えば、５つオフロード可能な機能ブロックがあり、１回目測定の結果、２番と4番のオフロードが高速化できた場合は、２番と４番両方をオフロードするパターンで２回目測定を行い、２番と４番単独でオフロードする場合より高速となっている場合は、解として選択する。

［ハードウェア構成］
　第１および第２の実施形態に係るオフロードサーバは、例えば図２２に示すような構成の物理装置であるコンピュータ９００によって実現される。
　図２２は、オフロードサーバ１，２００の機能を実現するコンピュータの一例を示すハードウェア構成図である。コンピュータ９００は、ＣＰＵ（Central Processing Unit）９０１、ＲＯＭ（Read Only Memory）９０２、ＲＡＭ９０３、ＨＤＤ（Hard Disk Drive）９０４、入出力Ｉ／Ｆ（Interface）９０５、通信Ｉ／Ｆ９０６およびメディアＩ／Ｆ９０７を有する。

　ＣＰＵ９０１は、ＲＯＭ９０２またはＨＤＤ９０４に記憶されたプログラムに基づき作動し、図１、図１２に示すオフロードサーバ１，２００の各処理部による制御を行う。ＲＯＭ９０２は、コンピュータ９００の起動時にＣＰＵ９０１により実行されるブートプログラムや、コンピュータ９００のハードウェアに係るプログラム等を記憶する。

　ＣＰＵ９０１は、入出力Ｉ／Ｆ９０５を介して、マウスやキーボード等の入力装置９１０、および、ディスプレイ等の出力装置９１１を制御する。ＣＰＵ９０１は、入出力Ｉ／Ｆ９０５を介して、入力装置９１０からデータを取得するともに、生成したデータを出力装置９１１へ出力する。

　ＨＤＤ９０４は、ＣＰＵ９０１により実行されるプログラムおよび当該プログラムによって使用されるデータ等を記憶する。通信Ｉ／Ｆ９０６は、通信網（例えば、ＮＷ（Network）９２０）を介して他の装置からデータを受信してＣＰＵ９０１へ出力し、また、ＣＰＵ９０１が生成したデータを、通信網を介して他の装置へ送信する。

　メディアＩ／Ｆ９０７は、記録媒体９１２に格納されたプログラムまたはデータを読み取り、ＲＡＭ９０３を介してＣＰＵ９０１へ出力する。ＣＰＵ９０１は、目的の処理に係るプログラムを、メディアＩ／Ｆ９０７を介して記録媒体９１２からＲＡＭ９０３上にロードし、ロードしたプログラムを実行する。記録媒体９１２は、ＤＶＤ（Digital Versatile Disc）、ＰＤ（Phase change rewritable Disk）等の光学記録媒体、ＭＯ（Magneto Optical disk）等の光磁気記録媒体、磁気記録媒体、導体メモリテープ媒体又は半導体メモリ等である。

　例えば、コンピュータ９００が第１および第２の実施形態に係るオフロードサーバ１，２００として機能する場合、コンピュータ９００のＣＰＵ９０１は、ＲＡＭ９０３上にロードされたプログラムを実行することによりオフロードサーバ１，２００の機能を実現する。また、ＨＤＤ９０４には、ＲＡＭ９０３内のデータが記憶される。ＣＰＵ９０１は、目的の処理に係るプログラムを記録媒体９１２から読み取って実行する。この他、ＣＰＵ９０１は、他の装置から通信網（ＮＷ９２０）を介して目的の処理に係るプログラムを読み込んでもよい。

［効果］
　以下、本発明に係るオフロードサーバ等の効果について説明する。
　第１の実施形態に係るオフロードサーバ１（図１参照）は、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、オフロードサーバ１は、アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部１１２と、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部１１３と、アプリケーションプログラムのループ文を特定し、特定した各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部１１４と、コンパイルエラーが出る繰り返し文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部１１５と、を備える。また、並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン１４に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部１１６と、性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、を備える。

　このようにすることにより、オフロードサーバ１は、移行元言語がＣ言語、Python、Javaを含む多様な言語の場合でも共通的な方式でＧＰＵに自動オフロードすることができる。これにより、移行元言語に合わせて、処理を検討したり実装する必要がなくなり、コストダウンを図ることができる。

　さらに、オフロードサーバ１は、移行元言語がＣ言語、Python、またはJavaのいずれの場合であっても、ＣＰＵ－ＧＰＵ間のデータ転送回数を低減しつつ、アプリケーションプログラムの特定処理をアクセラレータに自動でオフロードすることで、全体の処理能力を向上させることができる。これにより、ＣＵＤＡ等のスキルが無いユーザでもＧＰＵを使い高性能処理ができる。また、従来ＧＰＵでの高性能化が検討されていない汎用的なＣＰＵ向けアプリケーションを高性能化できる。また、高性能計算用サーバでない汎用的マシンのＧＰＵにオフロードすることができる。

　第２の実施形態に係るオフロードサーバ２００（図１２参照）は、アプリケーションプログラムの特定処理をＧＰＵまたはＰＬＤにオフロードするオフロードサーバであって、アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、オフロードサーバ２００は、ＧＰＵまたはＰＬＤにオフロード可能なライブラリおよびＩＰコアを記憶するコードパターンＤＢ２３０と、アプリケーションプログラムのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するアプリケーションコード分析部１１２と、検出された外部ライブラリ呼び出しをキーにして、コードパターンＤＢ２３０からライブラリおよびＩＰコアを取得する置換機能検出部２１３と、アプリケーションプログラムのソースコードの置換元の処理記述を、置換機能検出部２１３が取得した置換先のライブラリおよびＩＰコアの置換先の処理記述として置換するとともに、置換したライブラリおよびＩＰコアの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵまたはＰＬＤにオフロードする置換処理部２１４と、ホストプログラムとのインタフェースを作成するオフロードパターン作成部２１５と、作成されたＧＰＵまたはＰＬＤ処理パターンの前記アプリケーションをコンパイルして、実行ファイルを作成する実行ファイル作成部１１７と、作成された実行ファイルをアクセラレータ検証用装置に配置し、ＧＰＵまたはＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部１１６と、を備え、実行ファイル作成部１１７は、性能測定用処理による性能測定結果をもとに、複数のＧＰＵまたはＰＬＤ処理パターンから最高処理性能のＧＰＵまたはＰＬＤ処理パターンを選択し、最高処理性能のＧＰＵまたはＰＬＤ処理パターンをコンパイルして、最終実行ファイルを作成する。

　このようにすることにより、オフロードサーバ２００は、移行元言語がＣ言語、Python、Javaを含む多様な言語の場合でも共通的な方式でＧＰＵに自動オフロードすることができる。これにより、移行元言語に合わせて、処理を検討したり実装する必要をなり、コストダウンを図ることができる。

　さらに、オフロードサーバ２００は、移行元言語がＣ言語、Python、またはJavaのいずれの場合であっても、アプリケーションコードの置換元の処理記述を、置換先のライブラリおよびＩＰコア処理記述に置換して、オフロード可能な機能ブロックとして、ＧＰＵやＰＬＤ（ＦＰＧＡ等）にオフロードする。すなわち、個々のループ文でなく、行列積算やフーリエ変換等のより大きな単位で、ＦＰＧＡやＧＰＵ等ハードウェア向けのアルゴリズム含めて実装された機能ブロックをオフロードする。これにより、ＧＰＵやＰＬＤ（ＦＰＧＡ等）への自動オフロードにおいて、機能ブロックの単位でオフロードすることで、オフロード処理の高速化を図ることができる。その結果、ＧＰＵ、ＦＰＧＡ、ＩｏＴデバイス等環境が多様になる中で、アプリケーションを環境に合わせて適応させることが可能になり、高性能にアプリケーションを動作させることができる。

　第１および第２の実施形態に係るオフロードサーバ１，２００において、アプリケーションプログラムが、Ｃ言語の場合、ループ文のＧＰＵ処理をOpenACC文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてＣ言語コンパイラを用いてＧＰＵオフロードすることを特徴とする。

　このようにすることにより、移行元言語がＣ言語の場合に共通的な方式で、ＧＰＵに自動オフロードすることができる。

　第１および第２の実施形態に係るオフロードサーバ１，２００において、アプリケーションプログラムが、Pythonの場合、ループ文のＧＰＵ処理をＣＵＤＡ文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてpyCUDAを用いてＧＰＵオフロードすることを特徴とする。

　このようにすることにより、移行元言語がPythonの場合に共通的な方式で、ＧＰＵに自動オフロードすることができる。

　第１および第２の実施形態に係るオフロードサーバ１，２００において、アプリケーションプログラムが、Pythonの場合、ループ文のＧＰＵ処理をOpenACC文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてpyACCを用いてＧＰＵオフロードすることを特徴とする。

　このようにすることにより、移行元言語がPythonの場合にＣ言語の場合と同様に、共通的な方式でＧＰＵに自動オフロードすることができる。

　第１および第２の実施形態に係るオフロードサーバ１，２００において、アプリケーションプログラムが、Javaの場合、ループ文のＧＰＵ処理をJavaのラムダ文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてJava仮想マシンを用いてＧＰＵオフロードすることを特徴とする。

　このようにすることにより、移行元言語がJavaの場合に共通的な方式でＧＰＵに自動オフロードすることができる。

　本発明は、コンピュータを、上記オフロードサーバとして機能させるためのオフロードプログラムとした。

　このようにすることにより、一般的なコンピュータを用いて、上記オフロードサーバ２００の各機能を実現させることができる。

［変形例］
　第１の実施形態に係るオフロードサーバ１と第２の実施形態に係るオフロードサーバ２００とを組み合わせて、データ転送指定部１１３（図１参照）は、アプリケーションコード分析部１１２が行うコード解析結果をもとに、機能ブロックオフロード、ループ文オフロードの順に試行するようにデータ転送を行うととともに、機能ブロックオフロードが可能であった場合は、オフロード可能であった機能ブロック部分を抜いたコードに対して、ループ文オフロードを試行するデータ転送を行う構成としてもよい。

　この構成により、オフロードサーバは、まず、アプリケーションプログラムの利用依頼があると、構文解析ライブラリを用いてコード解析を行い、次に、機能ブロックオフロード、ループ文オフロードの順に試行を行う。機能ブロックオフロードが可能であった場合は、オフロード可能であった機能ブロック部分を抜いたコードに対して、ループ文オフロードを試行し、性能測定の結果、最高性能のパターンを解とする。これにより、ループ文と機能ブロックに関しては、アルゴリズム含めて処理内容に合わせてオフロードする機能ブロックオフロードの方が高速化できる。機能ブロックオフロード、ループ文オフロードの順に試行を行うことで、処理の高速化を図ることができ、全体の処理能力を向上させることができる。

　また、上記各実施形態において説明した各処理のうち、自動的に行われるものとして説明した処理の全部又は一部を手作業で行うこともでき、あるいは、手作業で行われるものとして説明した処理の全部又は一部を公知の方法で自動的に行うこともできる。この他、上述文書中や図面中に示した処理手順、制御手順、具体的名称、各種のデータやパラメータを含む情報については、特記する場合を除いて任意に変更することができる。
　また、図示した各装置の各構成要素は機能概念的なものであり、必ずしも物理的に図示の如く構成されていることを要しない。すなわち、各装置の分散・統合の具体的形態は図示のものに限られず、その全部又は一部を、各種の負荷や使用状況などに応じて、任意の単位で機能的又は物理的に分散・統合して構成することができる。

　１，２００　オフロードサーバ
　１１，２１０　制御部
　１２　入出力部
　１３，１３０　記憶部
　１４　検証用マシン (アクセラレータ検証用装置)
　１５　商用環境
　１１１　アプリケーションコード指定部
　１１２　アプリケーションコード分析部
　１１３　データ転送指定部
　１１４　並列処理指定部
　１１４ａ　オフロード範囲抽出部
　１１４ｂ　中間言語ファイル出力部
　１１５　並列処理パターン作成部
　１１６　性能測定部
　１１６ａ　バイナリファイル配置部
　１１７　実行ファイル作成部
　１１８　本番環境配置部
　１１９　性能測定テスト抽出実行部
　１２０　ユーザ提供部
　１２５　アプリケーションコード
　１３１　テストケースＤＢ
　１３２　中間言語ファイル
　１５１　各種デバイス
　１５２　ＣＰＵ-ＧＰＵを有する装置
　１５３　ＣＰＵ-ＦＰＧＡを有する装置
　１５４　ＣＰＵを有する装置
　２１３　置換機能検出部
　２１４　置換処理部
　２１５　オフロードパターン作成部
　２３０　コードパターンＤＢ

Claims

　アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、
　前記アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、
　前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、
　前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、
　前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部と、
　コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、
　前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、
　性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成する実行ファイル作成部と、
　を備えることを特徴とするオフロードサーバ。
　アプリケーションプログラムの特定処理をＧＰＵ（Graphics Processing Unit）またはＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバであって、
　前記アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、
　前記ＧＰＵまたは前記ＰＬＤにオフロード可能なライブラリおよびＩＰコアを記憶する記憶部と、
　前記アプリケーションプログラムのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するアプリケーションコード分析部と、
　検出された前記外部ライブラリ呼び出しをキーにして、前記記憶部から前記ライブラリおよび前記ＩＰコアを取得する置換機能検出部と、
　前記アプリケーションプログラムのソースコードの置換元の処理記述を、前記置換機能検出部が取得した前記ライブラリおよび前記ＩＰコアの置換先の処理記述として置換するとともに、
　置換した前記ライブラリおよび前記ＩＰコアの処理記述を、オフロード対象の機能ブロックとして、前記ＧＰＵまたは前記ＰＬＤにオフロードする置換処理部と、
　ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出するオフロードパターン作成部と、
　ホストプログラムとのインタフェースを作成するオフロードパターン作成部と、
　作成されたＧＰＵまたはＰＬＤ処理パターンの前記アプリケーションプログラムをコンパイルして、実行ファイルを作成する実行ファイル作成部と、
　作成された前記実行ファイルをアクセラレータ検証用装置に配置し、前記ＧＰＵまたは前記ＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部と、を備え、
　前記実行ファイル作成部は、前記性能測定用処理による性能測定結果をもとに、複数の前記ＧＰＵまたはＰＬＤ処理パターンから最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンを選択し、最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンをコンパイルして、最終実行ファイルを作成する
　ことを特徴とするオフロードサーバ。
　前記アプリケーションプログラムが、Ｃ言語の場合、ループ文のＧＰＵ処理をOpenACC文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてＣ言語コンパイラを用いてＧＰＵオフロードする
　ことを特徴とする請求項１または請求項２に記載のオフロードサーバ。
　前記アプリケーションプログラムが、Pythonの場合、ループ文のＧＰＵ処理をＣＵＤＡ文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてpyCUDAを用いてＧＰＵオフロードする
　ことを特徴とする請求項１または請求項２に記載のオフロードサーバ。
　前記アプリケーションプログラムが、Pythonの場合、ループ文のＧＰＵ処理をOpenACC文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてpyACCを用いてＧＰＵオフロードする
　ことを特徴とする請求項１または請求項２に記載のオフロードサーバ。
　前記アプリケーションプログラムが、Javaの場合、ループ文のＧＰＵ処理をJavaのラムダ文法で指定し、機能ブロックのＧＰＵ処理をＣＵＤＡライブラリ呼び出すようにしてJava仮想マシンを用いてＧＰＵオフロードする
　ことを特徴とする請求項１または請求項２に記載のオフロードサーバ。
　アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバのオフロード制御方法であって、
　前記アプリケーションプログラムは、Ｃ言語、Python、およびJavaより選択される少なくとも一つであり、
　前記オフロードサーバは、
　前記アプリケーションプログラムのソースコードを分析するステップと、
　前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うステップと、
　前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルするステップと、
　コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成するステップと、
　前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行するステップと、
　性能測定結果をもとに、複数の前記並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の前記並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成して、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の前記並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の前記並列処理パターンをコンパイルして実行ファイルを作成するステップと、を実行する
　ことを特徴とするオフロード制御方法。
　コンピュータを、請求項１から請求項６のいずれか一項に記載のオフロードサーバとして機能させるためのオフロードプログラム。