JP6249360B2

JP6249360B2 - ハードウェア設計装置，及びハードウェア設計用プログラム

Info

Publication number: JP6249360B2
Application number: JP2013105024A
Authority: JP
Inventors: 伸一山際
Original assignee: University of Tsukuba NUC
Current assignee: University of Tsukuba NUC
Priority date: 2013-05-17
Filing date: 2013-05-17
Publication date: 2017-12-20
Anticipated expiration: 2033-05-17
Also published as: JP2014225194A

Description

本発明は、ハードウェア設計装置，及びハードウェア設計用プログラムに関する。

現在、ネットワークには、様々なセンサ及びデバイスが接続され、これらのセンサ及びデバイスから刻々と出力されるデータがネットワーク上でデータストリームを形成する。データストリームを形成するデータ（ストリームデータと呼ばれる）を滞りなく処理する（リアルタイムに処理する）手法として、ストリームコンピューティングがある。

ストリームコンピューティングでは、ストリームデータに関して所定のデータ単位が入力として決定され、このデータ単位に対する複数の処理（演算）が直列に実行されるパイプライン処理が行われる。

パイプライン処理を実行する典型的な環境として、ＧＰＵ（Graphical Processing Unit）を用いたソフトウェア処理がある。ＧＰＵは、並列処理を実行可能な複数のプロセッ
サを含んでおり、パイプライン処理の手順を記述したプログラムに従って、各プロセッサにパイプライン処理に係る複数の処理を実行させることで、パイプライン処理結果を得ることができる。

OpenCLのような、ＧＰＵ向けのプログラミング言語では、ＧＰＵが並列処理を実行可能な複数のプロセッサを有することを考慮した、パイプライン処理手順を記述することができる。例えば、プログラム上で、パイプライン処理に係る複数の演算を、複数のプロセッサに割り当てることができる。

特開２０１２−１７４２６８号公報

"OpenCL規格を用いたFPGAデザインの導入"、[online]、アルテラ、[平成２５年５月１３日検索]、インターネット＜URL: http://www.altera.co.jp/literature/wp/wp-01173-opencl_j.pdf＞ Mencer, O., ASC, "a stream compiler for computing with FPGAs, Computer-Aided Design of Integrated Circuits and Systems," IEEE Transactions on (Volume:25 , Issue: 9 ) , pp. 603 - 1617, Sept. 2006. Shinichi Yamagiwa, leonel Sousa, "Caravela: A Novel Stream-Based Distributed Computing Environment," Computer , vol.40, no.5, pp.70-77, May 2007.

しかしながら、ＧＰＵ及びＧＰＵを用いたシステムは、回路規模が大きく、またコストの上昇を招来する。このため、パイプライン処理を行うために、ＧＰＵ乃至ＧＰＵシステムを導入することが困難、又は非現実的である場合が少なくない。

このため、ストリームデータ，ストリームデータに対して行う処理（演算）の内容などに応じたパイプライン処理を実行するハードウェア（パイプライン処理用のディジタル回
路：パイプラインハードウェアと呼ぶ）の開発が求められている。

パイプラインハードウェアを設計する際には、パイプライン処理で実行される複数の演算（四則演算等）をそれぞれ実行する複数の演算コンポーネント（「ハードウェアコンポーネント」と呼ぶ。以下単に「コンポーネント」と表記）が用意され、演算の順序に従って複数のコンポーネントが配線により接続される。パイプラインハードウェアにはクロックが入力され、クロック毎に演算が進められる。演算に係る遅延（ディレイ）はクロック数として定義される。

従来におけるハードウェア記述言語（hardware description language：HDL）を用いたパイプラインハードウェアの設計では、設計者がパイプライン処理における演算内容に応じて複数のコンポーネントを選択し、選択したコンポーネント間の配線接続によってパイプラインハードウェアを設計する。このとき、設計者がコンポーネント間のデータ到達タイミングを計りながら、最終的な演算結果が正しく出力されるための遅延（ディレイ）をマニュアル操作で挿入する。

コンポーネントは様々な遅延を持つ。例えば、加算器は、実装技術（ＬＳＩのプロセス）の違いにより遅延が異なる。さらに、積算回路については繰り返し処理が伴うことで遅延が可変となることがある。このように、パイプラインハードウェアの設計では、各コンポーネントの遅延を考慮することにより、全体の時間バランスを考慮することが要求される。

しかしながら、例えば、コンポーネントの入れ替え（選択変更）により、コンポーネント自体の遅延が変化した場合には、この変化をコンポーネント間に挿入されたディレイで吸収できない場合が起こり得る。この影響がパイプライン全体に及ぶ場合には、パイプラインハードウェアの設計を最初からやり直すことが要求される虞があった。このように、設計者に対して、コンポーネントの選択と、挿入する遅延の考慮とを慎重に行うことが要求され、これは設計者にとって大きな負担であった。

現在のところ、パイプライン処理の演算内容に応じたパイプラインハードウェアのひな形となるモデルを自動的に生成する機構はない。このため、コンポーネントの特性と、コンポーネントのパイプラインへの適用を、目的とするパイプラインハードウェアに応じて一意に決定可能な機構、特に、コンポーネントを考慮した遅延を自動的に挿入可能な機構もない。

本発明は、上記の事情に鑑みなされたものであり、パイプライン処理の演算内容に応じたパイプラインハードウェアのモデルを自動的に生成可能な技術を提供することを目的とする。

本発明は、上記課題を解決するために以下の手段を採用する。すなわち、本発明は、ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、前記パイプライン処理で行われる演算が記述されたプログラムと、前記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、前記演算の内容及び前記定義情報に応じた２以上のハードウェアコンポーネントを前記コンポーネント情報に基づき特定し、前記２以上のハードウェアコンポーネントが前記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成する制御装置を含むハードウェア設計装置である。

また、本発明は、上記した制御装置における処理をコンピュータに実行させるハードウ
ェア設計用プログラム，このようなプログラムを記録したコンピュータ読み取り可能な記録媒体，及び上記制御装置によるパイプラインハードウェア記述の生成方法としても特定することができる。

本発明によれば、パイプライン処理の演算内容に応じたパイプラインハードウェアのモデルを自動的に生成可能な技術を提供することができる。

図１は、ハードウェア設計用プログラムの実行によってハードウェア設計装置として機能する情報処理装置（コンピュータ）の構成例を示す図である。図２は、ＳＰＣの仕組みを説明する図である。図３は、ＸＭＬで記述されたフローモデルの例を示す。図４は、カウンタ生成に係る処理を説明する図である。図５は、コンポーネントの定義ファイルの例を示す。図６は、コンポーネントの選択に係る説明図である。図７は、抽象構文木の例を示す図である。図８は、演算式“a[id] = b[id] + c[id] * d[id] - e[id]”に応じて作成されたパイプラインハードウェアのモデル（ＨＡＭ）における遅延を説明する図である。図９は、遅延挿入の例を示す説明図である。図１０は、ＣＰＵによって実行されるＳＰＣにおける処理を大略して示すフローチャートである。図１１は、ＨＤＬによるパイプラインハードウェアの記述例を示す。

以下、図面を参照して本発明の実施形態について説明する。実施形態の構成及び設定は例示であり、本発明は実施形態の構成及び設定に限定されない。

以下、実施形態に係るハードウェア設計用プログラム、及びハードウェア設計用プログラムの実行によりハードウェア設計装置として機能する情報処理装置（コンピュータ）について説明する。

ハードウェア設計用プログラムは、ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、上記パイプライン処理で行われる演算が記述されたプログラムと、上記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、上記演算の内容及び上記定義情報に応じた２以上のハードウェアコンポーネントを上記コンポーネント情報に基づき特定する。

さらに、ハードウェア設計用プログラムは、上記２以上のハードウェアコンポーネントが上記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成する。このようなハードウェア設計用プログラムを“ストリーム・パイプライン・コンパイラ（ＳＰＣ）”と呼ぶ。

＜情報処理装置（ハードウェア設計装置）＞
図１は、ハードウェア設計用プログラムの実行によってハードウェア設計装置として機能する情報処理装置（コンピュータ）の構成例を示す図である。情報処理装置１０として、例えば、パーソナルコンピュータ（ＰＣ），ワークステーション，専用又は汎用のサーバマシンを適用することができる。

図１において、情報処理装置１０は、例として、バスＢを介して相互に接続された、ＣＰＵ１１と、主記憶装置１２と、補助記憶装置１３と、入力装置１４と、出力装置１５と、通信インタフェース回路（通信Ｉ／Ｆ）１６とを備える。

主記憶装置１２は、ＣＰＵ１１の作業領域として使用されるメインメモリとして機能する。メインメモリは、例えば、ＲＡＭ（Random Access Memory）及びＲＯＭ（Read Only Memory）によって形成される。

補助記憶装置１３は、制御装置に相当するＣＰＵ１１によって実行される、ハードウェア設計用プログラムを含む各種のプログラム，及び各プログラムの実行時に使用されるデータを記憶する。補助記憶装置１３は、例えば、不揮発性記録媒体であり、例えば、ハードディスク，フラッシュメモリ，ＥＥＰＲＯＭ（Electrically Erasable Programmable Read-Only Memory）の少なくとも１つを用いて形成することができる。主記憶装置１２及
び補助記憶装置１３のそれぞれは、記億装置，記録媒体の一例である。

入力装置１４は、キーボード，マウスやタッチパネルのようなポインティングデバイスを含み、情報（データ）の入力に使用される。出力装置１５は、例えば、ディスプレイ装置であり、情報を画面に表示する。通信Ｉ／Ｆ１６は、ネットワークとの通信処理を司る。

ＳＰＣは、補助記憶装置１３にインストールされており、ＣＰＵ１１がＳＰＣを主記憶装置１２に読み出し、ロードして実行することによって、情報処理装置１０は、ハードウェア設計装置として機能することができる。

＜ＳＰＣ＞
図２は、ＳＰＣの仕組みを説明する図である。図２に示すように、ＳＰＣに対する入力として、“Flow-model（フローモデル）”と呼ばれる、パイプライン処理のモデルを定義したモデル定義情報と、パイプライン処理で行われる演算が記述されたプログラムとが所定の記述言語で記述されたファイル（パイプライン処理モデルのファイル）が用意される。各ファイルは、補助記憶装置１３に記憶され、ＣＰＵ１１によるＳＰＣの実行に際して使用される。

ＳＰＣは、フローモデルのファイルを入力として、フローモデルにおいて定義されたパイプライン処理を行うパイプラインハードウェアの抽象化モデル（Hardware Abstraction
Model：ＨＡＭ）を生成する（図２（１））。抽象化モデル（ＨＡＭ）は、フローモデルにおける演算内容に従った演算を行う演算子である２以上のコンポーネントを特定し、２以上のコンポーネントを合成することによって生成される。さらに、ＳＰＣは、ＨＡＭをハードウェア記述言語（ＨＤＬ）で記述したパイプラインハードウェアの記述（Hardware
Description）を生成して出力する（図２（２））。

このように、ＳＰＣは、パイプラインハードウェアを設計するためのひな形として利用可能なＨＡＭ及びパイプラインハードウェアの記述を自動的に生成することができる。この点で、設計者の負担を減らすことができる。

図２において、パイプラインハードウェア１は、所定の複数の入力データを得て、パイプライン処理により所望の出力データを出力するハードウェアである。パイプラインハードウェア１は、パイプラインを形成する複数のコンポーネント２（ハードウェアコンポーネントとして機能する複数の部分ハードウェア）を含み、コンポーネント２間の配線接続によって合成される。

＜＜フローモデル＞＞
フローモデルの記述言語として、例えば、ＸＭＬ（Extensible Markup Language）を含む様々なマークアップ言語を適用可能である。

パイプライン処理で行われる演算が記述されたプログラムの一例として、ストリームコンピューティング向けのプログラム言語であるOpenCL（Open Computing Language）で記
述されたプログラムを使用することができる。但し、パイプライン処理で行われる演算内容が特定可能に記述される限り、OpenCL以外の他のプログラム言語を適用することもできる。

OpenCLは、ＧＰＵ（Graphics Processing Unit）用のプログラムを記述することができる。一般に、ＧＰＵは、複数のプロセッサを有し、これらのプロセッサを用いた並列処理を行うことができる。OpenCLは、「プロセッサインデックス」，「プロセッサＩＤ」と呼ばれる“id”値を指定する構文（例えば、id=get_grobal_id(0)：括弧内の数字"０"はid
値（id番号））の記述と、例えば、“c[id]=+a[id]-100+b[id]”のような、idの指定を含む演算内容の記述とにより、並列処理において個々の処理を行うプロセッサを指定することができる。

“フローモデル”のファイルは、パイプライン処理で行われる演算が記述されたプログラム（カーネルプログラムと呼ばれる）として、OpenCLを用いたカーネルプログラムの記述を含むことができる。カーネルプログラムにおいて、“id”を指定した演算内容が記述される。

OpenCLのような、ＧＰＵのようなアクセラレータで実行されるストリームコンピューティング向けのプログラミング言語におけるパイプライン処理の手順に係る記述は、ストリームデータを順次、処理していくハードウェア構成の記述と等価である。このため、当該プログラムがカーネルプログラムとしてＳＰＣに入力され、ＳＰＣがカーネルプログラムをコンパイルすることで、パイプラインハードウェアのコンポーネント割り出し乃至選択を行うことが可能となる。

定義情報は、パイプライン処理に使用される入力データ，パイプライン処理の結果として出力される出力データの定義情報を含む。入力データ及び出力データの定義情報として、カーネルプログラムの関数引数を入力データ又は出力データとする定義情報が記述される。

図３は、ＸＭＬで記述されたフローモデルの例を示す。図３において、カーネルのタグ（<kernel> </kernel>）で挟まれた部分が、OpenCLのプログラム、すなわちカーネルプログラムの記述である。そして、カーネルプログラムに記述された引数“ａ”及び“ｂ”を入力（input）とし、演算結果の“ｃ”を出力（output）とするモデル定義情報が記述さ
れている。図２の例では、“ａ”，“ｂ”，“ｃ”のそれぞれの定義（名称，データタイプ，データ長等）が記述されている。

なお、上記説明では、カーネルプログラムと定義情報とがフローモデルのファイルにおいて一つにまとめられているが、カーネルプログラムと定義情報とは個別のファイルであっても良い。

＜＜パイプラインハードウェアの抽象化モデル（ＨＡＭ）＞＞
次に、パイプラインハードウェアの抽象化モデル（ＨＡＭ）の詳細について説明する。ＳＰＣによって生成されるＨＡＭは、以下のように規定（設定）される。
・入力は、“（入力）データ”，“（入力）データのアドレス”，“クロック”，及び“
リセット”を含む。
・出力は、“（出力）データ”，“（出力）データの有効性を示すValid信号”，“（出
力）データのアドレス”を含む。
・パイプラインハードウェアに対する入力は、メモリ（図１のメモリ３を参照）からの読み出しによって行われる。
・出力データは、メモリ（図１のメモリ４を参照）に書き込まれる。
・入力時における“（入力）データのアドレス”から読み出された“（入力）データ”のメモリからの読み出し時間“T_ad”は、ＳＰＣの合成変数として与えられる。
・パイプラインハードウェアへの入力データを“入力ポート”，パイプラインハードウェアからの出力データを“出力ポート”と呼ぶ。

[入出力に係る規定]
パイプラインハードウェアに対する入力データ（入力ポート）及び出力データ（出力ポート）に関して、例えば、以下のように規定される。
・パイプラインハードウェアを構成するための演算式（例えば、図２のフローモデルにおける“c[id]=+a[id]-100+b[id]”）に現れる入力ポート（ａ及びｂ）は必ず演算式の右辺になくてはならない。
・上記演算式において、出力ポート（ｃ）は左辺にだけ現れなくてはならない。

パイプライン構造を作る際に、或る演算の出力データが他の演算の入力データとされると、パイプラインが再帰構造を有する状態となり、過去のデータ（或る演算の出力データ）を再利用するためのメモリが必要となる。そこで、本実施形態では、メモリの使用を回避すべく、データが入力データ（入力ポート）と、出力データ（出力ポート）とのどちらであるかを規定する。

上記規定が満たされない場合には、エラーが出力される。例えば、ＣＰＵ１１が、出力装置１５にエラーを表示する。これによって、プログラム中でのメモリ発生可能性（再帰構造）を検出でき、設計者（プログラマ）に修正を促すことで、メモリを使用しないＨＡＭを生成できる。また、再帰構造の排除によって、ＨＡＭ生成の際におけるコンポーネントの合成が容易になる。

[カウンタの生成]
入力とされるプログラム（上記したフローモデル中のカーネルプログラム）で使用される入力データ（入力ポート）及び出力データ（出力ポート）の各アドレスは、カーネルプログラムで記述されたプロセッサＩＤ値（id値：「識別子」に相当）に基づいて指定される。プログラムの構造において、プロセッサＩＤは連続的な値を有し、カウンタのインクリメントによってプロセッサＩＤが変更される。変更されたプロセッサＩＤが次のアドレスとして機能する。

図４は、カウンタ生成に係る処理を説明する図である。例えば、図４に示すようなカーネルプログラムを含むフローモデルのファイルがＳＰＣに供給されたと仮定する。この場合、ＳＰＣ（を実行するＣＰＵ１１）は、ファイルからカーネルプログラムを抽出し、構文解析を行う。

このとき、“id=get_global_id(0)”のようなプロセッサＩＤを指定する構文がある場
合には、ＳＰＣは、演算式（“c[id]=+a[id]-100+b[id]”）に基づき、図４の下段に図示するようなパイプラインハードウェアのモデルを生成する。

図４において、カウンタ値は、入力ポート“ａ”及び“ｂ”、並びに出力ポート“ｃ”に対するアドレスとして供給される。“ａ”のアドレス“a_addr”及び“ｂ”のアドレス
“b_addr”は、入力ポート側にあるメモリ３（図２）のアドレスとして機能する。そして、アドレスとして指定された箇所から読み出された“ａ”の値は、即値“１００”とともに、減算器（Subtract）に入力される。“ｂ”の値は、減算器の出力とともに、加算器（adder）に入力される。そして、加算器から出力ポートである“ｃ”の値が出力される。
“ｃ”の値は、アドレス“c_addr”として指定されたカウンタ値に対応する、出力ポート側のメモリ４（図２）の記憶領域に書き込まれる。そして、所定の契機（例えば、“ｃ”の出力“書き込み”）で、カウンタ値がインクリメント（現在のカウンタ値に１を加算）される。そして、次のカウンタ値に対応するメモリ３の記憶領域から、入力ポート“ａ”及び“ｂ”を取得し、次のカウンタ値に対応するメモリ４の記憶領域に出力ポート“ｂ”を書き込むことができる。

なお、カウンタ値としての“ａ”のアドレス“a_addr”及び“ｂ”のアドレス“b_addr”は、同じ値であるが、ポート毎に異なるメモリ空間（記憶領域）をアクセスするようにして、異なるデータが入力ポートとして得られるようにすることができる。例えば、入力ポート“ａ”の読み出し用メモリと入力ポート“ｂ”の読み出し用メモリとが物理的に別にされることが考えられる。或いは、ポート毎に異なるバンクを持ったデュアルポートメモリを適用することが考えられる。また、入力ポート“ａ”及び“ｂ”に関して、意図して同じメモリ空間にアクセスする場合もあり得る。

このように、実施形態は、プログラムが演算を行うプロセッサを識別子の記述により指定可能なストリームコンピューティング向けのプログラム言語で記述され、且つプログラム中の演算に関して識別子の指定が記述されている場合に、当該演算によって得られる出力データの書き込みアドレスを示すカウンタ値を出力するカウンタを含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成し、カウンタから出力されるカウンタ値は、演算の実行毎に変更される構成を含む。

なお、「インクリメント」の用語は、現在のカウンタ値に１を加算することであり、このような動作を行う場合に、カウンタの回路構成が最も簡易となる。但し、実施形態では、“id”番号の指定構文と、当該“id”番号を用いた演算式があるときに、その演算毎に異なるアドレスを生成及び出力するカウンタが生成されるようにすれば良い。このため、カウンタ値は、演算毎に、カウンタ値を変更するようにされていれば良い。このとき、カウンタ値は、例えば“規則的に増加又は減少”するように変更可能であり、１回のカウンタ値の変更において増加又は減少する数値は、１でも２以上であっても良い。また、id値として採り得る値は、本実施形態では連続する値としているが、離散値（規則的に増加、減少する値）であっても良い。

カウンタを利用した簡易なアドレス生成回路の作成によって、ＨＡＭ（パイプラインハードウェア）の回路構成を簡潔にすることができる。なお、カウンタ値は、採り得る値の最大値となったときには、次に採り得る値の最小値が出力されるようにすることができる。また、カウンタ値とメモリの記憶領域とが関連づけられ、或るカウンタ値に対して対応するメモリの記憶領域に対する読み出し／書き込みが行われるようにすることができる。

[コンポーネントの指定]
ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報として、例えば、コンポーネントの定義ファイルが、コンポーネント毎に用意される。コンポーネント定義ファイルは、以下のような情報を含む。
・コンポーネントの実装のファイル（例えば、ＨＤＬファイル）。
・クロック，イネーブル，リセットのそれぞれの入力と、各入力に合致する実装（コンポーネント）におけるポート（端子）の識別情報（例えばポート名）との組（関連）。
・Valid出力と、Valid出力に対応する実装（部分ハードウェア）上のポート識別情報との組（関連）
・演算において、演算式の左側（左辺）に位置すべき出力ポート，及び演算式の（右辺）に位置すべき入力ポートにそれぞれ対応する実装上のポート識別情報との組（関連）。
・演算に要求されるクロックサイクル数（遅延）
・演算タイプ（加算，減算，積算など）

コンポーネント定義ファイルは、例えば、フローモデルのＸＭＬファイルから独立したＸＭＬファイルとしてＳＰＣに供給される。但し、フローモデルのＸＭＬに含まれてＳＰＣに供給されるようにしても良い。或いは、コンポーネント定義ファイルは、データベース（ＤＢ）上で管理され、必要に応じてＳＰＣによりアクセス（参照）されるようにしても良い。コンポーネント定義ファイルとして、演算タイプが同一であるが仕様が異なる複数のコンポーネントに関する複数のファイルを含むことができる。ＤＢは、例えば、補助記憶装置１３に記憶される。

図５は、コンポーネント定義ファイルの例を示す図であり、減算器（Subtract）のコンポーネント定義ファイルが例示されている。当該コンポーネント定義ファイルの記述として、コンポーネント名“sub”，演算タイプ（Category：SUB），及びRTL（Register Transfer Level）ファイル名，が記述されている。コンポーネント定義ファイルには、各入力データ“ａ”及び“ｂ”の定義と、出力データ“ｃ”の定義と、イネーブル（enable），有効（Valid），クロック（ck），リセット（reset）及び遅延（delay）に係る記述が含
まれる。

図６は、コンポーネントの選択に係る説明図である。ＳＰＣは、実行時において、複数のコンポーネント定義ファイルを取得及び参照して、目的のパイプラインの形成に適合する定義ファイルを選択することができる。例えば、図６では、ＳＰＣは、カウンタ（Counter）、加算器(adder)，減算器(Subtract)，積算回路(Multiplier)，遅延（Delay），ネ
ゲート(negate)のような様々なコンポーネントの定義ファイルを受け取り、その中から、演算に適合したコンポーネントを選択（抽出）し（図６では、加算器（１），減算器，及び積算回路）、これらを合成（配線接続）したパイプラインハードウェアを生成する様子が図示されている。

このとき、或る演算タイプに関して複数の定義ファイルがある場合には、複数の定義ファイルから、目的に合致した定義ファイルを選択する。例えば、図６の例では、加算器に関して二つのコンポーネント定義ファイル（adder(1), adder(2)）が得られている。ここで、例えば、最小クロックサイクル数でパイプラインを作成する（遅延を最小にする）ことが目的とされる場合には、ＳＰＣは、クロックサイクル数が少ない順で複数の定義ファイルをソートし、最小クロックサイクル数のコンポーネントの定義ファイル（図６の例では、“adder(1)”）を選択する。これによって、最短の遅延時間でパイライン処理の演算結果を出力可能なパイプラインハードウェアを作成することができる。

このように、コンポーネントの定義ファイルにおいて、同タイプのコンポーネント間で目的に応じた適用の優先順位を決定可能な情報（優先情報）が含まれることで、目的に応じたコンポーネントを自動的に選択する（特定する）ことが可能となる。

[遅延挿入（出力タイミング調整）]
ＳＰＣは、入力プログラム（カーネルプログラム）に記述されたパイプライン処理で実行される演算の構文解析を行い、二分木構造を有する抽象構文木（Abstract Syntax Tree：ＡＳＴ）を作成する。ＡＳＴの作成は、ＳＰＣによるカーネルプログラムのコンパイル時に、bisonやyaccのようなバーサジェネレータを用いて行うことができる。

図７は、抽象構文木の例を示す図であり、プログラム中の演算式“a[id] = b[id] + c[id] * d[id] - e[id]”に対する抽象構文木を例示する。図７に示すように、ＡＳＴの幹
ノードには、代入（例えば等式における等号）が指定される。代入部分（幹ノード）の左側の葉ノードには出力ポート“ａ”が指定される。一方、代入部分の右側には、葉ノードと枝ノードとが配置される。なお、ノード間を結ぶ線（リンク）は、ハードウェアにおける信号線として宣言される。

葉ノードには、入力ポート，プログラム中のローカル変数，即値が指定される。図７の例では、入力ポート“ｂ”，“ｃ”，“ｄ”及び“ｅ”の葉ノードが配置されている。枝ノードには、演算式に従った演算子が配置され、演算子に応じたコンポーネントがマッピングされる。例えば、図６において、“ｃ”及び“ｄ”を入力とする演算子“*”には、
コンポーネントとして、積算回路がマッピングされる。

また、枝ノードには、葉ノード又は下位の枝ノードの出力（演算結果）が入力される。例えば、演算子“−”の枝ノード（減算器がマッピングされる）は、入力ポート“ｅ”の値と、演算子“*”の演算結果とが入力される。そして、演算子“+”（加算器がマッピングされる）には、入力ポート“ｂ”の値と、演算子“−”の演算結果とが入力される。

このようにして、図６の下側に示すような、パイプライン処理に応じた２以上のコンポーネントとしての積算回路（積算器），減算器，及び加算器が合成されたパイプラインハードウェアのモデルが生成される。

ＡＳＴの生成に当たっては、上記した入出力に係る規定に従い、幹ノードの左側に入力ポートが現れたり、右側に出力ポートが現れたりした場合には、パイプラインが再帰構造を有することになる。この場合、ＳＰＣを実行するＣＰＵ１１は、処理を停止して、エラーを出力装置１５に出力する。

このようにして、演算の代入部分を幹ノードとし、この幹ノードの左側にパイプライン処理の出力データを示す葉ノードが置かれ、且つ幹ノードの右側が前記パイプライン処理に係る複数の入力データのそれぞれを示す葉ノードと前記複数の入力データを用いて出力データを算出するために使用される演算子を示す枝ノードとを含む二分木構造で表された抽象構文木の生成が試行され、このような抽象構文木が生成されない場合に、エラーが出力される。

これに対し、構文解析によって、入出力の規定に従ったＡＳＴが生成された場合には、ＳＰＣを実行するＣＰＵ１１は、幹ノードの右側に関し、各深さにおいて、枝ノードからの出力タイミングが同じとなるように、遅延を挿入する処理を行う。

図７は、上記した演算式“a[id] = b[id] + c[id] * d[id] - e[id]”に応じて作成さ
れたパイプラインハードウェアのモデル（ＨＡＭ）における遅延を説明する図である。各入力ポート“ｂ”〜“ｅ”のメモリ３（図２）からの読み出し時間は、それぞれ遅延時間Ｔ_adで同じであると仮定する。

そして、時間Ｔ_adの経過後に積算器（演算子“*”）が積算結果を出力するまでの遅延
時間がＴ₁で、Ｔ₁から減算器（演算子“−”）が減算結果を出力するまでの遅延時間がＴ₂で、Ｔ₂から加算器（演算子“+”）が加算結果を出力するまでの遅延時間がＴ₃である。

このとき、積算器への入力タイミングは同じであるが、減算器及び加算器では、２つのデータの入力タイミングがそれぞれ異なる。このため、２つのデータの双方が入力される
まで、演算処理を開始することができない。この結果、減算器及び加算器の少なくとも一方から正確な演算結果が出力されない虞がある。

そこで、ＳＰＣは、以下のようにして、遅延（遅延時間）を自動的に挿入する処理を行う。図８は、遅延挿入（出力タイミング調整）の説明図である。図８の上側には、図７に示した各遅延時間をＡＳＴの各リンクにあてはめた状態を示す。

これに対し、各枝ノードに至る２つの入力に関する遅延量が同じになるように、遅延を挿入する。具体的に説明すると、ＳＰＣは、（ａ）枝ノードにおける二つの入力に係る遅延の差を算出し、（ｂ）差分が生じた場合には、その差分を、遅延が小さい側に挿入する。

例えば、図８の下側のＡＳＴにおいて、演算子“+”の枝ノードに係る２つの入力に対
する各遅延はそれぞれＴ_adであるので、遅延の挿入は行われない。これに対し、演算子“−”の枝ノードに注目すると、演算子“+”側の入力に係る遅延時間は、Ｔ_ad＋Ｔ₁であるのに対し、入力ポート“ｅ”側の入力に係る遅延時間はＴ_adである。このため、差分|Ｔ₁＋Ｔ_ad−Ｔ_ad|が生じる。そこで、当該差分を、演算子“*”の枝ノードによる遅延Ｄ（“*”）として、入力ポート“ｅ”側に挿入する。

これによって、演算子“−”の枝ノードに対するデータの入力タイミングを一致させることができるので、正確な演算結果を出力することが可能となる。同様に、演算子“+”
の枝ノードに着目した場合には、入力ポート“ｂ”側の入力に関して、遅延の差分Ｄ（“−”）＝|Ｔ₂＋Ｔ（“+”）−Ｔ_ad|が挿入される。なお、演算子“+”の枝ノードから演
算結果が出力されるまでの遅延は、Ｔ（“＋”）＝Ｔ₁＋Ｔ₂＋Ｔ₃＋Ｔ_adである。

以上のように、ＡＳＴの右側において、各枝ノードへの２つの入力のそれぞれにおける遅延が均等になるように遅延が挿入される。これによって、各枝ノード、すなわちコンポーネント（演算器）における出力タイミングを一致させることが可能となることで、二つの入力のタイミングを合わせることができる。これによって、正確な演算結果が適正なタイミングで出力されるようにすることができる。

＜ＳＰＣの処理フロー＞
図９は、ＣＰＵ１１によって実行されるＳＰＣにおける処理を大略して示すフローチャートである。図９において、ＣＰＵ１１は、例えば、入力装置１４を用いた操作に応じて、ＳＰＣの実行を開始すると、補助記憶装置１３に記憶されたパイプライン処理のプログラム及びモデル定義情報（フローモデルのＸＭＬファイル）を取得する（０１）。

次に、ＣＰＵ１１は、補助記憶装置１３に記憶された、複数のコンポーネントの定義ファイル（ＸＭＬファイル）を取得する（０２）。

次に、ＣＰＵ１１は、フローモデルのＸＭＬファイル中のカーネルプログラム（OpenCL）を取り出してコンパイルする。このとき、プログラムからプロセッサＩＤ（id）を指定する構文が見つかると（０２ＡのＹｅｓ）、ＣＰＵ１１は、図４を用いて説明した手法を用いてカウンタの生成を行う（０３）。すなわち、OpenCLの“get_global_id”関数を発
見したときに、カウンタを生成する。このとき、生成されたカウンタの値は、代入先の変数を使う部分、すなわち、出力データをメモリに書き込むポートに接続する。これによって、出力データの書き込みアドレスがカウンタ値によって制御される。また、カウンタ値は、入力ポートに対するアドレスとして使用することもできる。なお、プロセッサＩＤを含む構文が発見されない場合（０２ＡのＮｏ）には、ＣＰＵ１１は、処理を０４に進める。

次に、ＣＰＵ１１は、図７を用いて説明した手法で、ＡＳＴを作成する（０４）。このとき、作成されたＡＳＴが入出力に係る規定に合致するか否かを判定する（０５）。ＡＳＴが規定に合致しない場合（０５，ＮＧ）には、エラー出力が行われ（０６）、処理が終了する。これによって、設計者（プログラマ）に対し、パイプラインの見直しを行う機会を提供することができる。

これに対し、ＡＳＴが規定に合致する場合（０５，ＯＫ）には、ＣＰＵ１１は、コンポーネント及びポートのマッピングを行う（０７）。すなわち、ＣＰＵ１１は、図６に示したように、複数のコンポーネントの定義ファイルから、ＡＳＴの生成により得られた演算子と演算タイプが合致するコンポーネントを抽出する。このとき、ＣＰＵ１１は、或る演算子について複数のコンポーネントのファイルが抽出された場合には、所定の優先順位の決定ルール（例えば、クロックサイクル数の小さい順）に従って、クロックサイクル数（すなわち遅延）が最も小さいコンポーネントを選択する。そして、選択したコンポーネントを、その定義ファイル中の情報に基づいて演算子にマッピングする。また、ＣＰＵ１１は、或る演算子について１つだけコンポーネントが抽出された場合には、そのまま当該コンポーネントを演算子に対してマッピングする処理を行う。

次に、ＣＰＵ１１は、マッピングされたコンポーネント間の配線接続を、ＡＳＴ及びコンポーネント定義ファイル内の情報に基づき行うことで、複数のコンポーネントを合成する。これによって、ＨＡＭが作成される（０８）。

ＨＡＭが作成されると、次に、ＣＰＵ１１は、遅延（タイミング）調整を行う（０９）。すなわち、ＣＰＵ１１は、図８〜図１０を用いて説明した手法で、遅延の挿入を行う。このとき、遅延は、ＤＦＦを用いた遅延フリッププロップや、ゲートを複数段つないだ遅延器のような、所定の遅延回路の記述がＨＡＭに含められる。

そして、ＣＰＵ１１は、ＨＤＬの生成及び出力処理を行う（１０）。ＨＤＬとして、例えば、ＶＨＤＬや Verilog HDLを用いることができる。ここでは、ＶＨＤＬが使用されていると仮定する。

ＣＰＵ１１は、フローモデルのファイルで定義した入出力とプログラムの情報から、ＶＨＤＬにおけるentity宣言を出力する。また、ＣＰＵ１１は、ＨＡＭにおいて使用されているコンポーネントについて、ＶＨＤＬにおけるcomponent宣言を出力する。また、ＣＰ
Ｕ１１は、ＡＳＴにおけるノード間の接続（リンク）をsignal（信号線）として宣言し、さらに、architecture宣言の中で、ＡＳＴにおけるコンポーネント（ノード）の接続状態に基づき、パイプライン構造を作成する。このようにして、パイプラインハードウェアのＶＨＤＬによる記述（Hardware Description）が生成される。当該記述は、出力装置１５にて出力（表示又は印刷）されることができる。

図１１は、ＨＤＬによるパイプラインハードウェアの記述例を示す。図１１に示す記述例は、図４の下側に示したパイプラインハードウェアのモデルに対応する。当該記述から、回路図を表すことも可能である。

＜実施形態の効果＞
実施形態によれば、パイプライン処理における演算の内容が記述されたプログラム，パイプライン処理の入出力に係る定義情報，及びコンポーネント情報を用いて、パイプラインハードウェアのＨＤＬによる記述が自動的に作成される。このため、ＧＰＵを用いたパイプライン処理のプログラムをパイプラインハードウェアのＨＤＬ記述に自動的に変換することが可能となる。これによって、ハードウェア設計の作業負担の軽減を図ることがで
きる。

また、上記したように、カウンタの生成によって、少なくともパイプラインの出力に関して簡易なアドレス制御を行うことができ、パイプラインハードウェアの回路構成の複雑化を抑えることができる。また、同タイプの複数のコンポーネントの中から、目的に沿った優先順位が最も高いコンポーネントが選択されるようにすることで、例えば、パイプライン処理に要する時間を最短にすることができる。また、コンポーネントへの２つの入力間で遅延が均等になるように遅延の挿入が行われることで、正確な演算結果が得られるようにすることができる。

このように、実施形態に係るハードウェア設計装置を用いることで、ＧＰＵを用いたシステムよりも回路規模が小さい、すなわち、ＧＰＵシステムよりも小型化及び簡易化されたパイプラインハードウェアを容易に設計することが可能となる。従って、ストリームデータ，及びストリームデータに対する処理の内容に特化した、様々なパイプラインハードウェアを容易に得ることが可能となる。そして、設計されたパイプラインハードウェアの回路規模は小さいため、その適用に係るコストを抑えることができる。

なお、実施形態では、１つの演算式に基づくパイプライン処理を例示したが、２以上の演算式に基づくパイプライン処理に対しても、本実施形態に係るハードウェア設計装置を適用することができる。この場合、或る演算式の次の演算式において、或る演算式の出力が次の演算式の入力として扱われる。このように、ハードウェア設計装置は、複数の演算式を含むパイプライン処理に適用可能である。もちろん、１つの演算式に対するパイプラインハードウェアを直列に接続して、複数の演算式に応じたパイプライン処理を行うパイプラインハードウェアを構築することも可能である。

１０・・・情報処理装置（コンピュータ）
１１・・・ＣＰＵ（制御装置）
１２・・・主記憶装置
１３・・・補助記憶装置
１４・・・入力装置
１５・・・出力装置

Claims

ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、前記パイプライン処理で行われる演算が記述されたプログラムと、前記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、前記演算の内容及び前記定義情報に応じた２以上のハードウェアコンポーネントを前記コンポーネント情報に基づき特定し、前記２以上のハードウェアコンポーネントが前記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成する制御装置を含み、
前記プログラムは、前記演算を行うプロセッサを識別子の記述により指定可能なストリームコンピューティング向けのプログラム言語で記述され、
前記制御装置は、前記プログラム中の演算に関して識別子の指定が記述されている場合には、当該演算によって得られる出力データの書き込みアドレスを示すカウンタ値を出力するカウンタを含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成し、
前記カウンタから出力されるカウンタ値は、前記演算の実行毎に変更される
ことを特徴とするハードウェア設計装置。
前記制御装置は、前記プログラムに記述された前記パイプライン処理で実行される演算の構文解析を行い、前記演算の代入部分を幹ノードとし、この幹ノードの左側に前記パイプライン処理の出力データを示す葉ノードが置かれ、且つ幹ノードの右側が前記パイプライン処理に係る複数の入力データのそれぞれを示す葉ノードと前記複数の入力データを用いて前記出力データを算出するために使用される演算子を示す枝ノードとを含む二分木構造で表された抽象構文木の生成を試行し、前記抽象構文木が生成されないときにエラーを出力する
請求項１に記載のハードウェア設計装置。
前記制御装置は、前記抽象構文木が生成されたときに、前記枝ノードに対する２つの入力のそれぞれにおける遅延が均等になるように遅延を挿入し、挿入した遅延を含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成する
請求項２に記載のハードウェア設計装置。
前記制御装置は、前記コンポーネント情報が同一の演算タイプを有する複数のハードウェアコンポーネントの定義を含むときに、演算による遅延が最も小さいハードウェアコンポーネントを前記２以上のハードウェアコンポーネントの１つとして選択する
請求項１から３のいずれか１項に記載のハードウェア設計装置。
ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、前記パイプライン処理で行われる演算が記述されたプログラムと、前記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、前記演算の内容及び前記定義情報に応じた２以上のハードウェアコンポーネントを前記コンポーネント情報に基づき特定するステップと、
特定された２以上のハードウェアコンポーネントが前記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成するステップとをコンピュータに実行させるハードウェア設計用プログラムにおいて、
前記プログラムは、前記演算を行うプロセッサを識別子の記述により指定可能なストリームコンピューティング向けのプログラム言語で記述され、
前記制御装置は、前記プログラム中の演算に関して識別子の指定が記述されている場合には、当該演算によって得られる出力データの書き込みアドレスを示すカウンタ値を出力するカウンタを含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成するステップを前記コンピュータに実行させ、
前記カウンタから出力されるカウンタ値は、前記演算の実行毎に変更される
ことを特徴とするハードウェア設計用プログラム。
前記プログラムに記述された前記パイプライン処理で実行される演算の構文解析を行うステップと、
前記演算の代入部分を幹ノードとし、この幹ノードの左側に前記パイプライン処理の出力データを示す葉ノードが置かれ、且つ幹ノードの右側が前記パイプライン処理に係る複数の入力データのそれぞれを示す葉ノードと前記複数の入力データを用いて前記出力データを算出するために使用される演算子を示す枝ノードとを含む二分木構造で表された抽象構文木の生成を試行するステップと、
前記抽象構文木が生成されないときにエラーを出力するステップと
を前記コンピュータに実行させる請求項５に記載のハードウェア設計用プログラム。
前記抽象構文木が生成されたときに、前記枝ノードに対する２つの入力のそれぞれにおける遅延が均等になるように遅延を挿入するステップと、
挿入した遅延を含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成するステップと
を前記コンピュータに実行させる請求項６に記載のハードウェア設計用プログラム。
前記コンポーネント情報が同一の演算タイプを有する複数のハードウェアコンポーネントの定義を含むときに、演算による遅延が最も小さいハードウェアコンポーネントを前記２以上のハードウェアコンポーネントの１つとして選択するステップ
を前記コンピュータに実行させる請求項５から７のいずれか１項に記載のハードウェア設計用プログラム。