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

Description

本発明は、通常のＣＰＵ（Central Processing Unit）で処理するソフトウェア処理プログラムを、ＧＰＵ（Graphics Processing Unit）、ＦＰＧＡ（Field Programmable Gate Array）、メニーコアＣＰＵ等のアクセラレータが混在する環境で高速に処理するようにオフロードするオフロードサーバ、オフロード制御方法およびオフロードプログラムに関する。

近年、ＣＰＵの半導体集積度が１．５年で２倍になるというムーアの法則が減速するのではないかと言われている。そのような状況から、少コアのＣＰＵだけでなく、ＦＰＧＡやＧＰＵ等のデバイスの活用が増えている。例えば、マイクロソフト社はＦＰＧＡを使って検索エンジンの効率を高める取り組みをしている。アマゾン社は、ＦＰＧＡやＧＰＵ等をクラウドのインスタンスとして提供している。

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

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

そこで発明者は、一度記述したコードを、配置先の環境に存在するＧＰＵやＦＰＧＡ、メニーコアＣＰＵ等が利用できるように、変換、リソース設定等を自動で行い、アプリケーションを高性能に動作させることを目的とした環境適応ソフトウェアを非特許文献１に提案した。
発明者は、環境適応ソフトウェアの要素として、ソースコードのループ文及び機能ブロックを、ＦＰＧＡ、ＧＰＵに自動オフロードする方式を非特許文献２，３，４にて提案している。

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 ＦＰＧＡ 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.

近年、ＧＰＵの並列計算パワーを画像処理でないものにも使うＧＰＧＰＵ（General Purpose GPU）を行うための環境としてＣＵＤＡが普及している。ＣＵＤＡは、ＮＶＩＤＩＡ社が開発したＧＰＧＰＵ向け環境である。

更にＦＰＧＡ、メニーコアＣＰＵ、ＧＰＵ等のヘテロなデバイスを同じように扱うための仕様としてOpenCLが出ており、その開発環境も出てきている。ＣＵＤＡ、OpenCLは、ＦＰＧＡ等のカーネルとＣＰＵのホストとの間のメモリデータのコピーや解放の記述を明示的に行う等、Ｃ言語の拡張によりプログラムを行うものであり、プログラムの難度は高い。

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

ＣＵＤＡ、OpenCL、OpenACC、OpenMP等の技術仕様を用いることで、ＦＰＧＡやＧＰＵやメニーコアＣＰＵへオフロードすることは可能である。しかしアクセラレータ処理自体は行えるようになっても、高速化することには課題がある。

例えば、マルチコアＣＰＵ向けに自動並列化機能を持つコンパイラとして、Intelコンパイラ等がある。これらは、自動並列化時に、コードの中のループ文中で並列処理可能な部分を抽出して、並列化している。しかし、メモリ処理等の影響で、単に並列化可能ループ文を並列化しても性能が出ないことも多い。なお、ここでループ文とは、C/C++ソースコードにおいてはfor/while/do-whileである。

具体的にいうと、ＦＰＧＡやＧＰＵ等で高速化する際、OpenCLやＣＵＤＡの技術者は、チューニングを繰り返したり、OpenACCコンパイラ等を用いて適切な並列処理範囲を探索して試行している。このため、技術スキルが乏しいプログラマが、ＦＰＧＡやＧＰＵ、メニーコアＣＰＵを活用してアプリケーションを高速化することは難しい。また、自動並列化技術等を使う場合、並列処理箇所探索の試行錯誤等の稼働が必要であった。

現状、ヘテロなデバイスに対するオフロードは手作業での取組みが主流である。発明者は、環境適応ソフトウェアのコンセプトを提案し、自動オフロードを検討している。しかし、ＧＰＵ単体へのオフロードである等、ＧＰＵやＦＰＧＡやメニーコアＣＰＵ等の多様なアクセラレータが混在している移行先環境は想定されていない。

そこで、本発明は、オフロードサーバ、オフロード制御方法およびオフロードプログラムにおいて、複数種類のアクセラレータが混在する環境が移行先であっても、自動で高性能化できることを課題とする。

前記した課題を解決するため、本発明のオフロードサーバは、ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析するコード分析部と、前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを作成し、前記ソフトウェアプログラムのループ文を各前記アクセラレータの何れかにオフロードするパターンを作成する処理パターン作成部と、前記機能ブロックまたは前記ループ文を前記アクセラレータにオフロードする各前記パターンを検証環境にデプロイして性能を測定する性能測定部と、前記性能測定部が前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを測定した性能が所望のものでなかった場合、前記処理パターン作成部により、前記ソフトウェアプログラムの前記ループ文を各前記アクセラレータの何れかにオフロードするパターンを作成して、前記性能測定部により、各前記パターンを前記検証環境にデプロイして性能を測定する制御部と、を備えることを特徴とする。
その他の手段については、発明を実施するための形態のなかで説明する。

本発明によれば、複数種類のアクセラレータが混在する環境が移行先であっても、自動で高性能化できる。更に移行先が同一ノードに複数種類のアクセラレータを持つ場合は、複数種類のアクセラレータに対して同時にオフロードすることで、単一アクセラレータの利用の場合よりも高速なオフロードを実現する。

オフロードサーバの構成例を示す機能ブロック図である。 オフロードサーバの機能を実現するコンピュータの一例を示すハードウェア構成図である。 アクセラレータにオフロードするパターンの選択処理のフローチャートである。 アクセラレータにオフロードするパターンの選択処理のフローチャートである。 アクセラレータにオフロードするパターンの選択処理のフローチャートである。 機能ブロックを１ノードにオフロードする処理のフローチャートである。 ループ文を１ノードにオフロードする処理のフローチャートである。 機能ブロックをアクセラレータへオフロードするパターンの選択処理のフローチャートである。 機能ブロックをアクセラレータへオフロードするパターンの選択処理のフローチャートである。 オフロードサーバの自動オフロード処理を示す図である。 オフロードサーバの単純遺伝的アルゴリズムによる制御部の探索イメージを示す図である。 ループ文のメニーコアＣＰＵへの遺伝子配列マッピングを示す図である。 ループ文をメニーコアＣＰＵへオフロードするパターンの選択処理のフローチャートである。 ループ文をメニーコアＣＰＵへオフロードするパターンの選択処理のフローチャートである。 ループ文のＧＰＵへの遺伝子配列マッピングを示す図である。 ループ文をＧＰＵへオフロードするパターンの選択処理のフローチャートである。 ループ文をＧＰＵへオフロードするパターンの選択処理のフローチャートである。 ループ文をＦＰＧＡへオフロードするパターンの選択処理のフローチャートである。 ループ文をＦＰＧＡへオフロードするパターンの選択処理のフローチャートである。

以降、本発明を実施するための形態を、各図を参照して詳細に説明する。
発明者は、環境適応ソフトウェアのコンセプトを具体化するために、これまでに、プログラムのループ文のＧＰＵ自動オフロード、ＦＰＧＡ自動オフロード、プログラムの機能ブロックの自動オフロードの方式を提案してきた。これらの要素技術検討も踏まえて、移行先環境について対象を定義し、対象が多様化した場合でも踏襲する基本的考えを記述する、更に、個々の移行先環境への自動オフロードと、移行先環境が多様化した際の自動オフロード方式について提案する。

《多様移行先環境時の基本的考え》
本実施形態で対象とする多様な移行先環境としては、ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵの３つとする。ＧＰＵ、ＦＰＧＡについてはＣＰＵとは異なるヘテロジニアスなハードウェアとして歴史も深く、ＣＵＤＡやOpenCLを用いた手動でのオフロードによる高速化事例も多く、市場も大きい。また、メニーコアＣＰＵに関しては、近年１６コア以上の多数のコアを搭載したＣＰＵが千ドルから数千ドルの低価格でも市場に出るようになり、OpenMP等の技術仕様を用いて並列化を行い、手作業でチューニングすることで、高速化する事例が出てきている。

移行先環境が単なるＣＰＵだけでない場合で、自動で高速にオフロードするため、検証環境の実機で性能測定し、進化計算手法等の手法と組み合わせて、徐々に高速なオフロードパターンを見つけるアプローチをとる。

これは、今まで発明者が提案したＧＰＵオフロード等の場合と同様である。理由として、性能に関しては、コード構造だけでなく、処理するハードウェアのスペック、データサイズ、ループ回数等の実際に処理する内容によって大きく変わるため、静的に予測する事が困難であり、動的な測定が必要だからと考えるからである。市中には、ループ文を見つけコンパイル段階で並列化する自動並列化コンパイラがあるが、並列化可能ループ文の並列化だけでは性能を測定すると低速になる場合も多いため、性能測定は必要と考える。

また、オフロードする対象については、プログラムのループ文と機能ブロックとする。これは、今まで検討したＧＰＵやＦＰＧＡオフロード等の場合と同様である。ループ文については、処理時間がかかるプログラムの処理の大半はループで費やされているという現状から、オフロードのターゲットとしてループ文が考えられる。

一方、機能ブロックについては、特定処理を高速化する際に、処理内容や処理ハードウェアに適したアルゴリズムを用いることが多いため、個々のループ文の並列処理等に比べ、大きく高速化できる場合がある。本実施形態では、行列積算やフーリエ変換等の頻繁に使われる機能ブロック単位で、メニーコアＣＰＵやＦＰＧＡやＧＰＵ等のアクセラレータに応じたアルゴリズムで実装された処理（ＩＰコアやＣＵＤＡライブラリ等）に置換することで高速化する。

図１は、オフロードサーバ１の構成例を示す機能ブロック図である。
オフロードサーバ１は、アプリケーションの特定処理を、ＧＰＵやＦＰＧＡやメニーコアＣＰＵ等のアクセラレータに自動的にオフロードする装置である。
図１に示すように、オフロードサーバ１は、制御部１１と、入出力部１２と、記憶部１３と、検証用マシン１４（アクセラレータ検証用装置)と、を含んで構成される。

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

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

テストケースデータベース１３１には、検証対象ソフトに対応した試験項目のデータを格納する。例えば、MySQL等のデータベースシステムなら、TPC-C等のトランザクション試験等である。コードパターンデータベース１３３は、ライブラリ名に対応した、置換可能なアクセラレータライブラリまたはアクセラレータ用ＩＰコアを格納する。
検証用マシン１４は、環境適応ソフトウェアシステムの検証用環境として、ＧＰＵやＦＰＧＡやメニーコアＣＰＵ等のアクセラレータを備える。

制御部１１は、オフロードサーバ１の全体の制御を司る自動オフロード機能部である。制御部１１は、例えば、記憶部１３に格納されたプログラム（オフロードプログラム）を図２のＣＰＵ９１０が、実行することにより実現される。

制御部１１は、コード指定部１１１と、コード分析部１１２と、処理指定部１１４と、処理パターン作成部１１５と、性能測定部１１６と、実行ファイル作成部１１７と、本番環境配置部１１８と、性能測定テスト抽出実行部１１９と、ユーザ提供部１２０とを備える。

《コード指定部１１１》
コード指定部１１１は、入力されたソースコードの指定を行う。具体的には、コード指定部１１１は、受信したファイルに記載されたソースコードを、コード分析部１１２に渡す。

《コード分析部１１２》
コード分析部１１２は、ソフトウェアプログラムのソースコードを分析し、for/do-while/whileなどのループ文や、ＦＦＴライブラリや行列演算や乱数生成等の機能ブロックを把握する。これらループ文や機能ブロックは、アクセラレータによる実行が可能なものである。

《処理指定部１１４》
処理指定部１１４は、各機能ブロックに対して、アクセラレータへオフロードする処理に置換してコンパイルし、各ループ文に対して、アクセラレータにオフロードする並列処理指定文を指定してコンパイルする。処理指定部１１４は、オフロード範囲抽出部１１４ａと中間言語ファイル出力部１１４ｂとを含んでいる。

オフロード範囲抽出部１１４ａは、for/do-while/whileなどのループ文や、ＦＦＴや行列演算や乱数生成等の機能ブロックのように、メニーコアＣＰＵやＧＰＵやＦＰＧＡにオフロード可能な処理を特定して、オフロード範囲を抽出する。中間言語ファイル出力部１１４ｂは、抽出したオフロード可能処理の中間言語ファイル１３２を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。

《処理パターン作成部１１５》
処理パターン作成部１１５は、抽出したオフロード範囲に基づき、機能ブロックやループ文をアクセラレータにオフロードするパターンを作成する。

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

《性能測定部１１６》
性能測定部１１６は、処理パターンのソースコードをコンパイルして、検証用マシン１４に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する。
性能測定部１１６は、バイナリファイル配置部１１６ａを備える。バイナリファイル配置部１１６ａは、メニーコアＣＰＵとＧＰＵとＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイ(配置)する。

性能測定部１１６は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、処理パターン作成部１１５に戻す。この場合、処理パターン作成部１１５は、別のオフロードパターンを作成する。性能測定部１１６は、作成された別のオフロードパターンの性能測定を試行する。

《実行ファイル作成部１１７》
実行ファイル作成部１１７は、所定回数繰り返された性能測定結果をもとに、複数のオフロードパターンから最高処理性能のオフロードパターンを選択し、最高処理性能のオフロードパターンをコンパイルして実行ファイルを作成する。

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

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

《ユーザ提供部１２０》
ユーザ提供部１２０は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する（「価格・性能等の情報のユーザへの提供」）。テストケースデータベース１３１には、検証対象ソフトに対応した試験項目のデータが格納されている。ユーザ提供部１２０は、テストケースデータベース１３１に格納された試験項目を実施して得られた性能等のデータと、ユーザソフトを配置するリソース（仮想マシンやＧＰＵ等）の価格情報等のデータを、ユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。

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

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

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

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

図３は、オフロードサーバの機能を実現するコンピュータの一例を示すハードウェア構成図である。
コンピュータ９００は、ＣＰＵ９１０、ＲＡＭ９２０、ＲＯＭ９３０、ＨＤＤ９４０、通信インタフェース９５０、入出力インタフェース９６０、およびメディアインタフェース９７０を有する。

ＣＰＵ９１０は、ＲＯＭ９３０またはＨＤＤ９４０に格納されたプログラムに基づいて動作し、各部の制御を行う。ＲＯＭ９３０は、コンピュータ９００の起動時にＣＰＵ９１０によって実行されるブートプログラムや、コンピュータ９００のハードウェアに依存するプログラム等を格納する。

ＨＤＤ９４０は、ＣＰＵ９１０によって実行されるプログラム、および、かかるプログラムによって使用されるデータ等を格納する。通信インタフェース９５０は、通信網８０を介して他の機器からデータを受信してＣＰＵ９１０へ送り、ＣＰＵ９１０が生成したデータを通信網８０を介して他の機器へ送信する。

ＣＰＵ９１０は、入出力インタフェース９６０を介して、ディスプレイやプリンタ等の出力装置、および、キーボードやマウス等の入力装置を制御する。ＣＰＵ９１０は、入出力インタフェース９６０を介して、入力装置からデータを取得する。また、ＣＰＵ９１０は、生成したデータを入出力インタフェース９６０を介して出力装置へ出力する。

メディアインタフェース９７０は、記録媒体９８０に格納されたプログラムまたはデータを読み取り、ＲＡＭ９２０を介してＣＰＵ９１０に提供する。ＣＰＵ９１０は、かかるプログラムを、メディアインタフェース９７０を介して記録媒体９８０からＲＡＭ９２０上にロードし、ロードしたプログラムを実行する。記録媒体９８０は、例えばＤＶＤ（Digital Versatile Disc）等の光学記録媒体、ＭＯ（Magneto Optical disk）等の光磁気記録媒体、テープ媒体、磁気記録媒体、または半導体メモリ等である。

例えば、コンピュータ９００が本実施形態に係るオフロードサーバ１として機能する場合、コンピュータ９００のＣＰＵ９１０は、ＲＡＭ９２０上にロードされたプログラムを実行することにより、オフロードサーバ１の各部の機能を実現する。また、ＨＤＤ９４０には、オフロードサーバ１の各部内のデータが格納される。コンピュータ９００のＣＰＵ９１０は、これらのプログラムを記録媒体９８０から読み取って実行するが、他の例として、他の装置から通信網８０を介してこれらのプログラムを取得してもよい。

図３Ａと図３Ｂと図３Ｃは、アクセラレータにオフロードするパターンの選択処理のフローチャートである。
オフロードサーバ１は、ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵの３つの移行先環境に対して、機能ブロックとループ文の２つの方法で、自動オフロードを検討する。

制御部１１は、C/C++ソースコードが入力されると（Ｓ２００）、一連のオフロードパターンの選択処理を実行する。制御部１１は、図５Ａと図５Ｂの選択処理を呼び出して、メニーコアＣＰＵへの機能ブロックのオフロードをチェックし（Ｓ２０１）、所望の性能が得られたか否かを判定する（Ｓ２０２）。制御部１１は、所望の性能が得られたならば（Ｙｅｓ）、全測定から最高性能パターンを選択して（Ｓ２０８）、この図の処理を終了する。制御部１１は、所望の性能が得られていないならば（Ｎｏ）、ステップＳ２０３に進む。

ステップＳ２０３において、制御部１１は、図５Ａと図５Ｂの選択処理を呼び出して、ＧＰＵへの機能ブロックのオフロードをチェックする。そして制御部１１は、所望の性能が得られたか否かを判定する（Ｓ２０４）。制御部１１は、所望の性能が得られたならば（Ｙｅｓ）、全測定から最高性能パターンを選択して（Ｓ２０８）、この図の処理を終了する。制御部１１は、所望の性能が得られていないならば（Ｎｏ）、ステップＳ２０５に進む。

ステップＳ２０５において、制御部１１は、図５Ａと図５Ｂの選択処理を呼び出して、ＦＰＧＡへの機能ブロックのオフロードをチェックする。そして制御部１１は、所望の性能が得られたか否かを判定する（Ｓ２０６）。制御部１１は、所望の性能が得られたならば（Ｙｅｓ）、全測定から最高性能パターンを選択して（Ｓ２０８）、この図の処理を終了する。制御部１１は、所望の性能が得られていないならば（Ｎｏ）、コードから機能ブロックをオフロードできた部分を取り除き、図３ＢのステップＳ２０９に進む。
これらステップＳ２０１～Ｓ２０８の処理により、複数種類のアクセラレータが混在する環境が移行先であっても、機能ブロックを自動で高性能化できる。更に、制御部１１は、メニーコアＣＰＵ、ＧＰＵ、ＦＰＧＡの順に試行することにより、効率的に高速化可能なパターンを探索できる。

ステップＳ２０９において、制御部１１は、図９Ａと図９Ｂの選択処理を呼び出して、メニーコアＣＰＵへのループ文のオフロードをチェックする。そして制御部１１は、所望の性能が得られたか否かを判定する（Ｓ２１０）。制御部１１は、所望の性能が得られたならば（Ｙｅｓ）、全測定から最高性能パターンを選択して（Ｓ２０８）、この図の処理を終了する。制御部１１は、所望の性能が得られていないならば（Ｎｏ）、ステップＳ２１１に進む。

ステップＳ２１１において、制御部１１は、図１１Ａと図１１Ｂの選択処理を呼び出して、ＧＰＵへのループ文のオフロードをチェックする。そして制御部１１は、所望の性能が得られたか否かを判定する（Ｓ２１２）。制御部１１は、所望の性能が得られたならば（Ｙｅｓ）、全測定から最高性能パターンを選択して（Ｓ２０８）、この図の処理を終了する。制御部１１は、所望の性能が得られていないならば（Ｎｏ）、ステップＳ２１３に進む。

ステップＳ２１３において、制御部１１は、図１２Ａと図１２Ｂの選択処理を呼び出して、ＦＰＧＡへのループ文のオフロードをチェックする。そして制御部１１は、所望の性能が得られたか否かを判定する（Ｓ２１４）。制御部１１は、所望の性能が得られたならば（Ｙｅｓ）、全測定から最高性能パターンを選択して（Ｓ２０８）、この図の処理を終了する。制御部１１は、所望の性能が得られていないならば（Ｎｏ）、個々の移行先に対する６つの検証の中から最高性能のパターンを選択する（Ｓ２１５）。

これらステップＳ２０９～Ｓ２１５の処理により、複数種類のアクセラレータが混在する環境が移行先であっても、ループ文を自動で高性能化できる。更に、制御部１１は、メニーコアＣＰＵ、ＧＰＵ、ＦＰＧＡの順に試行することにより、効率的に高速化可能なパターンを探索できる。制御部は、アクセラレータへのオフロード対象として機能ブロックの試行の次にループ文を試行することより、効率的に高速化可能なパターンを探索できる。

そして制御部１１は、１ノードへの複数のオフロードが可能な候補があるか否かを判定する（Ｓ２１６）。
制御部１１は、１ノードへの複数のオフロードが可能な候補が無ければ（Ｎｏ）、この図の処理を終了する。制御部１１は、１ノードへの複数のオフロードが可能な候補があれば（Ｙｅｓ）、図３ＣのステップＳ２１７に進む。

図３ＣのステップＳ２１７において、制御部１１は、複数の機能ブロックをオフロード可能な候補が有るか否かを判定する。制御部１１は、複数の機能ブロックをオフロード可能な候補が有れば（Ｙｅｓ）、ステップＳ２１８に進み。複数の機能ブロックをオフロード可能な候補が無ければ（Ｎｏ）、ステップＳ２２３に進む。

ステップＳ２１８において、制御部１１は、複数の機能ブロックのオフロードの組み合わせパターンを作成する。制御部１１は、この組み合わせパターンの性能を測定する（Ｓ２１９）。そして制御部１１は、機能ブロックのオフロードの組み合わせで高速化できたか否かを判定する（Ｓ２２０）。

ステップＳ２２０において、制御部１１は、機能ブロックのオフロードを組み合わせた方が高速ならば、組み合わせた機能ブロックをソースコードから除去し（Ｓ２２１）、ステップＳ２２３の処理に進む。制御部１１は、機能ブロックのオフロードを組み合わせない方が高速ならば、より高速化できる一つの機能ブロックをソースコードのオフロード対象から除去し（Ｓ２２２）、ステップＳ２２３の処理に進む。

ステップＳ２２３において、制御部１１は、高速化できるループ文オフロードパターンが有るか否かと、それが単数のアクセラレータと複数のアクセラレータの何れであるかを判定する。制御部１１は、高速化できるループ文オフロードパターンが無いならば、ステップＳ２２７に進み、全測定から最高性能のパターンを選択すると、この選択処理を終了する。ここで、ループ文オフロードパターンＢはＧＰＵに３０個のfor文をオフロード、ループ文オフロードパターンＣはメニーコアＣＰＵに５０個のfor文をオフロードといった例の際に、Ｂだけ高速化できる場合を単数のアクセラレータ、Ｂ、Ｃの両方が高速化できる場合を複数のアクセラレータと記載する。

制御部１１は、ステップＳ２２３の判定結果が高速化できる単数のアクセラレータのループ文のオフロードパターンならば、除去した機能ブロックのオフロードと単数のアクセラレータのループ文のオフロードの組み合わせを作成する（Ｓ２２４）。そして制御部１１は、組み合わせパターンの性能を測定し（Ｓ２２６）、全測定から最高性能のパターンを選択すると（Ｓ２２７）、この選択処理を終了する。

制御部１１は、ステップＳ２２３の判定結果が高速化できる複数のアクセラレータのループ文のオフロードパターンならば、除去した機能ブロックのオフロードと複数のアクセラレータのループ文のオフロードの組み合わせを作成する（Ｓ２２５）。例えば、ＧＰＵへの機能ブロックオフロードＡ、ＧＰＵへのループ文オフロードパターンＢ、メニーコアＣＰＵへのループ文オフロードパターンＣが高速化できる場合は、Ａ＋Ｂ＋Ｃ、Ａ＋Ｂ、Ａ＋Ｃ、Ｂ＋Ｃの４種類の組み合わせが考えられる。
そして制御部１１は、組み合わせパターンの性能を測定し（Ｓ２２６）、全測定から最高性能のパターンを選択すると（Ｓ２２７）、この選択処理を終了する。
これらステップＳ２１７～Ｓ２２７の処理により、移行先が同一ノードに複数種類のアクセラレータを持つ場合は、複数種類のアクセラレータに対して同時にオフロードすることで、単一アクセラレータの利用の場合よりも高速なオフロードを実現する。

《オフロード試行の検証順番》
多様な移行先での自動オフロードについて検討する。
性能については、行う処理は同じでも、処理ハードウェアのスペック、処理内容（データサイズ、ループ回数等）により大きく変わるため、検証環境の実機で測定する事が必要である。移行先環境ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵの３つに対して、ループ文、機能ブロックの２つの方法でのオフロードがあるため、合計３×２の６つオフロードが考えられる。

まず、機能ブロックのオフロードとループ文のオフロードを比較すると、アルゴリズムを含めて処理内容に合わせてオフロードする機能ブロックのオフロードの方が高速化できる。

次に、仮想化等を行わない場合、ＧＰＵとＦＰＧＡとメニーコアＣＰＵで、中心価格帯はＧＰＵよりもメニーコアＣＰＵが高価で、ＦＰＧＡはそれらより更に高価である。また、１パターンの検証時間は、メニーコアＣＰＵとＧＰＵが略同等で、ＦＰＧＡがそれらよりも長くなる。ＦＰＧＡは、回路設定に数時間必要で、性能測定には時間がかかるのが現状である。

これらの現況を踏まえ、６つのオフロードで検証する順番として、メニーコアＣＰＵ向けの機能ブロックのオフロード、ＧＰＵ向けの機能ブロックのオフロード、ＦＰＧＡ向けの機能ブロックのオフロード、メニーコアＣＰＵ向けのループ文のオフロード、ＧＰＵ向けのループ文のオフロード、ＦＰＧＡ向けのループ文のオフロードとすることを提案する。この順番でオフロードを試行し、高性能となるパターンを探索していく。

前半の３つと後半の３つは、対象とするコードが変わってもよいとする。具体的には、前半の３つで機能ブロックのオフロードが可能であった場合、後半の３つのループ文オフロードはオフロード可能だった機能ブロック部分を抜いたコードに対して試行する。

なお、オフロード試行ではユーザが目標性能や価格を指定でき、ユーザが指定する範囲で十分に高速で低価格なオフロードパターンが６つ試行の前方側で見つかっていれば、以降の試行はしなくてもよい。このようにする理由として、機能ブロックオフロードの方がループ文オフロードに比べ、オフロードできる対象は少ないが、オフロードできる場合には、より高速化ができるためである。

また、自動オフロードする際は、できるだけ安価で短時間に高速なパターンを探索できることが望ましい。そこで、検証時間がかかるＦＰＧＡは最後とし、それよりも前の段階で十分にユーザ要件を満足するパターンが見つかっていれば、ＦＰＧＡ向けのオフロードパターンは検証しない事とする。

ＧＰＵとメニーコアＣＰＵに関しては、価格的にも検証時間的にも大きな差はないが、メモリも別空間となりデバイス自体が異なるＧＰＵに比して、メニーコアＣＰＵの方が、通常ＣＰＵとの差は小さい。そのため検証順は、メニーコアＣＰＵを先とし、メニーコアＣＰＵで十分にユーザ要件を満足するパターンが見つかっていれば、ＧＰＵ向けのオフロードパターンは検証しない事とする。
ＣＰＵとＧＰＵではデバイスが異なるため、丸め誤差の違い等で、正しくオフロードできても計算結果が異なる場合がある。

《複数オフロードの組み合わせによる高性能パターンの作成》
解となるパターンを作る際も試行錯誤が必要となる。６つのオフロード試行で、通常ＣＰＵよりも高速になるのが１つだけの場合は、そのパターンの実行ファイルを商用環境に配置すればよい。

なお、実行ファイル配置時は、最高性能のパターン選択だけでなく、利用するデバイスの価格に応じて、コストパフォーマンスが高いパターンを選択してもよい。なお、検証で複数が高速化できた場合には、解となるパターンは検討が必要である。

オフロードする際に、同一ノードのＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵにオフロードは高速化が容易だが、別ノードにオフロードする場合は、ＭＰＩ等の技術仕様が必要となり、また、通信オーバヘッドのため高速化は難しい。またサーバは、メニーコアＣＰＵとＧＰＵを備えた構成、メニーコアＣＰＵとＦＰＧＡを備えた構成が一般的だが、メニーコアＣＰＵとＧＰＵとＦＰＧＡを備えた構成は稀である。

複数アクセラレータでの高速化が可能な場合は、同一ノードにオフロードできることは最初の条件とする。同一ノードのメニーコアＣＰＵとＧＰＵは候補になり得るが、別ノードのＧＰＵとＦＰＧＡは候補にならない。

まず、機能ブロックオフロードに関して検討する。機能ブロックオフロードでは以下のルールとする。
（ａ） 同じ機能ブロックを複数種類のアクセラレータにオフロード可能な場合にはより高性能化効果が高いオフロード先を選択する。
（ｂ） 異なる機能ブロックが異なる種類のアクセラレータにオフロード可能で、かつ同一ノードにオフロードできない場合は、より高性能化効果が高い機能ブロックをオフロード対象にする。
（ｃ） 異なる機能ブロックが異なる種類のアクセラレータにオフロード可能で、かつ同一ノードにオフロードできる場合は、それら機能ブロックをオフロード対象にする。

ルール（ａ）の例として、シングルコアＣＰＵで１０秒を要するＦＦＴ処理を、ＧＰＵにオフロードすると２秒を要し、ＦＰＧＡにオフロードすると１秒を要し、メニーコアＣＰＵにオフロードすると５秒を要する場合を考える。この場合、オフロード先のアクセラレータとしてＦＰＧＡを選択するとよい。

ルール（ｂ）の例として、シングルコアＣＰＵで２０秒を要する乱数処理をＧＰＵにオフロードすると２秒を要し、シングルコアＣＰＵで１００秒を要するＦＦＴ処理をＦＰＧＡにオフロードすると１秒を要し、シングルコアＣＰＵで５０秒を要する行列計算をメニーコアＣＰＵにオフロードすると１０秒を要する場合を考える。この場合は、最も効果の高いＦＦＴ処理をＦＰＧＡにオフロードするパターンを選択するとよい。

ルール（ｃ）については、同一ノードの複数アクセラレータに異なる機能ブロックをオフロードして性能検証し、単体オフロードよりも高速の場合は解とする。
ルール（ｃ）の例として、シングルコアＣＰＵで２０秒を要する乱数処理をＧＰＵにオフロードすると２０秒を要し、シングルコアＣＰＵで１０秒を要するＦＦＴ処理をＦＰＧＡにオフロードすると１秒を要し、シングルコアＣＰＵで５０秒を要する行列計算をメニーコアＣＰＵにオフロードすると１０秒を要する場合を考える。更に、ＧＰＵとメニーコアＣＰＵの両方を搭載したノードがあるとする。この場合は、同一ノード上で、乱数処理はＧＰＵに、行列計算はメニーコアＣＰＵにオフロードするパターンを選択するとよい。

同一ノード上の複数種類のアクセラレータにオフロードする場合は、複数種類のアクセラレータに同時にオフロードすることで確実に高速化されるとも限らない。よって、オフロードサーバ１は、複数種類のアクセラレータに同時にオフロードして性能検証し、単体でオフロードするよりも高速化されていることを確認する。
このようにすることで、複数種類のアクセラレータが混在する環境に対して、機能ブロックを効率的に高速化可能なパターンを探索できる。

なお、利用頻度が高い処理である機能ブロックを、高い優先度でオフロードするようにしてもよい。その場合は、先行して行う６つのオフロード検証にて、性能測定するサンプルテストで、優先度が高い処理を中心的に行うサンプルテストとすればよい。

図４Ａは、機能ブロックを１ノードにオフロードする処理のフローチャートである。
オフロードサーバ１の制御部１１は、同じ機能ブロックを複数種類のアクセラレータにオフロード可能か否かを判定する（Ｓ５０）。制御部１１は、同じ機能ブロックを複数種類のアクセラレータにオフロード可能ならば（Ｙｅｓ）、より高性能化効果が高いオフロード先を選択し（Ｓ５１）、図４Ａの処理を終了する。制御部１１は、同じ機能ブロックを複数種類のアクセラレータにオフロードできないならば（Ｎｏ）、ステップＳ５２に進む。

ステップＳ５２において、制御部１１は、異なる機能ブロックが異なる種類のアクセラレータにオフロード可能か否かを判定する。制御部１１は、異なる機能ブロックが異なる種類のアクセラレータにオフロード可能ならば（Ｙｅｓ）、これら機能ブロックをオフロード対象にすると（Ｓ５５）、図４Ａの処理を終了する。
制御部１１は、異なる機能ブロックが異なる種類のアクセラレータにオフロードできないならば（Ｎｏ）、より高性能化効果が高い機能ブロックをオフロード対象にして（Ｓ５４）、図４Ａの処理を終了する。

次に、ループ文のオフロードに関して検討する。
ループ文のオフロードは、機能ブロックオフロードが可能な場合、オフロード可能な機能ブロック部分を除いたコードに対して行われる。ループ文オフロードでの高速化が、機能ブロックオフロードでのオフロード先と別ノードに対してのオフロードとなり、高性能化効果も小さい場合は、ループ文オフロードは行わない。

例として、シングルコアＣＰＵで１００秒を要するＦＦＴ処理をオフロードしてＦＰＧＡで１秒に処理できている場合、残りのループ文をＧＰＵにオフロードしたときに２０秒が５秒に高速化できることが分かったとしても、高性能化効果が小さいので、別ノードのＧＰＵに対してはオフロードは行わない。

ループ文のオフロードは、以下のルールとする。
（ｄ） 同じループ文が複数種類のアクセラレータにオフロードできる場合は、より高性能化効果が高いオフロード先を選択する。

（ｅ） 異なるループ文が異なる種類のアクセラレータにオフロード可能で、かつ同一ノードにオフロードできない場合は、より高性能化効果が高いループ文群をオフロード対象にする。

（ｆ） 異なるループ文が異なる種類のアクセラレータにオフロード可能で、かつ同一ノードにオフロードできる場合は、高性能化効果が高いループ文を優先してオフロードし、オフロードされていない残りのループ文群は高性能化効果が低いオフロード先にオフロードする。

ルール（ｄ）の例として、シングルコアＣＰＵで２０秒を要するＡＢＣＤＥのループ文をＧＰＵにオフロードすると２秒を要し、ＦＰＧＡにオフロードすると１秒を要し、メニーコアＣＰＵにオフロードすると５秒を要する場合を考える。この場合、ＡＢＣＤＥのループ文をＦＰＧＡにオフロードするパターンを選択するとよい。

ルール（ｅ）の例として、シングルコアＣＰＵで５０秒を要するＡＢＣのループ文をＧＰＵにオフロードすると２秒を要し、シングルコアＣＰＵで１００秒を要するＢＣＤのループをＦＰＧＡにオフロードすると１秒を要し、シングルコアＣＰＵで５０秒を要するＣＤＥのループ文をメニーコアＣＰＵにオフロードすると１０秒を要する場合を考える。この場合は、ＢＣＤのループ文をＦＰＧＡにオフロードするパターンを選択するとよい。

ルール（ｆ）については、同一ノードの複数種類のアクセラレータに異なるループをオフロードして性能検証し、単体よりも高速の場合はそのパターンを解として選択する。ただし、同一ファイルで定義された複数のループ文については別のアクセラレータに処理が分かれないように制限を行う。これは、関連が深い同一ファイル上の複数のループ文については、アクセラレータが分かれることでの複雑度向上を避けるためである。

ルール（ｆ）の例として、シングルコアＣＰＵで５０秒を要するＡＢＣのループ文をＧＰＵにオフロードすると２秒を要し、シングルコアＣＰＵで１０秒を要するＢＣＤのループ文をＦＰＧＡにオフロードすると１秒を要し、シングルコアＣＰＵで５０秒を要するＣＤＥのループ文をメニーコアＣＰＵにオフロードすると１０秒を要したと仮定し、かつＧＰＵとメニーコアＣＰＵの両方を搭載したノードがある場合を考える。この場合は、同一ノード上で、ＡＢＣのループ文はＧＰＵにオフロードし、ＤＥのループ文はメニーコアＣＰＵにオフロードすることを試行する。

図４Ｂは、ループ文を１ノードにオフロードする処理のフローチャートである。
オフロードサーバ１の制御部１１は、同じループ文を複数種類のアクセラレータにオフロード可能か否かを判定する（Ｓ６０）。制御部１１は、同じループ文を複数種類のアクセラレータにオフロード可能ならば（Ｙｅｓ）、より高性能化効果が高いオフロード先を選択し（Ｓ６１）、図４Ｂの処理を終了する。制御部１１は、同じループ文を複数種類のアクセラレータにオフロードできないならば（Ｎｏ）、ステップＳ６２に進む。

ステップＳ６２において、制御部１１は、異なるループ文が異なる種類のアクセラレータにオフロード可能か否かを判定する。制御部１１は、異なるループ文が異なる種類のアクセラレータにオフロード可能ならば（Ｙｅｓ）、これらループ文をオフロード対象にすると（Ｓ６５）、図４Ｂの処理を終了する。
制御部１１は、異なるループ文が異なる種類のアクセラレータにオフロードできないならば（Ｎｏ）、より高性能化効果が高いループ文をオフロード対象にして（Ｓ６４）、図４Ｂの処理を終了する。

同一ノード上の複数種類のアクセラレータにオフロードする場合は、確実に高速化されるとも限らないので、実際にオフロードして性能検証を行い、単体でオフロードするよりも高速化されていることを確認するとよい。
このようにすることで、複数種類のアクセラレータが混在する環境に対して、ループ文を効率的に高速化可能なパターンを探索できる。

このように、６つの性能測定の中で、高速化できているパターンを元に、最高性能となるパターンを作成する。同一ノードに対しては複数のアクセラレータにオフロードすることも可能であるが、複数のオフロードパターンを組み合わせる際は、その組み合わせで単体の場合よりも高性能化できているかを性能測定して確認することを通じて、最高性能のパターンを作成するため、複数回の測定が必要である。

図５Ａと図５Ｂは、機能ブロックをアクセラレータへオフロードするパターンの選択処理のフローチャートである。この選択処理は、非特許文献４の手法を用いている。
最初、コード分析部１１２は、C/C++ソースコードのオフロード分析を行う（Ｓ３０）。具体的には、コード分析部１１２は、Clang等の構文解析ツールを用いて、コードに含まれるライブラリ呼び出しや、機能処理を分析する。

オフロード範囲抽出部１１４ａは、このソースコードの外部ライブラリ呼び出しを検出し（Ｓ３１）、コードパターンデータベース１３３から、ライブラリ名をキーに、置換可能なアクセラレータライブラリまたはアクセラレータ用ＩＰコアを取得する（Ｓ３２）。ここでアクセラレータとは、メニーコアＣＰＵやＧＰＵやＦＰＧＡのうち何れかである。具体的にいうと、オフロード範囲抽出部１１４ａは、呼び出されているライブラリをキーに、コードパターンデータベース１３３に登録されているレコードから、高速化可能な実行ファイルやOpenCL等を取得する。

ここでオフロード範囲抽出部１１４ａは、把握した外部ライブラリ呼び出しについて、呼び出されているライブラリ名をキーとしてコードパターンデータベース１３３と照合することで、置換可能なアクセラレータライブラリまたはアクセラレータ用ＩＰコアを取得する。ここでアクセラレータライブラリとは、メニーコアＣＰＵやＧＰＵ向けの実行ファイルである。アクセラレータ用ＩＰコアとは、OpenCL等である。

なお、機能ブロックに対して、高速化するライブラリやＩＰコアがあるかの探索は、ライブラリ等の名前一致に加えて類似性検出ツールでの検出も行う。類似性検出ツールとは、Deckard等の、コピーコードやコピー後に変更したコードを検出するための、ソフト工学で利用されるツールである。類似性検出ツールは、abstract syntax tree類似性等を見て、コードの類似性を判定し、ＣＰＵで計算する場合は記述が同様になる処理や、他者の参照コードをコピーして変更した場合等を検出できる。

処理パターン作成部１１５は、ソースコードの置換元の処理記述を、置換先のアクセラレータに合わせて、メニーコアＣＰＵやＧＰＵ向けライブラリ、またはＦＰＧＡ向けのＩＰコアの処理記述に置換する（Ｓ３３）。更に処理パターン作成部１１５は、置換したアクセラレータライブラリまたはアクセラレータＩＰコアの処理記述を、オフロード対象の機能ブロックとして、アクセラレータにオフロードする（Ｓ３４）。この際に、オフロードできる処理が実際に高速化につながるか、コスト効果が十分かは分からないので、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを自動抽出する。

処理パターン作成部１１５は、アクセラレータライブラリやアクセラレータＩＰコア呼び出しのためのインタフェース処理を記述する（Ｓ３５）。このインタフェース処理は、例えばＦＰＧＡに対するOpenCL API等である。
次に実行ファイル作成部１１７は、作成したパターンをコンパイルする（Ｓ３６）。性能測定部１１６は、作成したパターンを検証環境で性能測定する（Ｓ３７）。これは、１回目の性能測定である。

１回目の性能測定後、実行ファイル作成部１１７は、１回目測定時に高速化できたパターンについて組み合わせパターンを作成し（Ｓ３８）、作成した組み合わせパターンをコンパイルする（Ｓ３９）。そして、性能測定部１１６は、作成した組み合わせパターンを検証環境で性能測定する（Ｓ４０）。これは、２回目の性能測定である。

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

非特許文献４で評価の通り、個々のループ文のオフロードに比べ、複数のループ文を含む機能ブロック単位で処理するアクセラレータに向けて、チューンした機能ブロックのオフロードは、高速化の度合いが高い。

《実装》
オフロードサーバ１は、コード解析にて、呼び出されているライブラリや定義されているクラス、構造体等のプログラム構造を把握する。

次にオフロードサーバ１は、呼び出されているライブラリを高速化できるＧＰＵ用ライブラリ、ＦＰＧＡ用ＩＰコア等の検出を行う。
オフロードサーバ１は、呼び出されているライブラリをキーに、コードパターンデータベース１３３に登録されているレコードから、高速化可能な実行ファイルやOpenCL等を取得する。高速化できる置換用機能が検出されたら、オフロードサーバ１は、その実行用ファイルを作成する。

メニーコアＣＰＵやＧＰＵ用ライブラリの場合は、置換用ライブラリ（ＣＵＤＡライブラリ等）を呼び出すよう、元の部分は削除して置換記述する。ＦＰＧＡ用ＩＰコアの場合は、取得したOpenCLコードを、元の部分をホストコードから削除してから、カーネルコードに置換記述する。それぞれ、置換記述が終わったら、メニーコアＣＰＵ向けにはgcc、ＧＰＵ向けにはPGIコンパイラ、ＦＰＧＡ向けにはIntel Acceleration Stackでコンパイルする。

ライブラリ呼び出しの場合について記載したが、類似性検知を用いる場合も並行して処理がされる。Deckardを用いて、検出されたクラス、構造体等の部分コードとＤＢに登録された比較用コードの類似性検知を行い、類似性が閾値越えの機能ブロックと該当するＧＰＵ用ライブラリやＦＰＧＡ用ＩＰコア等を発見する。特に置換元のコードと置換するライブラリやＩＰコアの引数や戻り値、型等のインタフェースが異なる場合、オフロードサーバ１は、オフロードを依頼したユーザに対して、置換先ライブラリやＩＰコアに合わせて、インタフェースを変更してよいか確認し、確認後に実行用ファイルを作成する。

ここで、オフロードサーバ１は、検証環境のＧＰＵやＦＰＧＡで性能測定できる実行用ファイルを作成する。
機能ブロックオフロードにおいて、オフロードサーバ１は、置換機能ブロック一つずつに対してオフロードするしないを性能測定して高速化できるか確認する。６つのオフロード試行で、通常ＣＰＵより高速のパターンが、１つだけの場合はそのパターンを選択するが、複数の際には解となるパターンを、以下のロジックで作成する。

［自動オフロード動作］
本実施形態のオフロードサーバ１は、環境適応ソフトウェアシステムの要素技術として、ユーザアプリケーションロジックをアクセラレータに自動オフロードする技術に適用した例である。

図６は、オフロードサーバ１の自動オフロード処理を示す図である。
図６に示すように、オフロードサーバ１は、環境適応ソフトウェアシステムの要素技術に適用される。オフロードサーバ１は、制御部１１と、テストケースデータベース１３１と、中間言語ファイル１３２と、検証用マシン１４と、を有している。
オフロードサーバ１は、ユーザが利用するソースコード１３０を取得する。

ユーザは、リソース１５を利用する。リソース１５は、例えば、装置１５１、ＣＰＵとＧＰＵを有する装置１５２、ＣＰＵとＦＰＧＡを有する装置１５３、ＣＰＵを有する装置１５４である。オフロードサーバ１は、機能処理をＣＰＵとＧＰＵを有する装置１５２や、ＣＰＵとＦＰＧＡを有する装置１５３のアクセラレータに自動オフロードする。

以下、図６のステップ番号を参照して各部の動作を説明する。
《ステップＳ１１：ソースコードの指定》
ステップＳ１１において、コード指定部１１１（図１参照）は、ユーザに提供しているサービスの処理機能（画像分析等）を特定する。具体的には、コード指定部１１１は、入力されたソースコードを指定する。

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

《ステップＳ１３：オフロード範囲の抽出》
ステップＳ１３において、処理指定部１１４（図１参照）は、アプリケーションのソースコードに含まれるループ文を特定し、各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする。具体的には、オフロード範囲抽出部１１４ａ（図１参照）は、ループ文やＦＦＴ等、ＧＰＵ・ＦＰＧＡにオフロード可能な範囲を抽出する。

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

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

《ステップＳ２１：実行ファイルの配置》
ステップＳ２１において、実行ファイル配置部１１６ａ（図１参照）は、メニーコアＣＰＵとＧＰＵとＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイする。

《ステップＳ２２：性能の測定》
ステップＳ２２において、性能測定部１１６（図１参照）は、配置したファイルを実行し、オフロードした際の性能を測定する。
オフロードする領域をより適切にするため、この性能測定結果は、処理パターン作成部１１５に戻され、処理パターン作成部１１５が、別のオフロードパターンを作成する。そして、性能測定部１１６は、別のオフロードパターンの性能を測定する（図６の符号ｅ参照）。

図６の矢印ｅに示すように、制御部１１は、ステップＳ１２からステップＳ２２の処理を繰り返し実行する。

制御部１１の自動オフロード機能をまとめると、下記である。
すなわち、処理指定部１１４は、ソースコードのループ文を特定し、各ループ文に対して、並列処理指定文を指定して、コンパイルする。そして、処理パターン作成部１１５は、コンパイルエラーが出るループ文を、オフロード対象外とし、コンパイルエラーが出ないループ文に対して、並列処理するか否かの指定を行う並列処理パターンを作成する。そして、性能測定部１１６は、該当並列処理パターンのソースコードをコンパイルして、検証用マシン１４に配置し、検証用マシン１４で性能測定用処理を実行する。実行ファイル作成部１１７は、所定回数繰り返された、性能測定結果をもとに、複数の並列処理パターンから最高処理性能のパターンを選択し、選択パターンをコンパイルして実行ファイルを作成する。

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

《ステップＳ２４：性能測定テストと実行》
ステップＳ２４において、性能測定テスト抽出実行部１１９は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースデータベース１３１から抽出し、抽出した性能試験を自動実行する。

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

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

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

［遺伝的アルゴリズムを用いたループ文自動オフロード］
ＧＰＵやメニーコアＣＰＵへのループ文自動オフロードは、図６のステップＳ１２～ステップＳ２２を繰り返し、最終的にステップＳ２３でデプロイするオフロードコードを得るための処理である。

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

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

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

そこで、本実施形態では、並列化を想定していない汎用プログラムから、自動で適切なオフロード領域を抽出する。このため、オフロードサーバ１は、機能ブロックのオフロードをメニーコアＣＰＵ→ＧＰＵ→ＦＰＧＡの順でチェックしたのち、ループ文のオフロードをメニーコアＣＰＵ→ＧＰＵ→ＦＰＧＡの順でチェックする。オフロード効果の高い機能ブロックをループ文よりも先にチェックすることと、検証コストが低い順であるメニーコアＣＰＵ→ＧＰＵ→ＦＰＧＡの順でチェックすることにより、効率的に高速化可能なパターンを探索できる。

［単純遺伝的アルゴリズムによる制御部１１の探索イメージ］
図７は、オフロードサーバの単純遺伝的アルゴリズムによる制御部１１の探索イメージを示す図である。
遺伝的アルゴリズムは、生物の進化過程を模倣した組み合わせ最適化手法の一つである。遺伝的アルゴリズムのフローチャートは、初期化→評価→選択→交叉→突然変異→終了判定となっている。

本実施形態では、遺伝的アルゴリズムの中で、処理を単純にした単純遺伝的アルゴリズムを用いる。単純遺伝的アルゴリズムは、遺伝子は１と０のみとし、ルーレット選択、一点交叉、突然変異は１箇所の遺伝子の値を逆にする等、単純化された遺伝的アルゴリズムである。

《初期化》
初期化では、ソースコードの全ループ文の並列可否をチェック後、並列可能ループ文を遺伝子配列にマッピングする。ＧＰＵ処理する場合は１、ＧＰＵ処理しない場合は０とする。遺伝子は、指定の個体数Ｍを準備し、１つのループ文にランダムに１と０を割り当てる。

具体的には、制御部１１（図１参照）は、ユーザが利用するソースコード１３０（図６参照）を取得し、ソースコード１３０のコードパターン１４１からループ文の並列可否をチェックする。図８に示すように、コードパターン１４１ｄから３つのループ文が検出された場合、各ループ文に対して１桁、ここでは３つのループ文に対し３桁の１または０をランダムに割り当てる。例えば、シングルコアＣＰＵで処理する場合には０、メニーコアＣＰＵに出す場合には１とする。ただし、この段階では１または０をランダムに割り当てる。
遺伝子長に該当するコードが３桁であり、３桁の遺伝子長のコードは２^３＝８パターン、例えば１００、１１０、…となる。なお、図７では、コードパターン１４１中の丸印（○印）をコードのイメージとして示している。

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

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

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

《突然変異》
突然変異では、一定の突然変異率Ｐｍで、個体の遺伝子の各値を０から１または１から０に変更する。
また、局所解を避けるため、突然変異を導入する。なお、演算量を削減するために突然変異を行わない態様でもよい。そして、ステップＳ３０２で示すように、これら交叉や突然変異を施した各コードを次世代のコードとする。

《終了判定》
クロスオーバーと突然変異後の次世代コードパターンを生成する。
終了判定では、指定の世代数Ｔ回、繰り返しを行った後に処理を終了し、最高適合度の遺伝子を解とする。
例えば、性能測定して、処理速度が速い個体を３つ、例えば遺伝子１１０，０１０，００１の個体を選ぶ。この３つを遺伝的アルゴリズムにより、次の世代は、組み換えをして、例えば新しいパターン（並列処理パターン）０１１（一例）を作っていく。このとき、組み換えをしたパターンに、勝手に０を１にするなどの突然変異を入れる。上記を繰り返して、一番処理速度が速いパターンを見付ける。指定世代（例えば、２０世代）などを決めて、最終世代で残ったパターンを、最後の解とする。

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

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

OpenACCには、ディレクティブ“#pragma acc kernels”で指定して、ＧＰＵ向けバイトコードを抽出し、実行によりＧＰＵオフロードを可能とするコンパイラがある。この“#pragma”に、ループ文のコマンドを書くことにより、そのループ文がＧＰＵで動くか否かを判定することができる。

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

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

検証用マシン１４でデプロイして、ベンチマーク、例えば画像処理であればその画像処理でベンチマークする、その処理時間が短い程、適応度が高いと評価する。例えば、処理時間の逆数、処理時間に１０秒かかる個体の適応度は１とする。処理時間に１００秒かかる個体の適応度は０．１とする。処理時間に１秒かかる個体の適応度は１０とする。

オフロードサーバ１は、適応度が高いものを選択する。オフロードサーバ１は、例えば１０個の個体のなかから、３～５個を選択して、それを組み替えて新しいコードパターンを作る。作成途中で、前と同じものができる場合がある。この場合、同じベンチマークを行う必要はないので、オフロードサーバ１は、前と同じ適応度のデータを使う。本実施形態のオフロードサーバ１は、コードパターンと、その処理時間を記憶部１３に保存している。

以上で、単純遺伝的アルゴリズムによる制御部１１の探索イメージについて説明した。

図８は、ループ文のメニーコアＣＰＵへの遺伝子配列マッピングを示す図である。
コードパターン１４１ｄは、３つのループ文を含んでいる。ここで、各ループ文に対してバイナリ１桁が割り振られ、３つのループ文に対して３桁の１または０がランダムに割り当てられる。

コードパターン１４１ｄの最初のforループ文には、ディレクティブ“#pragma omp parallel for”が付与されている。このとき、最初のforループ文は、メニーコアにオフロードするようにコンパイルされる。更に、コードパターン１４１ｄの左側には、コードパターン１４１ｄの遺伝子配列100が示されている。

図９Ａと図９Ｂは、ループ文をメニーコアＣＰＵへオフロードするパターンの選択処理のフローチャートである。
最初、コード分析部１１２は、C/C++ソースコードのループ文を分析する（Ｓ９０）。処理指定部１１４は、C/C++ソースコードのループ文、参照関係を特定する（Ｓ９１）。

処理指定部１１４は、ベンチマークツールを動作させて、各ループ文のループ回数を把握すると共に、各ループを閾値で振分けると（Ｓ９２）、各ループ文の並列処理可能性をチェックする（Ｓ９３）。

なお、ＧＰＵでの自動オフロードで用いていたPGIコンパイラは、並列化不能時はコンパイラがエラーを出力していた。しかし、gcc等のOpenMPコンパイラはそのようなエラーはプログラマの責任となる。そこで、自動化するために、OpenMP指示句で行う処理は、ループ文をメニーコアＣＰＵで並列処理するかどうかだけと単純化するとともに、並列処理した場合の最終計算結果が正しいかどうかのチェックも性能測定時に行うことで、正しい計算結果が出るパターンだけ進化計算の中で残る仕組みとする。

次に処理指定部１１４は、並列処理可能性を有するループ文の数をカウントして遺伝子長とし（Ｓ９４）、初期値として指定個体数の遺伝子配列を準備する（Ｓ９５）。ここで、処理指定部１１４は、０と１をランダムに割り当てた遺伝子配列を所定個体数だけ作成する。

処理指定部１１４は、C/C++ソースコードを、遺伝子にマッピングし、指定個体数パターンを準備する（Ｓ９６）。具体的には、処理指定部１１４は、準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをC/C++コードに挿入する（例えば図８の#pragmaディレクティブ参照）。

制御部１１は、ステップＳ１００～Ｓ１０８の処理を、指定世代数だけ繰り返す。
また、指定世代数の繰り返しにおいて、さらにステップＳ１０１～Ｓ１０４の処理について指定個体数だけ繰り返す。すなわち、指定世代数の繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。

指定個体数の繰り返しにおいて、処理パターン作成部１１５（図１参照）は、遺伝子パターンに応じてディレクティブ指定したC/C++ソースコード、すなわち、各オフロードパターンをgccコンパイラでコンパイルする（Ｓ１０２）。すなわち、処理パターン作成部１１５は、作成したC/C++コードを、メニーコアＣＰＵを備えた検証用マシン１４上のgccコンパイラでコンパイルする。
ここで、ネストした複数のループ文を並列処理を指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。

性能測定部１１６（図１参照）は、メニーコアＣＰＵ搭載の検証用マシン１４に、オフロードパターンをコンパイルした実行ファイル、すなわち、コンパイルしたオフロードパターンをデプロイして（Ｓ１０３）、配置した実行ファイルを実行し、オフロードした際のベンチマーク性能を測定する（Ｓ１０４）。

ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、遺伝的アルゴリズム処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。

ステップＳ１０５において、制御部１１は、指定個体数の繰り返しが完了していないならば、ステップＳ１０１の処理に戻り、指定個体数を全て繰り返したならば、ステップＳ１０６に進む。

次にステップＳ１０６で、実行ファイル作成部１１７（図１参照）は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。
実行ファイル作成部１１７（図１参照）は、各個体の計算結果とシングルコアＣＰＵの計算結果との差を評価し、許容範囲となる個体を選択する（Ｓ１０７）。実行ファイル作成部１１７は、性能測定の際に、最終計算結果が並列処理しない場合と同じ結果であることを、オリジナルコードを通常ＣＰＵで処理した場合と比較し、もし差分が許容できない程大きい場合はそのパターンの適応度は０として次世代には選ばれないようにする。実行ファイル作成部１１７は、並列処理した場合の最終計算結果が正しいかどうかのチェックも性能測定時に行うことで、進化計算の中で、正しい計算結果が出るパターンに絞り込むことができる。

実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する（Ｓ１０８）。これは、次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行うためである。

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

ステップＳ１０９において、制御部１１は、指定世代数の繰り返しが完了していないならば、ステップＳ１００の処理に戻り、指定世代数を全て繰り返したならば、ステップＳ１１０に進む。

ステップＳ１１０で、実行ファイル作成部１１７は、指定世代数の遺伝的アルゴリズム処理終了後、最高性能の遺伝子配列に該当するC/C++コード（最高性能の並列処理パターン）を解とすると、図９Ｂの処理を終了する。
このようにすることで、メニーコアＣＰＵの環境が移行先であっても、ループ文を効率的に高速化可能なパターンを探索できる。

メニーコアＣＰＵもＧＰＵと同様に、多数の計算コアを生かして、処理を並列化することで高速化する。ＧＰＵと異なる点は、メニーコアＣＰＵの場合、メモリは共通であるため、ＧＰＵへのオフロードでしばしば問題となったＣＰＵとＧＰＵのメモリ間のデータ転送によるオーバヘッドは考慮する必要がない。また、メニーコアＣＰＵでのプログラム処理の並列化には、OpenMP仕様が頻繁に利用される。OpenMPは、“#pragma omp parallel for”等の指示句でプログラムに対して、並列処理等を指定する仕様である。OpenMPでの処理並列化は、OpenMPプログラマが責任を持つこととなっており、並列化できない処理を並列化した場合には、コンパイラがエラー出力をするわけでなく、計算結果が誤って出力される。

これらを踏まえ、本実施形態では、ループ文のメニーコアＣＰＵ向け自動オフロードは、ループの並列処理可否を、OpenMPにおいて“#pragma”で指定するパターンを複数作成し、検証環境で実際の性能測定を繰り返すことで、徐々に高速化していく進化計算手法を採用する。

《実装》
オフロードサーバ１は、C/C++ソースコードを解析して、ループ文を発見するとともに、ループ文内で使われる変数データ、その変数の処理等の、プログラム構造を把握する。

並列処理自体が不可なループ文は排除する必要があるため、オフロードサーバ１は、各ループ文に対して、ＧＰＵで処理するディレクティブ挿入を試行し、エラーが出るループ文は遺伝的アルゴリズムの対象外とする。ここで、エラーが出ないループ文の数が遺伝子長となる。

次に、オフロードサーバ１は、初期値として、指定個体数の遺伝子配列を準備する。遺伝子の各値は、０と１をランダムに割当てて作成する。オフロードサーバ１は、準備された遺伝子配列に応じて、遺伝子の値が１の場合はＧＰＵやメニーコアＣＰＵ処理を指定するディレクティブをC/C++コードに挿入する。

オフロードサーバ１は、ディレクティブを挿入されたC/C++コードをgccでコンパイルし、コンパイルした実行ファイルをデプロイして性能を測定する。性能測定において、オフロードサーバ１は、処理時間とともに、例えば、PGIコンパイラのPCAST機能を用いて並列処理した場合の計算結果が、元のコードと大きく差分がないかチェックし、許容外の場合は、適応度を０とする。

オフロードサーバ１は、全個体に対して、性能測定後、処理時間に応じて、各個体の適合度を設定する。設定された適合度に応じて、残す個体を選択する。選択された個体に対して、交叉、突然変異、そのままコピーによる遺伝的アルゴリズム処理を行い、次世代の個体群を作成する。

オフロードサーバ１は、次世代の個体に対して更に、指示挿入、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行い、これを指定された世代数だけ繰り返す。指定世代数の遺伝的アルゴリズム処理終了後、オフロードサーバ１は、最高性能の遺伝子配列に該当する、ディレクティブ付きC/C++コードを解とする。

図１０は、ループ文のＧＰＵへの遺伝子配列マッピングを示す図である。
コードパターン１４１ｅの最初のforループ文には、“#pragma acc kernels”が付与されている。このとき、最初のforループ文は、ＧＰＵにオフロードするようにコンパイルされる。更に、コードパターン１４１ｅの左側には、コードパターン１４１ｅの遺伝子配列100が示されている。

図１１Ａと図１１Ｂは、ループ文をＧＰＵへオフロードするパターンの選択処理のフローチャートである。この選択処理は、非特許文献２の手法を用いている。
オフロードサーバ１は、C/C++向けOpenACCコンパイラを用いて以下の処理を行う。

コード分析部１１２（図１参照）は、C/C++ソースコードのコード解析によりループ文を分析する（Ｓ１２０）。そして、処理指定部１１４（図１参照）は、C/C++ソースコードのループ文とその参照関係を特定すると（Ｓ１２１）、ベンチマークツールを動作させ、各ループ文のループ回数を把握し、閾値で振分ける（Ｓ１２２）。そして、処理指定部１１４は、各ループ文の並列処理可能性をチェックする（Ｓ１２３）。

次に制御部１１は、ステップＳ１２４～Ｓ１２７の処理を、ループ文の数だけ繰り返す。
処理指定部１１４は、各ループ文に対して並列処理を指定してコンパイルする（Ｓ１２５）。図１０に示すように、OpenACCにおける並列処理の指定は、“#pragma acc kernels”である。
コンパイルのエラー時に処理指定部１１４は、並列処理の指定を削除し、当該ループ文をオフロード対象から除外する（Ｓ１２６）。そして制御部１１は、全てのループ文について処理を繰り返したならば、ステップＳ１２８に進む。

処理指定部１１４は、コンパイルエラーが出ないループ文の数をカウントして遺伝子長とすると（Ｓ１２８）、指定個体数の遺伝子配列を準備する。ここでは、初期値として０と１をランダムに割り当てた指定個体数の遺伝子配列を作成する。次に処理指定部１１４は、C/C++ソースコードを、遺伝子にマッピングして、指定個体数のパターンを準備する（Ｓ１２９）。具体的にいうと、処理指定部１１４は、準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをC/C++コードに挿入する（例えば図１０の#pragmaディレクティブ参照）。

制御部１１は、ステップＳ１３０～Ｓ１３７の処理を、指定世代数だけ繰り返す。
また、指定世代数の繰り返しにおいて、さらにステップＳ１３１～Ｓ１３４の処理について指定個体数だけ繰り返す。すなわち、指定世代数の繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。

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

性能測定部１１６（図１参照）は、ＣＰＵとＧＰＵを搭載した検証用マシン１４に、オフロードパターンをコンパイルした実行ファイル、すなわちコンパイルしたオフロードパターンをデプロイし（Ｓ１３３）、配置した実行ファイルを実行し、オフロードした際のベンチマーク性能を測定する（Ｓ１３４）。

ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、遺伝的アルゴリズム処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。

ステップＳ１３５において、制御部１１は、指定個体数の繰り返しが完了していないならば、ステップＳ１３１の処理に戻り、指定個体数を全て繰り返したならば、ステップＳ１３６に進む。

ステップＳ１３６で、実行ファイル作成部１１７（図１参照）は、処理時間が短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。そして実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する（Ｓ１３７）。これは、次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行うためである。
すなわち、実行ファイル作成部１１７は、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。そして実行ファイル作成部１１７は、設定された適合度に応じて、残す個体を選択する。そして実行ファイル作成部１１７は、選択された個体に対して、交叉、突然変異、そのままコピーの遺伝的アルゴリズム処理を行い、次世代の個体群を作成する。

ステップＳ１３８において、制御部１１は、指定世代数の繰り返しが完了していないならば、ステップＳ１３０の処理に戻り、指定世代数を全て繰り返したならば、ステップＳ１３９に進む。

指定世代数の遺伝的アルゴリズム処理終了後のステップＳ１３９において、実行ファイル作成部１１７は、最高性能の遺伝子配列に該当するC/C++コード（最高性能の並列処理パターン）を解として、図１１Ｂの処理を終了する。

基本的な課題として、コンパイラがこのループ文はＧＰＵで並列処理できないという制限を見つけることは可能だが、このループ文はＧＰＵの並列処理に適しているという適合性を見つけることは難しい。

一般的にループ回数が多い等のループの方が、ＧＰＵオフロードに適していると言われるが、実際にＧＰＵにオフロードすることでどの程度の性能になるかは、実測なしに予測は困難である。そのため、このループをＧＰＵにオフロードするという指示を手作業で行い、性能測定を試行錯誤することが頻繁に行われている。

非特許文献２では、ＧＰＵにオフロードする適切なループ文の発見を、進化計算手法の一つである遺伝的アルゴリズムで自動的に行うことを提案している。ＧＰＵ処理を想定していない通常ＣＰＵ向け汎用プログラムから、最初に並列可能ループ文のチェックを行い、次に並列可能ループ文群に対して、ＧＰＵ実行の際を１、ＣＰＵ実行の際を０と値を置いて遺伝子化して、検証環境で性能検証を反復し適切な領域を探索している。

並列可能ループ文に絞った上で、遺伝子の部分の形で、高速化可能な並列処理パターンを保持し組み換えていくことで、取り得る膨大な並列処理パターンから、効率的に高速化可能なパターンを探索している。

非特許文献２では、ループ文の適切な抽出に加えて、ネストループ文の中で利用される変数について、ループ文をＧＰＵにオフロードする際に、ネストの下位でＣＰＵ－ＧＰＵ転送が行われると下位のループの度に転送が行われ効率的でないため、上位でＣＰＵ－ＧＰＵ転送が行われても問題ない変数については、上位でまとめて転送を行うことを提案している。
ループ文のＧＰＵ自動オフロードについては、進化計算手法を用いた最適化と、ＣＰＵ－ＧＰＵ転送の低減により、自動でのオフロードを可能としている。

図１２Ａと図１２Ｂは、ループ文をＦＰＧＡへオフロードするパターンの選択処理のフローチャートである。この選択処理は、非特許文献３の手法を用いている。
最初、コード分析部１１２は、オフロードしたいソースコードを分析し（Ｓ１５０）、ソースコードの言語に合わせて、ループ文や変数の情報を分析する。

そして、オフロード範囲抽出部１１４ａは、ソースコードのループ文および参照関係を特定すると（Ｓ１５１）、特定したループ文に対して、ＦＰＧＡオフロードを試行するかどうか候補を絞る。ループ文に対してオフロード効果があるかどうかは、算術強度が一つの指標となる。

オフロード範囲抽出部１１４ａは、算術強度分析ツールを用いてアプリケーションのループ文の算術強度を算出する（Ｓ１５２）。算術強度は、計算数が多いと増加し、アクセス数が多いと減少する指標である。算術強度が高い処理は、プロセッサにとって重い処理となる。そこで、オフロード範囲抽出部１１４ａは、算術強度分析ツールでループ文の算術強度を分析し、算術強度が高いループ文をオフロード候補に絞る。

高算術強度のループ文であっても、それをＦＰＧＡで処理する際に、ＦＰＧＡリソースを過度に消費してしまうのは問題である。そこで、ここでは、高算術強度ループ文をＦＰＧＡ処理する際のリソース量の算出について述べる。

ＦＰＧＡにコンパイルする際の処理としては、OpenCL等の高位言語からハードウェア記述のＨＤＬ等のレベルに変換され、それに基づき実際の配線処理等がされる。この時、配線処理等は多大な時間がかかるが、ＨＤＬ等の途中状態の段階までは時間は分単位でしかかからない。ＨＤＬ等の途中状態の段階であっても、ＦＰＧＡで利用するFlip FlopやLook Up Table等のリソースは分かる。このため、ＨＤＬ等の途中状態の段階をみれば、利用するリソース量はコンパイルが終わらずとも短時間でわかる。

そこで、本実施形態では、処理パターン作成部１１５は、対象のループ文をOpenCL等の高位言語化し、まずリソース量を算出する。また、ループ文をオフロードした際の算術強度とリソース量が決まるため、算術強度／リソース量または算術強度×ループ回数／リソース量をリソース効率とする。そして、高リソース効率のループ文をオフロード候補として更に絞り込む。

図１２Ａのフローに戻って説明を続ける。処理パターン作成部１１５は、gcov、gprof等のプロファイリングツールを用いてアプリケーションのループ文のループ回数を測定し（Ｓ１５３）、ループ文のうち、高算術強度で高ループ回数のループ文を絞り込む（Ｓ１５４）。

処理パターン作成部１１５は、絞り込まれた各ループ文をＦＰＧＡにオフロードするためのOpenCLを作成する（Ｓ１５５）。

ここで、ループ文のOpenCL化について、補足して説明する。すなわち、ループ文をOpenCL等によって、高位言語化する際には、２つの処理が必要である。一つは、ＣＰＵ処理のプログラムを、カーネル（ＦＰＧＡ）とホスト（ＣＰＵ）に、OpenCL等の高位言語の文法に従って分割することである。もう一つは、分割する際に、高速化するための技法を盛り込むことである。一般に、ＦＰＧＡを用いて高速化するためには、ローカルメモリキャッシュ、ストリーム処理、複数インスタンス化、ループ文の展開処理、ネストループ文の統合、メモリインターリーブ等がある。これらは、ループ文によっては、絶対効果があるわけではないが、高速化するための手法として、よく利用されている。

次に、高リソース効率のループ文が幾つか選択されたので、それらを用いて性能を実測するオフロードパターンを実測する数だけ作成する。ＦＰＧＡでの高速化は、１個の処理だけＦＰＧＡリソース量を集中的にかけて高速化する形もあれば、複数の処理にＦＰＧＡリソースを分散して高速化する形もある。オフロードサーバ１は、選択された単ループ文のパターンを一定数作り、ＦＰＧＡ実機で動作する前段階としてプレコンパイルする。

処理パターン作成部１１５は、作成したOpenCLをプレコンパイルして利用するリソース量を算出し（Ｓ１５６）、高リソース効率のループ文を絞り込む（Ｓ１５７）。

そして、実行ファイル作成部１１７は、絞り込んだループ文をオフロードするためのOpenCLをコンパイルする（Ｓ１５８）。性能測定部１１６は、コンパイルされたプログラムの性能を測定する（Ｓ１５９）。これは、１回目の性能測定である。

処理パターン作成部１１５は、性能測定された中でＣＰＵに比べ高性能化されたループ文をリスト化し（Ｓ１６０）、リストのループ文を組み合わせてオフロードするOpenCLを作成する（Ｓ１６１）。処理パターン作成部１１５は、組み合わせたオフロードOpenCLでプレコンパイルして利用するリソース量を算出する（Ｓ１６２）。これは、２回目のリソース量の算出である。なお、処理パターン作成部１１５は、プレコンパイルせず、１回目測定前のプレコンパイルでのリソース量の和でもよい。このようにすれば、プレコンパイル回数を削減することができる。

実行ファイル作成部１１７は、組み合わせたオフロードOpenCLをコンパイルする（Ｓ１６３）。性能測定部１１６は、コンパイルされたプログラムの性能を測定する（Ｓ１６４）。これは、２回目の性能測定である。

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

このように、ループ文のＦＰＧＡ自動オフロードは、算術強度とループ回数が高くリソース効率が高いループ文に絞って、オフロードパターンを作り、検証環境で実測を通じて高速なパターンの探索を行う。

ＦＰＧＡで、処理時間が長時間かかる特定のループ文をＦＰＧＡにオフロードして高速化することを考えた際に、どのループをオフロードすれば高速になるかの予測は難しいため、ＧＰＵと同様に検証環境で性能測定を自動で行うことを提案している。しかし、ＦＰＧＡは、OpenCLをコンパイルして実機で動作させるまでに数時間以上かかるため、ＧＰＵ自動オフロードでの遺伝的アルゴリズムを用いて何回も反復して性能測定することは、処理時間が膨大となり現実的ではない。

そこで、本実施形態では、ＦＰＧＡにオフロードする候補のループ文を絞ってから、性能測定試行を行う形をとる。具体的には、まず、発見されたループ文に対して、ROSEフレームワーク等の算術強度分析ツールを用いて算術強度が高いループ文を抽出する。更に、gcovやgprof等のプロファイリングツールを用いてループ回数が多いループ文も抽出する。

そして、算術強度やループ回数が多いループ文を候補として、OpenCL化を行う。算術強度やループ回数が多いオフロード候補ループ文に対して、作成したOpenCLをプレコンパイルして、リソース効率が高いループ文を見つける。これは、コンパイルの途中で、作成するFlip FlopやLook Up Table等のリソースは分かるため、利用するリソース量が少ないループ文に更に絞り込む。

ループ文が幾つか選択されたため、それらを用いて性能を実測する。ＦＰＧＡでの高速化は、１個の処理だけＦＰＧＡリソース量を集中的にかけて高速化する形もあれば、複数の処理にＦＰＧＡリソースを分散して高速化する形もある。そこで、選択された単ループ文に対してＦＰＧＡ実機で動作するようコンパイルして性能測定し、更に高速化できた単ループ文に対してはその組み合わせのパターンも作り、２回目の性能測定をする。オフロードサーバ１は、検証環境で性能測定された複数パターンの中で、最高速のパターンを解として選択する。

ループ文のＦＰＧＡ自動オフロードについては、算術強度やループ回数、リソース量を用いて候補ループ文を絞り込んでから、検証環境での複数パターン性能測定を行い、自動でのオフロードを可能としている。

《実装》
オフロードサーバ１は、C/C++ソースコードを解析して、ループ文を発見するとともに、ループ文内で使われる変数データ等の、プログラム構造を把握する。
次に、オフロードサーバ１は、ROSEを実行して、各ループ文の算術強度を取得し、gcovを用いて、各ループ文のループ回数を取得する。

次に、オフロードサーバ１は、高算術強度、高ループ回数の個々のループ文をＦＰＧＡにオフロードするOpenCLコードを生成する。OpenCLコードは当該ループ文をＦＰＧＡカーネルとして、残りをＣＰＵホストプログラムとして分割したものである。

次に、オフロードサーバ１は、作成したOpenCLコードに対してプレコンパイルをして、利用するFlip Flop等のリソース量を算出する。使用リソース量は全体リソース量の割合で表示される。ここでオフロードサーバ１は、算術強度、ループ回数とリソース量から、高算術強度、高ループ回数かつ低リソース量（高リソース効率）のループ文を選定する。

次に、オフロードサーバ１は、選定したループ文を候補に、実測するパターンを作り、コンパイルする。オフロードサーバ１は、選定されたループ文一つずつに対してオフロードした場合の性能測定をして高速化できるかを確認する。その結果、複数が高速化できた場合、オフロードサーバ１は、その組み合わせのOpenCLも作成して性能測定を行い、単体よりも高速化できるか確認する。ただし、ＦＰＧＡはリソース量が限られるため、オフロードサーバ１は、複数オフロード時にリソース量が上限値に収まらない場合は作成しない。
最後に、オフロードサーバ１は、複数の測定パターンから高速なパターンを解として選択する。

《本発明とその作用効果》
請求項１に記載の発明では、ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析するコード分析部と、前記機能ブロックを前記アクセラレータにオフロードするパターンを作成する処理パターン作成部と、前記機能ブロックをオフロードする各前記パターンを検証環境にデプロイして性能を測定する性能測定部と、前記性能測定部が測定した性能が所望のものであった場合、前記処理パターン作成部により、前記ソフトウェアプログラムのうち、前記機能ブロックを除く残りのループ文をオフロードするパターンを作成して、前記性能測定部により、各前記パターンを前記検証環境にデプロイして性能を測定し、前記性能測定部が測定した性能が所望のものでなかった場合、前記処理パターン作成部により、前記ソフトウェアプログラムのループ文をオフロードするパターンを作成して、前記性能測定部により、各前記パターンを前記検証環境にデプロイして性能を測定する制御部と、を備えることを特徴とするオフロードサーバとした。

このようにすることで、複数種類のアクセラレータが混在する環境が移行先であっても、自動で高性能化できる。更に移行先が同一ノードに複数種類のアクセラレータを持つ場合は、複数種類のアクセラレータに対して同時にオフロードすることで、単一アクセラレータの利用の場合よりも高速なオフロードを実現する。

請求項２に記載の発明では、前記アクセラレータは複数種類であり、前記制御部は、各前記アクセラレータにオフロードするパターンのうち、最も高速なものを選択する、ことを特徴とする請求項１に記載のオフロードサーバとした。

このようにすることで、複数種類のアクセラレータが混在する環境が移行先であっても、自動で高性能化できる。

請求項３に記載の発明では、前記制御部は、前記アクセラレータへのオフロード対象として前記機能ブロックの試行の次に前記ループ文を試行し、前記機能ブロックまたは／および前記ループ文のオフロード先の前記アクセラレータとして、メニーコアＣＰＵ、ＧＰＵ、ＦＰＧＡの順に試行する、ことを特徴とする請求項１または２に記載のオフロードサーバとした。

このようにすることで、複数種類のアクセラレータが混在する環境に対して、効率的に高速化可能なパターンを探索できる。

請求項４に記載の発明では、前記制御部は、前記機能ブロックまたは前記ループ文のオフロードにおいて、元の処理よりも高速なパターンが複数のアクセラレータに対して見つかり、かつ同じ機能ブロックまたはループ文が前記複数のアクセラレータにオフロードできる場合、より性能が高くなるアクセラレータをオフロード先とし、異なる機能ブロックまたは異なるループ文が異なるアクセラレータにオフロードでき，かつ同一ノードのアクセラレータにオフロードできない場合、より性能が高くなるアクセラレータをオフロード先とし、異なる機能ブロックまたは異なるループ文が異なるアクセラレータにオフロードでき，かつ同一ノードのアクセラレータにオフロードできる場合、前記機能ブロックまたは前記ループ文のオフロード先をそれぞれ異なるアクセラレータとする、ことを特徴とする請求項１に記載のオフロードサーバとした。

このようにすることで、複数種類のアクセラレータが混在する環境に対して、機能ブロックまたはループ文を効率的に高速化可能なパターンを探索できる。

請求項５に記載の発明では、前記制御部は、前記機能ブロックと前記ループ文のオフロードにおいて、前記ループ文のオフロード先が，前記機能ブロックのオフロード先とは別ノードとなる場合，ループ文のオフロードによる性能の向上が前記機能ブロックのオフロードによる性能向上以下ならば、前記ループ文をオフロードするパターンを試行せず、かつ前記機能ブロックをオフロードするパターンを試行し、前記ループ文のオフロードによる性能の向上が前記機能ブロックのオフロードによる性能向上を超えたならば、前記機能ブロックをオフロードするパターンを試行せず、かつ前記ループ文をオフロードするパターンを試行し、前記ループ文のオフロード先が前記機能ブロックのオフロード先と同一ノードとなる場合，前記機能ブロックと前記ループ文とをオフロードしたパターンを試行する、ことを特徴とする請求項１に記載のオフロードサーバとした。

このようにすることで、複数種類のアクセラレータが混在する環境に対して、機能ブロックまたはループ文を効率的に高速化可能なパターンを探索できる。

請求項６に記載の発明では、ソフトウェアプログラムのソースコードが、メニーコアＣＰＵによる実行が可能なループ文を含むか否かを分析するコード分析部と、前記ループ文を前記メニーコアＣＰＵにオフロードするパターンを作成する処理パターン作成部と、各前記ループ文を遺伝子パターンにマッピングし、各前記ループ文をオフロードする各前記パターンを検証環境にデプロイして性能を測定し、高い性能を高い適応度に設定して、次世代のパターンを遺伝的アルゴリズムのエリート選択，交叉，突然変異の処理により作成することを繰り返す性能測定部と、を備え、前記性能測定部は、前記検証環境での性能の測定において、前記メニーコアＣＰＵが処理した場合の計算結果と、元の処理の計算結果を比較し、前記メニーコアＣＰＵが処理した場合の計算結果の誤差が許容閾値を超える場合、遺伝子パターンの適応度を低く設定する、ことを特徴とするオフロードサーバとした。

このようにすることで、メニーコアＣＰＵの環境が移行先であっても、ループ文を効率的に高速化可能なパターンを探索できる。

請求項７に記載の発明では、コード分析部が、ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析するステップと、処理パターン作成部が、前記機能ブロックを前記アクセラレータにオフロードするパターンを作成するステップと、性能測定部が、前記機能ブロックをオフロードする各前記パターンを検証環境にデプロイして性能を測定するステップと、を実行し、前記性能測定部が測定した性能が所望のものであった場合、前記処理パターン作成部が、前記ソフトウェアプログラムのうち、前記機能ブロックを除く残りのループ文をオフロードするパターンを作成するステップと、前記性能測定部が、各前記パターンを前記検証環境にデプロイして性能を測定するステップと、を実行し、前記性能測定部が測定した性能が所望のものでなかった場合、前記処理パターン作成部が、前記ソフトウェアプログラムのループ文をオフロードするパターンを作成するステップと、前記性能測定部が、各前記パターンを前記検証環境にデプロイして性能を測定するステップと、を実行する、ことを特徴とするオフロード制御方法とした。

このようにすることで、複数種類のアクセラレータが混在する環境が移行先であっても、自動で高性能化できる。更に移行先が同一ノードに複数種類のアクセラレータを持つ場合は、複数種類のアクセラレータに対して同時にオフロードすることで、単一アクセラレータの利用の場合よりも高速なオフロードを実現する。

請求項８に記載の発明では、コンピュータに、ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析させる手順、前記機能ブロックを前記アクセラレータにオフロードするパターンを作成させる手順、前記機能ブロックをオフロードする各前記パターンを検証環境にデプロイして性能を測定させる手順、測定した性能が所望のものであった場合、前記ソフトウェアプログラムのうち、前記機能ブロックを除く残りのループ文をオフロードするパターンを作成させ、各前記パターンを前記検証環境にデプロイして性能を測定させる手順、測定した性能が所望のものでなかった場合、前記ソフトウェアプログラムのループ文をオフロードするパターンを作成させ、各前記パターンを前記検証環境にデプロイして性能を測定させる手順、を実行させるためのオフロードプログラムとした。

このようにすることで、複数種類のアクセラレータが混在する環境が移行先であっても、自動で高性能化できる。更に移行先が同一ノードに複数種類のアクセラレータを持つ場合は、複数種類のアクセラレータに対して同時にオフロードすることで、単一アクセラレータの利用の場合よりも高速なオフロードを実現する。

１ オフロードサーバ
２ クラウドレイヤ
３ ネットワークレイヤ
４ デバイスレイヤ
１０ ゲートウェイ
２０ ネットワークエッジ
３０ データセンタ
１１ 制御部
１２ 入出力部
１３ 記憶部
１４ 検証用マシン
１５ リソース
１１１ コード指定部
１１２ コード分析部
１１４ 処理指定部
１１４ａ オフロード範囲抽出部
１１４ｂ 中間言語ファイル出力部
１１５ 処理パターン作成部
１１６ 性能測定部
１１６ａ 実行ファイル配置部
１１７ 実行ファイル作成部
１１８ 本番環境配置部
１１９ 性能測定テスト抽出実行部
１２０ ユーザ提供部
１３０ ソースコード
１３１ テストケースデータベース
１３２ 中間言語ファイル
１３３ コードパターンデータベース
１４１，１４２ コードパターン
１５１～１５４ 装置
８０ 通信網
９００ コンピュータ
９１０ ＣＰＵ
９２０ ＲＡＭ
９３０ ＲＯＭ
９４０ ＨＤＤ
９５０ 通信インタフェース
９６０ 入出力インタフェース
９７０ メディアインタフェース
９８０ 記録媒体

Claims (7)

1. ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析するコード分析部と、
   前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを作成し、前記ソフトウェアプログラムのループ文を各前記アクセラレータの何れかにオフロードするパターンを作成する処理パターン作成部と、
   前記機能ブロックまたは前記ループ文を前記アクセラレータにオフロードする各前記パターンを検証環境にデプロイして性能を測定する性能測定部と、
   前記性能測定部が前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを測定した性能が所望のものでなかった場合、前記処理パターン作成部により、前記ソフトウェアプログラムの前記ループ文を各前記アクセラレータの何れかにオフロードするパターンを作成して、前記性能測定部により、各前記パターンを前記検証環境にデプロイして性能を測定する制御部と、
   を備えることを特徴とするオフロードサーバ。
2. 前記アクセラレータは複数種類であり、
   前記制御部は、前記機能ブロックまたは前記ループ文を各前記アクセラレータの何れかにオフロードするパターンのうち、最も高速なものを選択する、
   ことを特徴とする請求項１に記載のオフロードサーバ。
3. 前記制御部は、前記アクセラレータへのオフロード対象として前記機能ブロックの試行の次に前記ループ文を試行し、前記機能ブロックまたは／および前記ループ文のオフロード先の前記アクセラレータとして、メニーコアＣＰＵ、ＧＰＵ、ＦＰＧＡの順に試行する、
   ことを特徴とする請求項１または２に記載のオフロードサーバ。
4. 前記制御部は、前記機能ブロックまたは前記ループ文のオフロードにおいて、元の処理よりも高速なパターンが複数のアクセラレータに対して見つかり、かつ同じ機能ブロックまたはループ文が前記複数のアクセラレータにオフロードできる場合、より性能が高くなるアクセラレータをオフロード先とし、
   異なる機能ブロックまたは異なるループ文が異なるアクセラレータにオフロードでき，かつ同一ノードのアクセラレータにオフロードできない場合、より性能が高くなるアクセラレータをオフロード先とし、
   異なる機能ブロックまたは異なるループ文が異なるアクセラレータにオフロードでき，かつ同一ノードのアクセラレータにオフロードできる場合、前記機能ブロックまたは前記ループ文のオフロード先をそれぞれ異なるアクセラレータとする、
   ことを特徴とする請求項１に記載のオフロードサーバ。
5. 前記制御部は、前記機能ブロックと前記ループ文のオフロードにおいて、前記ループ文のオフロード先が，前記機能ブロックのオフロード先とは別ノードとなる場合，前記ループ文のオフロードによる性能の向上が前記機能ブロックのオフロードによる性能向上以下ならば、前記ループ文をオフロードするパターンを試行せず、かつ前記機能ブロックをオフロードするパターンを試行し、
   前記ループ文のオフロードによる性能の向上が前記機能ブロックのオフロードによる性能向上を超えたならば、前記機能ブロックをオフロードするパターンを試行せず、かつ前記ループ文をオフロードするパターンを試行し、
   前記ループ文のオフロード先が前記機能ブロックのオフロード先と同一ノードとなる場合，前記機能ブロックと前記ループ文とをオフロードしたパターンを試行する、
   ことを特徴とする請求項１に記載のオフロードサーバ。
6. コード分析部が、ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析するステップと、
   処理パターン作成部が、前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを作成するステップと、
   性能測定部が、前記機能ブロックを各前記アクセラレータの何れかにオフロードする各前記パターンを検証環境にデプロイして性能を測定するステップと、を実行し、
   前記性能測定部が前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを測定した性能が所望のものでなかった場合、前記処理パターン作成部が、前記ソフトウェアプログラムのループ文を各前記アクセラレータの何れかにオフロードするパターンを作成するステップと、
   前記性能測定部が、各前記パターンを前記検証環境にデプロイして性能を測定するステップと、を実行する、
   ことを特徴とするオフロード制御方法。
7. コンピュータに、
   ソフトウェアプログラムのソースコードがアクセラレータによる実行が可能な機能ブロックを含むか否かを分析させる手順、
   前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを作成させる手順、
   前記機能ブロックを各前記アクセラレータの何れかにオフロードする各前記パターンを検証環境にデプロイして性能を測定させる手順、
   前記機能ブロックを各前記アクセラレータの何れかにオフロードするパターンを測定した性能が所望のものでなかった場合、前記ソフトウェアプログラムのループ文を各前記アクセラレータの何れかにオフロードするパターンを作成させ、各前記パターンを前記検証環境にデプロイして性能を測定させる手順、
   を実行させるためのオフロードプログラム。
   

Images