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

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

Info

Publication number
JP7716632B2
JP7716632B2 JP2023576454A JP2023576454A JP7716632B2 JP 7716632 B2 JP7716632 B2 JP 7716632B2 JP 2023576454 A JP2023576454 A JP 2023576454A JP 2023576454 A JP2023576454 A JP 2023576454A JP 7716632 B2 JP7716632 B2 JP 7716632B2
Authority
JP
Japan
Prior art keywords
application program
loop
offload
processing
server
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
JP2023576454A
Other languages
English (en)
Other versions
JPWO2023144926A1 (ja
Inventor
庸次 山登
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
NTT Inc
NTT Inc USA
Original Assignee
Nippon Telegraph and Telephone Corp
NTT Inc USA
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Nippon Telegraph and Telephone Corp, NTT Inc USA filed Critical Nippon Telegraph and Telephone Corp
Publication of JPWO2023144926A1 publication Critical patent/JPWO2023144926A1/ja
Application granted granted Critical
Publication of JP7716632B2 publication Critical patent/JP7716632B2/ja
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Classifications

    • G—PHYSICS
    • G06—COMPUTING OR CALCULATING; COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F11/00—Error detection; Error correction; Monitoring
    • G06F11/30—Monitoring
    • G06F11/34—Recording or statistical evaluation of computer activity, e.g. of down time, of input/output operation ; Recording or statistical evaluation of user activity, e.g. usability assessment
    • G—PHYSICS
    • G06—COMPUTING OR CALCULATING; COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G—PHYSICS
    • G06—COMPUTING OR CALCULATING; COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/60—Software deployment
    • G—PHYSICS
    • G06—COMPUTING OR CALCULATING; COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • General Engineering & Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Computer Hardware Design (AREA)
  • Quality & Reliability (AREA)
  • Debugging And Monitoring (AREA)

Description

本発明は、機能処理をＧＰＵ（Graphics Processing Unit）やＦＰＧＡ（Field Programmable Gate Array）等のアクセラレータに自動オフロードし、変換したアプリケーションプログラム（以下適宜、アプリケーションという）を適切な場所に配置するオフロードサーバ、オフロード制御方法およびオフロードプログラムに関する。
ＣＰＵ（Central Processing Unit）以外のヘテロな計算リソースを用いることが増えている。例えば、ＧＰＵ（アクセラレータ）を強化したサーバで画像処理を行ったり、ＦＰＧＡ（アクセラレータ）で信号処理をアクセラレートすることが始まっている。ＦＰＧＡは、製造後に設計者等が構成を設定できるプログラム可能なゲートアレイであり、ＰＬＤ（Programmable Logic Device）の一種である。Amazon Web Services (AWS)（登録商標）では、ＧＰＵインスタンス、ＦＰＧＡインスタンスが提供されており、オンデマンドにそれらリソースを使うこともできる。Microsoft（登録商標）は、ＦＰＧＡを用いて検索を効率化している。
サービス連携技術等を用いて、多彩なアプリケーションの創出が期待されるが、更に進歩したハードウェアを生かすことで、動作アプリケーションの高性能化が期待できる。しかし、そのためには、動作させるハードウェアに合わせたプログラミングや設定が必要である。例えば、ＣＵＤＡ（Compute Unified Device Architecture）、OpenＣＬ（Open Computing Language）といった多くの技術知識が求められ、ハードルは高い。OpenＣＬは、あらゆる計算資源（ＣＰＵやＧＰＵに限らない）を特定のハードに縛られず統一的に扱えるオープンなＡＰＩ（Application Programming Interface）である。
ＧＰＵやＦＰＧＡをユーザのアプリケーションで容易に利用できるようにするため下記が求められる。すなわち、動作させる画像処理、暗号処理等の汎用アプリケーションを環境にデプロイする際に、プラットフォームがアプリケーションロジックを分析し、ＧＰＵ、ＦＰＧＡに自動で処理をオフロードすることが望まれる。
ＧＰＵの計算能力を画像処理以外にも使うＧＰＧＰＵ（General Purpose GPU）のための開発環境ＣＵＤＡが発展している。ＣＵＤＡは、ＧＰＧＰＵ向けの開発環境である。また、ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵ等のヘテロハードウェアを統一的に扱うための標準規格としてOpenＣＬも登場している。
ＣＵＤＡやOpenＣＬでは、Ｃ言語の拡張によるプログラミングを行う。ただし、ＧＰＵ等のデバイスとＣＰＵの間のメモリコピー、解放等を記述する必要があり、記述の難度は高い。実際に、ＣＵＤＡやOpenＣＬを使いこなせる技術者は数多くはいない。
簡易にＧＰＧＰＵを行うため、ディレクティブベースで、ループ文等の並列処理すべき個所を指定し、ディレクティブに従いコンパイラがデバイス向けコードに変換する技術がある。技術仕様としてOpenACC（Open Accelerator）等、コンパイラとしてＰＧＩコンパイラ（登録商標）等がある。例えば、OpenACCを使った例では、ユーザはC/C++/Fortran言語で書かれたコードに、OpenACCディレクティブで並列処理させる等を指定する。ＰＧＩコンパイラは、コードの並列可能性をチェックして、ＧＰＵ用、ＣＰＵ用実行バイナリを生成し、実行モジュール化する。IBM JDK（登録商標）は、Java（登録商標）のlambda形式に従った並列処理指定を、ＧＰＵにオフロードする機能をサポートしている。これらの技術を用いることで、ＧＰＵメモリへのデータ割り当て等を、プログラマは意識する必要がない。
このように、OpenＣＬ、ＣＵＤＡ、OpenACC等の技術により、ＧＰＵやＦＰＧＡへのオフロード処理が可能になっている。
しかし、オフロード処理自体は行えるようになっても、適切なオフロードには課題が多い。例えば、Intelコンパイラ（登録商標）のように自動並列化機能を持つコンパイラがある。自動並列化する際は、プログラム上のfor文（繰り返し文）等の並列処理部を抽出する。ところが、ＧＰＵを用いて並列に動作させる場合は、ＣＰＵ-ＧＰＵメモリ間のデータやり取りによるオーバヘッドのため、性能が出ないことも多い。ＧＰＵを用いて高速化する際は、スキル保持者が、OpenＣＬやＣＵＤＡでのチューニングや、ＰＧＩコンパイラ等で適切な並列処理部を探索することが必要になっている。
このため、スキルが無いユーザがＧＰＵを使ってアプリケーションを高性能化することは難しいし、自動並列化技術を使う場合も、for文を並列するかしないかの試行錯誤チューニング等により、利用開始までに多くの時間がかかっている。
配置に関して、ネットワークリソースの最適利用として、ネットワーク上にあるサーバ 群に対してＶＮ（Virtual Network）の埋め込み位置を最適化する研究がある（非特許文献１参照）。非特許文献１では、通信トラヒックを考慮したＶＮの最適配置を決定する。しかし、単一リソースの仮想ネットワークが対象で、キャリアの設備コストや全体的応答時間の削減が目的で、個々に異なるアプリケーションの処理時間や、個々のユーザのコストや応答時間要求等の条件は考慮されていない。
並列処理箇所の試行錯誤を自動化する取り組みとして、非特許文献２が挙げられる。
非特許文献２は、一度記述したコードで、配置先の環境に存在するＧＰＵやＦＰＧＡ、メニーコアＣＰＵ等を利用できるように、変換、リソース設定等を自動で行い、アプリケーションを高性能で動作させることを目的とした、環境適応ソフトウェアを提案している。併せて、非特許文献２は、環境適応ソフトウェアの要素として、アプリケーションコードのループ文を、ＧＰＵに自動オフロードする方式を提案し性能向上を評価している。
非特許文献３は、環境適応ソフトウェアの要素として、アプリケーションコードのループ文を、ＦＰＧＡに自動オフロードする方式を提案し性能向上を評価している。
非特許文献４は、環境適応ソフトウェアの要素として、ＧＰＵ等向けに自動変換した後、アプリケーションを実行するリソース量（仮想マシンコアの数など）を、適正化する手法を評価している。
河島滉太，大歳達也，大下裕一，村田正幸, "分散型モデル予測制御にもとづくスケーラビリティを有する仮想ネットワーク埋め込み手法," 電子情報通信学会技術報告, vol.115, pp.19-24, 2015. Y. Yamato, "Study of parallel processing area extraction and data transfer number reduction for automatic GPU offloading of IoT applications," Journal of Intelligent Information Systems, Springer, DOI: 10.1007/s10844-019-00575-8, Aug. 2019. (Electronic Publishing) Vol.54, No.3, pp.567-584, May 2020. (Print Publishing) Y. Yamato, "Automatic Offloading Method of Loop Statements of Software to FPGA," International Journal of Parallel, Emergent and Distributed Systems, Taylor & Francis, DOI: 10.1080/17445760.2021.1916020, Apr. 2021. Y. Yamato, "Arranging resource amount after automatic GPU offloading," The Ninth International Symposium on Computing and Networking (CANDAR 2021), Nov. 2021.
非特許文献１～４では、自動オフロード時の処理時間の短縮を中心に評価している。
ＧＰＵ、ＦＰＧＡ等のヘテロジニアスなデバイスに処理をオフロードする際に、変換したアプリケーションをユーザ要望（価格、応答時間）を満たして動作させることについては提案されていないという課題がある。
このような点に鑑みて本発明がなされたのであり、ＧＰＵやＦＰＧＡ等のオフロードデバイスに配置できるよう自動変換した際に、変換したアプリケーションをユーザのコストまたは応答時間の要求を満たして最適に配置することを課題とする。
前記した課題を解決するため、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部と、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部と、を備えることを特徴とするオフロードサーバとした。
本発明によれば、ＧＰＵやＦＰＧＡ等のオフロードデバイスに配置できるよう自動変換した際に、変換したアプリケーションをユーザのコストまたは応答時間の要求を満たして最適に配置することができる。
本発明の第１実施形態に係るオフロードサーバの構成例を示す機能ブロック図である。 第１実施形態に係るオフロードサーバを用いた自動オフロード処理を示す図である。 第１実施形態に係るオフロードサーバのSimple GAによる制御部（自動オフロード機能部）の探索イメージを示す図である。 比較例の通常ＣＰＵプログラムの例を示す図である。 比較例の単純ＣＰＵプログラムを利用してＣＰＵからＧＰＵへデータ転送する場合のループ文の例を示す図である。 第１実施形態に係るオフロードサーバのネスト一体化をした場合のＣＰＵからＧＰＵへデータ転送する場合のループ文の例を示す図である。 第１実施形態に係るオフロードサーバの転送一体化をした場合のＣＰＵからＧＰＵへデータ転送する場合のループ文の例を示す図である。 第１実施形態に係るオフロードサーバの転送一体化をし、かつ一時領域を利用した場合のＣＰＵからＧＰＵへデータ転送する場合のループ文の例を示す図である。 第１実施形態に係るオフロードサーバの実装の動作概要を説明するフローチャートである。 第１実施形態に係るオフロードサーバの実装の動作概要を説明するフローチャートである。 第１実施形態に係るオフロードサーバのＧＰＵオフロード試行の後に追加されるリソース比とリソース量の設定および新規アプリケーションの配置を説明するフローチャートである。 第１実施形態に係るオフロードサーバの計算ノードのトポロジーの一例を示す図である。 第１実施形態に係るオフロードサーバの平均応答時間のアプリケーション配置数変化を示すグラフである。 本発明の第２実施形態に係るオフロードサーバの構成例を示す機能ブロック図である。 第２実施形態に係るオフロードサーバの実装の動作概要を説明するフローチャートである。 上記第２実施形態に係るオフロードサーバの性能測定部の性能測定処理を示すフローチャートである。 第２実施形態に係るオフロードサーバのＰＬＤ処理パターン作成部の探索イメージを示す図である。 第２実施形態に係るオフロードサーバのＣコードからOpenCL最終解の探索までの流れを説明する図である。 本発明の各実施形態に係るオフロードサーバの機能を実現するコンピュータの一例を示すハードウェア構成図である。
以下、図面を参照して本発明を実施するための形態（以下、「本実施形態」という）におけるオフロードサーバについて説明する。
（本発明の自動オフロードの基本的な考え方）
本発明者は、環境適応ソフトウェアのコンセプトを具体化するために、これまでに、プログラムのループ文のＧＰＵ自動オフロード、ＦＰＧＡ自動オフロード、変換アプリケーションの実行リソース適正化の方式を提案してきた（非特許文献２、３、４参照）。これら非特許文献２、３、４の要素技術の検討も踏まえて、本発明の基本的な考え方を述べる。
<ＣＰＵとオフロードデバイスのリソース比の適切化>
まず、デバイスにオフロードするプログラム変換ができた後の、ＣＰＵとオフロードデバイスのリソース比の適切化について説明する。
非特許文献２等の手法により、ＧＰＵやＦＰＧＡ等のオフロードデバイスに通常のプログラムを自動オフロードすることができる。
現在、マルチコアＣＰＵ、メニーコアＣＰＵは、仮想マシンやコンテナによる仮想化により、全コアの何割を割り当てる等が柔軟にできるようになっている。ＧＰＵについても、近年ＣＰＵ同様の仮想化が行われ、ＧＰＵの全コアの何割を割り当てる等の運用が可能になりつつある。ＦＰＧＡに関しては、リソース使用量は、Look Up TableやFlip Flopの設定数で表されることが多く、利用されていないゲートについては別用途に使うことができる。
このように、ＣＰＵ、ＧＰＵ、ＦＰＧＡとも全リソースの一部を使う運用が可能であり、ＣＰＵとオフロードデバイスのリソースを用途に応じて適切化することはコストパフォーマンスを高める上で重要である。
また、非特許文献２等の手法を用いて、アプリケーションをＣＰＵとＧＰＵ処理のコードに変換することはできる。しかし、コード自体は、適切であっても、ＣＰＵとＧＰＵとのリソース量が適切なバランスでない場合には、性能が出ない。例えば、ある処理を行う際に、ＣＰＵの処理時間が１０００秒、ＧＰＵの処理時間が１秒では、オフロードできる処理をＧＰＵである程度高速化しても、全体的にはＣＰＵがボトルネックとなっている。
さらに、非特許文献５の「K. Shirahata, H. Sato and S. Matsuoka, "Hybrid Map Task Scheduling for GPU-Based Heterogeneous Clusters,"IEEE Second International Conference on Cloud Computing Technology and Science (CloudCom), pp.733-740, Dec. 2010.」では、ＣＰＵとＧＰＵを使ってMapReduce（登録商標）フレームワークでタスク処理している際に、ＣＰＵとＧＰＵの実行時間が同じになるようMapタスクを配分することで、全体の高性能化を図っている。
本発明者は、ＣＰＵとオフロードデバイスのリソース比を下記のように決めることを想到した。すなわち、何れかのデバイスでの処理がボトルネックとなることを避けるため、上記非特許文献等も参考に、テストケースの処理時間から、ＣＰＵとオフロードデバイスの処理時間が同等オーダになるように、ＣＰＵとオフロードデバイスのリソース比（以下、「リソース比」という）を決定する。
また、本発明者は、非特許文献２の手法のように、自動オフロードの際、検証環境での性能測定結果に基づいて徐々に高速化していく手法を採る。理由としては、性能に関しては、コード構造だけでなく、実際に処理するハードウェアのスペック、データサイズ、ループ回数等の実際に処理する内容によって大きく変わるためである。また、性能は、静的に予測することが困難であり、動的な測定が必要だからである。そのため、コード変換の際に、既に検証環境での性能測定結果があるので、その結果を用いてリソース比を定める。
性能測定の際には、テストケースを指定して測定を行う。例えば、検証環境でのテストケースの処理時間が、ＣＰＵ処理：１０秒、ＧＰＵ処理：５秒の場合では、ＣＰＵ側のリソースは２倍で同等の処理時間程度と考えられる。このため、リソース比は２：１となる。なお、特にある処理をオフロードで高速化したいといったユーザ要望については、その処理を含むテストケースを準備して、そのテストケースに対して非特許文献２等の手法で高速化することでユーザ要望が反映される。
<ＣＰＵとオフロードデバイスのリソース量の決定と自動検証>
次に、ＣＰＵとオフロードデバイスのリソース量（以下、「リソース量」という）の決定と自動検証について説明する。
上記<ＣＰＵとオフロードデバイスのリソース比の適切化>により、リソース比が定まった場合、次に商用環境へのアプリケーションの配置を行う。
商用環境への配置の際は、ユーザが指定したコスト要求を満たすように、リソース比は可能な限りキープ（維持）したまま、リソース量を決定する。例えば、ＣＰＵに関して、１ＶＭは１０００円／月、ＧＰＵは４０００円／月、リソース比は２：１が適切であるとする。そして、ユーザの予算は、月１００００円以内であると想定する。この場合には、リソース比を２：１としても、ユーザの予算内である月１００００円以内に収まるので、適切なリソース比２：１をキープしたリソース量、すなわちＣＰＵは「２」、ＧＰＵは「１」を確保して商用環境に配置することになる。また、ユーザの予算が、月５０００円以内であった場合には、適切なリソース比２：１はキープできない。この場合、リソース量として、ＣＰＵは「１」、ＧＰＵは「１」を確保して配置する。
商用環境にリソースを確保してプログラムを配置した後は、ユーザが利用する前に動作することを確認するため、自動検証が行われる。自動検証では、性能検証テストケースやリグレッションテストケースが実行される。性能検証テストケースは、ユーザが指定した想定テストケースをJenkins（登録商標）等の試験自動実行ツールを用いて行い、処理時間やスループット等を測定する。リグレッションテストケースは、システムにインストールされるミドルウェアやＯＳ等のソフトウェアの情報を取得して、それらに対応するリグレッションテストをJenkins等を用いて実行する。これらの自動検証を、少ないテストケースの準備で行うための検討は非特許文献６（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, Feb. 2015.）等でなされており、この非特許文献６の技術を用いる。
性能検証テストケースでは、オフロードした場合でも計算結果が不正でないかをチェックする。また、性能検証テストケースでは、オフロードしない場合との計算結果差分もチェックする。例えば、ＧＰＵを処理するＰＧＩコンパイラは、PCAST（登録商標）という機能のＰＧＩ_compare（登録商標）やacc_compare（登録商標）というＡＰＩ（Application Programming Interface）で、ＧＰＵを使う場合使わない場合の計算結果差分を確認できる。
なお、ＧＰＵとＣＰＵでは丸め誤差が異なる等、並列処理等を正しくオフロードしても完全に計算結果が一致しない場合もある。そのため、例えばIEEE 754仕様による確認等を行い、許容できる差分かをユーザに提示し、ユーザに確認をしてもらう。
自動検証の結果として、性能検証テストケースの処理時間やそのスループット、計算結果差分およびリグレッションテストの実行結果の情報が、ユーザに提示される。ユーザには、さらに確保したリソース（ＶＭの数やスペック等）とその価格が提示されており、ユーザはそれら情報を参照して運用開始を判断する。
<リソース、リソース比、テストケース処理時間>
本実施形態におけるリソース、リソース比、テストケース処理時間について述べる。
・リソースについて
ＣＰＵ、ＧＰＵ、ＦＰＧＡ等は仮想資源のインスタンスとして提供されるようになってきている。
リソースとして、ＣＰＵのコア数、クロック、メモリ量、ディスクサイズ、ＧＰＵのコア数、クロック、メモリ量、ＦＰＧＡのゲート規模（Intel（登録商標）の場合はLE（登録商標）、Xilinx（登録商標）の場合LC（登録商標）が単位となる）がある。クラウド等の事業者は、それらをパッケージ化して、small sizeの仮想マシンやＧＰＵインスタンスといった形で提供している。仮想化する場合は、利用するインスタンスの数が利用するリソース量といえる。
・リソース比について
ＣＰＵ、ＧＰＵ、ＦＰＧＡのインスタンス数の比がリソース比となる。インスタンス数が１つ、２つ、３つであれば、リソース比は１：２：３である。
・テストケース処理時間について
本実施形態は、ユーザが指定するテストケースを高速化するオフロードパターンを探索して発見する。テストケースは、ＤＢ（データベース）であればTPC-C（登録商標）のようなトランザクション処理数であり、ＦＦＴであればサンプルデータでのフーリエ変換処理の実行である。処理時間は、そのサンプル処理を実行した際の実行時間である。例えば、処理Ａの処理時間は、オフロード前は１０秒であったものが、オフロード後は２秒になるといった形で、ＣＰＵで実行した場合と、オフロードデバイスで実行した場合との実行時間がそれぞれ取得される。
<ループ文の発見>
コンパイラが、このループ文はＧＰＵの並列処理に適しているという適合性を見つけることは難しいのが現状である。ＧＰＵにオフロードすることでどの程度の性能、電力消費量になるかは、実測してみないと予測は難しい。そのため、このループ文をＧＰＵにオフロードするという指示を手動で行い、測定の試行錯誤が行われている。
本発明は、ＧＰＵにオフロードする適切なループ文の発見を、進化計算手法である遺伝的アルゴリズム（ＧＡ：Genetic Algorithm）を用いて自動的に行う。すなわち、並列可能ループ文群に対して、ＧＰＵ実行の際を１、ＣＰＵ実行の際を０に値を置いて遺伝子化し、検証環境で反復測定し適切なパターンを探索する。
（第１の実施形態）
次に、本発明を実施するための形態（以下、「本実施形態」と称する。）における、オフロードサーバ１等について説明する。
［ループ文のＧＰＵ自動オフロード］
図１は、本発明の第１の実施形態に係るオフロードサーバ１の構成例を示す機能ブロック図である。
オフロードサーバ１は、アプリケーションの特定処理をアクセラレータに自動的にオフロードする装置である。
図１に示すように、オフロードサーバ１は、制御部１１と、入出力部１２と、記憶部１３と、検証用マシン１４（Verification machine）(アクセラレータ検証用装置)と、を含んで構成される。
入出力部１２は、各デバイス等との間で情報の送受信を行うための通信インターフェイスと、タッチパネルやキーボード等の入力装置や、モニタ等の出力装置との間で情報の送受信を行うための入出力インターフェイスとから構成される。
記憶部１３は、ハードディスクやフラッシュメモリ、ＲＡＭ（Random Access Memory）等により構成され、制御部１１の各機能を実行させるためのプログラム（オフロードプログラム）や、制御部１１の処理に必要な情報（例えば、中間言語ファイル(Intermediate file)１３３）が一時的に記憶される。
記憶部１３は、テストケースＤＢ（Test case database）１３１、設備リソースＤＢ１３２、中間言語ファイル(Intermediate file)１３３を備える。
テストケースＤＢ１３１は、検証対象ソフトに対応した試験項目のデータを格納する。試験項目のデータは、例えばMySQL等のデータベースシステムの場合、TPC-C等のトランザクション試験のデータである。
設備リソースＤＢ１３２は、事業者が保持するサーバ等のリソースと価格等の事前に準備された情報と、それらがどの程度使われているかの情報を保持する。例えばＧＰＵインスタンスを３収容できるサーバが１０台あり、１ＧＰＵインスタンスは月５０００円であり、１０台のうち、Ａ，Ｂの２台はフルに使われており、Ｃの１台はインスタンスが１つだけ使われている等の情報である。この情報は、ユーザが運用条件（コスト、性能等の条件）を指定した際に、確保するリソース量を決定するために利用される。ユーザ運用条件は、ユーザがオフロード依頼時に指定したコスト条件（例えば、月１００００円以内の予算等）と性能条件（例えば、ＴＰＣ－Ｃ等のトランザクションスループットが何以上やサンプルフーリエ変換処理が1スレッドで何秒以内等）である。
中間言語ファイル１３３は、高水準言語と機械語の中間に介在するプログラミング言語の形で制御部１１の処理に必要な情報を一時的に記憶する。
検証用マシン１４は、環境適応ソフトウェアの検証用環境として、ＣＰＵ、ＧＰＵ、ＦＰＧＡを備える。
制御部１１は、オフロードサーバ１全体の制御を司る自動オフロード機能部（Automatic Offloading function）である。制御部１１は、例えば、記憶部１３に格納されたアプリケーションプログラム（オフロードプログラム）を不図示のＣＰＵ（Central Processing Unit）が、ＲＡＭに展開し実行することにより実現される。
制御部１１は、アプリケーションコード指定部（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）１２２と、を備える。
<アプリケーションコード指定部１１１>
アプリケーションコード指定部１１１は、入力されたアプリケーションコードの指定を行う。具体的には、アプリケーションコード指定部１１１は、受信したファイルに記載されたアプリケーションコードを、アプリケーションコード分析部１１２に渡す。
<アプリケーションコード分析部１１２>
アプリケーションコード分析部１１２は、処理機能のソースコードを分析し、ループ文やＦＦＴライブラリ呼び出し等の構造を把握する。
<データ転送指定部１１３>
データ転送指定部１１３は、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行（後記する#pragma acc kernels、#pragma acc data copyin(a，ｂ)、#pragma acc data copyout(a，ｂ)、#prama acc parallel loop、#prama acc parallel loop vectorなど）を用いたデータ転送指定を行う。
<並列処理指定部１１４>
並列処理指定部１１４は、アプリケーションプログラムのループ文（繰り返し文）を特定し、各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする。
並列処理指定部１１４は、オフロード範囲抽出部（Extract offload able area）１１４ａと、中間言語ファイル出力部（Output intermediate file）１１４ｂと、を備える。
オフロード範囲抽出部１１４ａは、ループ文やＦＦＴ等、ＧＰＵ・ＦＰＧＡにオフロード可能な処理を特定し、オフロード処理に応じた中間言語を抽出する。
中間言語ファイル出力部１１４ｂは、抽出した中間言語ファイル１３３を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。
<リソース比決定部１１５>
リソース比決定部１１５は、性能測定結果をもとに、ＣＰＵとオフロードデバイスの処理時間（テストケースＣＰＵ処理時間とオフロードデバイス処理時間）を、リソース比として決定する（後記）。具体的には、リソース比決定部１１５は、ＣＰＵとオフロードデバイスの処理時間が同等オーダになるように、リソース比を決定する。また、リソース比決定部１１５は、ＣＰＵとオフロードデバイスの処理時間の差分が所定閾値以上の場合、リソース比を所定の上限値に設定する。
<リソース量設定部１１６>
リソース量設定部１１６は、決定したリソース比をもとに、所定のコスト条件を満たすように、ＣＰＵおよびオフロードデバイスのリソース量を設定する（後記）。具体的には、リソース量設定部１１６は、決定したリソース比を維持して、所定のコスト条件を満たす最大のリソース量を設定する。また、リソース量設定部１１６は、決定したリソース比を維持した最小リソース量の設定で所定のコスト条件を満たさない場合は、リソース比を崩してＣＰＵとオフロードデバイスのリソース量をコスト条件を満たすより小さい値（例えば、最小）で設定する。
<配置設定部１７０>
配置設定部１７０は、変換したアプリケーションを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションの配置場所を計算して設定する。具体的には、配置設定部１７０は、設備リソースＤＢ１３２のサーバ、リンクのスペック情報、既存アプリケーションの配置情報に基づいて、線形計画手法で、新規アプリケーションの配置先（ＡＰＬの配置場所）を計算して設定する。線形計画手法では、例えば、後記［数１］［数２］に示す線形計画式の目的関数および制約条件を用いる。後記［数１］［数２］に示す線形計画式は、設備リソースＤＢ１３２に保存されており、配置設定部１７０が、設備リソースＤＢ１３２から読み出し、配置設定部１７０が処理するメモリ上で展開される。
<並列処理パターン作成部１１７>
並列処理パターン作成部１１７は、コンパイルエラーが出るループ文（繰り返し文）に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する。
<性能測定部１１８>
性能測定部１１８は、並列処理パターンのアプリケーションプログラムをコンパイルして、検証用マシン１４に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する。
性能測定部１１８は、バイナリファイル配置部（Deploy binary files）１１８ａを備える。バイナリファイル配置部１１８ａは、ＧＰＵやＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイ(配置)する。
性能測定部１１８は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、オフロード範囲抽出部１１４ａに戻す。この場合、オフロード範囲抽出部１１４ａは、別の並列処理パターン抽出を行い、中間言語ファイル出力部１１４ｂは、抽出された中間言語をもとに、性能測定を試行する（後記図２の符号ａ参照）。
<実行ファイル作成部１１９>
実行ファイル作成部１１９は、所定回数繰り返された、性能測定結果をもとに、複数の並列処理パターンから高処理性能の並列処理パターンを複数選択し、高処理性能の並列処理パターンを交叉、突然変異処理により別の複数の並列処理パターンを作成する。そして、実行ファイル作成部１１９は、新たに性能測定までを行い、指定回数の性能測定後に、性能測定結果をもとに、複数の並列処理パターンから最高処理性能の並列処理パターンを選択し、最高処理性能の並列処理パターンをコンパイルして実行ファイルを作成する。
<本番環境配置部１２０>
本番環境配置部１２０は、作成した実行ファイルを、ユーザ向けの本番環境に配置する（「最終バイナリファイルの本番環境への配置」）。本番環境配置部１２０は、最終的なオフロード領域を指定したパターンを決定し、ユーザ向けの本番環境にデプロイする。
<性能測定テスト抽出実行部１２１>
性能測定テスト抽出実行部１２１は、実行ファイル配置後、テストケースＤＢ１３１から性能試験項目を抽出し、性能試験を実行する（「最終バイナリファイルの本番環境への配置」）。
性能測定テスト抽出実行部１２１は、実行ファイル配置後、ユーザに性能を示すため、性能試験項目をテストケースＤＢ１３１から抽出し、抽出した性能試験を自動実行する。
<ユーザ提供部１２２>
ユーザ提供部１２２は、性能試験結果を踏まえた、価格・性能等の情報をユーザに提示する（「価格・性能等の情報のユーザへの提供」）。テストケースＤＢ１３１には、性能試験項目が格納されている。ユーザ提供部１２２は、テストケースＤＢ１３１に格納された試験項目に対応した性能試験の実施結果に基づいて、価格、性能等のデータを、上記性能試験結果と共にユーザに提示する。ユーザは、提示された価格・性能等の情報をもとに、サービスの課金利用開始を判断する。ここで、本番環境への一括デプロイには、非特許文献7（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.１２０6/s13677-014-0007-3, 12 pages, June 2014.）の技術を、また、性能自動試験には、前述の非特許文献６の技術を用いればよい。
［遺伝的アルゴリズムの適用］
オフロードサーバ１は、オフロードの最適化にＧＡ（Genetic Algorithms）を用いることができる。ＧＡを用いた場合のオフロードサーバ１の構成は下記の通りである。
すなわち、並列処理指定部１１４は、遺伝的アルゴリズムに基づき、コンパイルエラーが出ないループ文（繰り返し文）の数を遺伝子長とする。並列処理パターン作成部１１７は、アクセラレータ処理をする場合を１または０のいずれか一方、しない場合を他方の０または１として、アクセラレータ処理可否を遺伝子パターンにマッピングする。
並列処理パターン作成部１１７は、遺伝子の各値を１か０にランダムに作成した指定個体数の遺伝子パターンを準備する。性能測定部１１８は、各個体に応じて、アクセラレータにおける並列処理指定文を指定したアプリケーションコードをコンパイルして、検証用マシン１４に配置する。性能測定部１１８は、検証用マシン１４において性能測定用処理を実行する。
ここで、性能測定部１１８は、途中世代で、以前と同じ並列処理パターンの遺伝子が生じた場合は、当該並列処理パターンに該当するアプリケーションコードのコンパイル、および、性能測定はせずに、性能測定値としては同じ値を使う。
また、性能測定部１１８は、コンパイルエラーが生じるアプリケーションコード、および、性能測定が所定時間で終了しないアプリケーションコードについては、タイムアウトの扱いとして、性能測定値を所定の時間（長時間）に設定する。
実行ファイル作成部１１９は、全個体に対して、性能測定を行い、処理時間の短い個体ほど適合度が高くなるように評価する。実行ファイル作成部１１９は、全個体から、適合度が所定値（例えば、全個数の上位ｎ％、または全個数の上位ｍ個 ｎ，ｍは自然数）より高いものを性能の高い個体として選択し、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。実行ファイル作成部１１９は、指定世代数の処理終了後、最高性能の並列処理パターンを解として選択する。
以下、上述のように構成されたオフロードサーバ１の自動オフロード動作について説明する。
［自動オフロード動作］
図２は、オフロードサーバ１を用いた自動オフロード処理を示す図である。
図２に示すように、オフロードサーバ１は、環境適応ソフトウェアの要素技術に適用される。オフロードサーバ１は、制御部（自動オフロード機能部）１１と、テストケースＤＢ１３１と、設備リソースＤＢ１３２と、中間言語ファイル１３３と、検証用マシン１４と、を有している。
オフロードサーバ１は、ユーザが利用するアプリケーションコード（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文が２００個あっても、ＧＰＵにオフロードできるものは３０個くらいである。ここでは、エラーになるものを除外し、この３０個について、ＧＡを行う。
OpenＡＣＣには、ディレクティブ #pragma acc kernelsで指定して、ＧＰＵ向けバイトコードを抽出し、実行によりＧＰＵオフロードを可能とするコンパイラがある。この#pragmaに、for文のコマンドを書くことにより、そのfor文がＧＰＵで動くか否かを判定することができる。
例えばC/C++を使った場合、C/C++のコードを分析し、for文を見付ける。for文を見付けると、OpenＡＣＣで並列処理の文法である#pragma acc kernels、#prama acc parallel loopや#prama acc parallel loop vectorを使ってfor文に対して書き込む。詳細には、#pragma acc kernels、#prama acc parallel loopや#prama acc parallel loop vectorに、一つ一つfor文を入れてコンパイルして、エラーであれば、そのfor文はそもそも、ＧＰＵ処理できないので、除外する。
このようにして、残るfor文を見付ける。そして、エラーが出ないものを、長さ（遺伝子長）とする。エラーのないfor文が５つであれば、遺伝子長は５であり、エラーのないfor文が１０であれば、遺伝子長は１０である。なお、並列処理できないものは、前の処理を次の処理に使うようなデータに依存がある場合である。
以上が準備段階である。次にＧＡ処理を行う。
for文の数に対応する遺伝子長を有するコードパターンが得られている。始めはランダムに並列処理パターン１００１０、０１００１、００１０１、…を割り当てる。ＧＡ処理を行い、コンパイルする。その時に、オフロードできるfor文であるにもかかわらず、エラーがでることがある。それは、for文が階層になっている（どちらか指定すればＧＰＵ処理できる）場合である。この場合は、エラーとなったfor文は、残してもよい。具体的には、処理時間が多くなった形にして、タイムアウトさせる方法がある。
検証用マシン１４でデプロイして、ベンチマーク、例えば画像処理であればその画像処理でベンチマークする、その処理時間が短い程、適応度が高いと評価する。例えば、処理時間の-1/2乗で、処理時間１秒かかるものは１、１００秒かかるものは０．１、０．０１秒かかるものは１０とする。
適応度が高いものを選択して、例えば１０個のなかから、３～５個を選択して、それを組み替えて新しいコードパターンを作る。このとき、作成途中で、前と同じものができる場合がある。その場合、同じベンチマークを行う必要はないので、前と同じデータを使う。本実施形態では、コードパターンと、その処理時間は記憶部１３に保存しておく。
以上で、Simple GAによる制御部（自動オフロード機能部）１１の探索イメージについて説明した。次に、データ転送の一括処理手法について述べる。
［データ転送の一括処理手法］
<基本的な考え方>
ＣＰＵ-ＧＰＵ転送の削減のため、ネストループの変数をできるだけ上位で転送することに加え、本発明は、多数の変数転送タイミングを一括化し、さらにコンパイラが自動転送してしまう転送を削減する。
転送の削減にあたり、ネスト単位だけでなく、ＧＰＵに転送するタイミングがまとめられる変数については一括化して転送する。例えば、ＧＰＵの処理結果をＣＰＵで加工してＧＰＵで再度処理させるなどの変数でなければ、複数のループ文で使われるＣＰＵで定義された変数を、ＧＰＵ処理が始まる前に一括してＧＰＵに送り、全ＧＰＵ処理が終わってからＣＰＵに戻すなどの対応も可能である。
コード分析時にループおよび変数の参照関係を把握するため、その結果から複数ファイルで定義された変数について、ＧＰＵ処理とＣＰＵ処理が入れ子にならず、ＣＰＵ処理とＧＰＵ処理が分けられる変数については、一括化して転送する指定をOpenACCのdata copy文を用いて指定する。
ＧＰＵ処理の始まる前に一括化して転送され、ループ文処理のタイミングで転送が不要な変数はdata presentを用いて転送不要であることを明示する。
ＣＰＵ-ＧＰＵのデータ転送時は、一時領域を作成し（#pragma acc declare create）、データは一時領域に格納後、一時領域を同期（#pragma acc update）することで転送を指示する。
<比較例>
まず、比較例について述べる。
比較例は、通常ＣＰＵプログラム（図４参照）、単純ＧＰＵ利用（図５参照）、ネスト一括化（非特許文献２）（図６参照）である。なお、以下の記載および図中のループ文の文頭の<1>～<4>等は、説明の便宜上で付したものである（他図およびその説明においても同様）。
図４に示す通常ＣＰＵプログラムのループ文は、ＣＰＵプログラム側で記述され、
<1> ループ〔for(i=0; i<10; i++)〕{
}
の中に、
<2> ループ〔for(j=0; j<20; j++〕 {
がある。図４の符号ｆは、上記 <2>ループにおける、変数ａ，ｂの設定である。
また、
<3> ループ〔for(k=0; k<30; k++)〕{
}
と、
<4> ループ〔for(l=0; l<40; l++)〕{
}
と、が続く。図４の符号ｇは、上記<3>ループにおける変数ｃ，ｄの設定であり、図４の符号ｈは、上記<4>ループにおける変数ｅ，ｆの設定である。
図４に示す通常ＣＰＵプログラムは、ＣＰＵで実行される（ＧＰＵ利用しない）。
図５は、図４に示す通常ＣＰＵプログラムを、単純ＧＰＵ利用して、ＣＰＵからＧＰＵへのデータ転送する場合のループ文を示す図である。データ転送の種類は、ＣＰＵからＧＰＵへのデータ転送、および、ＧＰＵからＣＰＵへのデータ転送がある。以下、ＣＰＵからＧＰＵへのデータ転送を例にとる。
図５に示す単純ＧＰＵ利用のループ文は、ＣＰＵプログラム側で記述され、
<1> ループ〔for(i=0; i<10; i++)〕{
}
の中に、
<2> ループ〔for(j=0; j<20; j++〕 {
がある。
さらに、図５の符号ｉに示すように、 <1> ループ〔for(i=0; i<10; i++)〕{
}の上部に、ＰＧＩコンパイラによるfor文等の並列処理可能処理部を、OpenＡＣＣのディレクティブ #pragma acc kernels（並列処理指定文）で指定している。
図５の符号ｉを含む破線枠囲みに示すように、#pragma acc kernelsによって、ＣＰＵからＧＰＵへデータ転送される。ここでは、このタイミングでａ，ｂが転送されるため１０回転送される。
また、図５の符号ｊに示すように、 <3> ループ〔for(k=0; k<30; k++)〕{
}の上部に、ＰＧＩコンパイラによるfor文等の並列処理可能処理部を、OpenＡＣＣのディレクティブ #pragma acc kernelsで指定している。
図５の符号ｊを含む破線枠囲みに示すように、#pragma acc kernelsによって、このタイミングでｃ，ｄが転送される。
ここで、 <4> ループ〔for(l=0; l<40; l++)〕{
}の上部には、#pragma acc kernelsを指定しない。このループは、ＧＰＵ処理しても効率が悪いのでＧＰＵ処理しない。
図６は、ネスト一括化（非特許文献２）による、ＣＰＵからＧＰＵおよびＧＰＵからＣＰＵへのデータ転送する場合のループ文を示す図である。
図６に示すループ文では、図６の符号ｋに示す位置に、ＣＰＵからＧＰＵへのデータ転送指示行、ここでは変数ａ，ｂの copyin 節の #pragma acc data copyin(a，ｂ)を挿入する。なお、本明細書では表記の関係でcopyin(a，ｂ)について、括弧()を付している。後記copyout(a，ｂ)、datacopyin(a，ｂ，ｃ，ｄ)についても同様の表記方法を採る。
上記 #pragma acc data copyin(a，ｂ)は、変数ａの設定、定義を含まない最上位のループ（ここでは、 <1> ループ〔for(i=0; i<10; i++)〕{
}の上部）に指定される。
図６の符号ｋを含む一点鎖線枠囲みに示すタイミングでａ，ｂが転送されるため１回転送が発生する。
また、図６に示すループ文では、図６の符号ｌに示す位置に、ＧＰＵからＣＰＵへのデータ転送指示行、ここでは変数ａ，ｂの copyout 節の #pragma acc data copyout(a，ｂ)を挿入する。
上記 #pragma acc data copyout(a，ｂ)は、 <1> ループ〔for(i=0; i<10; i++)〕{
}の下部に指定される。
このように、ＣＰＵからＧＰＵへのデータ転送において、変数ａの copyin 節の #pragma acc data copyin(a，ｂ)を、上述した位置に挿入することによりデータ転送を明示的に指示する。これにより、できるだけ上位のループでデータ転送を一括して行うことができ、図５に示す単純ＧＰＵ利用のループ文のようにループ毎に毎回データを転送する非効率な転送を避けることができる。
<実施形態>
次に、本実施形態について述べる。
《転送不要な変数をdata presentを用いて明示》
本実施形態では、複数ファイルで定義された変数について、ＧＰＵ処理とＣＰＵ処理が入れ子にならず、ＣＰＵ処理とＧＰＵ処理が分けられる変数については、一括化して転送する指定をOpenACCのdata copy文を用いて指定する。併せて、一括化して転送され、そのタイミングで転送が不要な変数はdata presentを用いて明示する。
図７は、本実施形態のＣＰＵ-ＧＰＵのデータ転送時の転送一括化によるループ文を示す図である。図７は、比較例の図６のネスト一括化に対応する。
図７に示すループ文では、図７の符号ｍに示す位置に、ＣＰＵからＧＰＵへのデータ転送指示行、ここでは変数ａ，ｂ，ｃ，ｄの copyin 節の #pragma acc datacopyin(a，ｂ，ｃ，ｄ)を挿入する。
上記 #pragma acc data copyin(a，ｂ，ｃ，ｄ)は、変数ａの設定、定義を含まない最上位のループ（ここでは、 <1> ループ〔for(i=0; i<10; i++)〕{
}の上部）に指定される。
このように、複数ファイルで定義された変数について、ＧＰＵ処理とＣＰＵ処理が入れ子にならず、ＣＰＵ処理とＧＰＵ処理が分けられる変数については、一括化して転送する指定をOpenACCのdata copy文#pragma acc data copyin(a，ｂ，ｃ，ｄ)を用いて指定する。
図７の符号ｍを含む一点鎖線枠囲みに示すタイミングでａ，ｂ，ｃ，ｄが転送されるため１回転送が発生する。
そして、上記#pragma acc data copyin(a，ｂ，ｃ，ｄ)を用いて一括化して転送され、そのタイミングで転送が不要な変数は、図７の符号ｎを含む二点鎖線枠囲みに示すタイミングで既にＧＰＵに変数があることを明示するdata present文#pragma acc data present (a，ｂ)を用いて指定する。
上記#pragma acc data copyin(a，ｂ，ｃ，ｄ)を用いて一括化して転送され、そのタイミングで転送が不要な変数は、図７の符号ｏを含む二点鎖線枠囲みに示すタイミングで既にＧＰＵに変数があることを明示するdata present文#pragma acc data present(ｃ，ｄ)を用いて指定する。
<1>、<3>のループがＧＰＵ処理されＧＰＵ処理が終了したタイミングで、ＧＰＵからＣＰＵへのデータ転送指示行、ここでは変数ａ，ｂ，ｃ，ｄの copyout 節の #pragma acc datacopyout(a，ｂ, c, d)を、図７の<3>ループが終了した位置ｐに挿入する。
一括化して転送する指定により一括化して転送できる変数は一括転送し、既に転送され転送が不要な変数はdata presentを用いて明示することで、転送を削減して、オフロード手段のさらなる効率化を図ることができる。しかし、OpenACCで転送を指示してもコンパイラによっては、コンパイラが自動判断して転送してしまう場合がある。コンパイラによる自動転送とは、OpenACCの指示と異なり、本来はＣＰＵ-ＧＰＵ間の転送が不要であるにもかかわらずコンパイラ依存で自動転送されてしまう事象のことである。
《データの一時領域格納》
図８は、本実施形態のＣＰＵ-ＧＰＵのデータ転送時の転送一括化によるループ文を示す図である。図８は、図７のネスト一括化および転送不要な変数明示に対応する。
図８に示すループ文では、図８の符号ｑに示す位置に、ＣＰＵ-ＧＰＵのデータ転送時、一時領域を作成するOpenACCのdeclare create文#pragma acc declare createを指定する。これにより、ＣＰＵ-ＧＰＵのデータ転送時は、一時領域を作成し（#pragma acc declare create）、データは一時領域に格納される。
また、図８の符号ｒに示す位置に、一時領域を同期するためのOpenACCのdeclare create文#pragma acc updateを指定することで転送を指示する。
このように、一時領域を作成し、一時領域でパラメータを初期化して、ＣＰＵ-ＧＰＵ転送に用いることで、不要なＣＰＵ-ＧＰＵ転送を遮断する。OpenACCの指示では意図しないが性能を劣化する転送を削減することができる。
［ＧＰＵオフロード処理］
上述したデータ転送の一括処理手法により、オフロードに適切なループ文を抽出し、非効率なデータ転送を避けることができる。
ただし、上記データ転送の一括処理手法を用いても、ＧＰＵオフロードに向いていないプログラムも存在する。効果的なＧＰＵオフロードには、オフロードする処理のループ回数が多いことが必要である。
そこで、本実施形態では、本格的なオフロード処理探索の前段階として、プロファイリングツールを用いて、ループ回数を調査する。プロファイリングツールを用いると、各行の実行回数を調査できるため、例えば、５０００万回以上のループを持つプログラムをオフロード処理探索の対象とする等、事前に振り分けることができる。以下、具体的に説明する（図２で述べた内容と一部重複する）。
本実施形態では、まず、アプリケーションコード分析部１１２（図１）がアプリケーションを分析し、for，do，while等のループ文を把握する。次に、サンプル処理を実行し、プロファイリングツールを用いて、各ループ文のループ回数を調査し、一定の値以上のループがあるか否かで、探索を本格的に行うか否かの判定を行う。
探索を本格的に行うと決まった場合は、ＧＡの処理に入る（図２参照）。初期化ステップでは、アプリケーションコードの全ループ文の並列可否をチェックした後、並列可能ループ文をＧＰＵ処理する場合は１、しない場合は０として遺伝子配列にマッピングする。遺伝子は、指定の個体数が準備されるが、遺伝子の各値にはランダムに１，０の割り当てをする。
ここで、遺伝子に該当するコードでは、ＧＰＵ処理すると指定されたループ文内の変数データ参照関係から、データ転送の明示的指示（#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の構文解析ライブラリ等を使用する。
実装では、最初に、そのアプリケーションがＧＰＵオフロード効果があるかの見込みを得るため、ベンチマークを実行し、上記構文解析で把握したfor文のループ回数を把握する。ループ回数把握には、GNUカバレッジのgcov等を用いる。プロファイリングツールとしては、「GNUプロファイラ(gprof)」、「GNUカバレッジ(gcov)」が知られている。双方とも各行の実行回数を調査できるため、どちらを用いてもよい。実行回数は、例えば、１０００万回以上のループ回数を持つアプリケーションのみ対象とするようにできるが、この値は変更可能である。
ＣＰＵ向け汎用アプリケーションは、並列化を想定して実装されているわけではない。そのため、まず、ＧＰＵ処理自体が不可なfor文は排除する必要がある。そこで、各for文一つずつに対して、ＧＰＵ処理の#pragma acc kernelsや#prama acc parallel loopや#prama acc parallel loop vectorディレクティブ挿入を試行し、コンパイル時にエラーが出るかの判定を行う。コンパイルエラーに関しては、幾つかの種類がある。for文の中で外部ルーチンが呼ばれている場合、ネストfor文で異なる階層が重複指定されている場合、break等でfor文を途中で抜ける処理がある場合、for文のデータにデータ依存性がある場合等がある。アプリケーションによって、コンパイル時エラーの種類は多彩であり、これ以外の場合もあるが、コンパイルエラーは処理対象外とし、#pragmaディレクティブは挿入しない。
コンパイルエラーは自動対処が難しく、また対処しても効果が出ないことも多い。外部ルーチンコールの場合は、#pragma acc routineにより回避できる場合があるが、多くの外部コールはライブラリであり、それを含めてＧＰＵ処理してもそのコールがネックとなり性能が出ない。for文一つずつを試行するため、ネストのエラーに関しては、コンパイルエラーは生じない。また、break等により途中で抜ける場合は、並列処理にはループ回数を固定化する必要があり、プログラム改造が必要となる。データ依存が有る場合はそもそも並列処理自体ができない。
ここで、並列処理してもエラーが出ないループ文の数がａの場合、ａが遺伝子長となる。遺伝子の１は並列処理ディレクティブ有、０は無に対応させ、長さａの遺伝子に、アプリケーションコードをマッピングする。
次に、初期値として，指定個体数の遺伝子配列を準備する。遺伝子の各値は、図３で説明したように、０と１をランダムに割当てて作成する。準備された遺伝子配列に応じて、遺伝子の値が１の場合はＧＰＵ処理を指定するディレクティブ ＼#pragma acc kernels，＼#pragma acc parallel loop，＼#pragma acc parallel loop vectorをC/C++コードに挿入する。single loop等はparallelにしない理由としては、同じ処理であればkernelsの方が、ＰＧＩコンパイラとしては性能が良いためである。この段階で、ある遺伝子に該当するコードの中で、ＧＰＵで処理させる部分が決まる。
並列処理およびデータ転送のディレクティブを挿入されたC/C++コードを、ＧＰＵを備えたマシン上のＰＧＩコンパイラでコンパイルを行う。コンパイルした実行ファイルをデプロイし、ベンチマークツールで性能と電力使用量を測定する。
全個体数に対して、ベンチマーク性能測定後、ベンチマーク処理時間と電力使用量に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。
次世代の個体に対して、ディレクティブ挿入、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。ここで、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイル、性能測定をせず、以前と同じ測定値を用いる。
指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当する、ディレクティブ付きC/C++コードを解とする。
この中で、個体数、世代数、交叉率、突然変異率、適合度設定、選択方法は、ＧＡのパラメータであり、別途指定する。提案技術は、上記処理を自動化することで、従来、専門技術者の時間とスキルが必要だった、ＧＰＵオフロードの自動化を可能にする。
図９Ａ－Ｂは、上述した実装の動作概要を説明するフローチャートであり、図９Ａと図９Ｂは、結合子で繋がれる。
C/C++向けOpenＡＣＣコンパイラを用いて以下の処理を行う。
<コード解析>
ステップＳ１０１で、アプリケーションコード分析部１１２（図１参照）は、C/C++アプリのコード解析を行う。
<ループ文特定>
ステップＳ１０２で、並列処理指定部１１４（図１参照）は、C/C++アプリのループ文、参照関係を特定する。
<ループ文の並列処理可能性>
ステップＳ１０３で、並列処理指定部１１４は、各ループ文のＧＰＵ処理可能性をチェックする（#pragma acc kernels）。
<ループ文の繰り返し>
制御部（自動オフロード機能部）１１は、ステップＳ１０４のループ始端とステップＳ１１７のループ終端間で、ステップＳ１０５－Ｓ１１６の処理についてループ文の数だけ繰り返す。
<ループの数の繰り返し（その１）>
制御部（自動オフロード機能部）１１は、ステップＳ１０５のループ始端とステップＳ１０８のループ終端間で、ステップＳ１０６－Ｓ１０７の処理についてループ文の数だけ繰り返す。
ステップＳ１０６で、並列処理指定部１１４は、各ループ文に対して、OpenACCでＧＰＵ処理（#pragma acc kernels）を指定してコンパイルする。
ステップＳ１０７で、並列処理指定部１１４は、エラー時は、次の指示句でＧＰＵ処理可能性をチェックする（#pragma acc parallel loop）。
<ループの数の繰り返し（その２）>
制御部（自動オフロード機能部）１１は、ステップＳ１０９のループ始端とステップＳ１１２のループ終端間で、ステップＳ１１０－Ｓ１１１の処理についてループ文の数だけ繰り返す。
ステップＳ１１０で、並列処理指定部１１４は、各ループ文に対して、OpenACCでＧＰＵ処理（#pragma acc parallel loop）を指定してコンパイルする。
ステップＳ１１１で、並列処理指定部１１４は、エラー時は、次の指示句でＧＰＵ処理可能性をチェックする（#pragma acc parallel loop vector）。
<ループの数の繰り返し（その３）>
制御部（自動オフロード機能部）１１は、ステップＳ１１３のループ始端とステップＳ１１６のループ終端間で、ステップＳ１１４－Ｓ１１５の処理についてループ文の数だけ繰り返す。
ステップＳ１１４で、並列処理指定部１１４は、各ループ文に対して、OpenACCでＧＰＵ処理（#pragma acc parallel loop vector）を指定してコンパイルする。
ステップＳ１１５で、並列処理指定部１１４は、エラー時は、当該ループ文からはＧＰＵ処理指示句を除去する。
<for文の数カウント>
ステップＳ１１８で、並列処理指定部１１４は、コンパイルエラーが出ないループ文（ここではfor文）の数をカウントし、遺伝子長とする。
<指定個体数パターン準備>
次に、初期値として、並列処理指定部１１４は、指定個体数の遺伝子配列を準備する。ここでは、０と１をランダムに割当てて作成する。
ステップＳ１１９で、並列処理指定部１１４は、C/C++アプリコードを、遺伝子にマッピングし、指定個体数パターン準備を行う。
準備された遺伝子配列に応じて、遺伝子の値が１の場合は並列処理を指定するディレクティブをC/C++コードに挿入する（例えば図３の#pragmaディレクティブ参照）。
制御部（自動オフロード機能部）１１は、図９ＢのステップＳ１２０のループ始端とステップＳ１３１のループ終端間で、ステップＳ１２１－Ｓ１３０の処理について指定世代数繰り返す。
また、上記指定世代数繰り返しにおいて、さらにステップＳ１２１のループ始端とステップＳ１２６のループ終端間で、ステップＳ１２２－Ｓ１２５の処理について指定個体数繰り返す。すなわち、指定世代数繰り返しの中で、指定個体数の繰り返しが入れ子状態で処理される。
<データ転送指定>
ステップＳ１２２で、データ転送指定部１１３は、変数参照関係をもとに、明示的指示行（#pragma acc data copy/copyin/copyout/presentおよび#pragam acc declarecreate, #pragma acc update）を用いたデータ転送指定を行う。
<コンパイル>
ステップＳ１２３で、並列処理パターン作成部１１７（図１参照）は、遺伝子パターンに応じてディレクティブ指定したC/C++コードをＰＧＩコンパイラでコンパイルする。すなわち、並列処理パターン作成部１１７は、作成したC/C++コードを、ＧＰＵを備えた検証用マシン１４上のＰＧＩコンパイラでコンパイルを行う。
ここで、ネストfor文を複数並列指定する場合等でコンパイルエラーとなることがある。この場合は、性能測定時の処理時間がタイムアウトした場合と同様に扱う。
ステップＳ１２４で、性能測定部１１８（図１参照）は、ＣＰＵ-ＧＰＵ搭載の検証用マシン１４に、実行ファイルをデプロイする。
ステップＳ１２５で、性能測定部１１８は、配置したバイナリファイルを実行し、オフロードした際のベンチマーク性能を測定する。
ここで、途中世代で、以前と同じパターンの遺伝子については測定せず、同じ値を使う。つまり、ＧＡ処理の中で、以前と同じパターンの遺伝子が生じた場合は、その個体についてはコンパイルや性能測定をせず、以前と同じ測定値を用いる。
ステップＳ１２７で、性能測定部１１８（図１参照）は、処理時間を測定する。
ステップＳ１２８で、性能測定部１１８は、測定した処理時間をもとに評価値を設定する。
ステップＳ１２９で、実行ファイル作成部１１９（図１参照）は、処理時間の短い個体ほど適合度が高くなるように評価し、性能の高い個体を選択する。実行ファイル作成部１１９は、測定された複数パターンの中で、短時間かつ低電力使用量のパターンを解として選択する。
ステップＳ１３０で、実行ファイル作成部１１９は、選択された個体に対して、交叉、突然変異の処理を行い、次世代の個体を作成する。実行ファイル作成部１１９は、次世代の個体に対して、コンパイル、性能測定、適合度設定、選択、交叉、突然変異処理を行う。
すなわち、全個体に対して、ベンチマーク性能測定後、ベンチマーク処理時間に応じて、各遺伝子配列の適合度を設定する。設定された適合度に応じて、残す個体の選択を行う。実行ファイル作成部１１９は、選択された個体に対して、交叉処理、突然変異処理、そのままコピー処理のＧＡ処理を行い、次世代の個体群を作成する。
ステップＳ１３２で、実行ファイル作成部１１９は、指定世代数のＧＡ処理終了後、最高性能の遺伝子配列に該当するC/C++コード（最高性能の並列処理パターン）を解とする。
<ＧＡのパラメータ>
上記、個体数、世代数、交叉率、突然変異率、適合度設定、選択方法は、ＧＡのパラメータである。ＧＡのパラメータは、例えば、以下のように設定してもよい。
実行するSimple GAの、パラメータ、条件は例えば以下のようにできる。
遺伝子長：並列可能ループ文数
個体数Ｍ：遺伝子長以下
世代数Ｔ：遺伝子長以下
適合度：(処理時間)(-1/2)
この設定により、ベンチマーク処理時間が短い程、高適合度になる。また、適合度を、処理時間の(-1/2)乗を含む形とすることで、処理時間が短い特定の個体の適合度が高くなり過ぎて、探索範囲が狭くなるのを防ぐことができる。また、性能測定が一定時間で終わらない場合は、タイムアウトさせ、処理時間１０００秒等の時間（長時間）であるとして、適合度を計算する。このタイムアウト時間は、性能測定特性に応じて変更させればよい。
選択：ルーレット選択
ただし、世代での最高適合度遺伝子は交叉も突然変異もせず次世代に保存するエリート保存も合わせて行う。
交叉率Ｐｃ：０．９
突然変異率Ｐｍ：０．０５
<コストパフォーマンス>
自動オフロード機能のコストパフォーマンスについて述べる。
NVIDIA Tesla等の、ＧＰＵボードのハードウェアの価格だけを見ると、ＧＰＵを搭載したマシンの価格は、通常のＣＰＵのみのマシンの約２倍となる。しかし、一般にデータセンタ等のコストでは、ハードウェアやシステム開発のコストが１／３以下であり、電気代や保守・運用体制等の運用費が１／３超であり、サービスオーダ等のその他費用が１／３程度である。本実施形態では、暗号処理や画像処理等動作させるアプリケーションで時間がかかる処理を２倍以上高性能化できる。このため、サーバハードウェア価格自体は２倍となっても、コスト効果が十分に期待できる。
本実施形態では、gcov，gprof等を用いて、ループが多く実行時間がかかっているアプリケーションを事前に特定して、オフロード試行をする。これにより、効率的に高速化できるアプリケーションを見つけることができる。
<本番サービス利用開始までの時間>
本番サービス利用開始までの時間について述べる。
コンパイルから性能測定1回は３分程度とすると、２０の個体数、２０の世代数のＧＡで最大２０時間程度解探索にかかるが、以前と同じ遺伝子パターンのコンパイル、測定は省略されるため、８時間以下で終了する。多くのクラウドやホスティング、ネットワークサービスではサービス利用開始に半日程度かかるのが実情である。本実施形態では、例えば半日以内の自動オフロードが可能である。このため、半日以内の自動オフロードであれば、最初は試し利用ができるとすれば、ユーザ満足度を十分に高めることが期待できる。
より短時間でオフロード部分を探索するためには、複数の検証用マシンにより個体数分並列で性能測定することが考えられる。アプリケーションに応じて、タイムアウト時間を調整することも短時間化に繋がる。例えば、オフロード処理がＣＰＵでの実行時間の２倍かかる場合はタイムアウトとする等である。また、個体数、世代数が多い方が、高性能な解を発見できる可能性が高まる。しかし、各パラメータを最大にする場合、個体数×世代数だけコンパイル、および性能ベンチマークを行う必要がある。このため、本番サービス利用開始までの時間がかかる。本実施形態では、ＧＡとしては少ない個体数、世代数で行っているが、交叉率Ｐｃを０．９と高い値にして広範囲を探索することで、ある程度の性能の解を早く発見するようにしている。
［指示句の拡大］
本実施形態では、適用できるアプリケーション増加のため、指示句の拡大を行う。具体的には、ＧＰＵ処理を指定する指示句として、kernels指示句に加えて，parallel loop指示句、parallel loop vector指示句にも拡大する。
OpenACC標準では、kernelsは、single loopやtightly nested loopに使われる。また、parallel loopは、non-tightly nested loopも含めたループに使われる。parallel loop vectorは、parallelizeはできないがvectorizeはできるループに使われる。ここで、tightly nested loopとは、ネストループにて、例えば、ｉとjをインクリメントする二つのループが入れ子になっている時、下位のループでｉとｊを使った処理がされ、上位ではされないような単純なループである。また、ＰＧＩコンパイラ等の実装においては、kernelsは、並列化の判断はコンパイラが行い、parallelは並列化の判断はプログラマが行うという違いがある。
そこで、本実施形態では、single、tightly nested loopにはkernelsを使い、non-tightly nested loopにはparallel loopを使う。また、parallelizeできないがvectorizeできるループにはparallel loop vectorを使う。
ここで、parallel指示句にすることで、結果がkernelsの場合より信頼度が下がる懸念がある。しかし、最終的なオフロードプログラムに対して、サンプルテストを行い、ＣＰＵとの結果差分をチェックしその結果をユーザに見せて、ユーザに確認してもらうことを想定している。そもそも、ＣＰＵとＧＰＵではハードが異なるため，有効数字桁数や丸め誤差の違い等があり、kernelsだけでもＣＰＵとの結果差分のチェックは必要である。
［リソース比とリソース量の設定、および新規アプリケーションの配置フローチャート］
図１０は、ＧＰＵオフロード試行の後に追加されるリソース比とリソース量の設定および新規アプリケーションの配置を説明するフローチャートである。図１０に示すフローチャートは、図９Ａ－Ｂに示すＧＰＵオフロード試行後に実行される。
ステップＳ５１でリソース比決定部１１５は、ユーザ運用条件、テストケースＣＰＵ処理時間、オフロードデバイス処理時間を取得する。ユーザ運用条件は、ユーザがオフロードしたいコードを指定する際に合わせてユーザに指定してもらう。ユーザ運用条件は、リソース量設定部１１６が、設備リソースＤＢ１３２の情報を参照してリソース量を決定する際に利用する。
<リソース比決定部１１５の処理>
ステップＳ５２でリソース比決定部１１５は、性能測定結果をもとに、ＣＰＵとオフロードデバイスの処理時間（テストケースＣＰＵ処理時間とオフロードデバイス処理時間）の比を、リソース比として決定する。
本自動オフロードによって、コード変換の際には、既に検証環境での性能測定結果が得られている。この性能測定結果を用いて、リソース比決定部１１５は、ＣＰＵとオフロードデバイスのリソース比を決定する。具体的には、検証環境でのＣＰＵとオフロードデバイスの処理時間の比に対して、適正なリソース比を決定する。例えば、検証環境でのテストケース処理時間が、ＣＰＵ処理：１０秒、ＧＰＵ処理：５秒の場合は、リソース比は、ＣＰＵ：ＧＰＵ＝２：１となる。
リソース比決定部１１５は、ＣＰＵとオフロードデバイスの処理時間が同等オーダになるように、リソース比を決定する。ＣＰＵとオフロードデバイスの処理時間が同等オーダになるように、リソース比を決定することで、ＣＰＵとオフロードデバイスの処理時間を揃え、ＣＰＵとアクセラレータがＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵ等の混在環境であってもリソース量を適切に設定することができる。
リソース比決定部１１５は、ＣＰＵとオフロードデバイスの処理時間の差分が所定閾値以上の場合、リソース比を所定の上限値に設定する。すなわち、検証環境でのＣＰＵとオフロードデバイスの処理時間が、例えば１０倍以上差分がある場合にリソース比を１０倍以上にしてしまうと、コストパフォーマンス悪化につながる。この場合は、例えば、５：１等のリソース比を上限にする（上限値は、処理時間の５：１のリソース比）。リソース比に上限を設けることで、ＶＭ数の大幅増加を防ぐことができる。
<リソース量設定部１１６の処理>
ステップＳ５３でリソース量設定部１１６は、ユーザ運用条件と適切リソース比をもとに、リソース量を設定する。すなわち、リソース量設定部１１６は、ユーザが指定したコスト条件を満たすように、リソース比はできるだけキープして、リソース量を決定する。
リソース量設定部１１６は、適切リソース比を維持して、ユーザ運用条件を満たす最大のリソース量を設定する。具体例を挙げると、ＣＰＵ１ＶＭは１０００円／月、ＧＰＵは４０００円／月、リソース比は２：１が適切であるとし、ユーザは月１００００円以内の予算であるとする。この場合には、ＣＰＵは２、ＧＰＵは１を確保して商用環境に配置する。
リソース量設定部１１６は、リソース比を維持した最小リソース量でもユーザ運用条件を満たさない場合は、リソース比を崩してＣＰＵとオフロードデバイスのリソース量をコスト条件を満たすよう最小で設定する。具体例を挙げると、ＣＰＵ１ＶＭは１０００円／月、ＧＰＵは４０００円／月、リソース比は２：１が適切であるとし、ユーザは月５０００円以内の予算であるとする。この場合には、ユーザ予算が足りないため、リソース比はキープできないが、ＣＰＵとオフロードデバイスのリソース量をより小さく設定、すなわちＣＰＵは１、ＧＰＵは１を確保して配置する。
上記ステップＳ５３の処理を終え、商用環境にリソースを確保して配置した後は、ユーザが利用する前に、性能およびコストを確認するため、図２で述べた自動検証を実行する。これにより、商用環境でリソースを確保して、自動検証後、性能とコストをユーザに提示することができる。
<リソース比とリソース量の設定のまとめ>
リソース比を適切化するため、オフロードパターンの解を決める際の性能測定結果を用いる。実装は、テストケースの処理時間から、ＣＰＵとＧＰＵの処理時間が同等オーダになるようリソース比を定める。例えば、テストケースの処理時間が、ＣＰＵ処理：１０秒、ＧＰＵ処理：５秒の場合では、ＣＰＵ側のリソースは２倍で同等の処理時間程度と考えられるため、リソース比は２：１となる。なお、仮想マシン等の数は整数となるため、リソース比は処理時間から計算する際に、整数比となるように四捨五入する。
リソース比が決定されると、次に、商用環境へのアプリケーション配置を行う際のリソース量の設定を行う。実装は、リソース量決定には、ユーザがオフロード依頼時に指定したコスト要求を満たすように、リソース比はできるだけキープして、ＶＭ等の数を定める。具体的には、コスト範囲内で、リソース比をキープする中では、ＶＭ等の数は最大値を選択する。
例えば、ＣＰＵに関して１ＶＭは１０００円／月、ＧＰＵは４０００円／月、リソース比は２：１が適切であり、ユーザは月１００００円以内の予算であった場合には、ＣＰＵは２、ＧＰＵは１を確保する。また、コスト範囲内で、リソース比をキープできない場合は、ＣＰＵ１単位、ＧＰＵ１単位から始めてできるだけ適切なリソース比に近くなるよう、リソース量を設定する。例えば、月５０００円以内の予算であった場合には、リソース比はキープできないが、ＣＰＵは１、ＧＰＵは１を確保する。
リソース量を設定すると、実装では、例えばXen Serverの仮想化機能を用いて、ＣＰＵやＧＰＵのリソースを割り当てる。
ステップＳ５４で配置設定部１７０は、設備リソースＤＢ１３２のサーバ、リンクのスペック情報、既存アプリケーションの配置情報に基づいて、線形計画手法を用いて、新規アプリケーションの配置先（ＡＰＬの配置場所）を計算して設定する。
［変換したアプリケーションの最適配置］
本実施形態のオフロードサーバ１は、ＣＰＵ向けプログラムを、ＧＰＵ等のデバイスにオフロードした際に、アプリケーションをユーザのコスト等要求を満たして、応答時間等を短く動作するように、配置先を適正化する。
<アプリケーション配置場所の適切化>
本実施形態では、アプリケーションはクラウドだけでなく、ネットワークエッジやユーザエッジに配置できることを前提とする。ただし、ネットワークエッジやユーザエッジは、クラウドに比べサーバの集約度が低く分散している。このため、計算リソースのコストは、クラウドに比べ割高となる。すなわち、一般にＣＰＵやＧＰＵ等のハードウェアの価格は配置場所によらず一定であるものの、クラウドを運用するデータセンタでは集約されたサーバをまとめて監視や空調制御等できるため、運用費が割安となる。
例えば、計算ノードリンクの簡単なトポロジーとしては、図１１が挙げられる。
図１１は、計算ノードのトポロジーの一例を示す図である。図１１は、ＩｏＴシステムのように、ユーザ環境でデータを収集するＩｏＴデバイス等から、ユーザエッジにデータが送られ、ネットワークエッジを介してクラウドにデータが送られ、分析結果を会社の幹部が見る等で使われるトポロジーである。
図１１に示すように、アプリケーションを配置するトポロジーは、３層で構成され、クラウドレイヤー（例えば、データセンタ）の拠点数は「２」（ｎ１３，ｎ１４）、キャリアエッジレイヤー（例えば、局舎）は「３」、ユーザエッジレイヤー（例えば、ユーザ環境）は「４」（ｎ６－ｎ９）、インプットノードは「５」（ｎ１－ｎ５）とする。
ＩｏＴ等のアプリケーションを想定してインプットノードからＩｏＴデータ（ＩｏＴデバイスの一つである花粉センサや体温センサ等）がユーザエッジに収集され、アプリケーションの特性（応答時間の要求条件等）に応じて、ユーザエッジ、キャリアエッジで分析処理がされたり、クラウドまでデータをあげてから分析処理されたりされる。アウトプットノードは「１」（ｎ１５）であり、分析結果を会社の幹部が見る。例えば、インプットノードがＩｏＴデータ（花粉センサ）の場合は、アウトプットノードの統計・分析結果は気象庁の責任者が確認する。
図１１に示す配置トポロジー３層は、一例であり、例えば５層であってもよい。また、ユーザエッジ、キャリアエッジの数は、実際には数十～数百の場合もある。
計算ノードは、ＣＰＵ、ＧＰＵ、ＦＰＧＡの３種に分けられる。ＧＰＵやＦＰＧＡを備えるノードには、ＣＰＵも搭載されているが、仮想化技術（例えば、NVIDIA vＧＰＵ）により、ＧＰＵインスタンス、ＦＰＧＡインスタンスとして、ＣＰＵリソースも含む形で分割して提供される。
アプリケーションは、クラウド、キャリアエッジ、ユーザエッジに配置され、ユーザ環境に近い側程、応答時間を低減することが可能になる代わりに、計算リソースのコストが高くなる。本実施形態では、ＧＰＵやＦＰＧＡ向けに変換したアプリケーションを配置することになるが、配置する際に、ユーザは２種類のリクエストを発出できる。
一つ目は、コスト要求であり、アプリケーションを動作させるために許容できる計算リソースのコストを指定する形で、例えば月５０００円以内で動作させる等である。二つ目は、応答時間要求であり、アプリケーションを動作させる際の許容応答時間を指定する形で、例えば１０秒以内に応答を返す等である。従来から行われている設備設計では、例えば仮想ネットワークを収容するサーバを配置する場所を、トラフィック増加量等の長期的傾向を見て、計画的に設計している。
本実施形態では、下記（１），（２）の特徴がある。（１）配置されるアプリケーションは静的に定まっているのではなく、ＧＰＵやＦＰＧＡ向けに自動変換され、ＧＡ等を通じて利用形態に適したパターンが実測を通じて抽出される。このため、アプリケーションのコードや性能は動的に変わり得る。
（２）キャリアの設備コストや全体的応答時間だけを低減すればよいのではなく、計算リソースのコストや応答時間に対する個々のユーザ要求を満たす必要がある。また、アプリケーションの配置ポリシーも動的に変わり得る。
上記（１），（２）の特徴も踏まえ、本実施形態のアプリケーション配置は、ユーザからの配置依頼があった場合、変換を行い、変換したアプリケーションをその時点で適切なサーバに順次配置していく形とする。アプリケーションを変換しても、コストパフォーマンスが向上しない場合は、変換前のアプリケーション配置とする。例えば、ＧＰＵインスタンスはＣＰＵインスタンスの２倍のコストがかかる際に、変換しても２倍以上性能が改善されないならば、変換前を配置した方がよい。また、既に上限まで計算リソースや帯域が使われてしまっている場合はそのサーバには配置はできないことがある。
<アプリケーション適切配置のための線形計画式>
本実施形態では、アプリケーションの適切な配置場所を計算するための、線形計画手法の定式化を行う。線形計画手法は、具体的には、［式１］（以下の式（１）～式（４））、［式２］（以下の式（３）～式（６））に示す線形計画式のパラメータを用いる。
ここで、デバイスやリンクのコストや計算リソース上限、帯域上限等は、事業者が準備するサーバやネットワークに依存する。このため、それらのパラメータ値は事業者が事前に設定する。オフロードした際にアプリケーションが使用する計算リソース量、帯域、データ容量、処理時間は、自動変換する前の検証環境での試験での最終的に選択されたオフロードパターンでの計測値により決まり、環境適応機能により自動設定される。
ユーザ要求が計算リソースのコスト要求であるかまたは応答時間要求であるかで、線形計画式のパラメータにおける、目的関数と制約条件が変わる。
・コスト要求による、線形計画式のパラメータ
コスト要求により、一月幾ら以内での配置が必要な要求の場合は、下記［式１］に示す線形計画式のパラメータを用いる。
式（１）の応答時間の最小化が目的関数である。式（２）の計算リソースのコストがいくら以内であるかは、制約条件の一つである。さらに、式（３）（４）のサーバのリソース上限を超えていないかの制約条件も加わる。
・応答時間要求による、線形計画式のパラメータ
応答時間要求により、アプリケーションの応答時間が何秒以内での配置が必要な要求の場合は、下記［式２］に示す線形計画式のパラメータを用いる。
式（２）に対応する式（５）の計算リソースのコストの最小化が目的関数である。式（１）に対応する式（６）の応答時間が何秒以内であるかは、制約条件の一つである。さらに、式（３）（４）の制約条件も加わる。
・線形計画式のパラメータの説明
式（１）および式（６）は、アプリケーションｋの応答時間を計算するための式であり、式（１）の場合はＲｋが目的関数、式（６）の場合はＲｋがユーザが指定した上限を設定する制約条件である。
式（２）および式（５）は、アプリケーションｋを動作させるコスト（価格）Ｐｋを計算するための式であり、式（２）の場合はＰｋがユーザが指定した上限を設定する制約条件、式（５）の場合はＰｋが目的関数である。
式（３）および式（４）は、計算リソースおよび通信帯域の上限を設定する制約条件であり、他者が配置したアプリケーション含めて計算され、新規ユーザのアプリケーション配置によるリソース上限の超過を防ぐ。
式（１）乃至式（４）および、式（３）乃至式（６）の線形計画式を、ネットワークトポロジーや変換アプリケーションタイプ（ＣＰＵに対するコスト増と性能増等）、ユーザ要求、既配置アプリケーションの異なる条件に対して、ＧＬＰＫ（Gnu Linear Programming Kit）やＣＰＬＥＸ（IBM Decision Optimization）等の線形計画ソルバで解を導出することで、適切なアプリケーション配置を計算できる。適切配置計算後に実際の配置を、複数のユーザに対して、順次行っていくことで、複数のアプリケーションが各ユーザの要求に基づいて配置される。
以上のように、線形計画式に基づいて、新規にアプリケーションの配置依頼があった場合に計算し、順に配置することで、ユーザ要望を満たした配置が可能である。
ここで、アプリケーションプログラムの配置は、順次行われるため早い者勝ちと言えるが、アプリケーション１００個毎等、定期的に、既に配置済みのアプリケーションプログラム群の適正配置を再計算する。そして、ユーザの指定するコスト、応答時間に応じて、目的関数が極小化される配置を計算し、計算で定まった位置に、アプリケーションを再配置してもよい。
［評価］
線形計画手法の一態様である線形計画式に基づき、無償ソルバのＧＬＰＫ（登録商標）を用いて、複数のアプリケーションが適切に配置されていくことを、いくつかの条件を変更して確認した。
<評価条件>
・対象アプリケーション
配置対象のアプリケーションは、多くのユーザが利用すると想定されるフーリエ変換による画像処理をする。フーリエ変換処理（FFT）は、振動周波数の分析等、ＩｏＴでのモニタリングの様々な場面で利用されている。
NAS.FT（https://www.nas.nasa.gov/publications/npb.html）（登録商標）は、FFT処理のオープンソースアプリケーションの一つである。備え付けのサンプルテストの２０４８×２０４８サイズの計算を行う。ＩｏＴで、デバイスからデータをネットワーク転送するアプリケーションについて想定した際に、ネットワークコストを下げるため、デバイス側でFFT 処理等の一次分析をして送ることが想定される。
MRI-Q（http://impact.crhc.illinois.edu/parboil/）（登録商標）は、非デカルト空間の３次元ＭＲＩ再構成アルゴリズムで使用されるキャリブレーション用のスキャナー構成を表す行列Ｑを計算する。ＩｏＴ環境では、カメラビデオからの自動監視のために画像処理が必要になることが多く、画像処理の自動オフロードへのニーズはある。MRI-QはＣ言語アプリケーションで、パフォーマンス測定中に３次元ＭＲＩ画像処理を実行し、Large の６４×６４×６４サイズのサンプルデータを使用して処理時間を測定する。ＣＰＵ処理はＣ言語で、ＦＰＧＡ処理はOpenCL（登録商標）に基づき処理される。
本実施形態のＧＰＵ、ＦＰＧＡ自動オフロード技術により、NAS.FTはＧＰＵで高速化でき、MRI-QはＦＰＧＡで高速化でき、それぞれ、ＣＰＵに比べて５倍、７倍の高速化ができる。
・評価手法
アプリケーションを配置するトポロジーは、図１１に示すように３層で構成され、クラウドレイヤーの拠点数は「５」、キャリアエッジレイヤーは「２０」、ユーザエッジレイヤーは「６０」、インプットノードは「３００」とする。ＩｏＴ等のアプリケーションを想定してインプットノードからＩｏＴデータ等がユーザエッジに収集され、アプリケーションの特性（応答時間の要求条件等）に応じて、ユーザエッジ、キャリアエッジで分析処理がされたり、クラウドまでデータをあげてから分析処理されたりされる。
［式１］［式２］に示す線形計画式のパラメータを元に、ユーザ要求条件に基づいて、例えば１０００個のアプリケーションを配置する。アプリケーションは、ＩｏＴアプリケーションで、インプットノードから生じるデータを分析する想定である。インプットノード（「３００」あるとする）から配置依頼をランダムに生じさせる。
例えば、配置依頼数として、NAS.FT：MRI-Q=３:１の割合で１０００回アプリを配置依頼する。また、ユーザ要求として、配置依頼する際にアプリ毎に価格条件か応答時間条件が選ばれる。NAS.FTの場合、価格に関しては月７０００円上限か８５００円上限か１００００円上限、応答時間に関しては６秒上限か７秒条件か１０秒上限が選択される。MRI-Qの場合、価格に関しては月１２５００円上限か２００００円上限、応答時間に関しては、４秒上限か８秒上限が選択される。
ユーザ要求のバリエーションとして、３パターンがある。
パターン１：NAS.FTでは６種のリクエストを１／６ずつ、MRI-Qでは４種のリクエストを１／４ずつ選択する。
パターン２：リクエストは最低価格が上限の条件を選択（最初は７０００円、１２５００円）し、空きがない場合は次に安い価格条件とする。
パターン３：リクエストは最低応答時間が上限の条件を選択（最初は６秒、４秒）し、空きがない場合は次に速い応答時間条件とする。
・配置のシミュレーション
配置は、評価ツールとしてソルバＧＬＰＫ５.０（登録商標）を用いてシミュレーション実験により行う。規模のあるネットワーク配置の模擬のため、評価ツールを用いたシミュレーションになる。実利用の際は、アプリケーションのオフロード依頼が来た場合、検証環境を用いた繰返し性能試験でオフロードパターンを作成し、検証環境での性能試験結果に基づいて適切なリソース量を決める（図１０参照）。そして、ユーザ要望に応じてＧＬＰＫ等を用いて適切な配置を定め、実際にデプロイした際の正常確認試験や性能試験を自動で行い、その結果と価格をユーザに提示して、ユーザ判断後利用を開始する。
図１２は、平均応答時間のアプリケーション配置数変化を示すグラフである。図１２は、平均応答時間とアプリケーション配置数を、上記３パターンに対して取る。
パターン２ではクラウドから順に、パターン３ではエッジから順に埋まっていくことが確認できた。パターン１では、多様な依頼が来た場合に、ユーザ要求条件を満たして配置される。
図１２に示すように、パターン２では、４００配置位までは全てクラウドに配置され平均応答時間は最遅のままであるが、クラウドが埋まると段々下がっていくことが分かる。
パターン３では、NAS.FTはユーザエッジから、またMRI-Qはキャリアエッジから配置される。このため、平均応答時間については最短となる。しかし、数が増えるとクラウドにも配置されるため平均応答時間は遅くなる。パターン２では、平均応答時間は、パターン１やパターン３の中間であり、ユーザ要求に応じて配置される。このため、パターン２では、最初はクラウドに全て入るパターン１に比べて平均応答時間は適切に低減されている。
このように、ソフトウェアを配置先環境に合わせて自動適応させ、ＧＰＵ等に自動オフロードした際に、ユーザのコスト要求、応答時間要求に応える。すなわち、ＧＰＵ等のデバイスで処理できるよう、プログラムを変換し、アサインするリソース量が定まった後に変換したアプリケーションの最適配置を行う。
まとめると、まず、プログラム変換する際に検証環境で行った性能試験のデータから、アプリケーションの利用データ容量、計算リソース量、帯域、処理時間を設定する。変換アプリケーション毎に設定される値と、事前に設定されるサーバやリンクのコスト等の値から、線形計画式に基づき、アプリケーションの適切な配置が計算される。アプリケーション配置の際は、ユーザが指定する価格や応答時間のリクエストに基づき、一方が制約条件にもう一方が目的関数となる。線形計画ソルバにより適切な配置が計算され、提案方式は計算された場所にリソースを配置した際の、価格等をユーザに提示し、ユーザ承諾後に利用が開始される。
ＧＰＵ、ＦＰＧＡに自動オフロードしたアプリケーションに対して、ユーザのリクエストする価格条件や応答時間条件、アプリケーションの配置数等を変更して、適正配置を計算する。これにより、ユーザ要望に従った配置が可能になる。
（第２実施形態）
次に、本発明の第２実施形態における、オフロードサーバ１Ａ等について説明する。
第２実施形態は、ループ文のＦＰＧＡ自動オフロードに適用した例である。
本実施形態は、ＰＬＤ（Programmable Logic Device）として、ＦＰＧＡ（Field Programmable Gate Array）に適用した例について説明する。本発明は、プログラマブルロジックデバイス全般に適用可能である。
（原理説明）
ＦＰＧＡで、どのループをオフロードすれば高速になるかの予測は難しいため、ＧＰＵ同様検証環境で自動測定することを提案している。しかし、ＦＰＧＡは、OpenCLをコンパイルして実機で動作させるまで数時間以上かかるため、ＧＰＵ自動オフロードでのＧＡを用いて何回も反復して測定することは、処理時間が膨大となり行うことはできない。そこで、ＦＰＧＡにオフロードする候補のループ文を絞ってから、測定を行う形をとる。具体的には、発見されたループ文に対して、ＲＯＳＥ（登録商標）等の算術強度分析ツールを用いて算術強度が高いループ文を抽出する。更に、gcov（登録商標）等のプロファイリングツールを用いてループ回数が多いループ文も抽出する。
算術強度やループ回数が多いループ文を候補として、OpenCL 化を行う。OpenCL 化時には、ＣＰＵ処理プログラムを、カーネル（ＦＰＧＡ）とホスト（ＣＰＵ）に、OpenCL の文法に従って分割する。候補ループ文に対して、作成したOpenCL をプレコンパイルして、リソース効率が高いループ文を見つける。これは、コンパイルの途中で、作成するリソースは分かるため、利用するリソース量が十分少ないループ文に更に絞り込む。
候補ループ文が幾つか残るため、それらを用いて性能や電力使用量を実測する。選択された単ループ文に対してコンパイルして測定し、更に高速化できた単ループ文に対してはその組み合わせパターンも作り２回目の測定をする。測定された複数パターンの中で、短時間かつ低電力使用量のパターンを解として選択する。
ループ文のＦＰＧＡ オフロードについては、算術強度等を用いて絞り込んでから、測定を行い、低電力パターンの評価値を高めることで、自動での高速化、低電力化を行う。
図１３は、本発明の第２実施形態に係るオフロードサーバ１Ａの構成例を示す機能ブロック図である。本実施形態の説明に当たり、図１と同一構成部分には同一符号を付して重複箇所の説明を省略する。
オフロードサーバ１Ａは、アプリケーションの特定処理をアクセラレータに自動的にオフロードする装置である。
また、オフロードサーバ１Ａは、エミュレータに接続可能である。
図１３に示すように、オフロードサーバ１Ａは、制御部２１と、入出力部１２と、記憶部１３と、検証用マシン１４（Verification machine）(アクセラレータ検証用装置)と、を含んで構成される。
制御部２１は、オフロードサーバ１Ａ全体の制御を司る自動オフロード機能部（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）１２２と、を備える。
<ＰＬＤ処理指定部２１３>
ＰＬＤ処理指定部２１３は、アプリケーションのループ文（繰り返し文）を特定し、特定した各ループ文に対して、ＰＬＤにおけるパイプライン処理、並列処理をOpenＣＬで指定した複数のオフロード処理パターンを作成してコンパイルする。
ＰＬＤ処理指定部２１３は、オフロード範囲抽出部（Extract offload able area）２１３ａと、中間言語ファイル出力部（Output intermediate file）２１３ｂと、を備える。
オフロード範囲抽出部２１３ａは、ループ文やＦＦＴ等、ＦＰＧＡにオフロード可能な処理を特定し、オフロード処理に応じた中間言語を抽出する。
中間言語ファイル出力部２１３ｂは、抽出した中間言語ファイル１３３を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。
<算術強度算出部２１４>
算術強度算出部２１４は、例えばROSEフレームワーク（登録商標）等の算術強度（Arithmetic Intensity）分析ツールを用いて、アプリケーションのループ文の算術強度を算出する。算術強度は、プログラムの稼働中に実行した浮動小数点演算（floating point number，ＦＮ）の数を、主メモリへのアクセスしたbyte数で割った値（ＦＮ演算／メモリアクセス）である。
算術強度は、計算回数が多いと増加し、アクセス数が多いと減少する指標であり、算術強度が高い処理はプロセッサにとって重い処理となる。そこで、算術強度分析ツールで、ループ文の算術強度を分析する。ＰＬＤ処理パターン作成部２１５は、算術強度が高いループ文をオフロード候補に絞る。
算術強度の計算例について述べる。
１回のループの中での浮動小数点計算処理が１０回（１０FLOP）行われ、ループの中で使われるデータが２byteであるとする。ループ毎に同じサイズのデータが使われる際は、１０／２＝５ [FLOP/byte]が算術強度となる。
なお、算術強度では、ループ回数が考慮されないため、本実施形態では、算術強度に加えて、ループ回数も考慮して絞り込む。
<ＰＬＤ処理パターン作成部２１５>
ＰＬＤ処理パターン作成部２１５は、算術強度算出部２１４が算出した算術強度をもとに、算術強度が所定の閾値より高い（以下、適宜、高算術強度という）ループ文をオフロード候補として絞り込み、ＰＬＤ処理パターンを作成する。
また、ＰＬＤ処理パターン作成部２１５は、基本動作として、コンパイルエラーが出るループ文（繰り返し文）に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、ＰＬＤ処理するかしないかの指定を行うＰＬＤ処理パターンを作成する。
・ループ回数測定機能
ＰＬＤ処理パターン作成部２１５は、ループ回数測定機能として、プロファイリングツールを用いてアプリケーションのループ文のループ回数を測定し、ループ文のうち、高算術強度で、ループ回数が所定の回数より多い（以下、適宜、高ループ回数という）ループ文を絞り込む。ループ回数把握には、GNUカバレッジのgcov等を用いる。プロファイリングツールとしては、「GNUプロファイラ(gprof)」、「GNUカバレッジ(gcov)」が知られている。双方とも各ループの実行回数を調査できるため、どちらを用いてもよい。
また、算術強度分析では、ループ回数は特に見えないため、ループ回数が多く負荷が高いループを検出するため、プロファイリングツールを用いて、ループ回数を測定する。ここで、算術強度の高さは、ＦＰＧＡへのオフロードに向いた処理かどうかを表わし、ループ回数×算術強度は、ＦＰＧＡへのオフロードに関連する負荷が高いかどうかを表わす。
・OpenＣＬ（中間言語）作成機能
ＰＬＤ処理パターン作成部２１５は、OpenＣＬ作成機能として、絞り込まれた各ループ文をＦＰＧＡにオフロードするためのOpenCLを作成（OpenCL化）する。すなわち、ＰＬＤ処理パターン作成部２１５は、絞り込んだループ文をオフロードするOpenCLをコンパイルする。また、ＰＬＤ処理パターン作成部２１５は、性能測定された中でＣＰＵに比べ高性能化されたループ文をリスト化し、リストのループ文を組み合わせてオフロードするOpenCLを作成する。
OpenCL化について述べる。
ＰＬＤ処理パターン作成部２１５は、ループ文をOpenCL等の高位言語化する。まず、ＣＰＵ処理のプログラムを、カーネル（ＦＰＧＡ）とホスト（ＣＰＵ）に、OpenCL等の高位言語の文法に従って分割する。例えば、１０個のfor文の内一つのfor文をＦＰＧＡで処理する場合は、その一つをカーネルプログラムとして切り出し、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 Ｃ言語に沿って作成されたカーネルの、デバイス実行について説明した。
・リソース量算出機能
ＰＬＤ処理パターン作成部２１５は、リソース量算出機能として、作成したOpenCLをプレコンパイルして利用するリソース量を算出する（「１回目のリソース量算出」）。ＰＬＤ処理パターン作成部２１５は、算出した算術強度およびリソース量に基づいてリソース効率を算出し、算出したリソース効率をもとに、各ループ文で、リソース効率が所定の値より高いｃ個のループ文を選ぶ。
ＰＬＤ処理パターン作成部２１５は、組み合わせたオフロードOpenCLでプレコンパイルして利用するリソース量を算出する（「２回目のリソース量算出」）。ここで、プレコンパイルせず、１回目測定前のプレコンパイルでのリソース量の和でもよい。
<性能測定部１１８>
性能測定部１１８は、作成されたＰＬＤ処理パターンのアプリケーションをコンパイルして、検証用マシン１４に配置し、ＰＬＤにオフロードした際の性能測定用処理を実行する。
性能測定部１１８は、配置したバイナリファイルを実行し、オフロードした際の性能を測定するとともに、性能測定結果を、オフロード範囲抽出部２１３ａに戻す。この場合、オフロード範囲抽出部２１３ａは、別のＰＬＤ処理パターン抽出を行い、中間言語ファイル出力部２１３ｂは、抽出された中間言語をもとに、性能測定を試行する（図２の符号ａ参照）。
性能測定部１１８は、バイナリファイル配置部（Deploy binary files）１１８ａを備える。バイナリファイル配置部１１８ａは、ＧＰＵを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイ(配置)する。
性能測定の具体例について述べる。
ＰＬＤ処理パターン作成部２１５は、高リソース効率のループ文を絞り込み、実行ファイル作成部１１９が絞り込んだループ文をオフロードするOpenCLをコンパイルする。性能測定部１１８は、コンパイルされたプログラムの性能を測定する（「１回目の性能測定」）。
そして、ＰＬＤ処理パターン作成部２１５は、性能測定された中でＣＰＵに比べ高性能化されたループ文をリスト化する。ＰＬＤ処理パターン作成部２１５は、リストのループ文を組み合わせてオフロードするOpenCLを作成する。ＰＬＤ処理パターン作成部２１５は、組み合わせたオフロードOpenCLでプレコンパイルして利用するリソース量を算出する。
なお、プレコンパイルせず、１回目測定前のプレコンパイルでのリソース量の和でもよい。実行ファイル作成部１１９は、組み合わせたオフロードOpenCLをコンパイルし、性能測定部１１８は、コンパイルされたプログラムの性能を測定する（「２回目の性能測定」）。
<実行ファイル作成部１１９>
実行ファイル作成部１１９は、所定回数繰り返された、処理時間の測定結果をもとに、複数のＰＬＤ処理パターンから最高評価値のＰＬＤ処理パターンを選択し、最高評価値のＰＬＤ処理パターンをコンパイルして実行ファイルを作成する。
以下、上述のように構成されたオフロードサーバ１Ａの自動オフロード動作について説明する。
［自動オフロード動作］
本実施形態のオフロードサーバ１Ａは、環境適応ソフトウェアの要素技術としてユーザアプリケーションロジックのＦＰＧＡ自動オフロードに適用した例である。
図２に示すオフロードサーバ１Ａの自動オフロード処理を参照して説明する。
図２に示すように、オフロードサーバ１Ａは、環境適応ソフトウェアの要素技術に適用される。オフロードサーバ１Ａは、制御部（自動オフロード機能部）１１と、テストケースＤＢ１３１と、中間言語ファイル１３３と、検証用マシン１４と、を有している。
オフロードサーバ１Ａは、ユーザが利用するアプリケーションコード（Application code）１２５を取得する。
ユーザは、例えば、各種デバイス（Device）１５１、ＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３、ＣＰＵを有する装置１５４を利用する。オフロードサーバ１Ａは、機能処理をＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３のアクセラレータに自動オフロードする。
以下、図２のステップ番号を参照して各部の動作を説明する。
<ステップＳ２１：Specify application code>
ステップＳ２１において、アプリケーションコード指定部１１１（図１３参照）は、ユーザに提供しているサービスの処理機能（画像分析等）を特定する。具体的には、アプリケーションコード指定部１１１は、入力されたアプリケーションコードの指定を行う。
<ステップＳ１２：Analyze application code>
ステップＳ１２において、アプリケーションコード分析部１１２（図１３参照）は、処理機能のソースコードを分析し、ループ文やＦＦＴライブラリ呼び出し等の特定ライブラリ利用の構造を把握する。
<ステップＳ１３：Extract offload able area>
ステップＳ１３において、ＰＬＤ処理指定部２１３（図１３参照）は、アプリケーションのループ文（繰り返し文）を特定し、各繰り返し文に対して、ＦＰＧＡにおける並列処理またはパイプライン処理を指定して、高位合成ツールでコンパイルする。具体的には、オフロード範囲抽出部２１３ａ（図１３参照）は、ループ文等、ＦＰＧＡにオフロード可能な処理を特定し、オフロード処理に応じた中間言語としてOpenCLを抽出する。
<ステップＳ１４：Output intermediate file>
ステップＳ１４において、中間言語ファイル出力部２１３ｂ（図１３参照）は、中間言語ファイル１３３を出力する。中間言語抽出は、一度で終わりでなく、適切なオフロード領域探索のため、実行を試行して最適化するため反復される。
<ステップＳ１５：Compile error>
ステップＳ１５において、ＰＬＤ処理パターン作成部２１５（図１３参照）は、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ない繰り返し文に対して、ＦＰＧＡ処理するかしないかの指定を行うＰＬＤ処理パターンを作成する。
<ステップＳ２１：Deploy binary files>
ステップＳ２１において、バイナリファイル配置部１１８ａ（図１３参照）は、ＦＰＧＡを備えた検証用マシン１４に、中間言語から導かれる実行ファイルをデプロイする。バイナリファイル配置部１１８ａは、配置したファイルを起動し、想定するテストケースを実行して、オフロードした際の性能を測定する。
<ステップＳ２２：Measure performances>
ステップＳ２２において、性能測定部１１８（図１３参照）は、配置したファイルを実行し、オフロードした際の性能と電力使用量を測定する。
オフロードする領域をより適切にするため、この性能測定結果は、オフロード範囲抽出部２１３ａに戻され、オフロード範囲抽出部２１３ａが、別パターンの抽出を行う。そして、中間言語ファイル出力部２１３ｂは、抽出された中間言語をもとに、性能測定を試行する（図２の符号ａ参照）。性能測定部１１８は、検証環境での性能・電力使用量測定を繰り返し、最終的にデプロイするコードパターンを決定する。
図２の符号ａに示すように、制御部２１は、上記ステップＳ１２乃至ステップＳ２２を繰り返し実行する。制御部２１の自動オフロード機能をまとめると、下記である。すなわち、ＰＬＤ処理指定部２１３は、アプリケーションのループ文（繰り返し文）を特定し、各繰返し文に対して、ＦＰＧＡにおける並列処理またはパイプライン処理をOpenCL（中間言語）で指定して、高位合成ツールでコンパイルする。そして、ＰＬＤ処理パターン作成部２１５は、コンパイルエラーが出るループ文を、オフロード対象外とし、コンパイルエラーが出ないループ文に対して、ＰＬＤ処理するかしないかの指定を行うＰＬＤ処理パターンを作成する。そして、バイナリファイル配置部１１８ａは、該当ＰＬＤ処理パターンのアプリケーションをコンパイルして、検証用マシン１４に配置し、性能測定部１１８が、検証用マシン１４で性能測定用処理を実行する。実行ファイル作成部１１９は、所定回数繰り返された、性能測定結果をもとに、複数のＰＬＤ処理パターンから最高評価値（例えば、評価値＝（処理時間）－１／２が最も高いもの）のパターンを選択し、選択パターンをコンパイルして実行ファイルを作成する。
<ステップＳ２３：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等の高位言語からハードウェア記述のＨＤＬ等のレベルに変換され、それに基づき実際の配線処理等がされる。この時、配線処理等は多大な時間がかかるが、ＨＤＬ等の途中状態の段階までは時間は分単位でしかかからない。ＨＤＬ等の途中状態の段階であっても、ＦＰＧＡで利用するFlip FlopやLook Up Table等のリソースは分かる。このため、ＨＤＬ等の途中状態の段階をみれば、利用するリソース量はコンパイルが終わらずとも短時間でわかる。
そこで、本実施形態では、ＰＬＤ処理パターン作成部２１５は、対象のループ文をOpenCL等の高位言語化し、まずリソース量を算出する。また、ループ文をオフロードした際の算術強度とリソース量が決まるため、算術強度／リソース量または算術強度×ループ回数／リソース量をリソース効率とする。そして、高リソース効率のループ文をオフロード候補として更に絞り込む。
図１４のフローに戻って、ステップＳ２０４でＰＬＤ処理パターン作成部２１５は、gcov、gprof等のプロファイリングツールを用いてアプリケーションのループ文のループ回数を測定する。
ステップＳ２０５でＰＬＤ処理パターン作成部２１５は、ループ文のうち、高算術強度で高ループ回数のループ文を絞り込む。
ステップＳ２０６でＰＬＤ処理パターン作成部２１５は、絞り込まれた各ループ文をＦＰＧＡにオフロードするためのOpenCLを作成する。
ここで、ループ文のOpenCL化（OpenCLの作成）について、補足して説明する。すなわち、ループ文をOpenCL等によって、高位言語化する際には、２つの処理が必要である。一つは、ＣＰＵ処理のプログラムを、カーネル（ＦＰＧＡ）とホスト（ＣＰＵ）に、OpenCL等の高位言語の文法に従って分割することである。もう一つは、分割する際に、高速化するための技法を盛り込むことである。一般に、ＦＰＧＡを用いて高速化するためには、ローカルメモリキャッシュ、ストリーム処理、複数インスタンス化、ループ文の展開処理、ネストループ文の統合、メモリインターリーブ等がある。これらは、ループ文によっては、絶対効果があるわけではないが、高速化するための手法として、よく利用されている。
次に、高リソース効率のループ文が幾つか選択されたので、それらを用いて性能を実測するオフロードパターンを実測する数だけ作成する。ＦＰＧＡでの高速化は、１個の処理だけＦＰＧＡリソース量を集中的にかけて高速化する形もあれば、複数の処理にＦＰＧＡリソースを分散して高速化する形もある。選択された単ループ文のパターンを一定数作り、ＦＰＧＡ実機で動作する前段階としてプレコンパイルする。
ステップＳ２０７でＰＬＤ処理パターン作成部２１５は、作成したOpenCLをプレコンパイルして利用するリソース量を算出する（「１回目のリソース量算出」）。
ステップＳ２０８でＰＬＤ処理パターン作成部２１５は、高リソース効率のループ文を絞り込む。
ステップＳ２０９で実行ファイル作成部１１９は、絞り込んだループ文をオフロードするOpenCLをコンパイルする。
ステップＳ２１０で性能測定部１１８は、コンパイルされたプログラムの性能を測定する（「１回目の性能測定」）。候補ループ文が幾つか残るため、性能測定部１１８は、それらを用いて性能を実測する（詳細については、図１５のサブルーチン参照）。
ステップＳ２１１でＰＬＤ処理パターン作成部２１５は、性能測定された中でＣＰＵに比べ高性能化されたループ文をリスト化する。
ステップＳ２１２でＰＬＤ処理パターン作成部２１５は、リストのループ文を組み合わせてオフロードするOpenCLを作成する。
ステップＳ２１３でＰＬＤ処理パターン作成部２１５は、組み合わせたオフロードOpenCLでプレコンパイルして利用するリソース量を算出する（「２回目のリソース量算出」）。なお、プレコンパイルせず、１回目測定前のプレコンパイルでのリソース量の和でもよい。このようにすれば、プレコンパイル回数を削減することができる。
ステップＳ２１４で実行ファイル作成部１１９は、組み合わせたオフロードOpenCLをコンパイルする。
ステップＳ２１５で性能測定部１１８は、コンパイルされたプログラムの性能を測定する（「２回目の性能測定」）。性能測定部１１８は、選択された単ループ文に対してコンパイルして測定し、更に高速化できた単ループ文に対してはその組み合わせパターンも作り２回目の性能測定を行う（詳細については、図１５のサブルーチン参照）。
ステップＳ２１６で本番環境配置部１２０は、１回目と２回目の測定の中で最高性能のパターンを選択して本フローの処理を終了する。測定された複数パターンの中で、短時間のパターンを解として選択する。
このように、ループ文のＦＰＧＡ自動オフロードは、算術強度とループ回数が高くリソース効率が高いループ文に絞って、オフロードパターンを作り、検証環境で実測を通じて高速なパターンの探索を行う（図１４参照）。
図１５は、性能測定部１１８の性能・電力使用量測定処理を示すフローチャートである。本フローは、図１４のステップＳ２１０またはステップＳ２１５のサブルーチンコールにより呼び出され、実行される。
ステップＳ３０１で、性能測定部１１８は、ＦＰＧＡオフロード時に必要となる処理時間を測定する。
ステップＳ３０２で、性能測定部１１８は、測定した処理時間をもとに評価値を設定する。
ステップＳ３０３で、性能測定部１１８は、評価値が高い個体ほど適合度が高くなるように評価された評価値の高いパターンの性能を測定し、図１４のステップＳ２１０またはステップＳ２１５に戻る。
［オフロードパターンの作成例］
図１６は、ＰＬＤ処理パターン作成部２１５の探索イメージを示す図である。
制御部（自動オフロード機能部）２１（図１３参照）は、ユーザが利用するアプリケーションコード（Application code）１２５（図２参照）を分析し、図１６に示すように、アプリケーションコード１２５のコードパターン（Code patterns）２４１からfor文の並列可否をチェックする。図１６の符号ｒに示すように、コードパターン２４１から４つのfor文が見つかった場合、各for文に対してそれぞれ１桁、ここでは４つのfor文に対し４桁の１または０を割り当てる。ここでは、ＦＰＧＡ処理する場合は１、ＦＰＧＡ処理しない場合（すなわちＣＰＵで処理する場合）は０とする。
［ＣコードからOpenCL最終解の探索までの流れ］
図１７の手順Ａ－Ｆは、ＣコードからOpenCL最終解の探索までの流れを説明する図である。
アプリケーションコード分析部１１２（図１３参照）は、図１７の手順Ａに示す「Ｃコード」を構文解析し（<構文解析>：図１７の符号ｓ参照）、ＰＬＤ処理指定部２１３（図１３参照）は、図１７の手順Ｂに示す「ループ文、変数情報」を特定する（図１７の符号ｔ参照）。
算術強度算出部２１４（図１３参照）は、特定した「ループ文、変数情報」に対して、算術強度分析ツールを用いて算術強度分析（Arithmetic Intensity analysis）する（図１７の符号ｕ参照）。ＰＬＤ処理パターン作成部２１５は、算術強度が高いループ文をオフロード候補に絞る。さらに、ＰＬＤ処理パターン作成部２１５は、プロファイリングツールを用いてプロファイリング分析（Profiling analysis）を行って、高算術強度で高ループ回数のループ文をさらに絞り込む。
そして、ＰＬＤ処理パターン作成部２１５は、絞り込まれた各ループ文をＦＰＧＡにオフロードするためのOpenCLを作成（OpenCL化）する（図１７の符号ｖ参照）。
さらに、OpenCL化時にコード分割と共に展開等の高速化手法を導入する（後記）。
<「高算術強度，OpenCL化」具体例（その１）：手順Ｃ>
例えば、アプリケーションコード１３０のコードパターン２４１（図１６参照）から４つのfor文（４桁の１または０の割り当て）が見つかった場合、算術強度分析で３つが絞り込まれる（選ばれる）。すなわち、図１７の符号ｕに示すように、４つのfor文から、３つのfor文のオフロードパターン「1000」「0010」「0001」が絞り込まれる。
<OpenCL化時にコード分割と共に実行する「展開」例>
ＦＰＧＡからＣＰＵへのデータ転送する場合の、ＣＰＵプログラム側で記述されるループ文〔k=０; k<１０; k++〕 {
}
において、このループ文の上部に、＼pragma unrollを指示する。すなわち、
＼pragma unroll
for(k=０; k<１０; k++){
}
と記述する。
＼pragma unroll等のIntelやXilinx（登録商標）のツールに合った文法でunrollを指示すると、上記展開例であれば、i=０,i=１,i=２と展開してパイプライン実行することができる。このため、リソース量は１０倍使うことになるが、高速になる場合がある。
また、unrollで展開する数は全ループ回数個でなく５個に展開等の指定もでき、その場合は、ループ２回ずつが、５つに展開される。
以上で、「展開」例についての説明を終える。
次に、ＰＬＤ処理パターン作成部２１５は、オフロード候補として絞り込まれた高算術強度のループ文を、リソース量を用いてさらに絞り込む。すなわち、ＰＬＤ処理パターン作成部２１５は、リソース量を算出し、ＰＬＤ処理パターン作成部２１５は、高算術強度のループ文のオフロード候補の中から、リソース効率（＝算術強度／ＦＰＧＡ処理時のリソース量、または、算術強度×ループ回数／ＦＰＧＡ処理時のリソース量）分析して、リソース効率の高いループ文を抽出する。
図１７の符号ｖでは、ＰＬＤ処理パターン作成部２１５は、絞り込んだループ文をオフロードするためのOpenCLをコンパイル（<プレコンパイル>）する。
<「高算術強度，OpenCL化」具体例（その２）>
図１７の符号ｕに示すように、算術強度分析で絞り込まれた４つのオフロードパターン「1000」「0100」「0010」「0001」の中から、上記リソース効率分析により３つのオフロードパターン「1000」「0010」「0001」に絞り込む。
以上、図１７の手順Ｃに示す「高算術強度，OpenCL化」について説明した。
図１７の手順Ｄに示す「リソース効率の高いループ文」に対して、性能測定部１１８は、コンパイルされたプログラムの性能を測定する（「１回目の性能測定」）。
そして、ＰＬＤ処理パターン作成部２１５は、性能測定された中でＣＰＵに比べ高性能化されたループ文をリスト化する。以下、同様に、リソース量を算出、オフロードOpenCLコンパイル、コンパイルされたプログラムの性能を測定する。
<「高算術強度，OpenCL化」具体例（その３）>
図１７の符号ｗに示すように、３つのオフロードパターン「1000」「0010」「0001」について１回目測定を行う。その３つの測定の中で、「1000」「0010」の２つの性能が高くなったとすると、「1000」と「0010」の組合せについて２回目測定を行う。
図１７の符号ｘでは、実行ファイル作成部１１９は、絞り込んだループ文をオフロードするためのOpenCLをコンパイル（<本コンパイル>）する。
図１７の手順Ｅに示す「組合せパターン実測」は、候補ループ文単体、その後、その組合せで検証パターン測定することをいう。
<「高算術強度，OpenCL化」具体例（その４）>
図１７の符号ｙに示すように、「1000」と「0010」の組合せである「1010」について２回目測定する。２回測定し、その結果、１回目測定と２回目測定の中で最高速度の「0010」が選択された。このような場合、「0010」が最終の解となる。ここで、組合せパターンがリソース量制限のため測定できない場合がある。この場合、組合せについてはスキップして、単体の結果から最高速度のものを選ぶだけでもよい。
図１７の符号ｚでは、性能測定部１１８は、１回目測定と２回目測定の中で最高速度の良い「0010」を選択（<選択>）する。
以上により、図１７の手順Ｆに示す「OpenCL最終解」の「0010」（図１７の符号ａａ参照）が選択された。
<デプロイ（配置）>
OpenCL最終解の、最高処理性能のＰＬＤ処理パターンで、本番環境に改めてデプロイして、ユーザに提供する。
［実装例］
実装例を説明する。
ＦＰＧＡはIntel PAC with Intel Arria10 GX ＦＰＧＡ等が利用できる。
ＦＰＧＡ処理は、Intel Acceleration Stack（Intel FPGA SDK for OpenCL、Quartus Prime Version）等が利用できる。
Intel FPGA SDK for OpenCLは、標準OpenCLに加え、Intel向けの#pragma等を解釈する高位合成ツール（HLS）である。
実装例では、ＦＰＧＡで処理するカーネルとＣＰＵで処理するホストプログラムを記述したOpenCLコードを解釈し、リソース量等の情報を出力し、ＦＰＧＡの配線作業等を行い、ＦＰＧＡで動作できるようにする。ＦＰＧＡ実機で動作できるようにするには、１００行程度の小プログラムでも３時間程の長時間がかかる。ただし、リソース量オーバーの際は、早めにエラーとなる。また、ＦＰＧＡで処理できないOpenCLコードの際は、数時間後にエラーを出力する。
実装例では、C/C++アプリケーションの利用依頼があると、まず、C/C++アプリケーションのコードを解析して、for文を発見するとともに、for文内で使われる変数データ等のプログラム構造を把握する。構文解析には、LLVM/Clangの構文解析ライブラリ等が利用できる。
実装例では、次に、各ループ文のＦＰＧＡオフロード効果があるかの見込みを得るため、算術強度分析ツールを実行し、計算数、アクセス数等で定まる算術強度の指標を取得する。算術強度分析には、ROSEフレームワーク等が利用できる。算術強度上位個のループ文のみ対象とするようにする。
次に、gcov等のプロファイリングツールを用いて、各ループのループ回数を取得する。算術強度×ループ回数が上位a個のループ文を候補に絞る。
実装例では、次に、高算術強度の個々のループ文に対して、ＦＰＧＡオフロードするOpenCLコードを生成する。OpenCLコードは、該当ループ文をＦＰＧＡカーネルとして、残りをＣＰＵホストプログラムとして分割したものである。ＦＰＧＡカーネルコードとする際に、高速化の技法としてループ文の展開処理を一定数ｂだけ行ってもよい。ループ文展開処理は、リソース量は増えるが、高速化に効果がある。そこで、展開する数は、一定数ｂに制限してリソース量が膨大にならない範囲で行う。
実装例では、次に、ａ個のOpenCLコードに対して、Intel ＦＰＧＡ SDK for OpenCLを用いて、プレコンパイルをして、利用するFlip Flop、Look Up Table等のリソース量を算出する。使用リソース量は、全体リソース量の割合で表示される。ここで、算術強度とリソース量または算術強度とループ回数とリソース量から、各ループ文のリソース効率を計算する。例えば、算術強度が１０、リソース量が０.５のループ文は、１０／０.５＝２０、算術強度が３、リソース量が０.３のループ文は３／０.３＝１０がリソース効率となり、前者が高い。また、ループ回数をかけた値をリソース効率としてもよい。各ループ文で、リソース効率が高いｃ個を選定する。
実装例では、次に、ｃ個のループ文を候補に、実測するパターンを作る。例えば、１番目と３番目のループが高リソース効率であった場合、１番をオフロード、３番をオフロードする各OpenCLパターンを作成して、コンパイルして性能測定する。複数の単ループ文のオフロードパターンで高速化できている場合（例えば、１番と３番両方が高速化できている場合）は、その組合せでのOpenCLパターンを作成して、コンパイルして性能測定する（例えば１番と３番両方をオフロードするパターン）。
なお、単ループの組み合わせを作る際は、利用リソース量も組み合わせになる。このため、上限値に納まらない場合は、その組合せパターンは作らない。組合せも含めてｄ個のパターンを作成した場合、検証環境のＦＰＧＡを備えたサーバで性能測定を行う。性能測定には、高速化したいアプリケーションで指定されたサンプル処理を行う。例えば、フーリエ変換のアプリケーションであれば、サンプルデータでの変換処理をベンチマークに性能測定をする。
実装例では、最後に、複数の測定パターンの高速なパターンを解として選択する。
第２実施形態でも第１実施形態で述べたと同様の「リソース量決定と配置決定」を実行する（説明省略）。
［評価］
評価を説明する。
第２実施形態の［ループ文のＦＰＧＡ自動オフロード］では、第１実施形態の［ループ文のＧＰＵ自動オフロード］と同様に評価できる。
<評価対象>
評価対象は、第２実施形態の［ループ文のＦＰＧＡ自動オフロード］では、ＭＲＩ（Magnetic Resonance Imaging）画像処理のMRI-Qとする。
MRI-Qは、非デカルト空間の３次元ＭＲＩ再構成アルゴリズムで使用されるスキャナー構成を表す行列Ｑを計算する。MRI-Qは、Ｃ言語で記述されており、性能測定中に３次元ＭＲＩ画像処理を実行し、Large(最大)の６４×６４×６４サイズのデータで処理時間を測定する。ＣＰＵ処理は、Ｃ言語を用い、ＦＰＧＡ処理はOpenCL に基づき処理される。
<評価手法>
対象となるアプリケーションのコードを入力し、移行先のＧＰＵやＦＰＧＡに対して、Clang等で認識されたループ文オフロードを試行してオフロードパターンを決める。この際に、処理時間と電力使用量を測定する。最終オフロードパターンについて、電力使用量の時間変化を取得し、全てＣＰＵで処理する場合に比べた低電力化を確認する。
第２実施形態の［ループ文のＦＰＧＡ自動オフロード］では、ＧＡは行わず、算術強度等を用いて、測定パターンが４パターンとなるまで絞り込む。
オフロード対象ループ文： MRI-Q 16
パターン適合度：処理時間が低い程、評価値が高くなり、高適合度になる。第２の実施形態のMRI-Qでも前述の図１２のような形で、単純に安さ優先や応答時間優先の配置に比べて、コストや応答時間が改善できる。
［ハードウェア構成］
第１および第２の実施形態に係るオフロードサーバは、例えば図１８に示すような構成の物理装置であるコンピュータ９００によって実現される。
図１８は、オフロードサーバ１，１Ａの機能を実現するコンピュータの一例を示すハードウェア構成図である。コンピュータ９００は、ＣＰＵ９０１、ＲＡＭ９０２、ＲＯＭ９０３、ＨＤＤ９０４、アクセラレータ９０５、入出力インターフェイス（Ｉ／Ｆ）９０６、メディアインターフェイス（Ｉ／Ｆ）９０７、および通信インターフェイス（Ｉ／Ｆ：Interface）９０８を有する。
アクセラレータ９０５は、通信Ｉ／Ｆ９０８からのデータ、または、ＲＡＭ９０２からのデータの少なくとも一方のデータを高速に処理するアクセラレータ（デバイス）である。例えば、アクセラレータ９０５は、図２の各種デバイス（Device）１５１、ＣＰＵ-ＧＰＵを有する装置１５２、ＣＰＵ-ＦＰＧＡを有する装置１５３、ＣＰＵを有する装置１５４のアクセラレータである。
なお、アクセラレータ９０５として、ＣＰＵ９０１またはＲＡＭ９０２からの処理を実行した後にＣＰＵ９０１またはＲＡＭ９０２に実行結果を戻すタイプ（look-aside型）を用いてもよい。一方、アクセラレータ９０５として、通信Ｉ／Ｆ９０８とＣＰＵ９０１またはＲＡＭ９０２との間に入って、処理を行うタイプ（in-line型）を用いてもよい。
アクセラレータ９０５は、通信Ｉ／Ｆ９０８を介して外部装置９１５と接続される。入出力Ｉ／Ｆ９０６は、入出力装置９１６と接続される。メディアＩ／Ｆ９０７は、記録媒体９１７からデータを読み書きする。
ＣＰＵ９０１は、ＲＯＭ９０３またはＨＤＤ９０４に記憶されたプログラムに基づき作動し、ＲＡＭ９０２に読み込んだプログラム（アプリケーションや、その略のアプリとも呼ばれる）を実行することにより、図１、図１３に示すオフロードサーバ１，１Ａの各処理部による制御を行う。そして、このプログラムは、通信回線を介して配布したり、ＣＤ－ＲＯＭ等の記録媒体９１７に記録して配布したりすることも可能である。
ＲＯＭ９０３は、コンピュータ９００の起動時にＣＰＵ９０１によって実行されるブートプログラムや、コンピュータ９００のハードウェアに依存するプログラム等を格納する。
ＣＰＵ９０１は、入出力Ｉ／Ｆ９０６を介して、マウスやキーボード等の入力部、および、ディスプレイやプリンタ等の出力部からなる入出力装置９１６を制御する。ＣＰＵ９０１は、入出力Ｉ／Ｆ９０６を介して、入出力装置９１６からデータを取得するともに、生成したデータを入出力装置９１６へ出力する。なお、プロセッサとしてＣＰＵ９０１とともに、ＧＰＵ（Graphics Processing Unit）等を用いてもよい。
ＨＤＤ９０４は、ＣＰＵ９０１により実行されるプログラムおよび当該プログラムによって使用されるデータ等を記憶する。通信Ｉ／Ｆ９０８は、通信網（例えば、ＮＷ（Network））を介して他の装置からデータを受信してＣＰＵ９０１へ出力し、また、ＣＰＵ９０１が生成したデータを、通信網を介して他の装置へ送信する。
メディアＩ／Ｆ９０７は、記録媒体９１７に格納されたプログラムまたはデータを読み取り、ＲＡＭ９０２を介してＣＰＵ９０１へ出力する。ＣＰＵ９０１は、目的の処理に係るプログラムを、メディアＩ／Ｆ９０７を介して記録媒体９１７からＲＡＭ９０２上にロードし、ロードしたプログラムを実行する。記録媒体９１７は、ＤＶＤ（Digital Versatile Disc）、ＰＤ（Phase change rewritable Disk）等の光学記録媒体、ＭＯ（Magneto Optical disk）等の光磁気記録媒体、磁気記録媒体、導体メモリテープ媒体又は半導体メモリ等である。
例えば、コンピュータ９００が第１および第２の実施形態に係るオフロードサーバ１，１Ａとして機能する場合、コンピュータ９００のＣＰＵ９０１は、ＲＡＭ９０２上にロードされたプログラムを実行することによりオフロードサーバ１，１Ａの機能を実現する。また、ＨＤＤ９０４には、ＲＡＭ９０２内のデータが記憶される。ＣＰＵ９０１は、目的の処理に係るプログラムを記録媒体９１２から読み取って実行する。この他、ＣＰＵ９０１は、他の装置から通信網を介して目的の処理に係るプログラムを読み込んでもよい。
［効果］
以上説明したように、第１実施形態に係るオフロードサーバ１（図１参照）は、アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部１１２と、アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部１１３と、アプリケーションプログラムのループ文を特定し、特定した各ループ文に対して、アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部１１４と、コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部１１７と、並列処理パターンのアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部１１８と、変換したアプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部１７０と、を備える。
このようにすることにより、ＧＰＵ、ＦＰＧＡ等のアクセラレータに自動オフロードしたアプリケーションに対して、ユーザのリクエストする価格条件や応答時間条件、アプリケーションの配置数等を変更して、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する適正配置を計算する。これにより、変換したアプリケーションを計算リソースのコストまたは応答時間の要求を満たして、ユーザ要望に従った最適配置を実現することができる。
第２実施形態に係るオフロードサーバ１Ａ（図１３参照）は、アプリケーションプログラムの特定処理をＰＬＤにオフロードするオフロードサーバであって、アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部１１２と、アプリケーションプログラムのループ文を特定し、特定した各ループ文に対して、ＰＬＤにおけるパイプライン処理および並列処理をOpenCLで指定した複数のオフロード処理パターンにより作成してコンパイルするＰＬＤ処理指定部２１３と、アプリケーションプログラムのループ文の算術強度を算出する算術強度算出部２１４と、算術強度算出部２１４が算出した算術強度をもとに、算術強度が所定の閾値より高いループ文をオフロード候補として絞り込み、ＰＬＤ処理パターンを作成するＰＬＤ処理パターン作成部２１５と、作成されたＰＬＤ処理パターンのアプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、ＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部１１８と、変換したアプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部１７０と、を備える。
このようにすることにより、実際に性能測定するパターンを絞ってから検証環境に配置し、コンパイルしてＰＬＤ（例えば、ＦＰＧＡ）実機で性能測定することで、性能測定する回数を減らすことができる。これにより、ＰＬＤへの自動オフロードにおいて、アプリケーションのループ文の自動オフロードを高速で行うことができる。そして、変換したアプリケーションに対して、ユーザのリクエストする価格条件や応答時間条件、アプリケーションの配置数等を変更して、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する適正配置を計算する。これにより、変換したアプリケーションについて、計算リソースのコストまたは応答時間の要求を満たして、ユーザ要望に従った最適配置を実現することができる。
第１および第２実施形態に係るオフロードサーバ１，１Ａにおいて、配置設定部１７０は、サーバにアプリケーションプログラムを配置した際に、計算リソースのコストを極小化する配置、または、応答時間を極小化する配置を計算することを特徴とする。
このようにすることにより、変換したアプリケーションを計算リソースのコストまたは応答時間の要求を満たして最適に配置することができる。
第１および第２実施形態に係るオフロードサーバ１，１Ａにおいて、配置設定部１７０は、サーバにアプリケーションプログラムを配置した際に、計算リソースのコストを極小化する配置を［数１］に示す線形計画式に従って計算することを特徴とする。
このようにすることにより、例えば、一月いくら以内での配置が必要な要求の場合は、［数１］の式（１）の応答時間の最小化が目的関数となり、［数１］の式（２）の計算リソースのコストがいくら以内であるかが制約条件の一つとなる。よって、変換したアプリケーションを計算リソースのコストの要求を満たして最適に配置することができる。
第１および第２実施形態に係るオフロードサーバ１，１Ａにおいて、配置設定部１７０は、サーバにアプリケーションプログラムを配置した際に、応答時間を極小化する配置を［数２］に示す線形計画式に従って計算することを特徴とする。
このようにすることにより、例えば、アプリケーションの応答時間が何秒以内での配置が必要な要求の場合は、［数２］の式（５）のコストの最小化が目的関数となり、［数２］の式（６）の応答時間が何秒以内であるかが制約条件の一つとなる。よって、変換したアプリケーションをユーザの応答時間の要求を満たして最適に配置することができる。
本発明は、コンピュータを、上記オフロードサーバとして機能させるためのオフロードプログラムとした。
このようにすることにより、一般的なコンピュータを用いて、上記オフロードサーバ１，１Ａの各機能を実現させることができる。
また、上記各実施形態において説明した各処理のうち、自動的に行われるものとして説明した処理の全部又は一部を手作業で行うこともでき、あるいは、手作業で行われるものとして説明した処理の全部又は一部を公知の方法で自動的に行うこともできる。この他、上述文書中や図面中に示した処理手順、制御手順、具体的名称、各種のデータやパラメータを含む情報については、特記する場合を除いて任意に変更することができる。
また、図示した各装置の各構成要素は機能概念的なものであり、必ずしも物理的に図示の如く構成されていることを要しない。すなわち、各装置の分散・統合の具体的形態は図示のものに限られず、その全部又は一部を、各種の負荷や使用状況などに応じて、任意の単位で機能的又は物理的に分散・統合して構成することができる。
また、上記の各構成、機能、処理部、処理手段等は、それらの一部又は全部を、例えば集積回路で設計する等によりハードウェアで実現してもよい。また、上記の各構成、機能等は、プロセッサがそれぞれの機能を実現するプログラムを解釈し、実行するためのソフトウェアで実現してもよい。各機能を実現するプログラム、テーブル、ファイル等の情報は、メモリや、ハードディスク、ＳＳＤ（Solid State Drive）等の記録装置、又は、ＩＣ（Integrated Circuit）カード、ＳＤ（Secure Digital）カード、光ディスク等の記録媒体に保持することができる。
また、本実施形態では、組合せ最適化問題を、限られた最適化期間中に解を発見できるようにするため、遺伝的アルゴリズム（ＧＡ）の手法を用いているが、最適化の手法はどのようなものでもよい。例えば、local search（局所探索法）、Dynamic Programming（動的計画法）、これらの組み合わせでもよい。
また、本実施形態では、C/C++向けOpenＡＣＣコンパイラを用いているが、ＧＰＵ処理をオフロードできるものであればどのようなものでもよい。例えば、Java lambda（登録商標） GPU処理、IBM Java 9 SDK（登録商標）でもよい。なお、並列処理指定文は、これらの開発環境に依存する。
例えば、Java（登録商標）では、Java 8よりlambda形式での並列処理記述が可能である。IBM（登録商標）は、lambda形式の並列処理記述を、ＧＰＵにオフロードするJITコンパイラを提供している。Javaでは、これらを用いて、ループ処理をlambda形式にするか否かのチューニングをＧＡで行うことで、同様のオフロードが可能である。
また、本実施形態では、繰り返し文（ループ文）として、for文を例示したが、for文以外のwhile文やdo-while文も含まれる。ただし、ループの継続条件等を指定するfor文がより適している。
１，１Ａ オフロードサーバ
１１，２１ 制御部
１２ 入出力部
１３ 記憶部
１４ 検証用マシン (アクセラレータ検証用装置)
１１１ アプリケーションコード指定部
１１２ アプリケーションコード分析部
１１３ データ転送指定部
１１４ 並列処理指定部
１１４ａ，２１３ａ オフロード範囲抽出部
１１４ｂ，２１３ｂ 中間言語ファイル出力部
１１５ リソース比決定部
１１６ リソース量設定部
１１７ 並列処理パターン作成部
１１８ 性能測定部
１１８ａ バイナリファイル配置部
１１９ 実行ファイル作成部
１２０ 本番環境配置部
１２１ 性能測定テスト抽出実行部
１２２ ユーザ提供部
１２５ アプリケーションコード
１３１ テストケースＤＢ
１３２ 設備リソースＤＢ
１３３ 中間言語ファイル
１５１ 各種デバイス
１５２ ＣＰＵ-ＧＰＵを有する装置
１５３ ＣＰＵ-ＦＰＧＡを有する装置
１５４ ＣＰＵを有する装置
１７０ 配置設定部
２１３ ＰＬＤ処理指定部
２１４ 算術強度算出部
２１５ ＰＬＤ処理パターン作成部
９０５ アクセラレータ

Claims (9)

  1. アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、
    前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、
    前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部と、
    コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、
    前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部と、を備え、
    前記配置設定部は、サーバにアプリケーションプログラムを配置した際に、応答時間を極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロードサーバ。
  2. アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバであって、
    前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、
    前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うデータ転送指定部と、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルする並列処理指定部と、
    コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成する並列処理パターン作成部と、
    前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行する性能測定部と、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部と、を備え、
    前記配置設定部は、サーバにアプリケーションプログラムを配置した際に、計算リソースのコストを極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロードサーバ。
  3. アプリケーションプログラムの特定処理をＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバであって、
    前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記ＰＬＤにおけるパイプライン処理および並列処理をOpenCLで指定した複数のオフロード処理パターンにより作成してコンパイルするＰＬＤ処理指定部と、
    前記アプリケーションプログラムのループ文の算術強度を算出する算術強度算出部と、
    前記算術強度算出部が算出した算術強度をもとに、前記算術強度が所定の閾値より高いループ文をオフロード候補として絞り込み、ＰＬＤ処理パターンを作成するＰＬＤ処理パターン作成部と、
    作成された前記ＰＬＤ処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記ＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部と、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部と、を備え、
    前記配置設定部は、サーバにアプリケーションプログラムを配置した際に、応答時間を極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロードサーバ。
  4. アプリケーションプログラムの特定処理をＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバであって、
    前記アプリケーションプログラムのソースコードを分析するアプリケーションコード分析部と、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記ＰＬＤにおけるパイプライン処理および並列処理をOpenCLで指定した複数のオフロード処理パターンにより作成してコンパイルするＰＬＤ処理指定部と、
    前記アプリケーションプログラムのループ文の算術強度を算出する算術強度算出部と、
    前記算術強度算出部が算出した算術強度をもとに、前記算術強度が所定の閾値より高いループ文をオフロード候補として絞り込み、ＰＬＤ処理パターンを作成するＰＬＤ処理パターン作成部と、
    作成された前記ＰＬＤ処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記ＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部と、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定する配置設定部と、を備え、
    前記配置設定部は、サーバにアプリケーションプログラムを配置した際に、計算リソースのコストを極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロードサーバ。
  5. アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバのオフロード制御方法であって、
    前記オフロードサーバは、
    前記アプリケーションプログラムのソースコードを分析するステップと、
    前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うステップと、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルするステップと、
    コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成するステップと、
    前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行するステップと、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定するステップと、を実行し、
    前記アプリケーションプログラムの配置場所を計算して設定するステップにおいて、サーバにアプリケーションプログラムを配置した際に、応答時間を極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロード制御方法。
  6. アプリケーションプログラムの特定処理をアクセラレータにオフロードするオフロードサーバのオフロード制御方法であって、
    前記オフロードサーバは、
    前記アプリケーションプログラムのソースコードを分析するステップと、
    前記アプリケーションプログラムのループ文の中で用いられる変数の参照関係を分析し、ループ外でデータ転送してよいデータについては、ループ外でのデータ転送を明示的に指定する明示的指定行を用いたデータ転送指定を行うステップと、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記アクセラレータにおける並列処理指定文を指定してコンパイルするステップと、
    コンパイルエラーが出るループ文に対して、オフロード対象外とするとともに、コンパイルエラーが出ないループ文に対して、並列処理するかしないかの指定を行う並列処理パターンを作成するステップと、
    前記並列処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記アクセラレータにオフロードした際の性能測定用処理を実行するステップと、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定するステップと、を実行し、
    前記アプリケーションプログラムの配置場所を計算して設定するステップにおいて、サーバにアプリケーションプログラムを配置した際に、計算リソースのコストを極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロード制御方法。
  7. アプリケーションプログラムの特定処理をＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバのオフロード制御方法であって、
    前記オフロードサーバは、
    前記アプリケーションプログラムのソースコードを分析するステップと、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記ＰＬＤにおけるパイプライン処理および並列処理をOpenCLで指定した複数のオフロード処理パターンにより作成してコンパイルするステップと、
    前記アプリケーションプログラムのループ文の算術強度を算出するステップと、
    算出した算術強度をもとに、前記算術強度が所定の閾値より高いループ文をオフロード候補として絞り込み、ＰＬＤ処理パターンを作成するステップと、
    作成された前記ＰＬＤ処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記ＰＬＤにオフロードした際の性能測定用処理を実行するステップと、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定するステップと、を実行し、
    前記アプリケーションプログラムの配置場所を計算して設定するステップにおいて、サーバにアプリケーションプログラムを配置した際に、応答時間を極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロード制御方法。
  8. アプリケーションプログラムの特定処理をＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバのオフロード制御方法であって、
    前記オフロードサーバは、
    前記アプリケーションプログラムのソースコードを分析するステップと、
    前記アプリケーションプログラムのループ文を特定し、特定した各前記ループ文に対して、前記ＰＬＤにおけるパイプライン処理および並列処理をOpenCLで指定した複数のオフロード処理パターンにより作成してコンパイルするステップと、
    前記アプリケーションプログラムのループ文の算術強度を算出するステップと、
    算出した算術強度をもとに、前記算術強度が所定の閾値より高いループ文をオフロード候補として絞り込み、ＰＬＤ処理パターンを作成するステップと、
    作成された前記ＰＬＤ処理パターンの前記アプリケーションプログラムをコンパイルして、アクセラレータ検証用装置に配置し、前記ＰＬＤにオフロードした際の性能測定用処理を実行するステップと、
    変換した前記アプリケーションプログラムを、ユーザの指定するコストまたは応答時間の条件に応じて、ネットワーク上の、クラウドサーバ、キャリアエッジサーバ、ユーザエッジサーバのいずれかに配置する際、デバイスおよびリンクのコスト、計算リソース上限、帯域上限を制約条件とし、かつ計算リソースのコストまたは応答時間を目的関数とした線形計画式に基づいて、アプリケーションプログラムの配置場所を計算して設定するステップと、を実行し、
    前記アプリケーションプログラムの配置場所を計算して設定するステップにおいて、サーバにアプリケーションプログラムを配置した際に、計算リソースのコストを極小化する配置を次式に示す線形計画式に従って計算する
    ことを特徴とするオフロード制御方法。
  9. コンピュータを、請求項１乃至請求項４のいずれか一項に記載のオフロードサーバとして機能させるためのオフロードプログラム。
JP2023576454A 2022-01-26 2022-01-26 オフロードサーバ、オフロード制御方法およびオフロードプログラム Active JP7716632B2 (ja)

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
PCT/JP2022/002880 WO2023144926A1 (ja) 2022-01-26 2022-01-26 オフロードサーバ、オフロード制御方法およびオフロードプログラム

Publications (2)

Publication Number Publication Date
JPWO2023144926A1 JPWO2023144926A1 (ja) 2023-08-03
JP7716632B2 true JP7716632B2 (ja) 2025-08-01

Family

ID=87471201

Family Applications (1)

Application Number Title Priority Date Filing Date
JP2023576454A Active JP7716632B2 (ja) 2022-01-26 2022-01-26 オフロードサーバ、オフロード制御方法およびオフロードプログラム

Country Status (2)

Country Link
JP (1) JP7716632B2 (ja)
WO (1) WO2023144926A1 (ja)

Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111405569A (zh) 2020-03-19 2020-07-10 三峡大学 基于深度强化学习的计算卸载和资源分配方法及装置
WO2020175411A1 (ja) 2019-02-25 2020-09-03 日本電信電話株式会社 アプリケーション配置装置及びアプリケーション配置プログラム
WO2021156955A1 (ja) 2020-02-04 2021-08-12 日本電信電話株式会社 オフロードサーバ、オフロード制御方法およびオフロードプログラム
WO2021156956A1 (ja) 2020-02-04 2021-08-12 日本電信電話株式会社 オフロードサーバ、オフロード制御方法およびオフロードプログラム

Patent Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
WO2020175411A1 (ja) 2019-02-25 2020-09-03 日本電信電話株式会社 アプリケーション配置装置及びアプリケーション配置プログラム
WO2021156955A1 (ja) 2020-02-04 2021-08-12 日本電信電話株式会社 オフロードサーバ、オフロード制御方法およびオフロードプログラム
WO2021156956A1 (ja) 2020-02-04 2021-08-12 日本電信電話株式会社 オフロードサーバ、オフロード制御方法およびオフロードプログラム
CN111405569A (zh) 2020-03-19 2020-07-10 三峡大学 基于深度强化学习的计算卸载和资源分配方法及装置

Non-Patent Citations (1)

* Cited by examiner, † Cited by third party
Title
山登庸次，「環境適応アプリケーション配置適正化の初期検討」，engrxiv.org [online]，2021年07月11日，[取得日 2025年2月28日], 取得先 <https://engrxiv.org/preprint/view/1770/version/2691>，<DOI: https://doi.org/10.31224/osf.io/vdn8z>

Also Published As

Publication number Publication date
WO2023144926A1 (ja) 2023-08-03
JPWO2023144926A1 (ja) 2023-08-03

Similar Documents

Publication Publication Date Title
JP7063289B2 (ja) オフロードサーバのソフトウェア最適配置方法およびプログラム
JP6927424B2 (ja) オフロードサーバおよびオフロードプログラム
JP7632712B2 (ja) オフロードサーバ、および、オフロード制御方法
US12530296B2 (en) Systems, apparatus, articles of manufacture, and methods for improved data transfer for heterogeneous programs
Yamato Proposal and evaluation of GPU offloading parts reconfiguration during applications operations for environment adaptation
JP7322978B2 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
US12056475B2 (en) Offload server, offload control method, and offload program
US11947975B2 (en) Offload server, offload control method, and offload program
JP7716632B2 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
JP7806893B2 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
JP7544142B2 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
US12033235B2 (en) Offload server, offload control method, and offload program
JP7662037B2 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
JP7521597B2 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
JP7184180B2 (ja) オフロードサーバおよびオフロードプログラム
WO2024147197A1 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム
WO2024079886A1 (ja) オフロードサーバ、オフロード制御方法およびオフロードプログラム

Legal Events

Date Code Title Description
A621 Written request for application examination

Free format text: JAPANESE INTERMEDIATE CODE: A621

Effective date: 20240528

A131 Notification of reasons for refusal

Free format text: JAPANESE INTERMEDIATE CODE: A131

Effective date: 20250311

A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20250508

TRDD Decision of grant or rejection written
A01 Written decision to grant a patent or to grant a registration (utility model)

Free format text: JAPANESE INTERMEDIATE CODE: A01

Effective date: 20250617

A61 First payment of annual fees (during grant procedure)

Free format text: JAPANESE INTERMEDIATE CODE: A61

Effective date: 20250630

R150 Certificate of patent or registration of utility model

Ref document number: 7716632

Country of ref document: JP

Free format text: JAPANESE INTERMEDIATE CODE: R150