TOP > 国内特許検索 > ハードウェア設計装置,及びハードウェア設計用プログラム > 明細書

明細書 :ハードウェア設計装置,及びハードウェア設計用プログラム

発行国 日本国特許庁(JP)
公報種別 特許公報(B2)
特許番号 特許第6249360号 (P6249360)
公開番号 特開2014-225194 (P2014-225194A)
登録日 平成29年12月1日(2017.12.1)
発行日 平成29年12月20日(2017.12.20)
公開日 平成26年12月4日(2014.12.4)
発明の名称または考案の名称 ハードウェア設計装置,及びハードウェア設計用プログラム
国際特許分類 G06F  17/50        (2006.01)
H01L  21/82        (2006.01)
FI G06F 17/50 654A
H01L 21/82 C
請求項の数または発明の数 8
全頁数 18
出願番号 特願2013-105024 (P2013-105024)
出願日 平成25年5月17日(2013.5.17)
審査請求日 平成28年5月16日(2016.5.16)
特許権者または実用新案権者 【識別番号】504171134
【氏名又は名称】国立大学法人 筑波大学
発明者または考案者 【氏名】山際 伸一
個別代理人の代理人 【識別番号】100100549、【弁理士】、【氏名又は名称】川口 嘉之
【識別番号】100123319、【弁理士】、【氏名又は名称】関根 武彦
【識別番号】100105407、【弁理士】、【氏名又は名称】高田 大輔
審査官 【審査官】合田 幸裕
参考文献・文献 特表2006-505061(JP,A)
特開2010-218217(JP,A)
特開2005-018626(JP,A)
横山 正幸, 戸川 望, 柳澤 政生, 大附 辰夫,制御処理を主体としたハードウェア記述生成手法,DAシンポジウム’99,日本,社団法人情報処理学会,1999年 7月15日,第99巻/第8号,第195-200頁
山際 伸一,ストリームコンピューティング方式のプログラミングを利用したハードウェア合成,電子情報通信学会2011年基礎・境界ソサイエティ大会講演論文集,日本,社団法人電子情報通信学会,2011年 8月30日,SS-64~SS-65
調査した分野 G06F 17/50
H01L 21/82
IEEE Xplore
JSTPlus(JDreamIII)
特許請求の範囲 【請求項1】
ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、前記パイプライン処理で行われる演算が記述されたプログラムと、前記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、前記演算の内容及び前記定義情報に応じた2以上のハードウェアコンポーネントを前記コンポーネント情報に基づき特定し、前記2以上のハードウェアコンポーネントが前記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成する制御装置を含み、
前記プログラムは、前記演算を行うプロセッサを識別子の記述により指定可能なストリームコンピューティング向けのプログラム言語で記述され、
前記制御装置は、前記プログラム中の演算に関して識別子の指定が記述されている場合には、当該演算によって得られる出力データの書き込みアドレスを示すカウンタ値を出力するカウンタを含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成し、
前記カウンタから出力されるカウンタ値は、前記演算の実行毎に変更される
ことを特徴とするハードウェア設計装置。
【請求項2】
前記制御装置は、前記プログラムに記述された前記パイプライン処理で実行される演算の構文解析を行い、前記演算の代入部分を幹ノードとし、この幹ノードの左側に前記パイプライン処理の出力データを示す葉ノードが置かれ、且つ幹ノードの右側が前記パイプライン処理に係る複数の入力データのそれぞれを示す葉ノードと前記複数の入力データを用いて前記出力データを算出するために使用される演算子を示す枝ノードとを含む二分木構造で表された抽象構文木の生成を試行し、前記抽象構文木が生成されないときにエラーを出力する
請求項1に記載のハードウェア設計装置。
【請求項3】
前記制御装置は、前記抽象構文木が生成されたときに、前記枝ノードに対する2つの入力のそれぞれにおける遅延が均等になるように遅延を挿入し、挿入した遅延を含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成する
請求項に記載のハードウェア設計装置。
【請求項4】
前記制御装置は、前記コンポーネント情報が同一の演算タイプを有する複数のハードウェアコンポーネントの定義を含むときに、演算による遅延が最も小さいハードウェアコンポーネントを前記2以上のハードウェアコンポーネントの1つとして選択する
請求項1からのいずれか1項に記載のハードウェア設計装置。
【請求項5】
ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、前記パイプライン処理で行われる演算が記述されたプログラムと、前記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、前記演算の内容及び前記定義情報に応じた2以上のハードウェアコンポーネントを前記コンポーネント情報に基づき特定するステップと、
特定された2以上のハードウェアコンポーネントが前記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成するステップとをコンピュータに実行させるハードウェア設計用プログラムにおいて、
前記プログラムは、前記演算を行うプロセッサを識別子の記述により指定可能なストリームコンピューティング向けのプログラム言語で記述され、
前記制御装置は、前記プログラム中の演算に関して識別子の指定が記述されている場合には、当該演算によって得られる出力データの書き込みアドレスを示すカウンタ値を出力するカウンタを含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成するステップを前記コンピュータに実行させ、
前記カウンタから出力されるカウンタ値は、前記演算の実行毎に変更される
ことを特徴とするハードウェア設計用プログラム。
【請求項6】
前記プログラムに記述された前記パイプライン処理で実行される演算の構文解析を行うステップと、
前記演算の代入部分を幹ノードとし、この幹ノードの左側に前記パイプライン処理の出力データを示す葉ノードが置かれ、且つ幹ノードの右側が前記パイプライン処理に係る複数の入力データのそれぞれを示す葉ノードと前記複数の入力データを用いて前記出力データを算出するために使用される演算子を示す枝ノードとを含む二分木構造で表された抽象構文木の生成を試行するステップと、
前記抽象構文木が生成されないときにエラーを出力するステップと
を前記コンピュータに実行させる請求項に記載のハードウェア設計用プログラム。
【請求項7】
前記抽象構文木が生成されたときに、前記枝ノードに対する2つの入力のそれぞれにおける遅延が均等になるように遅延を挿入するステップと、
挿入した遅延を含む前記パイプラインハードウェアのハードウェア記述言語による記述を生成するステップと
を前記コンピュータに実行させる請求項に記載のハードウェア設計用プログラム。
【請求項8】
前記コンポーネント情報が同一の演算タイプを有する複数のハードウェアコンポーネントの定義を含むときに、演算による遅延が最も小さいハードウェアコンポーネントを前記2以上のハードウェアコンポーネントの1つとして選択するステップ
を前記コンピュータに実行させる請求項5から7のいずれか1項に記載のハードウェア設計用プログラム。
発明の詳細な説明 【技術分野】
【0001】
本発明は、ハードウェア設計装置,及びハードウェア設計用プログラムに関する。
【背景技術】
【0002】
現在、ネットワークには、様々なセンサ及びデバイスが接続され、これらのセンサ及びデバイスから刻々と出力されるデータがネットワーク上でデータストリームを形成する。データストリームを形成するデータ(ストリームデータと呼ばれる)を滞りなく処理する(リアルタイムに処理する)手法として、ストリームコンピューティングがある。
【0003】
ストリームコンピューティングでは、ストリームデータに関して所定のデータ単位が入力として決定され、このデータ単位に対する複数の処理(演算)が直列に実行されるパイプライン処理が行われる。
【0004】
パイプライン処理を実行する典型的な環境として、GPU(Graphical Processing Unit)を用いたソフトウェア処理がある。GPUは、並列処理を実行可能な複数のプロセッ
サを含んでおり、パイプライン処理の手順を記述したプログラムに従って、各プロセッサにパイプライン処理に係る複数の処理を実行させることで、パイプライン処理結果を得ることができる。
【0005】
OpenCLのような、GPU向けのプログラミング言語では、GPUが並列処理を実行可能な複数のプロセッサを有することを考慮した、パイプライン処理手順を記述することができる。例えば、プログラム上で、パイプライン処理に係る複数の演算を、複数のプロセッサに割り当てることができる。
【先行技術文献】
【0006】

【特許文献1】特開2012-174268号公報
【0007】

【非特許文献1】“OpenCL規格を用いたFPGAデザインの導入”、[online]、アルテラ、[平成25年5月13日検索]、インターネット<URL: http://www.altera.co.jp/literature/wp/wp-01173-opencl_j.pdf>
【非特許文献2】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.
【非特許文献3】Shinichi Yamagiwa, leonel Sousa, "Caravela: A Novel Stream-Based Distributed Computing Environment," Computer , vol.40, no.5, pp.70-77, May 2007.
【発明の概要】
【発明が解決しようとする課題】
【0008】
しかしながら、GPU及びGPUを用いたシステムは、回路規模が大きく、またコストの上昇を招来する。このため、パイプライン処理を行うために、GPU乃至GPUシステムを導入することが困難、又は非現実的である場合が少なくない。
【0009】
このため、ストリームデータ,ストリームデータに対して行う処理(演算)の内容などに応じたパイプライン処理を実行するハードウェア(パイプライン処理用のディジタル回
路:パイプラインハードウェアと呼ぶ)の開発が求められている。
【0010】
パイプラインハードウェアを設計する際には、パイプライン処理で実行される複数の演算(四則演算等)をそれぞれ実行する複数の演算コンポーネント(「ハードウェアコンポーネント」と呼ぶ。以下単に「コンポーネント」と表記)が用意され、演算の順序に従って複数のコンポーネントが配線により接続される。パイプラインハードウェアにはクロックが入力され、クロック毎に演算が進められる。演算に係る遅延(ディレイ)はクロック数として定義される。
【0011】
従来におけるハードウェア記述言語(hardware description language:HDL)を用いたパイプラインハードウェアの設計では、設計者がパイプライン処理における演算内容に応じて複数のコンポーネントを選択し、選択したコンポーネント間の配線接続によってパイプラインハードウェアを設計する。このとき、設計者がコンポーネント間のデータ到達タイミングを計りながら、最終的な演算結果が正しく出力されるための遅延(ディレイ)をマニュアル操作で挿入する。
【0012】
コンポーネントは様々な遅延を持つ。例えば、加算器は、実装技術(LSIのプロセス)の違いにより遅延が異なる。さらに、積算回路については繰り返し処理が伴うことで遅延が可変となることがある。このように、パイプラインハードウェアの設計では、各コンポーネントの遅延を考慮することにより、全体の時間バランスを考慮することが要求される。
【0013】
しかしながら、例えば、コンポーネントの入れ替え(選択変更)により、コンポーネント自体の遅延が変化した場合には、この変化をコンポーネント間に挿入されたディレイで吸収できない場合が起こり得る。この影響がパイプライン全体に及ぶ場合には、パイプラインハードウェアの設計を最初からやり直すことが要求される虞があった。このように、設計者に対して、コンポーネントの選択と、挿入する遅延の考慮とを慎重に行うことが要求され、これは設計者にとって大きな負担であった。
【0014】
現在のところ、パイプライン処理の演算内容に応じたパイプラインハードウェアのひな形となるモデルを自動的に生成する機構はない。このため、コンポーネントの特性と、コンポーネントのパイプラインへの適用を、目的とするパイプラインハードウェアに応じて一意に決定可能な機構、特に、コンポーネントを考慮した遅延を自動的に挿入可能な機構もない。
【0015】
本発明は、上記の事情に鑑みなされたものであり、パイプライン処理の演算内容に応じたパイプラインハードウェアのモデルを自動的に生成可能な技術を提供することを目的とする。
【課題を解決するための手段】
【0016】
本発明は、上記課題を解決するために以下の手段を採用する。すなわち、本発明は、ストリームデータを処理するためのパイプライン処理に使用可能な複数のハードウェアコンポーネントの定義を含むコンポーネント情報と、前記パイプライン処理で行われる演算が記述されたプログラムと、前記演算における入力及び出力の定義を少なくとも含む定義情報とを用いて、前記演算の内容及び前記定義情報に応じた2以上のハードウェアコンポーネントを前記コンポーネント情報に基づき特定し、前記2以上のハードウェアコンポーネントが前記パイプライン処理を行うように合成されたパイプラインハードウェアのハードウェア記述言語による記述を生成する制御装置を含むハードウェア設計装置である。
【0017】
また、本発明は、上記した制御装置における処理をコンピュータに実行させるハードウ
ェア設計用プログラム,このようなプログラムを記録したコンピュータ読み取り可能な記録媒体,及び上記制御装置によるパイプラインハードウェア記述の生成方法としても特定することができる。
【発明の効果】
【0018】
本発明によれば、パイプライン処理の演算内容に応じたパイプラインハードウェアのモデルを自動的に生成可能な技術を提供することができる。
【図面の簡単な説明】
【0019】
【図1】図1は、ハードウェア設計用プログラムの実行によってハードウェア設計装置として機能する情報処理装置(コンピュータ)の構成例を示す図である。
【図2】図2は、SPCの仕組みを説明する図である。
【図3】図3は、XMLで記述されたフローモデルの例を示す。
【図4】図4は、カウンタ生成に係る処理を説明する図である。
【図5】図5は、コンポーネントの定義ファイルの例を示す。
【図6】図6は、コンポーネントの選択に係る説明図である。
【図7】図7は、抽象構文木の例を示す図である。
【図8】図8は、演算式“a[id] = b[id] + c[id] * d[id] - e[id]”に応じて作成されたパイプラインハードウェアのモデル(HAM)における遅延を説明する図である。
【図9】図9は、遅延挿入の例を示す説明図である。
【図10】図10は、CPUによって実行されるSPCにおける処理を大略して示すフローチャートである。
【図11】図11は、HDLによるパイプラインハードウェアの記述例を示す。
【発明を実施するための形態】
【0020】
以下、図面を参照して本発明の実施形態について説明する。実施形態の構成及び設定は例示であり、本発明は実施形態の構成及び設定に限定されない。

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

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

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

【0024】
<情報処理装置(ハードウェア設計装置)>
図1は、ハードウェア設計用プログラムの実行によってハードウェア設計装置として機能する情報処理装置(コンピュータ)の構成例を示す図である。情報処理装置10として、例えば、パーソナルコンピュータ(PC),ワークステーション,専用又は汎用のサーバマシンを適用することができる。

【0025】
図1において、情報処理装置10は、例として、バスBを介して相互に接続された、CPU11と、主記憶装置12と、補助記憶装置13と、入力装置14と、出力装置15と、通信インタフェース回路(通信I/F)16とを備える。

【0026】
主記憶装置12は、CPU11の作業領域として使用されるメインメモリとして機能する。メインメモリは、例えば、RAM(Random Access Memory)及びROM(Read Only Memory)によって形成される。

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

【0028】
入力装置14は、キーボード,マウスやタッチパネルのようなポインティングデバイスを含み、情報(データ)の入力に使用される。出力装置15は、例えば、ディスプレイ装置であり、情報を画面に表示する。通信I/F16は、ネットワークとの通信処理を司る。

【0029】
SPCは、補助記憶装置13にインストールされており、CPU11がSPCを主記憶装置12に読み出し、ロードして実行することによって、情報処理装置10は、ハードウェア設計装置として機能することができる。

【0030】
<SPC>
図2は、SPCの仕組みを説明する図である。図2に示すように、SPCに対する入力として、“Flow-model(フローモデル)”と呼ばれる、パイプライン処理のモデルを定義したモデル定義情報と、パイプライン処理で行われる演算が記述されたプログラムとが所定の記述言語で記述されたファイル(パイプライン処理モデルのファイル)が用意される。各ファイルは、補助記憶装置13に記憶され、CPU11によるSPCの実行に際して使用される。

【0031】
SPCは、フローモデルのファイルを入力として、フローモデルにおいて定義されたパイプライン処理を行うパイプラインハードウェアの抽象化モデル(Hardware Abstraction
Model:HAM)を生成する(図2(1))。抽象化モデル(HAM)は、フローモデルにおける演算内容に従った演算を行う演算子である2以上のコンポーネントを特定し、2以上のコンポーネントを合成することによって生成される。さらに、SPCは、HAMをハードウェア記述言語(HDL)で記述したパイプラインハードウェアの記述(Hardware
Description)を生成して出力する(図2(2))。

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

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

【0034】
<<フローモデル>>
フローモデルの記述言語として、例えば、XML(Extensible Markup Language)を含む様々なマークアップ言語を適用可能である。

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

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

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

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

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

【0040】
図3は、XMLで記述されたフローモデルの例を示す。図3において、カーネルのタグ(<kernel> </kernel>)で挟まれた部分が、OpenCLのプログラム、すなわちカーネルプログラムの記述である。そして、カーネルプログラムに記述された引数“a”及び“b”を入力(input)とし、演算結果の“c”を出力(output)とするモデル定義情報が記述さ
れている。図2の例では、“a”,“b”,“c”のそれぞれの定義(名称,データタイプ,データ長等)が記述されている。

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

【0042】
<<パイプラインハードウェアの抽象化モデル(HAM)>>
次に、パイプラインハードウェアの抽象化モデル(HAM)の詳細について説明する。SPCによって生成されるHAMは、以下のように規定(設定)される。
・入力は、“(入力)データ”,“(入力)データのアドレス”,“クロック”,及び“
リセット”を含む。
・出力は、“(出力)データ”,“(出力)データの有効性を示すValid信号”,“(出
力)データのアドレス”を含む。
・パイプラインハードウェアに対する入力は、メモリ(図1のメモリ3を参照)からの読み出しによって行われる。
・出力データは、メモリ(図1のメモリ4を参照)に書き込まれる。
・入力時における“(入力)データのアドレス”から読み出された“(入力)データ”のメモリからの読み出し時間“Tad”は、SPCの合成変数として与えられる。
・パイプラインハードウェアへの入力データを“入力ポート”,パイプラインハードウェアからの出力データを“出力ポート”と呼ぶ。

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

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

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

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

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

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

【0049】
図4において、カウンタ値は、入力ポート“a”及び“b”、並びに出力ポート“c”に対するアドレスとして供給される。“a”のアドレス“a_addr”及び“b”のアドレス
“b_addr”は、入力ポート側にあるメモリ3(図2)のアドレスとして機能する。そして、アドレスとして指定された箇所から読み出された“a”の値は、即値“100”とともに、減算器(Subtract)に入力される。“b”の値は、減算器の出力とともに、加算器(adder)に入力される。そして、加算器から出力ポートである“c”の値が出力される。
“c”の値は、アドレス“c_addr”として指定されたカウンタ値に対応する、出力ポート側のメモリ4(図2)の記憶領域に書き込まれる。そして、所定の契機(例えば、“c”の出力“書き込み”)で、カウンタ値がインクリメント(現在のカウンタ値に1を加算)される。そして、次のカウンタ値に対応するメモリ3の記憶領域から、入力ポート“a”及び“b”を取得し、次のカウンタ値に対応するメモリ4の記憶領域に出力ポート“b”を書き込むことができる。

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

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

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

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

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

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

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

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

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

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

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

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

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

【0063】
また、枝ノードには、葉ノード又は下位の枝ノードの出力(演算結果)が入力される。例えば、演算子“-”の枝ノード(減算器がマッピングされる)は、入力ポート“e”の値と、演算子“*”の演算結果とが入力される。そして、演算子“+”(加算器がマッピングされる)には、入力ポート“b”の値と、演算子“-”の演算結果とが入力される。

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

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

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

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

【0068】
図7は、上記した演算式“a[id] = b[id] + c[id] * d[id] - e[id]”に応じて作成さ
れたパイプラインハードウェアのモデル(HAM)における遅延を説明する図である。各入力ポート“b”~“e”のメモリ3(図2)からの読み出し時間は、それぞれ遅延時間Tadで同じであると仮定する。

【0069】
そして、時間Tadの経過後に積算器(演算子“*”)が積算結果を出力するまでの遅延
時間がT1で、T1から減算器(演算子“-”)が減算結果を出力するまでの遅延時間がT2で、T2から加算器(演算子“+”)が加算結果を出力するまでの遅延時間がT3である。

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

【0071】
そこで、SPCは、以下のようにして、遅延(遅延時間)を自動的に挿入する処理を行う。図8は、遅延挿入(出力タイミング調整)の説明図である。図8の上側には、図7に示した各遅延時間をASTの各リンクにあてはめた状態を示す。

【0072】
これに対し、各枝ノードに至る2つの入力に関する遅延量が同じになるように、遅延を挿入する。具体的に説明すると、SPCは、(a)枝ノードにおける二つの入力に係る遅延の差を算出し、(b)差分が生じた場合には、その差分を、遅延が小さい側に挿入する。

【0073】
例えば、図8の下側のASTにおいて、演算子“+”の枝ノードに係る2つの入力に対
する各遅延はそれぞれTadであるので、遅延の挿入は行われない。これに対し、演算子“-”の枝ノードに注目すると、演算子“+”側の入力に係る遅延時間は、Tad+T1であるのに対し、入力ポート“e”側の入力に係る遅延時間はTadである。このため、差分|T1+Tad-Tad|が生じる。そこで、当該差分を、演算子“*”の枝ノードによる遅延D(“*”)として、入力ポート“e”側に挿入する。

【0074】
これによって、演算子“-”の枝ノードに対するデータの入力タイミングを一致させることができるので、正確な演算結果を出力することが可能となる。同様に、演算子“+”
の枝ノードに着目した場合には、入力ポート“b”側の入力に関して、遅延の差分D(“-”)=|T2+T(“+”)-Tad|が挿入される。なお、演算子“+”の枝ノードから演
算結果が出力されるまでの遅延は、T(“+”)=T1+T2+T3+Tadである。

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

【0076】
<SPCの処理フロー>
図9は、CPU11によって実行されるSPCにおける処理を大略して示すフローチャートである。図9において、CPU11は、例えば、入力装置14を用いた操作に応じて、SPCの実行を開始すると、補助記憶装置13に記憶されたパイプライン処理のプログラム及びモデル定義情報(フローモデルのXMLファイル)を取得する(01)。

【0077】
次に、CPU11は、補助記憶装置13に記憶された、複数のコンポーネントの定義ファイル(XMLファイル)を取得する(02)。

【0078】
次に、CPU11は、フローモデルのXMLファイル中のカーネルプログラム(OpenCL)を取り出してコンパイルする。このとき、プログラムからプロセッサID(id)を指定する構文が見つかると(02AのYes)、CPU11は、図4を用いて説明した手法を用いてカウンタの生成を行う(03)。すなわち、OpenCLの“get_global_id”関数を発
見したときに、カウンタを生成する。このとき、生成されたカウンタの値は、代入先の変数を使う部分、すなわち、出力データをメモリに書き込むポートに接続する。これによって、出力データの書き込みアドレスがカウンタ値によって制御される。また、カウンタ値は、入力ポートに対するアドレスとして使用することもできる。なお、プロセッサIDを含む構文が発見されない場合(02AのNo)には、CPU11は、処理を04に進める。

【0079】
次に、CPU11は、図7を用いて説明した手法で、ASTを作成する(04)。このとき、作成されたASTが入出力に係る規定に合致するか否かを判定する(05)。ASTが規定に合致しない場合(05,NG)には、エラー出力が行われ(06)、処理が終了する。これによって、設計者(プログラマ)に対し、パイプラインの見直しを行う機会を提供することができる。

【0080】
これに対し、ASTが規定に合致する場合(05,OK)には、CPU11は、コンポーネント及びポートのマッピングを行う(07)。すなわち、CPU11は、図6に示したように、複数のコンポーネントの定義ファイルから、ASTの生成により得られた演算子と演算タイプが合致するコンポーネントを抽出する。このとき、CPU11は、或る演算子について複数のコンポーネントのファイルが抽出された場合には、所定の優先順位の決定ルール(例えば、クロックサイクル数の小さい順)に従って、クロックサイクル数(すなわち遅延)が最も小さいコンポーネントを選択する。そして、選択したコンポーネントを、その定義ファイル中の情報に基づいて演算子にマッピングする。また、CPU11は、或る演算子について1つだけコンポーネントが抽出された場合には、そのまま当該コンポーネントを演算子に対してマッピングする処理を行う。

【0081】
次に、CPU11は、マッピングされたコンポーネント間の配線接続を、AST及びコンポーネント定義ファイル内の情報に基づき行うことで、複数のコンポーネントを合成する。これによって、HAMが作成される(08)。

【0082】
HAMが作成されると、次に、CPU11は、遅延(タイミング)調整を行う(09)。すなわち、CPU11は、図8~図10を用いて説明した手法で、遅延の挿入を行う。このとき、遅延は、DFFを用いた遅延フリッププロップや、ゲートを複数段つないだ遅延器のような、所定の遅延回路の記述がHAMに含められる。

【0083】
そして、CPU11は、HDLの生成及び出力処理を行う(10)。HDLとして、例えば、VHDLや Verilog HDLを用いることができる。ここでは、VHDLが使用されていると仮定する。

【0084】
CPU11は、フローモデルのファイルで定義した入出力とプログラムの情報から、VHDLにおけるentity宣言を出力する。また、CPU11は、HAMにおいて使用されているコンポーネントについて、VHDLにおけるcomponent宣言を出力する。また、CP
U11は、ASTにおけるノード間の接続(リンク)をsignal(信号線)として宣言し、さらに、architecture宣言の中で、ASTにおけるコンポーネント(ノード)の接続状態に基づき、パイプライン構造を作成する。このようにして、パイプラインハードウェアのVHDLによる記述(Hardware Description)が生成される。当該記述は、出力装置15にて出力(表示又は印刷)されることができる。

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

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

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

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

【0089】
なお、実施形態では、1つの演算式に基づくパイプライン処理を例示したが、2以上の演算式に基づくパイプライン処理に対しても、本実施形態に係るハードウェア設計装置を適用することができる。この場合、或る演算式の次の演算式において、或る演算式の出力が次の演算式の入力として扱われる。このように、ハードウェア設計装置は、複数の演算式を含むパイプライン処理に適用可能である。もちろん、1つの演算式に対するパイプラインハードウェアを直列に接続して、複数の演算式に応じたパイプライン処理を行うパイプラインハードウェアを構築することも可能である。
【符号の説明】
【0090】
10・・・情報処理装置(コンピュータ)
11・・・CPU(制御装置)
12・・・主記憶装置
13・・・補助記憶装置
14・・・入力装置
15・・・出力装置
図面
【図1】
0
【図2】
1
【図3】
2
【図4】
3
【図5】
4
【図6】
5
【図7】
6
【図8】
7
【図9】
8
【図10】
9
【図11】
10