JP5885481B2 - 情報処理装置、情報処理方法、及びプログラム - Google Patents

Description

本発明は、リダクション処理を行う並列計算機に対する命令を生成する情報処理装置、情報処理方法、及びプログラムに関する。

近年、ＣＰＵコアを複数用いて計算機の処理能力を向上させるアプローチが行われている。特に、ＧＰＵ（Ｇｒａｐｈｉｃｓ Ｐｒｏｃｅｓｓｉｎｇ Ｕｎｉｔ）にグラフィクス処理以外の処理を行わせる、ＧＰＧＰＵ（Ｇｅｎｅｒａｌ Ｐｕｒｐｏｓｅ ＧＰＵ）又はＧＰＵ Ｃｏｍｐｕｔｉｎｇが脚光を浴びている。ＧＰＵは数１０〜１０００個を超える多数の演算コアを持っており、全ての演算コアを動作させた場合のピーク性能は非常に高い。しかしながら、ＧＰＵの持つ高い性能を発揮するには、従来とは異なるプログラミング技法が必要になる。以下ではＧＰＧＰＵとしてはＮＶＩＤＩＡ社のＣＵＤＡを例にとり説明を行うが、詳細については非特許文献１で述べられているため、割愛する。

ＧＰＧＰＵは通常、ＳＰＭＤ（Ｓｉｎｇｌｅ Ｐｒｏｇｒａｍ−Ｍｕｌｔｉｐｌｅ Ｄａｔａ）で動作する。従って、同一プログラム（カーネル）が各々のスレッドで同時に実行される。ＧＰＧＰＵの演算能力は、多くの演算コアに休みなく処理をさせつづけることにより、より向上する。ところで多くのアプリケーションは、並列処理を行った後に各スレッドで行った演算結果をひとつにまとめ上げる処理が必要とする。こうした処理のうち、よく用いられるものとして並列リダクション処理が挙げられる。並列リダクション処理においては、複数のデータが徐々にまとめられて処理結果が得られる。このとき、データがまとめられるにつれて並列リダクション処理に参加するスレッドの数はだんだん減っていく。すなわち、何も行っていないスレッド（アイドルコア）が増えていくため、処理資源が無駄になる。並列リダクションの例としては非特許文献２に詳しいので、ここでは割愛する。

さらに並列リダクション処理においては、各スレッド間での通信が発生する。この通信を共有メモリを介して行う場合、複数のスレッドが同時に通信を行うため、アクセスのコンフリクトが発生する。コンフリクトしたアクセスは順に処理され、処理が完了するまで他のアクセスは待たされることになるため、大きく処理速度が低下してしまう。

特許文献１には、並列に動作可能な複数の演算器を備える計算機上で実行されるプログラムをコンパイルする方法が開示されている。特許文献１の方法によれば、ある命令を出す場合に見積もられた使用レジスタ数が、使用可能なレジスタ数よりも大きい場合に、同時にアクティブなレジスタの数を減らすために、その命令が別の命令に変更される。

特許第３３１１３８１号公報

ＮＶＩＤＩＡ ＣＵＤＡＴＭ ＮＶＩＤＩＡ ＣＵＤＡ Ｃ Ｐｒｏｇｒａｍｍｉｎｇ Ｇｕｉｄｅ Ｖｅｒｓｉｏｎ３．１．１ ７／２１／２０１０ ＣＵＤＡ Ｔｅｃｈｎｉｃａｌ Ｔｒａｉｎｉｎｇ Ｖｏｌｕｍｅ ＩＩ：ＣＵＤＡ Ｃａｓｅ Ｓｔｕｄｉｅｓ Ｑ２ ２００８

しかしながら特許文献１に記載の技術は、ＧＰＧＰＵのように複数のコアがＳＰＭＤで動作する場合が考慮されていない。すなわち、特許文献１に記載の技術によれば、複数のコアが異なる命令に従って動作するように指示される。しかしながらこのような動作を行わないＧＰＧＰＵにおいては、特許文献１に記載の技術に従うとかえって動作速度が低下することが考えられる。

本発明は、複数のスレッドが互いに通信しながら演算を行うシステムにおいて、メモリアクセスのコンフリクトを減らしながら演算コアの利用率を高め、演算速度を向上させることを目的とする。

本発明の目的を達成するために、例えば、本発明の情報処理装置は以下の構成を備える。すなわち、
複数のバンクで構成されるメモリに接続された複数の演算コアに対する命令を生成する情報処理装置であって、
前記複数の演算コアは、
前記命令によって指定された演算コアが互いに同期して、
初期データを前記演算コアが保持するレジスタへと読み込む読込サイクルと、
前記演算コアごとに予め対応付けられたバンク内の領域から読み込んだデータと、前記演算コアが保持するレジスタ内のデータとを用いて演算を行い、演算結果を前記演算コアが保持するレジスタに格納する演算サイクルと、
前記メモリへ前記演算コアが保持するレジスタ内のデータを書き込む書込サイクルとを用いるものであり、
前記情報処理装置は、
１回の前記書込サイクルにおいて演算コアがメモリへの書き込みアクセスを行う回数に少なくとも関連する値をそれぞれのバンクごとに保持する保持手段と、
前記複数の演算コアのうち、対応するバンクについての前記保持手段の保持する値が所定値未満である演算コアを、他の演算コアからデータを受信する継続演算コアとして選択する第１の選択手段と、
前記継続演算コアに対してデータを送信する送信演算コアを選択する第２の選択手段と、
前記送信演算コアに、前記書込サイクルにおいて、レジスタ内のデータを前記継続演算コアに対応付けられたバンク内の領域に書き込ませるための命令と、前記継続演算コアに、当該書込サイクルに続く前記演算サイクルにおいて、当該継続演算コアに対応付けられたバンク内の領域からデータを読み込ませるための命令と、を生成する生成手段と、
前記第１の選択手段が前記継続演算コアを選択する際に、選択された前記継続
演算コアに予め対応付けられたバンクについての前記保持手段の保持する値をインクリメントする更新手段と、
を備えることを特徴とする。

本発明によれば、複数のスレッドが互いに通信しながら演算を行うシステムにおいて、メモリアクセスのコンフリクトを減らしながら演算コアの利用率を高め、演算速度を向上させることができる。

実施例の結果を用いて動作するシステムの構成例。 実施例の結果を用いて動作するシステムを説明する図。 実施例の結果を用いて動作するシステムを説明する図。 実施例の結果を用いて動作するシステムを説明する図。 実施例の結果を用いて動作するシステムを説明する図。 実施例のスケジューリング方法を説明する図。 実施例を適用したシステムの構成例。 実施例の動作を説明する図。 実施例で用いられるデータテーブルの例。 実施例の動作を説明する図。 実施例の動作を説明する図。 実施例の動作を説明する図。 実施例の動作を説明する図。 実施例の動作を説明する図。 実施例の動作を説明する図。 実施例の動作を説明する図。 実施例の動作を実行しうるコンピュータの一例を示すブロック図。

以下、本発明の実施例を図面に基づいて説明する。ただし、本発明の範囲は以下の実施例に限定されるものではない。

図１に、並列リダクション処理を行うシステムの構成例を示す。図１に示すシステムは、以下で詳細に示すように、本実施例に係る情報処理装置が定めたスケジューリングに従って処理を行う。演算部１は、複数の処理エレメントＰＥ（演算コア）を備える。図１の例では、演算部１は８つの処理エレメントを備え、それぞれの処理エレメントをＰＥ０（１０）〜ＰＥ７（１７）と呼ぶことにする。また、各々の処理エレメントＰＥはクロスバスイッチ２を介して共有メモリ４へアクセスすることが可能である。

共有メモリ４は複数のバンクに分かれており、各バンクはそれぞれ異なるメモリコントローラを経由してクロスバスイッチ２と接続される。図１の例では、共有メモリ４は４つのバンクに分かれており、それぞれのバンクをバンク０（４０）〜バンク３（４３）と呼ぶことにする。また、図１のシステムはバンク０〜バンク３のそれぞれに対応するメモリコントローラを備え、それぞれメモリコントローラ０（３０）〜メモリコントローラ３（３３）と呼ぶことにする。

演算部１が有するそれぞれの処理エレメントＰＥが、共有メモリ４に対して同時にアクセス要求を発行する場合がある。異なるバンクへのアクセス要求が発行された場合であれば、演算部１は遅延なく共有メモリ４にアクセスすることができる。しかし、同一バンクへのアクセス要求が同時に発行された場合（これをバンクコンフリクトと呼ぶ）、メモリコントローラがこれらのアクセス要求を処理するためにはある程度の時間を要する。またバンクコンフリクトによる処理遅延は、同一のメモリコントローラに対するアクセス要求が多いほど、大きくなる。

図２は処理エレメント（ＰＥ０（１０）〜ＰＥ７（１７））のひとつの構成をさらに詳細に示した例である。処理ユニットＰＵ１００はあらゆる演算を行う。処理ソフトウェアには、処理動作の主体であるスレッドが、この処理ユニットＰＵ１００上で動作しているように見える。なお、以降の説明では、処理ユニットＰＥｎ内の処理ユニットＰＵ１００ＰＥで動作しているスレッドのことを、スレッドｎと呼ぶ。すなわち、処理ユニットＰＥ０〜ＰＥ７のそれぞれで動作しているスレッドを、スレッド１〜７と呼ぶ。本実施例において、全スレッドは同一プログラムの同一行を実行する。このため、いくつかのスレッドは必要のない処理を行うことがある。このような場合、不図示のマスクレジスタによって処理結果は破棄される。

リダクション処理は、処理エレメントが複数のデータから１つのデータを算出することを繰り返すことにより行われる。この処理を本明細書では統合処理と呼ぶことにする。例えば、多数のデータのうち２個ずつを統合することを繰り返すことにより、多数のデータから１つの処理結果（演算処理結果）を得ることができる。統合処理の例としては、複数のデータを加算すること、複数のデータの最大値を求めること、複数のデータの最小値を求めること、などがある。もっとも、その他の統合処理が本実施例において採用されてもよい。リダクション処理においては多数のデータからいくつかのデータを選択することと選択されたデータを統合することとが繰り返されるが、どのデータを先に選択しても通常は同じ結果が得られる。

統合処理においてＰＵ１００は、処理レジスタ内（１０１）の値と、データレジスタ内内（１０２）の値とを用いて演算を行い、演算結果を処理レジスタ１０１に書き戻す。データレジスタ１０２は、１以上のレジスタで構成されうる。図２の例では、データレジスタ１０２はデータレジスタ０（１０２０）及びデータレジスタ１（１０２１）を有する。また、リダクション処理を開始するにあたって、各処理エレメントのローカルレジスタ１０３は、予め初期データを格納している。この初期データは、リダクション処理を開始するにあたって不図示の記憶媒体からローカルレジスタ１０３へと読み込まれてもよい。本実施例において処理エレメントＰＥは２つのデータレジスタ１０２０，１０２１を有する。しかしながら、処理エレメントＰＥが有するデータレジスタの数は１つでもよいし、３つ以上でもよい。

図３は、共有メモリ４のメモリ領域がどのようにバンクにマップされているかを表す。図３において、１番目のデータはバンク０に、２番目のデータはバンク１に、３番目のデータはバンク２に、４番目のデータはバンク３に、それぞれマップされている。そして、５〜８番目のデータもまた、それぞれバンク０〜バンク３にマップされている。

図４は共有メモリ４に格納されるデータに対する、メモリ領域の割当て例である。それぞれの四角形はデータが格納される１つの領域を表し、四角形の内部に書かれている番号は、マップされているバンクの番号である。つまり、それぞれのメモリ領域にアクセスする場合には、番号で示されるバンク内へのアクセスが行われる。本実施例においては、１０グループの初期データ群のそれぞれについて、リダクション処理が行われるものとする。すなわち、１グループの初期データ群を構成するそれぞれのデータが各処理エレメントＰＥに格納されている。本実施例においては、１グループの初期データ群は８個の初期データで構成され、それぞれの初期データは各処理エレメントＰＥのローカルレジスタ１０３に格納されている。もっとも、１グループの初期データ群を構成する初期データの数が、処理エレメントＰＥの数に一致している必要はない。

各処理エレメントＰＥは、この１グループの初期データ群についてリダクション処理を行い、処理結果を得る。同様に各処理エレメントＰＥは、他の９グループについても、初期データ群を構成するデータを格納している。そして各処理エレメントＰＥは、他の９グループの初期データ群についてもリダクション処理を行い、９個の処理結果を得る。このように得られたリダクション処理結果は、共有メモリ内の結果出力領域５１に出力される。結果出力領域５１は１０グループのそれぞれに対応する１０個の領域（５１０〜５１９）を有し、１０個の出力結果は先頭から順に格納される。

また、共有メモリ４は通信用エリア５２を備える。この通信用エリア５２は、スレッド間の通信のために用いられる。各スレッドには、予め通信用エリアが割り当てられて（対応付けられて）いる。具体的には、通信領域５２０〜５２７は、それぞれスレッド０〜スレッド７に割り当てられている。そしてスレッド０〜スレッド７は、それぞれに割り当てられた通信領域５２０〜５２７内のデータを読み込むように構成されている。

また、それぞれの通信領域５２０〜５２７は、それぞれのデータレジスタ１０２０，１０２１に対応するデータ領域を有している。本実施例においては１つの処理エレメントＰＥは２つのデータレジスタ１０２０，１０２１を有するため、それぞれの通信領域５２０〜５２７は２つのデータ領域を有する。具体的には、通信領域５２０はデータ領域５２００と５２０１とを有し、通信領域５２７はデータ領域５２７０と５２７１とを有する。それぞれのデータレジスタは、対応するデータ領域内のデータを読み込むように構成されている。

例えばスレッド３からスレッド２へと通信を行う場合について、具体的に説明する。まずスレッド３は、データ領域５２２０又はデータ領域５２２１にデータを書き込む。次にスレッド２は、データ領域５２２０又はデータ領域５２２１内のデータを、データレジスタ１０２０又は１０２１に書き込む。この際、データ領域５２２０内のデータはデータレジスタ０（１０２０）へと、データ領域５２２１内のデータはデータレジスタ１（１０２１）へと、それぞれ読み込まれる。このように、各データレジスタが読み込むデータが格納されているデータ領域は、予め固定的に定められている。

図５は、並列リダクション処理における、処理ユニットＰＵ１００上で動作している１つのスレッドが行う統合処理を示すフローチャートである。ここで、この１つのスレッドをスレッドＡと呼ぶ。ステップＳ１１でスレッドＡは、ローカルレジスタ１０３内の初期データを処理レジスタ１０１に読み込む。ステップＳ１２でスレッドＡは、他のスレッドからリダクション対象データが送られてきている場合、通信用エリア５２からデータレジスタ０，１（１０２０，１０２１）にデータを読み込む。ステップＳ１３でスレッドＡは、処理レジスタ１０１内のデータとデータレジスタ０，１（１０２０，１０２１）内のデータとの統合処理を行う。例えば、統合処理として最大値を求める場合、処理レジスタ１０１内の値が“３”、データレジスタ０（１０２０）内の値が“５”であるならば、処理結果として”５”が得られる。そしてスレッドＡは、得られた処理結果を処理レジスタ１０１に書き戻す。この例では、統合処理により処理レジスタ１０１内の値は“５”に更新される。なお、初期データをローカルレジスタ１０３から読み出し、統合処理をすることなく読み出した値を他のスレッドに渡す場合は、ステップＳ１２およびＳ１３の処理は省略される。

ステップＳ１９でスレッドＡは、さらにリダクション処理を行うか否かを判定する。例えば、他のスレッドからリダクション対象データが送られてきている場合に、送られてきたデータを用いてリダクション処理を行うことができる。さらにリダクション処理を行う場合、処理はステップＳ１２に戻る。さらなるリダクション処理を行わずに処理結果を共有メモリ４に書き込む場合、処理はステップＳ１４に進む。

ステップＳ１４においてスレッドＡは、処理レジスタ１０１内の値を共有メモリ４に書き込む。例えば、処理結果を他のスレッドに渡す場合、スレッドＡは、通信用エリア５２内の相手スレッドについてのデータ領域に、処理レジスタ１０１内の値を書き込む。またもし１グループの初期データ群に対するリダクション処理が完了したのであれば、スレッドＡは結果出力領域５１内の対応する領域（５１０〜５１９）に処理レジスタ１０１内の値を書き込む。

並列リダクション処理においては、ステップＳ１２及びステップＳ１３における繰り返し回数は、それぞれのスレッドによって異なる。したがって、あるスレッドがステップＳ１２を実行しようとする際に、別のスレッドではステップＳ１４を実行しようとしていることがありえる。しかしながら上述のように、本実施例においては各スレッドは一度に同じ命令を実行するように構成されている。したがって、このように異なるスレッドが異なるステップを実行しようとすることは、実行速度に大きなペナルティをもたらしうる。

本実施例においては、図６のように、１回の統合処理が、初期データ設定フェーズ（ステップＳ２１、読込サイクル）、読込・統合処理フェーズ（ステップＳ２２、演算サイクル）、及び結果書込みフェーズ（ステップＳ２３、書込サイクル）の３つに分けられる。そして、命令によって指定された各スレッドが同期して各々のフェーズを開始するように、本実施例に係る情報処理装置はスケジューリングを行う。図５と図６とを比較すると、ステップＳ２１はステップＳ１１に、ステップＳ２２はステップＳ１２およびＳ１３に、ステップＳ２３はステップＳ１４に、それぞれ対応する。これらのフェーズを繰り返すことにより、リダクション処理の結果がメモリに書き込まれる。

本実施例によれば、あるスレッドで読込・統合処理フェーズＳ２２（共有メモリ読込Ｓ１２及び統合処理Ｓ１３）を実行している場合、他のスレッドも読込・統合処理フェーズＳ２２を実行するように、スケジューリングが行われる。読込・統合処理フェーズＳ２２を実行する必要のないスレッドによる処理結果はマスクレジスタによって破棄されるため、実質的にはこのようなスレッドは何も処理をしていない状態（アイドル状態）となる。本実施例の方法によれば、一見アイドル状態の時間が増えるように見える。しかしながら、異なる処理を複数のスレッドが同時に実行しようとすると、一方が処理をしている間、他方のスレッドはアイドル状態となる。もちろん逆も成り立つため、アイドル状態にある時間がかえって増えてしまう。そのため、本実施例のように複数のスレッドに同期して同じ処理を行わせることにより、アイドル状態にある時間が減り、より効率的な処理が実現されることが期待される。

スレッド間通信は共有メモリ４上の通信用エリア５２を介して行われ、リダクション結果は結果出力領域５１に書き込まれる。この際、複数のスレッドから共有メモリ４の同じバンクに対して同時にアクセスが行われると、バンクコンフリクトが起きるため、このことは速度的に大きなペナルティとなる。そこで本実施例では、バンクへの同時アクセス数は、許容できる限界の数（以下、許容アクセス数と呼ぶ）以下となるように制御される。なお、以下の説明において、「バンクコンフリクト数＝同時アクセス数−１」であり、「許容バンクコンフリクト数＝許容アクセス数−１」である。

本実施例に係る情報処理装置であるスケジューリングシステム６は、各スレッドに対する命令を順次生成することにより、各スレッドの動作をスケジューリングする。スケジューリングシステム６は、図１に示した実行環境に接続されていてもよいし、接続されていなくてもよい。例えば、実行環境とは独立なスケジューリングシステム６が、以下で説明する処理を行い、各スレッドに対する命令を生成してもよい。

生成された命令は命令出力バッファ６１に出力される。命令は例えば、実行環境におけるプログラム（たとえばＣＵＤＡのソースコード）という形で出力されうる。出力された命令は、例えば記憶媒体を介して、図１に示される実行環境へと入力されてもよい。そして、出力された命令に従って、各スレッドは動作を行う。こうして、それぞれのスレッドの動作がスケジューリングされる。スケジューリング処理を予め行っておくことは、実行速度を向上させる点で有利である。

また、命令出力バッファ６１に出力された命令は、図１に示される実行環境によって直接アクセスされてもよい。この場合命令出力バッファ６１は、それぞれの処理エレメントＰＥに対して（各スレッドに対して）備えられていてもよい。そしてそれぞれのスレッドは、対応する命令出力バッファ６１に格納された値に従って動作する。こうして、それぞれのスレッドの動作がスケジューリングされる。

このようなスケジューリングシステム６は、例えば通常のコンピュータを用いて実現することもできる。図１７は、このようなコンピュータの基本構成を示す図である。このコンピュータにおいて図７に示すスケジューリングシステム６の機能を実行するためには、各機能構成をプログラムにより表現し、このコンピュータに読み込ませればよい。こうして、このコンピュータでスケジューリングシステム６の全ての機能を実現することができる。この場合、図７をはじめとする構成要素の各々は関数、若しくはＣＰＵが実行するサブルーチンで機能させればよい。

また、コンピュータプログラムは通常、ＣＤ−ＲＯＭ等のコンピュータが読み取り可能な記憶媒体に格納されている。この記憶媒体を、コンピュータが有する読み取り装置（ＣＤ−ＲＯＭドライブ等）にセットし、システムにコピー若しくはインストールすることで実行可能になる。従って、係るコンピュータが読み取り可能な記憶媒体も本発明の範疇にあることは明らかである。

図１７においてＣＰＵ１７０１は、コンピュータ全体の動作をコントロールする。例えＣＰＵ１７０１は、一次記憶１７０２に格納されたプログラムの実行等を行う。一次記憶１７０２は、主にＲＡＭ等のメモリであり、二次記憶１７０３に記憶されたプログラム等を読み込んで格納する。二次記憶１７０３は、例えばハードディスク、ＣＤ−ＲＯＭ等がこれに該当する。プログラムは二次記憶１７０３に格納され、プログラム実行時に一次記憶１７０２に読み込んで、ＣＰＵ１７０１が実行処理を行う。入力デバイス１７０４とはコンピュータに情報を入力するデバイスであって、例えばマウスやキーボード等がこれに該当する。入力デバイス１７０４を用いることにより、ユーザがコンピュータに情報を入力することが可能であってもよい。出力デバイス１７０５とはコンピュータが情報を出力するデバイスであって、モニタ及びプリンタを含む。読込デバイス１７０６は、外部の情報を取得するためのデバイスである。読込デバイス１７０６は、メモリカードリーダ及びネットワークカードを含む。バス１７０８は、上述の各部を接続し、データのやりとりを可能とする。

図７に本実施例に係るスケジューリング方法を用いて処理エレメントＰＥに対する命令を生成する、スケジューリングシステム６の構成例を示す。スケジューリングシステム６は、以下の処理を行うことにより、命令出力バッファ６１に命令を書き込む。なお、ＧＰＧＰＵ上で動作するプログラムの開発においては、同一処理は、同一コードを用いてコーディングすることが通常である。しかしながら、こうしたプログラムを書くためのノウハウは既に知られていることから、本明細書においてはこうしたノウハウについては省略する。

以下では図７を参照しながら、本実施例におけるリダクション処理についてより詳しく説明する。本実施例では上述のように、１０グループの初期データ群のそれぞれについてリダクション処理が行われる。以下の説明では、それぞれの初期データ群（データグループ）をデータ０、データ１、・・・、データ９と呼ぶことにする。

リダクションスケジューリング部６２は、ワークメモリ６３を参照しながら、命令出力バッファ６１に命令を書き込む。リダクションスケジューリング部６２が行う処理の流れを図８に示す。ステップＳ３１においてリダクションスケジューリング部６２は、全てのデータグループについてリダクション処理が終了したか否かをチェックする。リダクション処理が終了している場合、リダクションスケジューリング部６２の処理は終了する。リダクション処理が終了していない場合、リダクションスケジューリング部６２は以下の１ループのリダクション処理（ステップＳ３２〜ステップＳ３６）を行う。

ステップＳ３２においてリダクションスケジューリング部６２は、初期データ設定フェーズＳ２１のスケジューリングを行う。ステップＳ３３においてリダクションスケジューリング部６２は、読込・統合処理フェーズＳ２２のスケジューリングを行う。またステップＳ３４及びＳ３５においてリダクションスケジューリング部６２は、結果書き込みフェーズＳ２３のスケジューリングを行う。リダクションスケジューリング部６２は、共有メモリ４のバンクコンフリクト数を含むリソース使用状況を、ワークメモリ６３上でシミュレーションしながら、これらのスケジューリングを行う。ステップＳ３６においてリダクションスケジューリング部６２は、次のループにおけるスケジューリングを行うのに先立って、処理対象データテーブル６３４を更新する。図９（Ａ）〜（Ｄ）は、ワークメモリ６３上でシミュレーションされたリソース使用状況を示すデータテーブルの例である。以下に、図９（Ａ）〜（Ｄ）を参照しながら、これらの処理についてより詳しく説明する。

ステップＳ３２において、リダクションスケジューリング部６２の初期設定スケジューリング部６２１は、初期データ設定フェーズＳ２１のスケジューリングを行う。本実施例においては、スレッド０からスレッド７までの全スレッドに対して、順番にスケジューリングを行う。具体的な例としては、初期設定スケジューリング部６２１はまず、命令出力バッファ６１にスレッド０についての命令を書き込む。その後初期設定スケジューリング部６２１は、スレッド１〜７についても、命令出力バッファ６１に順次命令を書き込む。こうして初期設定スケジューリング部６２１は、スケジューリングを行うことができる。

図１０（Ａ）は、ステップＳ３２の詳細なフローチャートである。ステップＳ３２においては、スレッド０〜７のそれぞれについて図１０（Ａ）に示される処理が行われ、それぞれのスレッドに対する処理は同様である。以下では、スレッド０についての処理について説明する。

ステップＳ４１において初期設定スケジューリング部６２１は、命令出力バッファ６１に同期命令を出力する。同期命令によって、各スレッドが初期データ設定フェーズＳ２１を同期して行うことができる。同期命令としては、例えばＣＵＤＡでは__syncthreads()という命令が用意されている。

ステップＳ４２において初期設定スケジューリング部６２１は、ローカルレジスタ１０３から処理レジスタ１０１へのデータの読み込みをスケジューリングする。具体的には初期設定スケジューリング部６２１は、データ０からデータ９までのそれぞれに順次着目する。そして、着目データグループについての初期データをローカルレジスタ１０３から処理レジスタ１０１へと読み込ませるか否かを判定する。

図１０（Ｂ）は、ステップＳ４２のより具体的な処理のフローチャートである。ステップＳ４２１で初期設定スケジューリング部６２１は、スレッド０がいずれかのデータグループに属するデータを処理するようにスケジューリングされているか否かを判定する。この判定は、処理対象データテーブル６３４を参照して行うことができる。

処理対象データテーブル６３４は、スレッド数分の要素（本実施例の場合８個）を持つ配列である。それぞれの要素は、各スレッドがどのデータグループを処理対象としているかを示す。各要素は無効を示す値“Ｆ”で初期化されている。スケジューリングシステム６が、あるデータグループについての初期データをあるスレッドに読み込ませる命令を生成した際に、そのスレッドに対応する要素には、そのデータグループを示す数値が設定される。例えば図９（Ａ）においては、グループ１についての初期データを読み込む命令が生成されたスレッドに対応する要素には、数値”１”が書き込まれている。また、スケジューリングシステム６が、あるスレッドに対してあるデータグループについてのリダクション処理を完了させる命令を生成した際には、そのスレッドに対応する要素は再度“Ｆ”で初期化される。例えばスケジューリングシステム６が、あるスレッドに処理レジスタ１０１内のデータを出力させる命令を生成した際に、そのスレッドに対応する要素は再度“Ｆ”で初期化される。

具体的には、処理対象データテーブル６３４において着目データグループに対応する要素が”Ｆ”ではなければ、スレッド０はいずれかのデータグループに属するデータを処理するようにスケジューリングされていると判定することができる。このような場合初期設定スケジューリング部６２１は、着目データグループについてのスケジューリングを終了し、次のデータグループに着目して図１０（Ｂ）の処理を行う。

一方で処理対象データテーブル６３４においてスレッド０に対応する要素が”Ｆ”である場合、処理はステップＳ４２２に進む。この場合、スレッド０はデータを処理するようにはスケジューリングされていない。そこで、次に初期設定スケジューリング部６２１は、着目データグループに属する初期データをスレッド０に読み込ませる命令を生成するか否かを判定する。具体的にはステップＳ４２２において初期設定スケジューリング部６２１は、着目データグループに属する初期データをスレッド０に読み込ませる命令を既に生成しているか否かを判定する。ステップＳ４２２の処理は、データタッチテーブル６３２を参照して行うことができる。

データタッチテーブル６３２はスレッド数×データグループ数（本実施例の場合、８×１０）の二次元配列であり、各要素は予め“Ｆ”で初期化されている。各要素はそれぞれのスレッド及びデータグループに対応している。あるスレッドがあるデータグループに属する初期データを読み込む命令が生成された場合、対応する要素に“Ｔ”が設定される。

初期設定スケジューリング部６２１は、スレッド０と着目データグループとに対応する要素に”Ｔ”が設定されている場合には、着目データグループに属する初期データをスレッド０に読み込ませる命令は既に生成されているものと判定することができる。この場合、初期設定スケジューリング部６２１は、着目データグループについてのスケジューリングを終了し、次のデータグループに着目して図１０（Ｂ）の処理を行う。

一方でスレッド０と着目データグループとに対応する要素に”Ｆ”が設定されている場合には、着目データグループに属する初期データをスレッド０に読み込ませる命令はまだ生成されていないものと判定することができる。この場合、処理はステップＳ４２３に進む。ステップＳ４２３において初期設定スケジューリング部６２１は、命令出力バッファ６１に、着目データグループについての初期データをローカルレジスタ１０３から処理レジスタ１０１へと読み込むことをスレッド０に指示する命令を書き込む。さらに初期設定スケジューリング部６２１は、データタッチテーブル６３２のスレッド０と着目データグループとに対応する要素を、”Ｔ”で更新する。さらに初期設定スケジューリング部６２１は、処理対象データテーブル６３４のスレッド０に対応する要素に、着目データグループを示す番号を設定する。

以上の図１０（Ａ）の処理をスレッド０〜７のそれぞれについて行うことにより、初期データ設定フェーズＳ２１の処理が完了する。

ステップＳ３３において、リダクションスケジューリング部６２の読込・統合スケジューリング部６２２は、読込・統合処理フェーズＳ２２のスケジューリングを行う。ステップＳ３３においても、ステップＳ３２と同様に、読込・統合スケジューリング部６２２はスレッド０からスレッド７までの全スレッドに対して順にスケジューリングを行う。図１１（Ａ）は、ステップＳ３３の詳細なフローチャートである。ステップＳ３３においても、ステップＳ３２と同様に、スレッド０〜７のそれぞれについて図１１（Ａ）に示される処理が行われ、それぞれのスレッドに対する処理は同様である。以下では、スレッド０についての処理について説明する。

ステップＳ５１において読込・統合スケジューリング部６２２は、ステップＳ４１と同様に、命令出力バッファ６１に同期命令を出力する。

ステップＳ５２において読込・統合スケジューリング部６２２は、通信用エリア５２からのデータの読み込み及び統合処理をスケジューリングする。

図１１（Ｂ）は、ステップＳ４２のより具体的な処理のフローチャートである。ステップＳ５２１において読込・統合スケジューリング部６２２は、スレッド０が継続スレッドであるか否かを判定する。この判定は、継続スレッドテーブル６３３を参照して行うことができる。

継続スレッド（継続演算コア）とは、前の次の結果書き込みフェーズＳ２３において、処理レジスタ１０１の値を通信用エリア５２に出力しないように命令されたスレッドである。継続スレッドは、読込・統合処理フェーズＳ２２において、他スレッドからの通信データを読み込み、統合処理を行う（ただし、データを読み込まない場合もある）。本実施例においては、後述する通信スレッド決定部６２４によって、各スレッドが継続スレッドであるか否かが判定されている。

継続スレッドテーブルは、スレッド数分の要素（本実施例の場合８個）を持つ配列であり、それぞれのスレッドに対していくつのスレッドからデータが送られてきているのかを示す。すなわち、継続スレッドテーブルの要素の最小値は０であり、最大値は各処理エレメントＰＥ内のデータレジスタの数（本実施例では２）である。また、継続スレッドではないスレッドに対応する要素は、値“Ｆ”を有する。

このように、スレッド０に対応する継続スレッドテーブル内の要素が”Ｆ”である場合、読込・統合スケジューリング部６２２は、スレッド０は継続スレッドではないと判定することができる。ステップＳ５２１で継続スレッドではないと判断された場合、スレッド０についてのスケジューリング処理は終了し、読込・統合スケジューリング部６２２は次のスレッドについて図１１（Ａ）の処理を行う。

ステップＳ５２１で継続スレッドであると判断された場合、処理はステップＳ５２２に進む。ステップＳ５２２において読込・統合スケジューリング部６２２は、「通信用エリア５２からデータをデータレジスタ１０２に読み込む」ことを示す命令を命令出力バッファ６１に出力する。それぞれのスレッドが通信用エリアから読み込むデータの数は、継続スレッドテーブル内の、それぞれのスレッドに対応する要素に格納されている。

本実施例においては、スレッド０に対応する継続スレッドテーブル内の要素が”１”である場合、読込・統合スケジューリング部６２２は、スレッド０にデータ領域５２００内の値をデータレジスタ０に読み込ませる命令を生成する。また、スレッド０に対応する継続スレッドテーブル内の要素が”２”である場合、読込・統合スケジューリング部６２２は、スレッド０にデータ領域５２００内の値をデータレジスタ０に読み込ませる命令を生成する。スレッド０に対応する継続スレッドテーブル内の要素が”２”である場合、読込・統合スケジューリング部６２２はさらに、スレッド０にデータ領域５２０１内の値をデータレジスタ１に読み込ませる命令を生成する。

このように本実施例においては、各スレッドに対応する継続スレッドテーブル内の要素が示す値がデータレジスタ１０２の数よりも少ない場合、この要素が示す数のデータを通信用エリアから読み込ませる命令が生成される。この場合読込・統合スケジューリング部６２２は、より先頭側に位置するデータを優先して読み込ませるように、命令を生成する。しかしながら別の実施例において読込・統合スケジューリング部６２２は、継続スレッドテーブル内の要素が”Ｆ”ではない場合に、全てのデータレジスタ１０２へと通信用エリア５２から値を読み込む命令を生成してもよい。

ステップＳ５２３において読込・統合スケジューリング部６２２は、「処理レジスタ１０１とデータレジスタ１０２の間で統合処理を行う」ことを示す命令を命令出力バッファ６１に出力する。この統合処理は、処理レジスタ１０１と各データレジスタ１０２の間で順次行われる。このような命令を受けたスレッドは、まず処理レジスタ１０１とデータレジスタ０（１０２０）との間での統合処理を行い、結果を処理レジスタ１０１に書き込む。さらにこのスレッドは、処理レジスタ１０１とデータレジスタ１（１０２１）との間で統合処理を行って結果を処理レジスタ１０１に書き込む。

この統合処理は、ステップＳ５２２においてデータレジスタ１０２へと読み込まれた値の数だけ繰り返されれば十分である。例えば、データレジスタ１０２０にのみ値が読み込まれた場合、すなわちスレッド０に対応する継続スレッドテーブル内の要素が”１”である場合、処理レジスタ１０１とデータレジスタ０（１０２０）との間でのみ統合が行われればよい。

しかしながら、データレジスタ１０２へと読み込まれた値の数に従って分岐処理を行うと、実行環境によってはかえって実行時間がかかってしまうことがある。そこで本実施例においては、ステップＳ５２２とステップＳ５２３との間で読込・統合スケジューリング部６２２は、「統合処理を行う必要がないデータレジスタ１０２のデータを書き換える」命令を命令出力バッファ６１に出力する。ここで読込・統合スケジューリング部６２２は、データレジスタ１０２の値を、統合処理を行っても結果に影響しない値に書き換えればよい。

具体的な例としては、統合処理として各要素の加算が行われ、スレッド０に対応する継続スレッドテーブル内の要素が”１”である場合、読込・統合スケジューリング部６２２は「“０”をデータレジスタ１（１０２１）にセットする」命令を出力すればよい。統合処理として最大値を求める場合も同様である。このような処理を行うことにより、ステップＳ５２３において読込・統合スケジューリング部６２２は、単純に処理レジスタ１０１と全てのデータレジスタ１０２との統合処理を行う命令を出力すればよい。この場合、実行時に条件分岐を行わなくてもよい。

ステップＳ３４及びＳ３５において、リダクションスケジューリング部６２の結果出力スケジューリング部６２３及び通信スレッド決定部６２４は、結果書込みフェーズＳ２３のスケジューリングを行う。結果書込みフェーズＳ２３において各スレッドは、結果出力領域５１にリダクション結果を出力するか、又は他のスレッドに対しデータを送信する。本実施例においては、どちらの場合にも共有メモリ４に対する書き込みが行われる。したがって、これらの動作は同一コードで実行可能である。また、１つのスレッドが１回の結果書き込みフェーズＳ２３において両方の動作を行うことはない。したがって本実施例においてそれぞれのスレッドは、１回の結果書込みフェーズＳ２３において、結果出力領域５１へリダクション結果を出力するか、他のスレッドに対してデータを送信するか、あるいはアイドル状態にあるようにスケジューリングされる。

もし、結果出力領域５１へのリダクション結果の出力と、他のスレッドに対するデータの送信とを、同一コードで実行することが困難であるときは、これらを異なるフェーズにおいて行えばよい。例えば、これらの間でメモリへの書込み手順が大きく異なる場合が挙げられる。この場合、各フェーズの先頭で同期命令出力を行う点を除いて、以下に示す本実施例と同様にスケジューリング処理を行えばよい。

まず、結果出力領域５１にリダクション処理の結果が出力されるステップＳ３４について説明する。ステップＳ３４の開始時に、結果出力スケジューリング部６２３は、バンクコンフリクトテーブル６３１と継続スレッドテーブル６３３とを初期化する。具体的には結果出力スケジューリング部６２３は、バンクコンフリクトテーブル６３１の各要素に値”０”を格納する。また結果出力スケジューリング部６２３は、継続スレッドテーブル６３３の各要素に値”Ｆ”を格納する。

ステップＳ３４においても、ステップＳ３２と同様に、結果出力スケジューリング部６２３はスレッド０からスレッド７までの全スレッドに対して順にスケジューリングを行う。図１２（Ａ）は、ステップＳ３４の詳細なフローチャートである。ステップＳ３４においても、ステップＳ３２と同様に、スレッド０〜７のそれぞれについて図１２（Ａ）に示される処理が行われ、それぞれのスレッドに対する処理は同様である。以下では、スレッド０についての処理について説明する。

ステップＳ６１において結果出力スケジューリング部６２３は、ステップＳ４１と同様に、命令出力バッファ６１に同期命令を出力する。ステップＳ６２において結果出力スケジューリング部６２３は、スレッド０の処理レジスタ１０１に格納されているデータを、結果出力領域５１に出力するか否かを判定する。

図１２（Ｂ）は、ステップＳ６２のより具体的な処理のフローチャートである。ステップＳ６２１において結果出力スケジューリング部６２３は、スレッド０が処理レジスタ１０１に格納しているデータが、結果出力領域５１に出力されるか否かを判定する。処理レジスタ１０１に格納されているデータが、各データグループについての最終的な処理結果である場合に、このデータは結果出力領域５１に出力される。

この判定は例えば、データタッチテーブル６３２及び処理対象データテーブル６３４を参照して行うことができる。ここで、スレッド０が処理しているデータグループを着目データグループとする。着目データグループについての処理を予め定められた数のスレッドが開始しており、かつ着目データグループを処理しているスレッドがスレッド０だけである場合に、スレッド０が処理レジスタ１０１に格納しているデータは結果出力領域５１に出力される。ここで、予め定められた数は通常、着目データグループに属する初期データをローカルレジスタ１０３に有しているスレッドの数に一致する。

着目データグループについての処理を予め定められた数のスレッドが開始している場合、データタッチテーブルにおいて着目データグループに対応する要素”Ｔ”の数は予め定められた数に一致する。本実施例においては１０スレッドのそれぞれが各データグループを処理するため、着目データグループに対応する要素”Ｔ”の数が１０である場合に、着目データグループについての処理を予め定められた数のスレッドが開始していると判定できる。また、処理対象データテーブル６３４において着目データグループに対応する値を有する要素の数が１つである場合、着目データグループを処理しているスレッドはスレッド０だけであると判定することができる。

スレッド０が処理レジスタ１０１に格納しているデータが、結果出力領域５１に出力されない場合、ステップＳ６２の処理は終了し、次のスレッドについて図１２（Ａ）の処理が行われる。スレッド０が処理レジスタ１０１に格納しているデータが、結果出力領域５１に出力される場合、処理はステップＳ６２２に進む。ステップＳ６２２において、結果出力スケジューリング部６２３は、着目データグループについての処理結果を書き込むバンクへのアクセス要求の数（アクセス回数）を確認する。本実施例においては上述のように、そこで同一バンクに対して同時にアクセスすることが可能なスレッド数の上限（許容アクセス数）を予め定めておき、この上限に従ってスケジューリングを行う。

各バンクに対するアクセス要求の数（管理情報）は、バンクコンフリクトテーブル６３１で管理され保持される。バンクコンフリクトテーブル６３１はバンク数分の要素（本実施例では４つ）を持つ配列である。バンクコンフリクトテーブル６３１の各要素は予め０に初期化されている。そして、結果出力スケジューリング部６２３が各バンクへのアクセスをスケジューリングするたびに、各バンクに対応する要素に１が加えられる。

そして結果出力スケジューリング部６２３は、着目データグループについての処理結果を書き込むバンクへのアクセス要求の数が、予め定められた許容アクセス数未満（所定回数未満）であるか否かを確認する。許容アクセス数未満である場合、処理はステップＳ６２３に進む。許容アクセス数以上であった場合、スレッド０についての図１２Ａの処理は終了する。すなわち、現在の結果書込みフェーズＳ２３においてスレッド０に結果出力領域５１へと結果を書き込ませる命令は生成されず、次回以降の結果書込みフェーズＳ２３においてスレッド０に結果を書き込ませる命令が生成される。

ステップＳ６２３で結果出力スケジューリング部６２３は、「処理レジスタ１０１内の値を、結果出力領域５１内の着目データグループに対応する領域に書き込む」ことを、スレッド０についての命令出力バッファ６１に出力する。さらに結果出力スケジューリング部６２３は、着目データグループについての処理結果を書き込むバンクに対応する、バンクコンフリクトテーブル６３１の要素の値に１を加える。

次に、他のスレッドに対しデータが送信されるステップＳ３５における、通信スレッド決定部６２４の処理について説明する。ステップＳ３５においては、通信スレッド決定部６２４はスレッド０〜スレッド７について順次スケジューリングを行うのではなく、以下のように処理を行う。

図１３Ａは、ステップＳ３５の詳細なフローチャートである。ステップＳ７１において通信スレッド決定部６２４は、全データグループのそれぞれについて、データを送信するスレッド（送信スレッド、送信演算コア）とデータを受信するスレッド（継続スレッド）とを決定する。本実施例では通信スレッド決定部６２４は、データ０からデータ９までのそれぞれに順次着目し、着目データグループについて送信スレッドと継続スレッドとを決定する。

図１３Ｂは、ステップＳ７１の詳細なフローチャートである。ステップＳ７１１で通信スレッド決定部６２４の継続スレッド決定部６２４１は、継続スレッドを１つ選択する（第１の選択）。もし継続スレッドがなければ、着目データグループについての図１３（Ｂ）の処理を終了し、次の着目データグループについて図１３（Ｂ）の処理を行う。ステップＳ７１２で通信スレッド決定部６２４の送信スレッド決定部６２４２は、ステップＳ７１１で選択された継続スレッドにデータを送信する送信スレッドを判定する（第２の選択）。ステップＳ７１１及びステップＳ７１２を繰り返すことにより、１組の結果書込フェーズと続く読込・統合処理フェーズとにおいてデータが送受信される継続スレッドと送信スレッドとのセットが繰り返し選択される。

ステップＳ７１１のより詳細なフローチャートを図１４に示す。ステップＳ８１で継続スレッド決定部６２４１は、着目データグループを処理しているスレッドから１つを選択する。ここで継続スレッド決定部６２４１は、通信用エリア５２内の対応するデータ領域が属するバンクへの同時アクセス数が最も少ないスレッドを選択する。

どのスレッドが着目データグループを処理しているのかは、処理対象データテーブル６３４を参照して判定することができる。例えば継続スレッド決定部６２４１は、処理対象データテーブル６３４における、各スレッドに対応する要素の値が、着目データグループの番号と一致するか否かを判定すればよい。一致する場合、そのスレッドは着目データグループを処理している。

また継続スレッド決定部６２４１は、バンクコンフリクトテーブル６３１を参照して、それぞれのバンクに対する同時アクセス数を判定することができる。上述のように、バンクコンフリクトテーブル６３１は、それぞれのバンクに対する同時アクセス数を要素として有している。

本実施例のように、対応するデータ領域が属するバンクへの同時アクセス数が最も少ないスレッドを継続スレッドとして選択することにより、それぞれのバンクへの同時アクセス数を平準化させることができる。なお、１つのスレッドが複数のデータ領域からデータを読み込む場合、このスレッドは複数のバンクからデータを読み込むかもしれない。このような場合、継続スレッド決定部６２４１は、それぞれのバンクについての同時アクセス数を判定する。そして継続スレッド決定部６２４１は、同時アクセス数の最大値を、スレッドに対応するデータ領域が属するバンクへの同時アクセス数として用いる。こうすることにより、全てのバンクのバンクコンフリクト数を許容アクセス数以下に抑えることができる。

ステップＳ８２において継続スレッド決定部６２４１は、ステップＳ７１１においてスレッドが選択されたか否かを判定する。スレッドが選択されなかった場合、継続スレッド決定部６２４１は、着目データグループについては継続スレッドがないものと判定する。そして、継続スレッド決定部６２４１は図１３（Ｂ）の処理を終了し、次の着目データグループについて送信スレッドと継続スレッドとを決定する。

ステップＳ８３において継続スレッド決定部６２４１は、ステップＳ７１１において選択されたスレッドについて、対応するデータ領域が属するバンクへの同時アクセス数が許容アクセス数未満であるかを否かを判定する。同時アクセス数が許容アクセス数以上である場合、継続スレッド決定部６２４１は、着目データグループについては更なる継続スレッドがないものと判定する。そして継続スレッド決定部６２４１は、図１３（Ｂ）の処理を終了し、次の着目データグループについて送信スレッドと継続スレッドとを決定する。

同時アクセス数が許容アクセス数未満である場合、継続スレッド決定部６２４１はステップＳ７１１において選択されたスレッドを継続スレッドとして判定する。ステップＳ８４において継続スレッド決定部６２４１は、継続スレッドテーブル６３３のうちステップＳ７１１において選択された継続スレッドに対応する要素を“０”で更新する。さらに継続スレッド決定部６２４１は、バンクコンフリクトテーブル６３１において、ステップＳ７１１において選択された継続スレッドに対応するデータ領域が属するバンクについての要素の値をインクリメントする（例えば１を加える）。こうして継続スレッド決定部６２４１は、バンクコンフリクトテーブル６３１を更新する。そして、処理はステップＳ７１２に進む。

ステップＳ７１２のより詳細なフローチャートを図１５に示す。ステップＳ９１で送信スレッド決定部６２４２は、ステップＳ７１１で選択された継続スレッドにデータを送信する送信スレッドを決定する。具体的には送信スレッド決定部６２４２は、着目データグループを処理対象としているスレッドのうち、まだ継続スレッド又は送信スレッドとして選択されていないスレッドを送信スレッドとして選択する。各スレッドが継続スレッド又は送信スレッドとして選択されているか否かは、継続スレッドテーブル６３３を参照して判定することができる。本実施例においては、スレッドに対応する継続スレッドテーブル６３３内の要素が”Ｆ”であり、かつ着目データグループを処理対象としているスレッドを、送信スレッド決定部６２４２は送信スレッドとして選択すればよい。

本実施例において送信スレッド決定部６２４２は、１つの継続スレッドに対応する送信スレッドを、継続スレッドが有するデータレジスタ１０２の数だけ選択する。この場合、各送信スレッドからのデータが、それぞれのデータレジスタ１０２へと読み込まれる。着目データグループを処理対象としており、かつ継続スレッド又は送信スレッドとして選択されていないスレッドの数がデータレジスタ１０２の数よりも多い場合、送信スレッド決定部６２４２は以下のようにして送信スレッドを選択することができる。

すなわち送信スレッド決定部６２４２は、各スレッドに対応するデータ領域が属するバンクに対する同時アクセス数がより多いスレッドを選択する。本実施例において、送信スレッドに対応するデータ領域には、データの書き込みは行われない。したがって、このように対応するデータ領域が属するバンクに対する同時アクセス数がより多いスレッドを送信スレッドとして選択する（継続スレッドとして選択しない）ことにより、各バンクへの同時アクセス数を平準化することができる。一方で、着目データグループを処理対象としており、かつ継続スレッド又は送信スレッドとして選択されていないスレッドの数が、継続スレッドが有するデータレジスタ１０２の数よりも少ない場合、送信スレッド決定部６２４２は全てのスレッドを選択できる。

ステップＳ９２において送信スレッド決定部６２４２は、継続スレッドテーブル６３３のうち、ステップＳ７１１で決定された継続スレッドに対応する要素を更新する。具体的には送信スレッド決定部６２４２は、継続スレッドにデータを送信する送信スレッドの数を、継続スレッドに対応する要素として継続スレッドテーブル６３３に格納する。

ステップＳ９３において送信スレッド決定部６２４２は、ステップＳ９１で選択された送信スレッドに対して、中間書込み処理を行うように命令を出力する。具体的には送信スレッド決定部６２４２は、「送信スレッドのそれぞれが、処理レジスタ１０１の内容を、継続スレッドが読み込むデータ領域にコピーする」ことを示す命令を、命令出力バッファ６１に出力する。

例えば継続スレッドがスレッド０、送信スレッドがスレッド２とスレッド３である場合、送信スレッド決定部６２４２は、「スレッド２が処理レジスタ１０１の内容をデータ領域５２００へコピーする」ことを示す命令を命令出力バッファ６１に出力する。また送信スレッド決定部６２４２は、「スレッド３が処理レジスタ１０１の内容をデータ領域５２０１へコピーする」ことを示す命令を命令出力バッファ６１に出力する。

図１６は、ステップＳ３６のより具体的な処理のフローチャートである。ステップＳ１０１でデータテーブル修正部６２５は、ステップＳ３４においていずれかのスレッドに結果出力領域５１に処理レジスタ１０１の内容を出力させる命令が生成されたか否かを判定する。またデータテーブル修正部６２５は、ステップＳ３５においていずれかのスレッドに通信用エリア５２に処理レジスタ１０１の内容を出力させる命令が生成されたか否かを判定する。どちらの命令も生成されていないのであれば、図１６の処理は終了する。どちらかの命令が生成された場合、処理はステップＳ１０２に進む。ステップＳ１０２においてデータテーブル修正部６２５は、処理対象データテーブル６３４内の、処理レジスタ１０１の内容を出力する命令の対象となるスレッドに対応する要素を、無効を示す値“Ｆ”で更新する。

上述の実施例では、共有メモリのバンク数およびメモリコントローラの数を４、処理エレメントの数を８、処理エレメント内のデータレジスタの数を２、リダクション処理を行うデータグループの数を１０とした。しかしながら、本発明はこれらの数に限定されない。また、上述の実施例に係るシステムにおいては、データレジスタ数及び許容アクセス数を変化させることにより、処理時間も変動する。また、最適なデータレジスタ数及び許容アクセス数は、ハードウェアの構成によって異なる。したがって、これらの値を適宜選択することにより、処理時間をより短くすることが可能である。

（他の実施形態）
本発明は、以下の処理を実行することによっても実現される。即ち、上述した実施形態の機能を実現するソフトウェア（プログラム）をネットワーク又は各種記憶媒体を介してシステム或いは装置に供給する。そして、そのシステム或いは装置のコンピュータ（又はＣＰＵやＭＰＵ等）がプログラムコードを読み出して実行する。この場合、そのプログラム、及び該プログラムを記憶した記憶媒体は本発明を構成することになる。

Claims (4)

1. 複数のバンクで構成されるメモリに接続された複数の演算コアに対する命令を生成する情報処理装置であって、
   前記複数の演算コアは、
   前記命令によって指定された演算コアが互いに同期して、
   初期データを前記演算コアが保持するレジスタへと読み込む読込サイクルと、
   前記演算コアごとに予め対応付けられたバンク内の領域から読み込んだデータと、前記演算コアが保持するレジスタ内のデータとを用いて演算を行い、演算結果を前記演算コアが保持するレジスタに格納する演算サイクルと、
   前記メモリへ前記演算コアが保持するレジスタ内のデータを書き込む書込サイクルとを用いるものであり、
   前記情報処理装置は、
   １回の前記書込サイクルにおいて演算コアがメモリへの書き込みアクセスを行う回数に少なくとも関連する値をそれぞれのバンクごとに保持する保持手段と、
   前記複数の演算コアのうち、対応するバンクについての前記保持手段の保持する値が所定値未満である演算コアを、他の演算コアからデータを受信する継続演算コアとして選択する第１の選択手段と、
   前記継続演算コアに対してデータを送信する送信演算コアを選択する第２の選択手段と、
   前記送信演算コアに、前記書込サイクルにおいて、レジスタ内のデータを前記継続演算コアに対応付けられたバンク内の領域に書き込ませるための命令と、前記継続演算コアに、当該書込サイクルに続く前記演算サイクルにおいて、当該継続演算コアに対応付けられたバンク内の領域からデータを読み込ませるための命令と、を生成する生成手段と、
   前記第１の選択手段が前記継続演算コアを選択する際に、選択された前記継続
   演算コアに予め対応付けられたバンクについての前記保持手段の保持する値をインクリメントする更新手段と、
   を備えることを特徴とする情報処理装置。
2. 前記第１の選択手段及び前記第２の選択手段は、複数の前記演算コアの中から、１組の書込サイクルと続く演算サイクルとにおいてデータが送受信される前記継続演算コアと前記送信演算コアとのセットを繰り返し選択し、
   前記第１の選択手段は、対応付けられたバンクについての前記保持手段の保持する値がより小さい演算コアを、前記継続演算コアとして選択し、
   前記第２の選択手段は、対応付けられたバンクについての前記保持手段の保持する値がより大きい演算コアを、前記送信演算コアとして選択する
   ことを特徴とする、請求項１に記載の情報処理装置。
3. 複数のバンクで構成されるメモリに接続された複数の演算コアに対する命令を生成する情報処理装置が行う情報処理方法であって、
   前記複数の演算コアは、
   前記命令によって指定された演算コアが互いに同期して、
   初期データを前記演算コアが保持するレジスタへと読み込む読込サイクルと、
   演算コアごとに予め対応付けられたバンク内の領域から読み込んだデータと、前記演算コアが保持するレジスタ内のデータとを用いて演算を行い、演算結果を前記演算コアが保持するレジスタに格納する演算サイクルと、
   前記メモリへ前記演算コアが保持するレジスタ内のデータを書き込む書込サイクルとを用いるものであり、
   前記情報処理装置は、１回の前記書込サイクルにおいて演算コアがメモリに書き込みアクセスを行う回数に少なくとも関連する値をそれぞれのバンクごとに保持する保持手段を備え、
   前記情報処理方法は、
   前記情報処理装置の第１の選択手段が、前記複数の演算コアのうち、対応するバンクについての前記保持手段の保持する値が所定値未満である演算コアを、他の演算コアからデータを受信する継続演算コアとして選択する第１の選択工程と、
   前記情報処理装置の第２の選択手段が、前記継続演算コアに対してデータを送信する送信演算コアを選択する第２の選択工程と、
   前記情報処理装置の生成手段が、前記送信演算コアに、前記書込サイクルにおいて、レジスタ内のデータを前記継続演算コアに対応付けられたバンク内の領域に書き込ませるための命令と、前記継続演算コアに、当該書込サイクルに続く前記演算サイクルにおいて、当該継続演算コアに対応付けられたバンク内の領域からデータを読み込ませるための命令と、を生成する生成工程と、
   前記第１の選択工程で前記継続演算コアを選択する際に、前記情報処理装置の更新手段が、選択された前記継続演算コアに予め対応付けられたバンクについての前記保持手段の保持する値をインクリメントする更新工程と、
   を備えることを特徴とする情報処理方法。
4. コンピュータを、請求項１又は２に記載の情報処理装置の各手段として機能させるための、コンピュータプログラム。

Images