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

Description

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

ＣＰＵ（Central Processing Unit）以外のヘテロな計算リソースを用いることが増えている。例えば、ＧＰＵ（Graphics Processing Unit）(アクセラレータ)を強化したサーバで画像処理を行ったり、ＦＰＧＡ（アクセラレータ）で信号処理をアクセラレートすることが始まっている。ＦＰＧＡは、製造後に設計者等が構成を設定できるプログラム可能なゲートアレイであり、ＰＬＤ（Programmable Logic Device）の一種である。Amazon Web Services (AWS)（登録商標）では、ＧＰＵインスタンス、ＦＰＧＡインスタンスが提供されており、オンデマンドにそれらリソースを使うこともできる。Microsoft（登録商標）は、ＦＰＧＡを用いて検索を効率化している。

OpenＩｏＴ（Internet of Things）環境では、サービス連携技術等を用いて、多彩なアプリケーションの創出が期待されるが、更に進歩したハードウェアを生かすことで、動作アプリケーションの高性能化が期待できる。しかし、そのためには、動作させるハードウェアに合わせたプログラミングや設定が必要である。例えば、ＣＵＤＡ（Compute Unified Device Architecture）、 OpenＣＬ（Open Computing Language）といった多くの技術知識が求められ、ハードルは高い。OpenＣＬは、あらゆる計算資源（ＣＰＵやＧＰＵに限らない）を特定のハードに縛られず統一的に扱えるオープンなＡＰＩ（Application Programming Interface）である。

ＧＰＵやＦＰＧＡをユーザのＩｏＴアプリケーションで容易に利用できるようにするため下記が求められる。すなわち、動作させる画像処理、暗号処理等の汎用アプリケーションをOpenＩｏＴ環境にデプロイする際に、OpenＩｏＴのプラットフォームがアプリケーションロジックを分析し、ＧＰＵ、ＦＰＧＡに自動で処理をオフロードすることが望まれる。

ＧＰＵの計算能力を画像処理以外にも使うＧＰＧＰＵ（General Purpose GPU）のための開発環境ＣＵＤＡが発展している。ＣＵＤＡは、ＧＰＧＰＵ向けの開発環境である。また、ＧＰＵ、ＦＰＧＡ、メニーコアＣＰＵ等のヘテロハードウェアを統一的に扱うための標準規格としてOpenＣＬも登場している。

ＣＵＤＡやOpenＣＬでは、Ｃ言語の拡張によるプログラミングを行う。ただし、ＧＰＵ等のデバイスとＣＰＵの間のメモリコピー、解放等を記述する必要があり、記述の難度は高い。実際に、ＣＵＤＡやOpenＣＬを使いこなせる技術者は数多くはいない。

簡易にＧＰＧＰＵを行うため、ディレクティブベースで、ループ文等の並列処理すべき個所を指定し、ディレクティブに従いコンパイラがデバイス向けコードに変換する技術がある。技術仕様としてOpenACC（Open Accelerator）等、コンパイラとしてＰＧＩ（The Portland Group, Inc.）コンパイラ（登録商標）等がある。例えば、OpenACCを使った例では、ユーザはC/C++/Fortran言語で書かれたコードに、OpenＡＣＣディレクティブで並列処理させる等を指定する。ＰＧＩコンパイラは、コードの並列可能性をチェックして、ＧＰＵ用、ＣＰＵ用実行バイナリを生成し、実行モジュール化する。IBM JDK（登録商標）は、Java（登録商標）のlambda形式に従った並列処理指定を、ＧＰＵにオフロードする機能をサポートしている。これらの技術を用いることで、ＧＰＵメモリへのデータ割り当て等を、プログラマは意識する必要がない。
このように、OpenＣＬ、ＣＵＤＡ、OpenACC等の技術により、ＧＰＵやＦＰＧＡへのオフロード処理が可能になっている。

しかし、オフロード処理自体は行えるようになっても、適切なオフロードには課題が多い。例えば、Intelコンパイラ（登録商標）のように自動並列化機能を持つコンパイラがある。自動並列化する際は、プログラム上のfor文（繰り返し文）等の並列処理部を抽出する。ところが、ＧＰＵを用いて並列に動作させる場合は、ＣＰＵ-ＧＰＵメモリ間のデータやり取りのオーバヘッドのため、性能が出ないことも多い。ＧＰＵを用いて高速化する際は、スキル保持者が、OpenＣＬやＣＵＤＡでのチューニングや、ＰＧＩコンパイラ等で適切な並列処理部を探索することが必要になっている。
このため、スキルが無いユーザがＧＰＵを使ってアプリケーションを高性能化することは難しいし、自動並列化技術を使う場合も、for文を並列するかしないかの試行錯誤のチューニング等、利用開始までに多くの時間がかかっている。

並列処理箇所の試行錯誤を自動化する取り組みとして、非特許文献１，２が挙げられる。非特許文献１，２は、ＧＰＵオフロードに適したループ文を、進化的計算手法を用いて検証環境での性能測定を繰り返すことで、適切に抽出し、ネストループ文内の変数をできるだけ上位のループでＣＰＵ-ＧＰＵ転送を一括化することで自動での高速化を行っている。

Y. Yamato, T. Demizu, H. Noguchi and M. Kataoka, "Automatic GPU Offloading Technology for Open IoT Environment, "IEEE Internet of Things Journal, Sep. 2018. 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. 山登庸次，"アプリケーションループ文のＦＰＧＡ自動オフロード手法の評価," 電子情報通信学会技術報告，SWIM2019-25, Feb. 2020. Numerical Recipes in C, https://www.cec.uchile.cl/cinetica/pcordero/MC_libros/NumericalRecipesinC.pdf https://www.intel.com/content/www/us/en/programmable/support/support-resources/design-examples/design-software/opencl/fft-2d.html

非特許文献１では、ＣＰＵ向けの汎用的コードから、ＧＰＵオフロードに向けて、適切な並列処理領域を自動抽出し、並列処理可能なループ文群に対してＧＡを用いて、より適切な並列処理領域を探索することで、ＧＰＵへの自動オフロードを実現している。しかし、特に、ＦＰＧＡで高速化する際は、ＣＰＵ向けのアルゴリズムからハードウェア処理に適したアルゴリズムに変更し、高速化していることが多いため、ループ文の単純なオフロードだけでは、手動でアルゴリズムから変えて高速化している場合に比べ、性能が不十分なことが多かった。

このような点に鑑みて本発明がなされたのであり、ＰＬＤ（例えば、ＦＰＧＡ）への自動オフロードにおいて、オフロード処理の高速化を図ることを課題とする。

前記した課題を解決するため、アプリケーションの特定処理をＧＰＵ（Graphics Processing Unit）またはＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバであって、前記ＧＰＵまたは前記ＰＬＤにオフロード可能なライブラリおよびＩＰコアを記憶する記憶部と、アプリケーションのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するアプリケーションコード分析部と、検出された前記外部ライブラリ呼び出しをキーにして、前記記憶部から前記ライブラリおよび前記ＩＰコアを取得する置換機能検出部と、前記アプリケーションのソースコードの置換元の処理記述を、前記置換機能検出部が取得した前記ライブラリおよび前記ＩＰコアの置換先の処理記述として置換するとともに、置換した前記ライブラリおよび前記ＩＰコアの処理記述を、オフロード対象の機能ブロックとして、前記ＧＰＵまたは前記ＰＬＤにオフロードする置換処理部と、ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出するオフロードパターン作成部と、作成されたＧＰＵまたはＰＬＤ処理パターンの前記アプリケーションをコンパイルして、実行ファイルを作成する実行ファイル作成部と、作成された前記実行ファイルをアクセラレータ検証用装置に配置し、前記ＧＰＵまたは前記ＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部と、を備え、前記実行ファイル作成部は、前記性能測定用処理による性能測定結果をもとに、複数の前記ＧＰＵまたはＰＬＤ処理パターンから最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンを選択し、最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンをコンパイルして、最終実行ファイルを作成することを特徴とするオフロードサーバとした。

本発明によれば、ＰＬＤへの自動オフロードにおいて、オフロード処理の高速化を図ることができる。

本発明の実施形態に係るオフロードサーバを含む環境適応ソフトウェアシステムを示す図である。 上記実施形態に係るオフロードサーバの構成例を示す機能ブロック図である。 上記実施形態に係るオフロードサーバの機能ブロックのオフロード処理を示す図である。 上記実施形態に係るオフロードサーバの制御部が、機能ブロックのオフロード処理において<処理Ａ－１>と<処理Ｂ－１>と<処理Ｃ－１>とを実行する場合のフローチャートである。 上記実施形態に係るオフロードサーバの制御部が、機能ブロックのオフロード処理において<処理Ａ－２>と<処理Ｂ－２>と<処理Ｃ－２>とを実行する場合のフローチャートである。 本発明の実施形態に係るオフロードサーバの機能を実現するコンピュータの一例を示すハードウェア構成図である。

次に、本発明を実施するための形態（以下、「本実施形態」と称する。）における、オフロードサーバ１等について説明する。
以下、明細書の説明において、ＰＬＤ（Programmable Logic Device）として、ＦＰＧＡ（Field Programmable Gate Array）に適用した例について説明する。本発明は、プログラマブルロジックデバイス全般に適用可能である。
（背景説明）
処理時間が長くかかる特定のループ文を、ＦＰＧＡにオフロードして高速化することを考えた際に、どのループをオフロードすれば高速になるかの予測は難しい。
このため、非特許文献３では、ＧＰＵ同様、検証環境で試行錯誤を自動で行うことを記載している。しかし、ＦＰＧＡは、ＧＰＵと異なり、コンパイルに数時間以上かかる。そこで、オフロード候補のループ文を絞ってから、実測試行を行うようにしている。

しかし、特に、ＦＰＧＡで高速化する際は、ＣＰＵ向けのアルゴリズムからハードウェア処理に適したアルゴリズムに変更し、高速化していることが多い。このため、手動でアルゴリズムから変えて高速化している場合に比べ、ループ文の単純なオフロードだけでは、性能が不十分なことが多かった。例えば、行列積算の場合を例に採ると、行列のすべてのデータをＦＰＧＡのローカルメモリで持つことは難しいため、まずデータＡを行方向に読み、次にデータＢを列方向に読み込むことで、容量制限があるローカルメモリを上手く活用する等のアルゴリズムで高速化している例がある。ＧＰＵの場合も、例えばフーリエ変換を、ＧＰＵ向けのアルゴリズムで実装してＧＰＵ向けに高速化したＣＵＤＡライブラリのcuFFT等がある。

しかしながら、このようなアプリケーションに応じた処理ハードウェア向けのアルゴリズム変更は、機械に自動で抽出させるのは現状では困難である。
そこで、本発明は、個々のループ文でなく、行列積算やフーリエ変換等のより大きな単位で、ＦＰＧＡやＧＰＵ等ハードウェア向けのアルゴリズム含めて実装された機能ブロックに置換する。これにより、人の既存ノウハウの活用し、オフロード処理の高速化を図る。

（実施形態）
図１は、本実施形態に係るオフロードサーバ１を含む環境適応ソフトウェアシステムを示す図である。
本実施形態に係る環境適応ソフトウェアシステムは、従来の環境適応ソフトウェアの構成に加え、オフロードサーバ１を含むことを特徴とする。オフロードサーバ１は、アプリケーションの特定処理をアクセラレータにオフロードするオフロードサーバである。また、オフロードサーバ１は、クラウドレイヤ２、ネットワークレイヤ３、デバイスレイヤ４の３層に位置する各装置と通信可能に接続される。クラウドレイヤ２にはデータセンタ３０が、ネットワークレイヤ３にはネットワークエッジ２０が、デバイスレイヤ４にはゲートウェイ１０が、それぞれ配設される。

そこで、本実施形態に係るオフロードサーバ１を含む環境適応ソフトウェアシステムでは、デバイスレイヤ、ネットワークレイヤ、クラウドレイヤのそれぞれのレイヤにおいて、機能配置や処理オフロードを適切に行うことによる効率化を実現する。主に、機能を３レイヤの適切な場所に配置し処理させる機能配置効率化と、画像分析等の機能処理をＧＰＵやＦＰＧＡ等のヘテロハードウェアにオフロードすることでの効率化を図る。クラウドレイヤでは、ＧＰＵやＦＰＧＡ等のヘテロジニアスなＨＷ（ハードウェア）（以下、「ヘテロデバイス」と称する。）を備えたサーバが増えてきている。例えば、Microsoft（登録商標）社のBing検索においても、ＦＰＧＡが利用されている。このように、ヘテロデバイスを活用し、例えば、行列計算等をＧＰＵにオフロードしたり、ＦＦＴ（Fast Fourier Transform）計算等の特定処理をＦＰＧＡにオフロードしたりすることで、高性能化を実現している。

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

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

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

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

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

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

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

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

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

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

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

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

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

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

<アプリケーションコード指定部１１１>
アプリケーションコード指定部１１１は、入力されたアプリケーションコードの指定を行う。具体的には、アプリケーションコード指定部１１１は、ユーザに提供しているサービスの処理機能（画像分析等）を特定する。

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

以下、図３のステップ番号を参照して各部の動作を説明する。
<ステップＳ１１：Specify application code>
ステップＳ１１において、アプリケーションコード指定部１１１（図２参照）は、ユーザに提供しているサービスの処理機能（画像分析等）を特定する。具体的には、アプリケーションコード指定部１１１は、入力されたアプリケーションコードの指定を行う。

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

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

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

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

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

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

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

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

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

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

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

なお、上記では、環境適応に必要な、コード変換、配置場所調整を一括して行う処理フローを説明したが、これに限らず、行いたい処理だけ切出すことも可能である。例えば、ＧＰＵ、ＦＰＧＡ向けにコード変換だけ行いたい場合は、上記ステップＳ１１～ステップＳ２２の、環境適応機能や検証環境等必要な部分だけ利用すればよい。

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

［実装例］
実装例を説明する。
<実装する際の利用ツール例>
ＧＰＵは、NVIDIA Quadro P4000(CUDA core: 1792, Memory: GDDR5 8GB)が利用できる。ＦＰＧＡは、Intel PAC with Intel Arria10 GX FPGA等が利用できる。
ＧＰＵ処理は、市中のＰＧＩコンパイラなどが利用できる。ＰＧＩコンパイラは、OpenACCを解釈するC/C++/Fortran向けコンパイラである。ＰＧＩコンパイラは、for文等のループ文を、OpenACCのディレクティブ #pragma acc kernels, #pragma acc parallel loopで指定することにより、ＧＰＵ向けバイトコードを生成し，実行によりGPUオフロードを可能としている。併せて、ＰＧＩコンパイラは、cuFFTやcuBLAS等のCUDAライブラリの呼び出しも処理が可能である。

ＦＰＧＡ処理は、Intel Acceleration Stack（Intel FPGA SDK for OpenCL，Quartus Prime Version）等が利用できる、Intel FPGA SDK for OpenCLは、標準OpenCLに加えIntel向けの#pragma等を解釈する高位合成ツール（HLS）である。高位合成ツールは、ＦＰＧＡで処理するカーネルとＣＰＵで処理するホストプログラムを記述したOpenCLコードを解釈するとともに、リソース量等の情報を出力し、ＦＰＧＡの配線作業等を行うことで、ＦＰＧＡで動作できるようにする。ＦＰＧＡ実機で動作できるようにするには、１００行程度の小プログラムでも３時間程の長時間がかかるが、リソース量オーバーの際は早めにエラーとなる。ＦＰＧＡの既存OpenCLコードをカーネルコードに組み込めば、OpenCLプログラム処理の中でオフロードが可能である。

コードパターンＤＢ１３０は、MySQL等が利用できる。コードパターンＤＢ１３０は、呼び出しているライブラリ名をキーに、高速化できるＧＰＵ用ライブラリやＦＰＧＡ用ＩＰコアを検索するためのレコードを保持する。ライブラリやＩＰコアには、それに紐づく名前やコードや実行ファイルが保持される。実行ファイルはその利用手法等も登録されている。併せて、コードパターンＤＢ１３０には、ライブラリやＩＰコアを類似性検出技術で検出するための、比較用コードとの対応関係も保持される。

類似性検出ツールには、Deckard等が利用できる。Deckardは、機能ブロックのオフロードの適用領域拡大のため、ライブラリ呼び出し以外にも、コードコピーし変更した機能等のオフロードを実現する。このため、Deckardは、照合対象となる部分コードと、コードパターンＤＢ１３０に登録されたコードの類似性を判定する。

実装の動作概要を説明する。
<実装の動作概要>
実装例では、Ｃ／Ｃ++アプリケーションの利用依頼があると、まず、Ｃ／Ｃ++アプリケーションのコードを解析して、ループ文をオフロードに使うためにループ文を検出する。また、<処理Ａ－１>で呼び出されているライブラリや<処理Ａ－２>で定義されているクラス、構造体等のプログラム構造を把握する。構文解析には、LLVM/Clangの構文解析ライブラリ等が利用できる。呼び出されている外部ライブラリがあるかどうかは、コードパターンＤＢ１３０の外部ライブラリリストと照合することで確認する。

実装例では、次に、<処理Ｂ－１>で呼び出されているライブラリを高速化できるＧＰＵ用ライブラリ、ＦＰＧＡ用ＩＰコアの検出を行う。呼び出されているライブラリをキーに、コードパターンＤＢ１３０に登録されているレコードから、高速化可能な実行ファイルやOpenCL等を取得する。高速化できる置換用機能が見つかった場合、<処理Ｃ－１>の実行用ファイルを作成する。ＧＰＵ用ライブラリの場合は、置換用ＣＵＤＡライブラリを呼び出すよう、Ｃ／Ｃ++コードに、元の部分は削除して置換記述する。ＦＰＧＡ用ＩＰコアの場合は、取得したOpenCLコードを、元の部分をホストコードから削除してから、カーネルコードに置換記述する。それぞれの置換記述が終わった場合、ＧＰＵ向けにはＰＧＩコンパイラでコンパイルし、ＦＰＧＡ向けにはIntel acceleration stackでコンパイルする。ＦＰＧＡに関しては、OpenCLコードに基づき、ＣＰＵとＦＰＧＡの接続が高位合成ツール（HLS）を介して行われる。

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

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

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

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

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

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

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

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

［効果］
以上説明したように、本実施形態に係るオフロードサーバ１は、アプリケーションの特定処理をＧＰＵまたはＰＬＤにオフロードするオフロードサーバであって、ＧＰＵまたはＰＬＤにオフロード可能なライブラリおよびＩＰコアを記憶するコードパターンＤＢ１３０と、アプリケーションのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するアプリケーションコード分析部１１２と、検出された外部ライブラリ呼び出しをキーにして、コードパターンＤＢ１３０からライブラリおよびＩＰコアを取得する置換機能検出部１１３と、アプリケーションのソースコードの置換元の処理記述を、置換機能検出部１１３が取得した置換先のライブラリおよびＩＰコアの置換先の処理記述として置換するとともに、置換したライブラリおよびＩＰコアの処理記述を、オフロード対象の機能ブロックとして、ＧＰＵまたはＰＬＤにオフロードする置換処理部１１４と、ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出するオフロードパターン作成部１１５と、作成されたＧＰＵまたはＰＬＤ処理パターンの前記アプリケーションをコンパイルして、実行ファイルを作成する実行ファイル作成部１１７と、作成された実行ファイルをアクセラレータ検証用装置に配置し、ＧＰＵまたはＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部１１６と、を備え、実行ファイル作成部１１７は、性能測定用処理による性能測定結果をもとに、複数のＧＰＵまたはＰＬＤ処理パターンから最高処理性能のＧＰＵまたはＰＬＤ処理パターンを選択し、最高処理性能のＧＰＵまたはＰＬＤ処理パターンをコンパイルして、最終実行ファイルを作成する。

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

本実施形態では、コードパターンＤＢ１３０は、ホストで計算する場合に記述が同様になる処理のクラスまたは構造体を記憶し、アプリケーションコード分析部１１２は、ソースコードからクラスまたは構造体の定義記述コードを検出し、置換機能検出部１１３は、検出されたクラスまたは構造体の定義記述コードをキーにして、コードパターンＤＢ１３０からライブラリおよびＩＰコアを取得し、置換処理部１１４は、アプリケーションのソースコードの置換元の処理記述を、取得したライブラリおよびＩＰコアの処理記述に置換するとともに、置換元と置換先で引数、戻り値の数または型が異なる場合に、その確認を通知することを特徴とする。

このようにすることにより、クラスまたは構造体の定義記述コードをキーに、ＧＰＵライブラリおよびＩＰコアを検出し、アプリケーションコードの置換元の処理記述を、置換先のＧＰＵライブラリ、ＩＰコア処理記述に置換する。また、置換元と置換先で引数、戻り値の数または型が異なる場合に、例えばユーザに確認する。これにより、登録されていないライブラリ呼び出し以外の機能処理を検出することができる。

本実施形態では、置換機能検出部１１３は、コピーコードやコピー後変更したコードを検出する類似性検出ツールを用いて、コードパターンＤＢ１３０からライブラリおよびＩＰコアを取得することを特徴とする。

このようにすることにより、類似性検出ツールを用いて、照合対象となる部分コードと、記憶部に登録された定義記述コードとの類似性を判定することができ、機能ブロックのオフロードの適用領域を拡大することができる。

本実施形態では、コードパターンＤＢ１３０は、ライブラリとして、ＣＵＤＡライブラリと当該ＣＵＤＡライブラリを利用するためのライブラリ利用手順とを記憶し、置換機能検出部１１３は、ライブラリ利用手順をもとに、コードパターンＤＢ１３０からＣＵＤＡライブラリを取得することを特徴とする。

このようにすることにより、ＣＵＤＡライブラリと利用手順も含めて記憶しておくことで、ソースコードからＣＵＤＡライブラリを検出することができる。これにより、アプリケーションコードの置換元の処理記述を、置換先のＣＵＤＡライブラリ処理記述に置換して、機能ブロックとして、ＧＰＵにオフロードすることができる。

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

このようにすることにより、OpenCLコードを、カーネルコードとして、ホストとＰＬＤとを接続することができ、ＰＬＤへのＩＰコアの実装を行うことができる。例えば、ＦＰＧＡ用ＩＰコアの場合は、取得したOpenCLコードを、元の部分をホストコードから削除してから、カーネルコードに置換記述する。また、ＦＰＧＡに関しては、OpenCLコードに基づき、ＣＰＵとＦＰＧＡの接続を、高位合成ツールを介して行うことができる。

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

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

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

また、上記の各構成、機能、処理部、処理手段等は、それらの一部又は全部を、例えば集積回路で設計する等によりハードウェアで実現してもよい。また、上記の各構成、機能等は、プロセッサがそれぞれの機能を実現するプログラムを解釈し、実行するためのソフトウェアで実現してもよい。各機能を実現するプログラム、テーブル、ファイル等の情報は、メモリや、ハードディスク、ＳＳＤ（Solid State Drive）等の記録装置、又は、ＩＣ（Integrated Circuit）カード、ＳＤ（Secure Digital）カード、光ディスク等の記録媒体に保持することができる。

また、本実施形態では、ＧＰＵ、ＦＰＧＡ処理をオフロードできるものであればどのようなものでもよい。

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

Claims (7)

1. アプリケーションの特定処理をＧＰＵ（Graphics Processing Unit）またはＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバであって、
   前記ＧＰＵまたは前記ＰＬＤにオフロード可能なライブラリおよびＩＰコアを記憶する記憶部と、
   アプリケーションのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するアプリケーションコード分析部と、
   検出された前記外部ライブラリ呼び出しをキーにして、前記記憶部から前記ライブラリおよび前記ＩＰコアを取得する置換機能検出部と、
   前記アプリケーションのソースコードの置換元の処理記述を、前記置換機能検出部が取得した前記ライブラリおよび前記ＩＰコアの置換先の処理記述として置換するとともに、
   置換した前記ライブラリおよび前記ＩＰコアの処理記述を、オフロード対象の機能ブロックとして、前記ＧＰＵまたは前記ＰＬＤにオフロードする置換処理部と、
   ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出するオフロードパターン作成部と、
   作成されたＧＰＵまたはＰＬＤ処理パターンの前記アプリケーションをコンパイルして、実行ファイルを作成する実行ファイル作成部と、
   作成された前記実行ファイルをアクセラレータ検証用装置に配置し、前記ＧＰＵまたは前記ＰＬＤにオフロードした際の性能測定用処理を実行する性能測定部と、を備え、
   前記実行ファイル作成部は、前記性能測定用処理による性能測定結果をもとに、複数の前記ＧＰＵまたはＰＬＤ処理パターンから最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンを選択し、最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンをコンパイルして、最終実行ファイルを作成する
   ことを特徴とするオフロードサーバ。
2. 前記記憶部は、ホストで計算する場合に記述が同様になる処理のクラスまたは構造体を記憶し、
   前記アプリケーションコード分析部は、前記ソースコードから前記クラスまたは前記構造体の定義記述コードを検出し、
   前記置換機能検出部は、検出された前記クラスまたは前記構造体の定義記述コードをキーにして、前記記憶部から前記ライブラリおよび前記ＩＰコアを取得し、
   前記置換処理部は、アプリケーションのソースコードの置換元の処理記述を、取得した前記ライブラリおよび前記ＩＰコアの処理記述に置換するとともに、置換元と置換先で引数、戻り値の数または型が異なる場合に、当該異なることを通知する
   ことを特徴とする請求項１に記載のオフロードサーバ。
3. 前記置換機能検出部は、コピーコードやコピー後変更したコードを検出する類似性検出ツールを用いて、前記記憶部から前記ライブラリおよび前記ＩＰコアを取得する
   ことを特徴とする請求項２に記載のオフロードサーバ。
4. 前記記憶部は、前記ライブラリとして、ＣＵＤＡライブラリと当該ＣＵＤＡライブラリを利用するためのライブラリ利用手順とを記憶し、
   前記置換機能検出部は、前記ライブラリ利用手順をもとに、前記記憶部から前記ＣＵＤＡライブラリを取得する
   ことを特徴とする請求項１に記載のオフロードサーバ。
5. 前記記憶部は、ＩＰコア関連の情報としてOpenCLコードを記憶し、
   前記オフロードパターン作成部は、前記OpenCLコードをもとにOpenCLインタフェースを用いてホストと前記ＰＬＤとを接続するとともに、前記OpenCLコードをもとに前記ＰＬＤへの前記ＩＰコアの実装を行う
   ことを特徴とする請求項１に記載のオフロードサーバ。
6. アプリケーションの特定処理をＧＰＵ（Graphics Processing Unit）またはＰＬＤ（Programmable Logic Device）にオフロードするオフロードサーバのオフロード制御方法であって、
   前記オフロードサーバは、
   前記ＧＰＵまたは前記ＰＬＤにオフロード可能なライブラリおよびＩＰコアを記憶する記憶部を備えており、
   アプリケーションのソースコードを分析して、当該ソースコードに含まれる外部ライブラリ呼び出しを検出するステップと、
   検出された前記外部ライブラリ呼び出しをキーにして、前記記憶部から前記ライブラリおよび前記ＩＰコアを取得するステップと、
   前記アプリケーションのソースコードの置換元の処理記述を、取得した前記ライブラリおよび前記ＩＰコアの置換先の処理記述として置換するステップと、
   置換した前記ライブラリおよび前記ＩＰコアの処理記述を、オフロード対象の機能ブロックとして、前記ＧＰＵまたはＰＬＤにオフロードするステップと、
   ホストプログラムとのインタフェースを作成し、検証環境での性能測定を通じて、オフロードするしないを試行することで、より高速となるオフロードパターンを抽出するステップと、
   作成されたＧＰＵまたはＰＬＤ処理パターンの前記アプリケーションをコンパイルして、実行ファイルを作成するとともに、
   性能測定用処理による性能測定結果をもとに、複数の前記ＧＰＵまたはＰＬＤ処理パターンから最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンを選択し、最高処理性能の前記ＧＰＵまたはＰＬＤ処理パターンをコンパイルして、最終実行ファイルを作成するステップと、
   作成された前記実行ファイルをアクセラレータ検証用装置に配置し、前記ＧＰＵまたは前記ＰＬＤにオフロードした際の前記性能測定用処理を実行するステップと、を実行する
   ことを特徴とするオフロード制御方法。
7. コンピュータを、請求項１乃至５のいずれか１項に記載のオフロードサーバとして機能させるためのオフロードプログラム。

Images