JP2008276740A5 - - Google Patents
Download PDFInfo
- Publication number
- JP2008276740A5 JP2008276740A5 JP2008015281A JP2008015281A JP2008276740A5 JP 2008276740 A5 JP2008276740 A5 JP 2008276740A5 JP 2008015281 A JP2008015281 A JP 2008015281A JP 2008015281 A JP2008015281 A JP 2008015281A JP 2008276740 A5 JP2008276740 A5 JP 2008276740A5
- Authority
- JP
- Japan
- Prior art keywords
- virtual
- thread
- threads
- cta
- program
- Prior art date
- Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
- Granted
Links
- 238000000034 method Methods 0.000 claims description 30
- 230000000875 corresponding Effects 0.000 claims description 17
- 230000011273 social behavior Effects 0.000 claims description 7
- 230000006399 behavior Effects 0.000 description 27
- 238000007667 floating Methods 0.000 description 14
- 238000004891 communication Methods 0.000 description 10
- UIIMBOGNXHQVGW-UHFFFAOYSA-M buffer Substances [Na+].OC([O-])=O UIIMBOGNXHQVGW-UHFFFAOYSA-M 0.000 description 8
- 230000004048 modification Effects 0.000 description 8
- 238000006011 modification reaction Methods 0.000 description 8
- 238000004422 calculation algorithm Methods 0.000 description 7
- 238000000354 decomposition reaction Methods 0.000 description 6
- 238000010586 diagram Methods 0.000 description 6
- 238000006243 chemical reaction Methods 0.000 description 5
- 230000000694 effects Effects 0.000 description 5
- 238000001914 filtration Methods 0.000 description 5
- 230000004044 response Effects 0.000 description 4
- 238000004364 calculation method Methods 0.000 description 3
- 125000004429 atoms Chemical group 0.000 description 2
- 230000004927 fusion Effects 0.000 description 2
- 239000003607 modifier Substances 0.000 description 2
- 230000003287 optical Effects 0.000 description 2
- 230000002093 peripheral Effects 0.000 description 2
- 238000000926 separation method Methods 0.000 description 2
- 238000000844 transformation Methods 0.000 description 2
- 230000001131 transforming Effects 0.000 description 2
- 230000005540 biological transmission Effects 0.000 description 1
- 239000000969 carrier Substances 0.000 description 1
- 230000015556 catabolic process Effects 0.000 description 1
- 239000003795 chemical substances by application Substances 0.000 description 1
- 238000004590 computer program Methods 0.000 description 1
- 230000001276 controlling effect Effects 0.000 description 1
- 238000007796 conventional method Methods 0.000 description 1
- 230000004059 degradation Effects 0.000 description 1
- 238000006731 degradation reaction Methods 0.000 description 1
- 238000005516 engineering process Methods 0.000 description 1
- 238000009434 installation Methods 0.000 description 1
- 239000011159 matrix material Substances 0.000 description 1
- 238000005293 physical law Methods 0.000 description 1
- 238000009877 rendering Methods 0.000 description 1
- 230000003068 static Effects 0.000 description 1
- 230000001360 synchronised Effects 0.000 description 1
Description
[0001]本発明は一般的には並列処理に関し、具体的には、並列スレッド・コンピューティングのための仮想アーキテクチャ及び命令セットに関する。
[0002]並列処理において、複数の処理装置(例えば、複数のプロセッサ・チップ又は単一チップ内の複数の処理コア)が同時に動作してデータを処理する。そのようなシステムは、複数部分への分解に適した問題を解決するために使用される。1つの例は画像フィルタリングである。この場合、1つの出力画像(又は複数の出力画像)の各画素は、1つの入力画像(又は複数の入力画像)の或る数の画素から計算される。各々の出力画素の計算は、一般的に全ての他の出力画素計算から独立しているので、異なる処理装置は異なる出力画素を並列に計算することができる。他の多くの問題も並列分解が容易である。一般的に、N個の路による並列実行は、そのような問題解決を、おおまかにNの因子だけ高速にする。
[0003]他の種の問題も、並列実行スレッドを相互に調整できれば、容易に並列処理することができる。1つの例は高速フーリエ変換(FFT)である。FFTは再帰的アルゴリズムであって、各々の段階で、前の段階の出力に計算が行われて新しい値が生成され、この新しい値は出力段階に到達するまで次の段階の入力として使用される。単一の実行スレッドは、前の段階から信頼できる出力データを確実に取得できるかぎり、複数段階行うことができる。もしタスクを複数のスレッドへ分割するのであれば、或る一定の調整機構を提供して、例えば、まだ書き込まれていない入力データをスレッドが読み出さないようにしなければならない。(この問題の1つの解決法は、2005年12月15日に出願された同一出願人による同時係属米国特許出願第11/303,780号で説明されている。)
[0004]しかし、並列処理システムのプログラミングは難しい。プログラマは、通常、処理ユニットが実際に実行できるコードを作り出すため、利用可能な処理装置の数及び能力(命令セット、データ・レジスタの数、相互接続など)を知っていなければならない。機械特有コンパイラは、この分野で相当の援助を提供できるが、異なるプロセッサへコードを移植する度に、コードを再コンパイルすることが依然として必要である。
[0005]更に、並列処理アーキテクチャの様々な様相が急速に発展している。例えば、新しいプラットフォーム・アーキテクチャ、命令セット、及びプログラミング・モデルが継続的に開発されている。並列アーキテクチャの様々な様相(例えば、プログラミング・モデル又は命令セット)は1つの世代から次の世代へと変化するので、それに従ってアプリケーション・プログラム、ソフトウェア・ライブラリ、コンパイラ、並びに他のソフトウェア及びツールも変化しなければならない。この非定常性は、並列処理コードの開発及び維持へ相当のオーバーヘッドを付加する。
[0006]スレッド間の調整が必要とされるとき、並列プログラミングは一層難しくなる。プログラマは、特定のプロセッサ又はコンピュータ・システムの中で、どのような機構がスレッド間通信をサポート(又はエミュレート)するために利用可能であるかを決定し、利用可能な機構を有効に使うコードを書かなければならない。異なるコンピュータ・システムでは利用可能及び/又は最適メカニズムは、一般的に異なることから、この種の並列コードは一般的に移植することができず、並列コードが実行される各々のハードウェア・プラットフォーム毎、コードを書き直さなければならない。
[0007]更に、プロセッサの実行可能なコードを提供することに加えて、プログラマは、更に、「マスタ」プロセッサの制御コードを提供しなければならない。制御コードは様々な処理装置の動作を調整し、例えば、どのようなプログラムを実行すべきか、どのような入力データを処理すべきかを各々の処理装置に命令する。そのような制御コードは、通常、特定のマスタ・プロセッサ及びプロセッサ間通信プロトコルに対して特有であり、もし異なるマスタ・プロセッサが代用されるのであれば、通常、書き直さなれければならない。
[0008]並列処理コードのコンパイル及び再コンパイルの困難性は、コンピューティング・テクノロジの発展と共にユーザがシステムをアップグレードしようとする意欲を低下させる。したがって、コンパイルされた並列処理コードを特定のハードウェア・プラットフォームから切り離して、目標とする並列アプリケーション及びツールのために安定した並列処理アーキテクチャ及び命令セットを提供することが望ましい。
[0009]本発明の実施形態は、並列スレッド・コンピューティングの仮想アーキテクチャ及び仮想命令セットを提供する。仮想並列アーキテクチャは、異なる仮想スレッド間で複数レベルのデータ共有及び調整(例えば、同期)を有する複数の仮想スレッドの並行実行をサポートする仮想プロセッサ、並びに仮想プロセッサを制御する仮想実行ドライバを定義する。仮想プロセッサの仮想命令セット・アーキテクチャは、仮想スレッドの振る舞いを定義するために使用され、並列スレッド振る舞い、例えば、データ共有及び同期に関系する命令を含む。仮想並列プラットフォームを使用して、プログラマは、仮想スレッドが並行して実行してデータを処理するアプリケーション・プログラムを開発することができる。アプリケーション・プログラムは、高度に移植可能な中間形態で、例えば、仮想並列プラットフォームを目標にしたプログラム・コードとして記憶及び配布される。導入時又は実行時に、ハードウェア特有仮想命令トランスレータ(translator)及び実行ドライバは、中間形態アプリケーション・コードを、このアプリケーション・コードが実行される特定のハードウェアへ適応させる。結果として、アプリケーション・プログラムは、一層移植可能となり、開発が容易になる。なぜなら、開発プロセスは特定の処理ハードウェアから独立しているからである。
[0010]本発明の1つの様相によれば、並列処理動作を定義する方法は、協調的仮想スレッドから成るアレイ内の多数の仮想スレッドの各々について実行される動作シーケンスを定義する第1のプログラム・コードを提供することを含む。第1のプログラム・コードは、アレイの代表的仮想スレッドについて実行されるスレッド単位命令のシーケンスを定義する仮想スレッド・プログラムへコンパイルされ、スレッド単位命令のシーケンスは、代表的仮想スレッドとアレイの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含む。仮想スレッド・プログラムは記憶され(例えば、メモリ又はディスクに)、後で目標プラットフォーム・アーキテクチャに合致する命令のシーケンスへ翻訳可能である。
[0011]更に、入力データセットを処理して出力データセットを生成するように適応させられた協調的仮想スレッドのアレイを定義ために、第2のプログラム・コードを提供することが可能である。アレイ内の各仮想スレッドは仮想スレッド・プログラムを並行して実行する。第2のプログラム・コードは、有利には、仮想関数ライブラリの関数呼び出しシーケンスへ変換される。ライブラリは、協調的仮想スレッドのアレイを初期化及び実行させる仮想関数を含む。この関数呼び出しシーケンスも記憶可能である。記憶された仮想スレッド・プログラム及び関数呼び出しシーケンスは、目標プラットフォーム・アーキテクチャの上で実行可能なプログラム・コードへ翻訳可能である。実行可能プログラム・コードは、協調的仮想スレッドのアレイを実行する1つ又は複数のプラットフォーム・スレッドを定義する。実行可能なプログラム・コードは目標プラットフォーム・アーキテクチャに合致したコンピュータ・システムの上で実行することができ、記憶メディア(例えば、コンピュータ・メモリ、ディスクなど)に記憶可能な出力データセットを生成することができる。
[0012]前述したように、仮想スレッド・プログラム・コード内のスレッド単位命令シーケンスは、有利には、代表的仮想スレッドとアレイの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含む。例えば、スレッド単位命令シーケンスは、1つ又は複数の他の仮想スレッドがシーケンス内の特定の点に達するまで、この特定の点における代表的仮想スレッドの動作の実行を一時停止する命令、1つ又は複数の他の仮想スレッドがアクセスを有する共有メモリの中に代表的仮想スレッドがデータを記憶する命令、1つ又は複数の他の仮想スレッドがアクセスを有する共有メモリの中に記憶されたデータを代表的仮想スレッドがアトミックに読み出し及び更新する命令などを含んでもよい。
[0013]仮想スレッド・プログラムは、更に、多数の仮想状態空間の1つで変数を定義する変数定義ステートメントを含んでもよい。異なる仮想状態空間は、仮想スレッド間のデータ共有の異なるモードに対応する。1つの実施形態において、少なくともスレッド単位非共有モード及びグローバル共有モードがサポートされる。他の実施形態では、追加モード、例えば、仮想スレッドの1つのアレイ内の共有モード、及び/又は仮想スレッドの複数のアレイ間の共有モードもサポートされてもよい。
[0014]本発明の他の様相によれば、目標プロセッサを動作させる方法は、入力プログラム・コードを提供することを含む。入力プログラム・コードは第1の部分及び第2の部分を含む。第1の部分は、入力データセットを処理して出力データセットを生成するように適応させられた仮想スレッドの1つのアレイにおける多数の仮想スレッドの各々について実行される動作シーケンスを定義し、第2の部分は上記仮想スレッドのアレイの次元を定義する。入力プログラム・コードの第1の部分は、アレイの代表的仮想スレッドのために実行されるスレッド単位命令のシーケンスを定義する仮想スレッド・プログラムへコンパイルされる。スレッド単位命令シーケンスは、代表的仮想スレッドとアレイの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含む。入力プログラム・コードの第2の部分は、仮想関数ライブラリへの関数呼び出しシーケンスへ変換される。ライブラリは、協調的仮想スレッドのアレイを初期化して実行させる仮想関数を含む。仮想スレッド・プログラム及び関数呼び出しシーケンスは、目標プラットフォーム・アーキテクチャの上で実行可能なプログラム・コードへ翻訳される。実行可能なプログラム・コードは、協調的仮想スレッドのアレイを実行する1つ又は複数の現実スレッドを定義する。実行可能なプログラム・コードは、目標プラットフォーム・アーキテクチャに合致するコンピュータ・システムの上で実行され、記憶メディアに記憶可能な出力データセットを生成する。
[0015]幾つかの実施形態において、仮想スレッドのアレイは、2つ以上の次元で定義可能である。更に、入力プログラム・コードの第2の部分も、仮想スレッドのアレイのグリッドの1つ又は複数の次元を定義する関数呼び出しを含んでもよく、グリッド内の各々のアレイが実行される。
[0016]任意の目標プラットフォーム・アーキテクチャが使用されてもよい。幾つかの実施形態において、目標プラットフォーム・アーキテクチャは、マスタ・プロセッサ及びコプロセッサを含む。翻訳中に、仮想スレッド・プログラムは、コプロセッサ上で定義された多数のスレッドによって並列に実行可能なプログラム・コードへ翻訳可能であり、関数呼び出しシーケンスは、マスタ・プロセッサ上で実行するコプロセッサ用ドライバ・プログラムへの呼び出しシーケンスへ翻訳される。他の実施形態において、目標プラットフォーム・アーキテクチャは中央処理装置(CPU)を含む。翻訳中に、仮想スレッド・プログラム及び関数呼び出しシーケンスの少なくとも一部分は、仮想スレッドの数よりも少ない或る数のCPUスレッドを使用して仮想スレッド・アレイを実行する目標プログラム・コードへ翻訳される。
[0017]本発明の更に他の実施形態によれば、目標プロセッサを動作させる方法は、入力データセットを処理して出力データセットを生成するように適応させられた仮想スレッド・アレイにおける多数の仮想スレッドの代表的仮想スレッドについて実行されるスレッド単位命令シーケンスを定義する仮想スレッド・プログラムを取得することを含む。スレッド単位命令シーケンスは、代表的仮想スレッドとアレイの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含む。仮想スレッド・アレイの次元を定義する追加のプログラム・コードも取得される。仮想スレッド・プログラム及び追加のプログラム・コードは、目標プラットフォーム・アーキテクチャの上で実行可能なプログラム・コードへ翻訳され、実行可能なプログラム・コードは、仮想スレッド・アレイを実行する1つ又は複数のプラットフォーム・スレッドを定義する。実行可能なプログラム・コードは、目標プラットフォーム・アーキテクチャに合致したコンピュータ・システム上で実行され、出力データセットを生成して、出力データセットをメモリ内に記憶する。
[0018]幾つかの実施形態において、仮想スレッド・プログラムは、高級プログラミング言語で書かれたソース・プログラム・コードを受け取り、コンパイルして仮想スレッド・プログラムを生成することによって取得されてもよい。代替として、仮想スレッド・プログラムは、記憶メディアから読み出されるか、ネットワークを介してリモート・コンピュータ・システムから受け取られてもよい。読み出され又は受け取られる仮想スレッド・コードは、事前に高級言語からコンパイルされているか、仮想命令セット・アーキテクチャに合致するコードとして直接作成されていることを理解すべきである。
[0019]下記の詳細な説明は、添付の図面と共に、本発明の性質及び利点の一層良好な理解を提供する。
[0030]本発明の実施形態は、並列スレッド・コンピューティングのための仮想アーキテクチャ及び命令セットを提供する。仮想アーキテクチャは、異なるスレッド間で複数レベルのデータ共有及び調整(例えば、同期)を有する複数スレッドの並行実行をサポートするプロセッサのモデル、及びモデル・プロセッサを制御する仮想実行ドライバを提供する。処理スレッドの振る舞いを定義するために使用される仮想命令セットは、並列スレッド振る舞いに関する命令、例えば、或るスレッド間でデータの共有を許す命令、及び異なるスレッドがプログラム内の或るプログラマ特定点で同期することを要求する命令を含む。仮想プラットフォームを使用して、プログラマは、データを処理するために並行及び協調的スレッドが実行されるアプリケーション・プログラムを開発することができる。ハードウェア特有仮想命令トランスレータ及び実行ドライバは、アプリケーション・コードが実行される特定のハードウェアへアプリケーション・コードを適応させる。結果として、アプリケーション・プログラムは移植が一層可能となり、開発が容易になる。なぜなら、開発プロセスは特定の処理ハードウェアから独立しているからである。
1.システムの概観
[0031]図1は、本発明の実施形態に従ったコンピュータ・システム100のブロック図である。コンピュータ・システム100は、メモリ・ブリッジ105を含むバス経路を介して通信する中央処理装置(CPU)102及びシステム・メモリ104を含む。例えば、ノースブリッジ・チップであってもよいメモリ・ブリッジ105は、バス又は他の通信経路106(例えば、HyperTransportリンク)を介してI/O(入力/出力)ブリッジ107へ接続される。例えば、サウスブリッジ・チップであってもよいI/Oブリッジ107は、ユーザ入力を1つ又は複数のユーザ入力デバイス108(例えば、キーボード、マウス)から受け取り、バス106及びメモリ・ブリッジ105を介して入力をCPU102へ回送する。並列処理サブシステム112は、バス又は他の通信経路113(例えば、PCI Express又はAccelerated Graphics
Portリンク)を介してメモリ・ブリッジ105へ結合される。1つの実施形態において、並列処理サブシステム112は、画素をディスプレイ・デバイス110(例えば、通常のCRT又はLCDベースのモニタ)へ引き渡すグラフィックス・サブシステムである。システム・ディスク114もI/Oブリッジ107へ接続される。スイッチ116は、I/Oブリッジ107と、他のコンポーネント、例えば、ネットワーク・アダプタ118及び様々なアドイン・カード120及び121との間の接続を提供する。USB又は他のポート接続、CDドライブ、DVDドライブなどを含む他のコンポーネント(明示的には示されない)も、I/Oブリッジ107へ接続されてもよい。図1の様々なコンポーネントを相互接続する通信経路は、任意の適切なプロトコル、例えば、PCI(Peripheral Component Interconnect)、PCI Express(PCI−E)、AGP(Accelerated Graphics Port)、HyperTransport、又は任意の他のバス又はポイントツーポイント通信プロトコルを使用して実現されてもよく、異なるデバイス間の接続は、当技術分野で公知の異なるプロトコルを使用してもよい。
[0031]図1は、本発明の実施形態に従ったコンピュータ・システム100のブロック図である。コンピュータ・システム100は、メモリ・ブリッジ105を含むバス経路を介して通信する中央処理装置(CPU)102及びシステム・メモリ104を含む。例えば、ノースブリッジ・チップであってもよいメモリ・ブリッジ105は、バス又は他の通信経路106(例えば、HyperTransportリンク)を介してI/O(入力/出力)ブリッジ107へ接続される。例えば、サウスブリッジ・チップであってもよいI/Oブリッジ107は、ユーザ入力を1つ又は複数のユーザ入力デバイス108(例えば、キーボード、マウス)から受け取り、バス106及びメモリ・ブリッジ105を介して入力をCPU102へ回送する。並列処理サブシステム112は、バス又は他の通信経路113(例えば、PCI Express又はAccelerated Graphics
Portリンク)を介してメモリ・ブリッジ105へ結合される。1つの実施形態において、並列処理サブシステム112は、画素をディスプレイ・デバイス110(例えば、通常のCRT又はLCDベースのモニタ)へ引き渡すグラフィックス・サブシステムである。システム・ディスク114もI/Oブリッジ107へ接続される。スイッチ116は、I/Oブリッジ107と、他のコンポーネント、例えば、ネットワーク・アダプタ118及び様々なアドイン・カード120及び121との間の接続を提供する。USB又は他のポート接続、CDドライブ、DVDドライブなどを含む他のコンポーネント(明示的には示されない)も、I/Oブリッジ107へ接続されてもよい。図1の様々なコンポーネントを相互接続する通信経路は、任意の適切なプロトコル、例えば、PCI(Peripheral Component Interconnect)、PCI Express(PCI−E)、AGP(Accelerated Graphics Port)、HyperTransport、又は任意の他のバス又はポイントツーポイント通信プロトコルを使用して実現されてもよく、異なるデバイス間の接続は、当技術分野で公知の異なるプロトコルを使用してもよい。
[0032]並列処理サブシステム112は、並列処理装置(PPU)122及び並列処理(PP)メモリ124を含む。これらは、例えば、1つ又は複数の集積回路デバイス、例えば、プログラム可能プロセッサ、アプリケーション特有集積回路(ASIC)、及びメモリ・デバイスを使用して実現されてもよい。PPU122は、有利には、1つ又は複数の処理コアを含む高度並列プロセッサを実現する。処理コアの各々は、多数の(例えば、数百の)スレッドを並行して実行することができる。PPU122は多種の計算を行うようにプログラム可能である。そのような計算の中には、線形及び非線形データ変換、ビデオ及び/又はオーディオ・データのフィルタリング、モデリング(例えば、物理法則を適用して物体の位置、速度、及び他の属性を決定すること)、画像のレンダリングなどが含まれる。PPU122は、システム・メモリ104及び/又はPPメモリ124から内部メモリへデータを転送し、データを処理し、結果のデータをシステム・メモリ104及び/又はPPメモリ124へ返してもよい。そのようなデータは、例えば、CPU102を含む他のシステム・コンポーネントからアクセス可能である。幾つかの実施形態において、PPU122はグラフィックス・プロセッサであり、メモリ・ブリッジ105及びバス113を介してCPU102及び/又はシステム・メモリ104によって供給されたグラフィックス・データから画素データを生成すること、PPメモリ124(このメモリは、例えば、従来のフレーム・バッファを含むグラフィックス・メモリとして使用可能である)と交信して画素データを記憶及び更新すること、画素データをディスプレイ110へ引き渡すことなどに関する様々なタスクを行うように構成可能である。幾つかの実施形態において、PPサブシステム112は、グラフィックス・プロセッサとして動作する1つのPPU122及び汎用計算に使用される別のPPU122を含んでもよい。PPUは同格のもの又は異なるものであってもよく、各々のPPUは自分自身の専用PPメモリ・デバイスを有してもよい。
[0033]CPU102はシステム100のマスタ・プロセッサとして動作し、他のシステム・コンポーネントの動作を制御及び調整する。具体的には、CPU102は、PPU122の動作を制御するコマンドを出す。幾つかの実施形態において、CPU102はPPU122のためにコマンドのストリームをコマンド・バッファへ書き込む。コマンド・バッファは、システム・メモリ104、PPメモリ124、又はCPU102及びPPU122の両者からアクセスできる他の記憶場所にあってもよい。PPU122はコマンド・バッファからコマンド・ストリームを読み出し、CPU102の動作とは同期しないでコマンドを実行する。
[0034]本明細書で示されるシステムは例示であって、変形及び修正が可能であることが理解される。ブリッジの数及び配列を含む接続トポロジは、所望のように修正されてもよい。例えば、幾つかの実施形態において、システム・メモリ104はブリッジを介することなくCPU102へ直接接続され、他のデバイスはメモリ・ブリッジ105及びCPU102を介してシステム・メモリ104と通信する。他の代替のトポロジにおいて、PPサブシステム112はメモリ・ブリッジ105ではなくI/Oブリッジ107へ接続される。更に他の実施形態において、I/Oブリッジ107及びメモリ・ブリッジ105は単一のチップへ集積されてもよい。本明細書で示された特定のコンポーネントは自由に選択でき、例えば、任意の数のアドイン・カード又は周辺デバイスがサポートされてもよい。幾つかの実施形態において、スイッチ116は削除され、ネットワーク・アダプタ118並びにアドイン・カード120及び121がI/Oブリッジ107へ直接接続する。
[0035]PPU122からシステム100の残部への接続も変更されてもよい。幾つかの実施形態において、PPシステム112は、システム100の拡張スロットへ挿入されるアドイン・カードとして実現される。他の実施形態において、PPUは、バス・ブリッジ、例えば、メモリ・ブリッジ105又はI/Oブリッジ107と一緒に、単一チップの上に集積可能である。更に他の実施形態において、PPU122の幾つか又は全部の要素がCPU102と一緒に集積されてもよい。
[0036]PPUは、ローカルメモリを含まないで、任意の量のローカルPPメモリを設けられてもよく、ローカルメモリ及びシステム・メモリを任意の組み合わせで使用してもよい。例えば、PPU122は、一体化メモリ・アーキテクチャ(UMA)実施形態としてグラフィックス・プロセッサであってもよい。そのような実施形態では、専用グラフィックス・メモリは殆ど又は全く設けられず、PPU122はシステム・メモリを排他的又は殆ど排他的に使用する。UMA実施形態において、PPUはブリッジ・チップの中に集積されるか、離散的チップとして提供されて、高速リンク(例えば、PCI−E)がPPUをブリッジ・チップ及びシステム・メモリへ接続する。
[0037]更に理解すべきは、任意の数のPPUをシステム内に含めることができることである。それは、例えば、複数のPPUを単一のアドイン・カードに含め、複数のアドイン・カードを経路113へ接続し、及び/又は1つ又は複数のPPUをシステム・マザーボードへ直接接続することによって行う。複数のPPUは並列に動作可能であり、単一のPPUよりも高いスループットでデータを処理する。
[0038]当業者は、更に、CPU及びPPUが単一のデバイスへ集積され得ること、CPU及びPPUが様々なリソース、例えば、命令論理、バッファ、キャッシュ、メモリ、処理エンジンなどを共有可能なこと、又は別々のリソースを並列処理及び他の動作のために設けてもよいことを認識する。したがって、本明細書の中でPPUと関連づけて説明される回路及び/又は機能の任意のもの及び全部が、適切に装備されたCPUの中で実現されてもよく、そのようなCPUによって達成されてもよい。
[0039]PPUを組み込んだシステムは、デスクトップ、ラップトップ、又は携帯用・パーソナル・コンピュータ、サーバ、ワークステーション、ゲームコンソール、組み込みシステムなどを含む多様な構成及び形態因子の中で実現されてもよい。
[0040]更に、本発明の1つの利点として、特定のコンピューティング・ハードウェアからの独立性が増加することを当業者は理解する。したがって、本発明の実施形態は、PPUを設けないシステムを含む任意のコンピュータ・システムを使用して、実施されてもよいことを理解すべきである。
2.仮想プログラミング・モデルの概観
[0041]本発明の実施形態において、コンピューティング・システムのPPU122又は他のプロセッサを使用して、スレッド・アレイを使用する汎用計算を実行することが望ましい。本明細書で使用されるように、「スレッド・アレイ」とは、入力データセット上で同じプログラムを並行して実行して出力データセットを生成する多数の(n0個の)スレッドのグループである。スレッド・アレイ内の各スレッドは一意のスレッド識別子(スレッドID)を割り当てられる。スレッドIDは、スレッドの実行中にスレッドへアクセス可能である。スレッドIDは、1次元又は多次元の数値(例えば、0からn0−1)として定義可能であり、スレッドの処理振る舞いの様々な様相を制御する。例えば、スレッドIDは、入力データセットのどの部分をスレッドが処理すべきかを決定し、及び/又は出力データセットのどの部分をスレッドが作成又は書き込むべきかを決定するために使用される。
[0041]本発明の実施形態において、コンピューティング・システムのPPU122又は他のプロセッサを使用して、スレッド・アレイを使用する汎用計算を実行することが望ましい。本明細書で使用されるように、「スレッド・アレイ」とは、入力データセット上で同じプログラムを並行して実行して出力データセットを生成する多数の(n0個の)スレッドのグループである。スレッド・アレイ内の各スレッドは一意のスレッド識別子(スレッドID)を割り当てられる。スレッドIDは、スレッドの実行中にスレッドへアクセス可能である。スレッドIDは、1次元又は多次元の数値(例えば、0からn0−1)として定義可能であり、スレッドの処理振る舞いの様々な様相を制御する。例えば、スレッドIDは、入力データセットのどの部分をスレッドが処理すべきかを決定し、及び/又は出力データセットのどの部分をスレッドが作成又は書き込むべきかを決定するために使用される。
[0042]幾つかの実施形態において、スレッド・アレイは、「協調的」スレッド・アレイ、すなわちCTAである。他のタイプのスレッド・アレイと同じように、CTAは、入力データセット上で同じプログラム(本明細書では「CTAプログラム」と呼ばれる)を並行して実行して、出力データセットを作成する複数スレッドのグループである。CTAにおいて、スレッドは、スレッドIDに依存する仕方でデータを相互に共有することによって協力することができる。例えば、CTAにおいて、データは1つのスレッドによって作成され、他のスレッドによって消費されることが可能である。幾つかの実施形態において、消費するスレッドがデータへのアクセスを試みる前に、作成するスレッドによってデータが実際に作成されたことを保証するため、データが共有される地点で、同期命令がCTAプログラム・コードの中に挿入されることが可能である。CTAスレッド間のデータ共有の範囲は、あるとすれば、CTAプログラムによって決定される。したがって、CTAを使用する特定のアプリケーションにおいて、CTAスレッドは、CTAプログラムに依存して、相互にデータを実際に共有しても共有しなくてもよく、「CTA」及び「スレッド・アレイ」の用語は、本明細書では同じ意味で使用されることを理解すべきである。
[0043]幾つかの実施形態において、CTA内のスレッドは、同じCTA内の他のスレッドと、入力データ及び/又は中間結果を共有する。例えば、CTAプログラムは、特定のデータが書き込まれる共有メモリ内のスレッドIDの関数であるアドレスを計算する命令を含んでもよい。各々のスレッドは自分自身のスレッドIDを使用して関数を計算し、対応する場所へ書き込む。アドレス関数は、有利には、異なるスレッドが、異なる場所に書き込むように定義される。関数が決定論的である限り、スレッドによって書き込まれる場所は予測可能である。CTAプログラムは、更に、データが読み出される共有メモリ内のスレッドIDの関数であるアドレスを計算する命令を含むことができる。適切な関数を定義し、同期手法を提供することによって、データはCTAの1つのスレッドによって共有メモリ内の所与の場所へ書き込まれ、同じCTAの異なるスレッドによって、予測可能な仕方で上記所与の場所から読み出されることが可能である。その結果、スレッド間データ共有の所望のパターンがサポートされ、CTA内の任意のスレッドが同一CTA内の他のスレッドとデータを共有することができる。
[0044]CTA(又は他のタイプのスレッド・アレイ)は、有利には、データ並列分解に適した計算を行うために使用される。本明細書で使用されるように、「データ並列分解」とは、入力データ上で同じアルゴリズムを並列に複数回実行して出力データを生成することによって、計算問題が解決される任意の場合を含む。例えば、データ並列分解の1つの例は、入力データセットの異なる部分へ同じ処理アルゴリズムを適用して出力データセットの異なる部分を生成することを引き起こす。データ並列分解に適した問題の例は、行列代数、任意数の次元における線形及び/又は非線形変換(例えば、高速フーリエ変換)、及び様々なフィルタリング・アルゴリズムを含む。フィルタリング・アルゴリズムには、任意数の次元における畳み込みフィルタ、複数次元における分離フィルタなどが含まれる。入力データセットの各部分に適用される処理アルゴリズムは、CTAプログラム中で特定され、CTA内の各スレッドは、入力データセットの一部分の上で同じCTAプログラムを実行する。CTAプログラムは、広範囲の数学及び論理演算を使用してアルゴリズムを実現することができ、プログラムは条件付き又は分岐実行経路、及び直接及び/又は間接メモリ・アクセスを含むことができる。
[0045]CTA、及びCTAの実行は、前述した出願第11/303,780号に詳細で説明される。
[0046]幾つかの場合において、関係のあるCTA(より一般的には、スレッド・アレイ)の「グリッド」を定義することが有用である。本明細書で使用されるように、CTAの「グリッド」とは或る数の(n1個の)CTAの集まりである。集まりの中の全CTAは同じサイズ(即ち、スレッド数)であり、同じCTAプログラムを実行する。グリッド内のn1個のCTAは、有利には、相互に独立である。これは、グリッド内のCTA実行が、グリッド内の他のCTA実行によって影響されないことを意味する。明らかなように、この特徴は、利用可能な処理コア間にCTAを配分するとき著しい柔軟性を提供する。
[0047]グリッド内の異なるCTAを区別するため、「CTA識別子」(又はCTA ID)が、有利には、グリッドの各CTAへ割り当てられる。スレッドIDのように、任意の一意の識別子(非限定的に、数値識別子を含む)をCTA IDとして使用することができる。1つの実施形態において、CTA IDは、0からn1−1までの単純に順次の(1次元の)索引値である。他の実施形態では、多次元索引スキームを使用することができる。CTA IDはCTAの全スレッドに共通であり、グリッド内の所与のCTAスレッドは、例えば、入力データを読み出すソース場所及び/又は出力データを書き込むデスティネーション場所を決定するために、スレッドIDと一緒にCTA IDを使用できる。このようにして、同じグリッドの異なるCTA内のスレッドは、同じデータセットの上で並行して動作してもよい。もっとも、幾つかの実施形態において、グリッド内の異なるCTA間でデータを共有することはサポートされていない。
[0048]CTAのグリッドを定義することは有用である。例えば、複数のCTAを使用して、単一の大きな問題の異なる部分を解くことが望まれる場合にそうである。例えば、フィルタリング・アルゴリズムを行って、高品位テレビ(HDTV)画像を生成することが望まれるかも知れない。当技術分野で知られるように、HDTV画像は2百万個を超える画素を含むことは可能である。もし各々のスレッドが1つの画素を生成するならば、実行されるスレッドの数は、単一のCTA内で処理可能なスレッド数を超過する(合理的なサイズ及びコストの処理プラットフォームが、通常の手法を使用して構築されると仮定する)。
[0049]この大きな処理タスクは、多数のCTA間で画像を分割することによって管理可能である。各々のCTAは出力画素の異なる部分(例えば、16×16タイル)を生成する。全てのCTAは同じプログラムを実行し、スレッドはCTA ID及びスレッドIDの組み合わせを使用して、入力データを読み出す場所及び出力データを書き込む場所を決定し、各々のCTAが入力データセットの正しい部分上で動作し、出力データセットの自己部分を正しい場所へ書き込む。
[0050]注意すべきことは、CTA内の(データを共有できる)スレッドとは異なり、グリッド内のCTAは、有利には、相互にデータを共有せず、又は相互に依存しないことである。即ち、同じグリッドの2つのCTAは順次に(どのような順序でも)又は並行して実行可能であり、依然として同じ結果を作成することができる。その結果、処理プラットフォーム(例えば、図1のシステム100)は、CTAのグリッドを実行して結果を取得することができる。この取得は、最初に1つのCTAを実行し、続いて次のCTAを実行し、以下同様にグリッドの全CTAを実行してしまうことによって行われる。代替として、もし十分なリソースが利用可能であれば、処理プラットフォームは同じグリッドを実行し、複数のCTAを並列に実行することによって、同じ結果を取得することができる。
[0051]幾つかの場合、CTAの複数(n2)のグリッドを定義することが望ましいかも知れない。この場合、各々のグリッドはデータ処理プログラム又はタスクの異なる部分を実行する。例えば、データ処理タスクは或る数の「解法ステップ」へ分割されてもよく、各々の解法ステップはCTAのグリッドを実行することによって行われる。他の例として、データ処理タスクは、連続した入力データセット(例えば、ビデオデータの連続フレーム)の上で同一又は類似の動作を実行することを含んでもよい。CTAグリッドは各入力データセットについて実行可能である。仮想プログラミング・モデルは、有利には、少なくともこれら3レベルの作業定義(即ち、スレッド、CTA、及びCTAグリッド)をサポートする。もし望まれるならば、追加のレベルもサポートされてもよい。
[0052]特定の問題を解決するために使用されるCTAサイズ(スレッドの数n0)、グリッドサイズ(CTAの数n1)、及びグリッド数(n2)は、問題分解を定義するプログラマ又は自動化エージェントの問題及び選好パラメータに依存することが理解される。したがって、幾つかの実施形態において、CTAサイズ、グリッドサイズ、及びグリッド数は、有利には、プログラマによって定義される。
[0053]CTAアプローチから利益を得る問題は、通常、並列に処理可能な多数のデータ要素の存在によって特徴づけられる。幾つかの場合、データ要素は出力要素であり、出力要素の各々は、入力データセットの異なる(可能性として重複する)部分に同じアルゴリズムを行うことによって生成される。他の場合、データ要素は入力要素であり、入力要素の各々は同じアルゴリズムを使用して処理される。
[0054]そのような問題は、常に、少なくとも2つのレベルへ分解され、前述したスレッド、CTA、及びグリッドへマップされる。例えば、各々のグリッドは、複雑なデータ処理タスクの中で1つの解法ステップの結果を表すことができる。各々のグリッドは、有利には、或る数の「ブロック」へ分割され、ブロックの各々は単一のCTAとして処理可能である。各々のブロックは、有利には、複数の「要素」、即ち、解かれる問題の要素部分(例えば、単一の入力データ点又は単一の出力データ点)を含有する。CTAの中で、各々のスレッドは1つ又は複数の要素を処理する。
[0055]図2A及び図2Bは、本発明の実施形態で使用される仮想プログラミング・モデルにおけるグリッド、CTA、及びスレッド間の関係を図示する。図2Aは或る数のグリッド200を示し、各々のグリッドはCTA202の2次元(2D)アレイから作られる。(ここで、類似オブジェクトの複数のインスタンスは、オブジェクト識別する参照数字、及び必要な場合にはインスタンスを識別する括弧内の数字で表される。)CTA202(0,0)について図2Bで示されるように、各々のCTA202はスレッド(Θ)204の2Dアレイを含む。各グリッド200の各CTA202内の各スレッド204について、形式I=[ig,ic,ii]の一意識別子が定義される。ここで、グリッド識別子igはグリッドを一意に識別し、CTA ID icはグリッド内のCTAを一意に識別し、スレッドID iiはCTA内のスレッドを一意に識別する。この実施形態において、識別子Iは1次元グリッド識別子ig、2次元CTA識別子ic、及び2次元スレッド識別子iiから構築されることが可能である。他の実施形態において、一意の識別子Iは、0≦ig<n2;0≦ic<n1;及び0≦ii<n0の三つ組み整数である。更に、他の実施形態において、グリッド識別子、CTA識別子、及びスレッド識別子の任意のもの又は全部が、1次元整数、2D座標対、3D三つ組みなどで表現されてもよい。一意のスレッド識別子Iは、例えば、1つのグリッド全体又は複数のグリッドについて入力データセットを包含するアレイ内で入力データのソース場所を決定するため、及び/又は1つのグリッド全体又は複数のグリッドについて出力データセットを包含するアレイ内で出力データを記憶する目標場所を決定するために使用されることが可能である。
[0056]例えば、HDTV画像の場合、各々のスレッド204は出力画像の画素に対応してもよい。CTA202のサイズ(スレッド204の数)は、問題分解での選択事項であり、単一のCTA202におけるスレッド最大数の制約(これはプロセッサ・リソースの有限性を反映する)のみによって限定される。グリッド200はHDTVデータのフレーム全体に対応可能であり、又は複数のグリッドが単一のフレームにマップされることができる。
[0057]幾つかの実施形態において、問題分解は均一である。これは、全てのグリッド200が同じ数及び同じ配列のCTA202を有し、全てのCTA202が同じ数及び同じ配列のスレッド204を有することを意味する。他の実施形態において、分解は非均一であってもよい。例えば、異なるグリッドは、異なる数のCTAを含んでもよく、異なるCTA(同じグリッド又は異なるグリッド内の)は、異なる数のスレッドを含んでもよい。
[0058]上記で定義されたCTAは、数十又は数百の並行スレッドを含むことができる。CTAが実行される並列処理システムは、そのような多数の並行スレッドをサポートしてもサポートしなくてもよい。1つの様相において、本発明はそのようなハードウェア制限からプログラマを解き放す。これは、実際のハードウェア能力がどのようなものであれ、プログラマがCTA及びCTAグリッドのモデルを使用して処理タスクを定義できるようにすることによってなされる。例えば、プログラマはコード(CTAプログラム)を書き、CTAの単一の代表的スレッドによって行われる処理タスクを定義し、CTAを或る数のそのようなスレッドとして定義し(スレッドの各々は一意の識別子を有する)、グリッドを或る数のCTAとして定義する(CTAの各々は一意の識別子を有する)ことができる。下記で説明するように、そのようなコードは、特定のプラットフォーム上で実行されるコードへ自動的に翻訳される。例えば、もしCTAが、或る数n0の並行スレッドを含むものと定義され、しかし目標プラットフォームが、1つのスレッドしかサポートしないならば、トランスレータは、n0個のスレッドの全てに割り当てられたタスクを行う1つの現実スレッドを定義することができる。もし目標プラットフォームが、2つ以上ではあるがn0個よりも少ない並行スレッドをサポートするならば、所望に応じてタスクを利用可能スレッド数へ分割することができる。
[0059]したがって、CTA及びグリッドのプログラミング・モデルは、仮想モデル、即ち、特定の物理的現実から切り離されたプログラマへの概念的援助であるモデルとして理解されるべきである。CTA及びグリッドの仮想モデルは、様々な程度の並列処理ハードウェア・サポートを有する多様な目標プラットフォームとして現実化可能である。具体的には、本明細書で使用される「CTAスレッド」の用語は、(可能性として、1つ又は複数の他の処理タスクと協力する)離散的処理タスクの仮想モデルを意味し、CTAスレッドは目標プラットフォーム上のスレッドに対して1対1でマップしてもマップしなくてもよいことを理解すべきである。
3.仮想アーキテクチャ
[0060]本発明の1つの様相によれば、CTA及びCTAグリッドを実行する仮想並列アーキテクチャが定義される。仮想並列アーキテクチャは、多数の並行CTAスレッドの実行をサポートする並列プロセッサ及び関連メモリ空間を表現する。多数の並行CTAスレッドは協調的に振る舞うことができ、例えば、所望の時点で相互にデータを共有し同期する。この仮想並列アーキテクチャは、多様な現実プロセッサ及び/又は処理システムの上にマップ可能である。そのような現実プロセッサ及び/又は処理システムには、例えば、図1のシステム100のPPU122が含まれる。仮想アーキテクチャは、有利には、異なるレベルのデータ共有及びアクセス・タイプをサポートする或る数の仮想メモリ空間、及び仮想プロセッサによって実行可能な関数の全部を識別する仮想命令セット・アーキテクチャ(ISA)を定義する。仮想アーキテクチャは、更に、有利には、仮想実行ドライバを定義する。仮想実行ドライバを、例えば、CTA又はCTAグリッドを定義及び起動することによって、CTA実行を制御するために使用することができる。
[0060]本発明の1つの様相によれば、CTA及びCTAグリッドを実行する仮想並列アーキテクチャが定義される。仮想並列アーキテクチャは、多数の並行CTAスレッドの実行をサポートする並列プロセッサ及び関連メモリ空間を表現する。多数の並行CTAスレッドは協調的に振る舞うことができ、例えば、所望の時点で相互にデータを共有し同期する。この仮想並列アーキテクチャは、多様な現実プロセッサ及び/又は処理システムの上にマップ可能である。そのような現実プロセッサ及び/又は処理システムには、例えば、図1のシステム100のPPU122が含まれる。仮想アーキテクチャは、有利には、異なるレベルのデータ共有及びアクセス・タイプをサポートする或る数の仮想メモリ空間、及び仮想プロセッサによって実行可能な関数の全部を識別する仮想命令セット・アーキテクチャ(ISA)を定義する。仮想アーキテクチャは、更に、有利には、仮想実行ドライバを定義する。仮想実行ドライバを、例えば、CTA又はCTAグリッドを定義及び起動することによって、CTA実行を制御するために使用することができる。
[0061]図3は、本発明の実施形態に従った仮想アーキテクチャ300のブロック図である。仮想アーキテクチャ300は仮想プロセッサ302を含み、仮想プロセッサ302は、多数のCTAスレッドを並列に実行するように構成された仮想コア308を有する。仮想アーキテクチャ300は、更に、仮想プロセッサ302へアクセス可能なグローバルメモリ304、及び仮想プロセッサ302の動作を制御するコマンドを供給する仮想ドライバ320を含む。仮想ドライバ320はグローバルメモリ304へのアクセスも有する。
[0062]仮想プロセッサ302は、仮想ドライバ320からコマンドを受け取って解釈するフロントエンド306、及び単一CTAのn0個の全スレッドを並行して実行できる実行コア308を含む。仮想コア308は多数の(n0個以上の)仮想処理エンジン310を含む。1つの実施形態において、各々の仮想処理エンジン310は1つのCTAスレッドを実行する。仮想処理エンジン310はそれぞれのCTAスレッドを並行して実行するが、必ずしも並列に実行することはない。1つの実施形態において、仮想アーキテクチャ300は仮想処理エンジン310の数T(例えば、384、500、768など)を特定する。この数はCTA内のスレッド数n0の上限を設定する。理解すべきは、仮想アーキテクチャ300の現実が、特定された数Tよりも少ない物理的処理エンジンを含んでもよく、単一の処理エンジンが幾つかのCTAスレッド、単一の「現実」(即ち、プラットフォームでサポートされた)スレッド又は並行の複数現実スレッドとして実行できることである。
[0063]仮想プロセッサ302は、更に、仮想命令装置312を含む。仮想命令装置312は、仮想処理エンジン310がそれぞれのCTAスレッドについて命令を供給されるように維持する。命令は、仮想アーキテクチャ300の一部分である仮想ISAによって定義される。並列スレッド・コンピューティングの仮想ISAの例は、下記で説明される。命令装置312は、命令を仮想処理エンジン310へ供給する過程で、CTAスレッド同期及びCTAスレッド振る舞いの他の協調的様相を管理する。
[0064]仮想コア308は、異なるレベルのアクセス可能性を有する内部データ記憶装置を提供する。特殊レジスタ311は読み出し可能であるが、仮想処理エンジン310のよる書き込みは不可能であり、図2の問題分解モデル内の各CTAスレッド「位置」を定義するパラメータを記憶するために使用される。1つの実施形態において、特殊レジスタ311は、CTAスレッドごとに(又は仮想処理エンジン310ごとに)、スレッドIDを記憶する1つのレジスタを含む。各々のスレッドIDレジスタは、仮想処理エンジン310のそれぞれの1つによってのみアクセス可能である。特殊レジスタ311は、更に、追加のレジスタを含んでもよく、これら追加のレジスタは全てのCTAスレッド(又は全ての仮想処理エンジン310)によって読み込み可能であり、CTA識別子、CTA次元、CTAが属するグリッドの次元、及びCTAが属するグリッドの識別子を記憶する。特殊レジスタ311は、仮想ドライバ320からフロントエンド306を介して受け取られたコマンドに応答して初期化中に書き込まれ、CTA実行の間に変化しない。
[0065]仮想ローカルレジスタ314は、各々のCTAスレッドによってスクラッチ空間として使用される。各々のレジスタは、1つのCTAスレッド(又は1つの仮想処理エンジン310)の排他的使用に割り振られ、ローカルレジスタ314内のデータは、ローカルレジスタが割り振られたCTAスレッドへのみアクセス可能である。共有メモリ316は、(単一のCTA内の)全てのCTAスレッドへアクセス可能である。共有メモリ316内の場所は、同じCTA内のCTAスレッド(又は仮想コア308内の仮想処理エンジン)へアクセス可能である。パラメータ・メモリ318はランタイムパラメータ(定数)を記憶する。ランタイムパラメータはCTAスレッド(又は仮想処理エンジン310)によって読み出し可能であるが、書き込みはできない。1つの実施形態において、仮想ドライバ320はパラメータをパラメータ・メモリ318に提供するが、この提供は、これらのパラメータを使用するCTAの実行開始を仮想プロセッサ302に命令する前に行われる。CTA内のCTAスレッド(又は仮想コア308内の仮想処理エンジン310)は、メモリ・インタフェース322を介してグローバルメモリ304にアクセスすることができる。
[0066]仮想アーキテクチャ300において、仮想プロセッサ302は仮想ドライバ320の制御のもとでコプロセッサとして動作する。仮想アーキテクチャの仕様は、有利には、仮想アプリケーション・プログラム・インタフェース(API)を含む。仮想APIは、仮想ドライバ320によって認識される関数呼び出し、及び各々の関数呼び出しが作成するように期待される振る舞いを識別する。並列スレッド・コンピューティングを目的とする仮想APIの例示的関数呼び出しは、下記で説明される。
[0067]仮想アーキテクチャ300は、様々なハードウェア・プラットフォームの上で現実化可能である。1つの実施形態において、仮想アーキテクチャ300は、図1のシステム100の中で現実化され、PPU122は仮想プロセッサ302を実現し、CPU102上で実行するPPUドライバ・プログラムは仮想ドライバ320を実現する。グローバルメモリ304は、システム・メモリ104及び/又はPPメモリ124の中で実現可能である。
[0068]1つの実施形態において、PPU122は1つ又は複数の処理コアを含む。処理コアは、単一命令複数データ(SIMD)及びマルチスレッド手法を使用して、(仮想命令装置312を実現する)単一の命令装置からの多数の(例えば、384又は768個の)スレッドの並行実行をサポートする。各々のコアはP個の(例えば、8個、16個などの)並列処理エンジン302のアレイを含む。処理エンジン302は命令装置からSIMD命令を受け取って実行するように構成され、P個のスレッドまでのグループが並列に処理されることを可能にする。コアはマルチスレッドされ、処理エンジンはG個(例えば、24個)までのスレッド・グループを並行して実行することができる。この並行実行は、例えば、各々のスレッドに関連した最新状態情報を維持し、処理エンジンが1つのスレッドから他のスレッドへ迅速に切り換わることによって行われる。こうして、コアは、並行して、各々P個のスレッドから成るG個のSIMDグループ、即ち、全部でP*G個の並行スレッドを並行して実行する。この現実化において、P*G≧n0である限り、(仮想の)CTAスレッドと、現実のPPU122の上で実行している並行スレッドとの間に、1対1の対応が存在し得る。
[0069]特殊レジスタ311はPPU122の中で実現可能である。この実現は、各々の処理コアにP*G項目レジスタ・ファイルを提供し、各々の項目がスレッドIDを記憶できるようにし、更にCTA ID、グリッドID、並びにCTA次元及びグリッド次元を記憶するためのグローバル読み出し可能レジスタの集合を提供することによって行われる。代替として、特殊レジスタ311は他の記憶場所を使用して実現可能である。
[0070]ローカルレジスタ314は、PPU122の中でローカルレジスタ・ファイルとして実現可能である。ローカルレジスタ・ファイルは物理的又は論理的にP個のレーン(lane)に分割され、各々のレーンは或る数の項目を有する(ここで、各々の項目は、例えば、32ビット・ワードを記憶できる)。P個の処理エンジンの各々へ1つのレーンを割り当て、異なるレーンの中の対応する項目に、同じプログラムを実行する別々のスレッドのデータを入れて、SIMD実行を容易にすることができる。レーンの異なる部分は、G個の並行スレッド・グループの異なるものへ割り振られ、したがってローカルレジスタ・ファイル内の所与の項目は、特定のスレッドへのみアクセス可能である。1つの実施形態において、ローカルレジスタ・ファイル内の或る項目は、スレッド識別子を記憶するために予約され、特殊レジスタ311の1つを実現する。
[0071]共有メモリ316は、PPU122の中で、共有レジスタ・ファイルとして実現されてもよく、処理エンジンが共有メモリ内の場所との間で読み出し又は書き込みを可能にする相互接続を有する共有オンチップ・キャッシュメモリとして実現されてもよい。パラメータ・メモリ318は、PPU122の中で、同じ共有レジスタ・ファイル又は共有メモリ316を実現する共有キャッシュメモリの中の指定セクションとして実現されてもよく、処理エンジンが読み出し専用アクセスを有する別個の共有レジスタ・ファイル又はオンチップ・キャッシュメモリとして実現されてもよい。1つの実施形態において、パラメータ・メモリを実現する区域は、更に、CTA ID及びグリッドID並びにCTA次元及びグリッド次元を記憶するために使用され、特殊レジスタ311の一部分を実現する。
[0072]1つの実施形態において、図1のCPU102で実行するPPUドライバ・プログラムは、メモリ(例えば、システム・メモリ104)内のプッシュバッファ(明示的には示されない)へコマンドを書き込むことによって、仮想API関数呼び出しに応答する。コマンドはプッシュバッファからPPU122によって読み出される。コマンドは、有利には、状態パラメータ、例えば、CTA内のスレッドの数、CTAを使用して処理される入力データセットのグローバルメモリ内の場所、実行されるCTAプログラムのグローバルメモリ内の場所、及び出力データが書き込まれるグローバルメモリ内の場所に関連づけられる。コマンド及び状態パラメータに応答して、PPU122は状態パラメータを自分のコアの1つへロードし、CTAパラメータ内で特定された或る数のスレッドが起動されてしまうまで、スレッドの起動を始める。1つの実施形態において、PPU122は制御論理を含み、この制御論理はスレッドが起動されるときスレッドIDをスレッドへ順次に割り当てる。スレッドIDは、例えば、ローカルレジスタ・ファイル内の指定場所、又はこの目的だけの特殊レジスタの中に記憶できる。
[0073]代替の実施形態において、仮想アーキテクチャ300は(例えば、幾つかのCPUにおける)単一スレッド処理コアで現実化される。処理コアは、n0個よりも少ない現実のスレッドを使用して全てのCTAスレッドを実行する。仮想プログラミング・モデルが、異なるCTAスレッド関連づける処理タスクは、例えば、1つのCTAスレッドのためにタスク(又はタスクの一部分)を実行し、続いて次のCTAスレッドのために実行し、以下同様に実行することによって単一のスレッドへ結合できる。ベクトル実行、SIMD実行、及び/又は機械内で利用可能な他の形式の並列実行は、複数CTAスレッドに関連づけられた処理タスクを並列に実行、又は同じCTAスレッドに関連づけられた複数の処理タスクを並列に実行するために活用できる。こうして、CTAが、単一のスレッド、n0個のスレッド、又は他の数のスレッドを使用して現実化できる。下記で説明されるように、仮想命令トランスレータは、有利には、目標仮想アーキテクチャ300へ書かれたコードを目標プラットフォームに特有の命令へ翻訳する。
[0074]本明細書で説明された仮想アーキテクチャは例であること、そして変形及び修正が可能であることが理解される。例えば、1つの代替の実施形態において、各々の仮想処理エンジンは、スレッドに割り当てられた一意のスレッドIDを記憶し、ローカル仮想レジスタ内の空間をこの目的に使用することのない専用のスレッドIDレジスタを有してもよい。
[0075]他の例として、仮想アーキテクチャは仮想コア308の内部構造の詳細を多かれ少なかれ特定してもよい。例えば、仮想コア308がP個のマルチスレッド仮想処理エンジンを含み、P路SIMDグループ内のCTAスレッドを実行するために処理エンジンが使用され、G個までのSIMDグループがコア308の中に共存し、P*GがT(CTA内のスレッドの最大数)を決定することを特定してもよい。異なるタイプのメモリ及び異なる共有レベルも特定可能である。
[0076]仮想アーキテクチャは、ハードウェア及び/又はソフトウェア要素の組み合わせを使用して各々のコンポーネントを定義及び制御する多様なコンピュータ・システムで現実化されてもよい。例として、ハードウェア・コンポーネントを使用する1つの現実化が説明されたが、本発明は特定のハードウェア現実からプログラミング・タスクの切り離しに関することを理解すべきである。
4.仮想アーキテクチャのプログラミング
[0077]図4は、本発明の実施形態に従って、仮想アーキテクチャ300を使用して目標プロセッサ又は目標プラットフォーム440を動作させる概念的モデル400である。モデル400が示すように、仮想アーキテクチャ300の存在は、目標プロセッサ又はプラットフォームのハードウェア実現から、コンパイルされたアプリケーション及びAPIを切り離す。
[0077]図4は、本発明の実施形態に従って、仮想アーキテクチャ300を使用して目標プロセッサ又は目標プラットフォーム440を動作させる概念的モデル400である。モデル400が示すように、仮想アーキテクチャ300の存在は、目標プロセッサ又はプラットフォームのハードウェア実現から、コンパイルされたアプリケーション及びAPIを切り離す。
[0078]アプリケーション・プログラム402は、前述した仮想プログラミング・モデルを利用するデータ処理アプリケーションを定義する。定義には、単一CTA及び/又はCTAグリッドの定義が含まれる。一般的に、アプリケーション・プログラム402は複数の様相を含む。第1に、プログラムは単一CTAスレッドの振る舞いを定義する。第2に、プログラムは(CTAスレッドの数で)CTAの次元を定義し、もしグリッドが使用されるのであれば、(CTAの数で)グリッドの次元を定義する。第3に、プログラムは、CTA(又はグリッド)によって処理される入力データセットを定義し、また出力データセットが記憶される場所を定義する。第4に、プログラムは、例えば、各々のCTA又はグリッドをいつ起動するかを含む、全体的な処理振る舞いを定義する。プログラムは、CTA又はグリッドの次元、新しいCTA又はグリッドの起動を控えるかどうかなどを動的に決定する追加のコードを含んでもよい。
[0079]アプリケーション・プログラム402は、高級プログラミング言語、例えば、C/C++、FORTRANなどで書かれてもよい。1つの実施形態において、アプリケーションC/C++プログラムは、1つの(仮想)CTAスレッドの振る舞いを直接特定する。他の実施形態において、アプリケーション・プログラムはデータ並列言語(例えば、Fortran 90、C*、又はData−Parallel C)を使用して書かれ、アレイ及び集合体データ構造へのデータ並列動作を特定する。そのようなプログラムは、1つの(仮想)CTAスレッドの振る舞いを特定する仮想ISAプログラム・コードにコンパイルできる。CTAスレッドの振る舞いを定義できるようにするため、言語拡張又は関数ライブラリが提供され、これらを介して、プログラマは並列CTAスレッド振る舞いを特定することができる。例えば、特殊記号又は変数は、スレッドID、CTA ID、及びグリッドIDに対応するように定義され、関数が提供される。関数を介して、プログラマは、CTAスレッドが他のCTAスレッドといつ同期すべきかを指示することができる。
[0080]アプリケーション・プログラム402がコンパイルされるとき、コンパイラ408は、CTAスレッド振る舞いを定義するアプリケーション・プログラム402の部分について、仮想ISAコード410を生成する。1つの実施形態において、仮想ISAコード410は図3の仮想アーキテクチャ300の仮想ISAとして表現される。仮想ISAコード410はプログラム・コードであるが、必ずしも特定の目標プラットフォーム上で実行できる形式ではない。仮想ISAコード410は他のプログラム・コードと同じように記憶及び/又は配布される。他の実施形態において、アプリケーション・プログラムは全体的又は部分的に仮想ISAコード410として特定されてもよく、コンパイラ408は全体的又は部分的にバイパスされてもよい。
[0081]仮想命令トランスレータ412は仮想ISAコード410を目標ISAコード414へ変換する。幾つかの実施形態において、目標ISAコード414は、目標プラットフォーム440によって直接実行できるコードである。例えば、図4の点線ボックスによって示されるように、1つの実施形態において、目標ISAコード414はPPU122内の命令装置430によって受け取られ、正しくデコードされる。目標プラットフォーム440の仕様に依存して、仮想ISAコード410は、目標プラットフォーム440上でn0個のスレッドの各々によって実行されるスレッド単位コードへ翻訳されてもよい。代替として、仮想ISAコード410はn0個よりも少ないスレッドとして実行されるプログラム・コードへ翻訳されてもよく、各々のスレッドが2つ以上のCTAスレッドに関する処理タスクを含む。
[0082]幾つかの実施形態において、CTA及び/又はグリッドの次元定義、並びに入力データセット及び出力データセットの定義は、仮想APIによって取り扱われる。アプリケーション・プログラム402は、仮想API関数のライブラリ404への呼び出しを含んでもよい。1つの実施形態において、仮想APIの仕様(例えば、関数名、入力、出力、及び効果を含むが、実現の詳細は含まない)がプログラマへ提供され、プログラマは仮想API呼び出しをアプリケーション・プログラム402の中へ直接組み込んで、仮想APIコード406を直接生成する。他の実施形態において、仮想APIコード406は、他の構文を使用してCTA及びグリッドを定義するアプリケーション・プログラム402をコンパイルすることによって生成される。
[0083]仮想APIコード406は、仮想実行ドライバ416を提供することによって部分的に現実化される。仮想実行ドライバ416はコード406の仮想APIコマンドを目標APIコマンド418へ翻訳し、目標APIコマンド418は目標プラットフォーム440によって処理されてもよい。例えば、図4の点線ボックスによって示されるように、1つの実施形態において、目標APIコマンド418はPPUドライバ432によって受け取られ、処理されてもよい。PPUドライバ432は、対応するコマンドをPPU122のフロントエンド434へ通信する。(この実施形態において、仮想実行ドライバ416は、PPUドライバ432の1つの様相又は部分であってもよい。)他の実施形態において、仮想実行ドライバはコプロセッサのドライバに対応しなくてもよい。仮想実行ドライバは単純に、仮想実行ドライバを実行する同じプロセッサ上で、他のプログラム又はスレッドを起動する制御プログラムであってもよく仮想実行ドライバを実行する同じプロセッサ上で、他のプログラム又はスレッドを起動する。
[0084]理解すべきは、仮想命令トランスレータ412及び仮想実行ドライバ416は、CTA実行をサポート可能な任意のプラットフォーム又はアーキテクチャのために作成できることである。異なるプラットフォーム又はアーキテクチャの仮想命令トランスレータ412が同じ仮想ISAから翻訳できる範囲で、同じ仮想ISAコード410を任意のプラットフォーム又はアーキテクチャと共に使用できる。こうして、アプリケーション・プログラム402は、各々の可能なプラットフォーム又はアーキテクチャのために再コンパイルされる必要はない。
[0085]更に、目標プラットフォーム440が、図4で示されるようなPPU及び/又はPPUドライバを含むことは必要でない。例えば、1つの代替の実施形態において、目標プラットフォームはCPUであり、このCPUはソフトウェア手法を使用して多数のスレッドの並行実行をエミュレートし、目標ISAコード及び目標APIコマンドは、目標CPUによって実行されるプログラム(又は相互に通信するプログラムのグループ)内の命令に対応する。目標CPUは、例えば、シングルコア又はマルチコアのCPUであってもよい。
5.仮想ISAの例
[0086]本発明の実施形態に従った仮想ISAの例を、今から説明する。前述したように、仮想ISAは、有利には、前述した仮想プログラミング・モデル(CTA及びグリッド)に対応する。したがって、この実施形態において、コンパイラ408によって生成された仮想ISAコード410は、図3の仮想コア308内で仮想処理エンジン310の1つによって実行される単一CTAスレッドの振る舞いを定義する。振る舞いは他のCTAスレッドとの協調的交信、例えば、同期及び/又はデータ共有を含んでもよい。
[0086]本発明の実施形態に従った仮想ISAの例を、今から説明する。前述したように、仮想ISAは、有利には、前述した仮想プログラミング・モデル(CTA及びグリッド)に対応する。したがって、この実施形態において、コンパイラ408によって生成された仮想ISAコード410は、図3の仮想コア308内で仮想処理エンジン310の1つによって実行される単一CTAスレッドの振る舞いを定義する。振る舞いは他のCTAスレッドとの協調的交信、例えば、同期及び/又はデータ共有を含んでもよい。
[0087]理解すべきは、本明細書で説明される仮想ISAは、単なる例であること、及び本明細書で説明される特定の要素、又は要素の組み合わせは、本発明の範囲を限定しないことである。幾つかの実施形態において、プログラマは仮想ISAでコードを書いてもよく、他の実施形態において、プログラマは他の高級言語(例えば、FORTRAN、C、C++)でコードを書き、コンパイラ408が仮想ISAコードを生成する。プログラマは、更に、或る部分は高級言語で書かれ、他の部分は仮想ISAで書かれる「混合」コードを書くことが可能である。
5.1.特殊変数
[0088]図5は、例示的仮想ISAによって定義された「特殊」変数を列挙する表500である(本明細書では、特殊変数を知らせるため「%」接頭辞が使用される)。これらの変数は図2のプログラミング・モデルに関する。このプログラミング・モデルにおいて、各々のスレッド204はCTA202内の位置によって識別され、転じてCTA202は、或る数のグリッド200の中にある特定の1つに存在する。幾つかの実施形態において、表500の特殊変数は、図3の仮想アーキテクチャ300の特殊レジスタ311に対応する。
[0088]図5は、例示的仮想ISAによって定義された「特殊」変数を列挙する表500である(本明細書では、特殊変数を知らせるため「%」接頭辞が使用される)。これらの変数は図2のプログラミング・モデルに関する。このプログラミング・モデルにおいて、各々のスレッド204はCTA202内の位置によって識別され、転じてCTA202は、或る数のグリッド200の中にある特定の1つに存在する。幾つかの実施形態において、表500の特殊変数は、図3の仮想アーキテクチャ300の特殊レジスタ311に対応する。
[0089]表500において、CTA及びグリッドの各々は3次元空間で定義されること、及び異なるグリッドは1次元空間で順次に番号を付けられることが想定される。仮想ISAは、CTAが起動されたとき図5の特殊変数が初期化されることを期待し、仮想ISAコードは、初期化なしに、これらの変数を単純に使用することができる。特殊変数の初期化は、仮想APIを参照して下記で説明される。
[0090]図5で示されるように、特殊変数%ntid=(%ntid.x,%ntid.y,%ntid.z)の最初の3ベクトルは、CTAの次元を(スレッドの数で)定義する。CTAの全てのスレッドは同じ%ntidベクトルを共有できる。仮想アーキテクチャ300において、%ntidベクトルの値は、後述するように、CTAの次元を確立する仮想API関数呼び出しを介して仮想プロセッサ302へ提供されることが期待される。
[0091]図5で示されるように、特殊変数%tid=(%tid.x,%tid.y,%tid.z)の2番目の3ベクトルは、CTAにおける所与のスレッドのスレッドIDを参照する。図3の仮想アーキテクチャ300において、CTAの各々のスレッドが起動されるとき、制約0≦%tid.x<%ntid.x、0≦%tid.y<%ntid.y、及び0≦%tid.z<ntid.zを満足させる一意の%tidベクトルを仮想プロセッサ302が割り当てることが期待される。1つの実施形態において、%tidベクトルはパック32ビット・ワードとして記憶されるように定義されてもよい(例えば、%tid.xには16ビット、%tid.yには10ビット、%tid.zには6ビット)。
[0092]図5で示されるように、特殊変数%nctaid=(%nctaid.x,%nctaid.y,%nctaid.z)の3番目の3ベクトルは、グリッドの次元を(CTAの数で)定義する。図3の仮想アーキテクチャ300において、CTAのグリッドの次元を確立する仮想API関数呼び出しを介して、%nctaidベクトルの値が仮想プロセッサ302へ提供されることが期待される。
[0093]図5で示されるように、特殊変数%ctaid=(%ctaid.x,%ctaid.y,%ctaid.z)の4番目の3ベクトルは、グリッドにおける所与のCTAのCTA IDを参照する。図3の仮想アーキテクチャ300において、CTAが起動されるとき、CTAについて制約0≦%ctaid.x<%nctaid.x、0≦%ctaid.y<%nctaid.y、及び0≦%ctaid.z<%nctaid.zを満足させる一意の%ctaidベクトルが仮想プロセッサ302へ提供されることが期待される。
[0094]特殊変数は、更に、CTAが属するグリッドのためにグリッド識別子を提供するスカラー%gridid変数を含む。図3の仮想アーキテクチャ300において、%gridid値が仮想プロセッサ302へ提供されて、現在のCTAを一部分とするグリッドを識別することが期待される。%gridid値は、有利には、例えば、複数のグリッドが使用されて大きな問題の異なる部分を解決するとき、仮想ISAコードの中で使用される。
5.2.プログラムで定義される変数及び仮想状態空間
[0095]仮想ISAによって、プログラマ(又はコンパイラ)は任意の数の変数を定義して、処理されているデータ項目を表すことができる。変数は、型及び「仮想状態空間」によって定義される。仮想状態空間は、変数がどのように使用されるか、また変数がどの範囲まで共有されるかを指示する。変数は、目標プラットフォームで利用可能なレジスタ又は他のメモリ構造を使用して現実化される。多くの目標プラットフォームにおいて、状態空間は特定の変数の現実化に使用されるメモリ構造の選択に影響を及ぼし得る。
[0095]仮想ISAによって、プログラマ(又はコンパイラ)は任意の数の変数を定義して、処理されているデータ項目を表すことができる。変数は、型及び「仮想状態空間」によって定義される。仮想状態空間は、変数がどのように使用されるか、また変数がどの範囲まで共有されるかを指示する。変数は、目標プラットフォームで利用可能なレジスタ又は他のメモリ構造を使用して現実化される。多くの目標プラットフォームにおいて、状態空間は特定の変数の現実化に使用されるメモリ構造の選択に影響を及ぼし得る。
[0096]図6は、例示的仮想ISA実施形態でサポートされる変数型を列挙する表600である。4つの型がサポートされる。即ち、型を有しないビット、符号付き整数、符号のない整数、及び浮動小数点である。型を有しない変数は、単純に単一のビット又は特定長ビットのグループである。符号付き及び符号のない整数形式、並びに浮動小数点形式は、通常の形式(例えば、IEEE754標準)に従って定義されてもよい。
[0097]この実施形態において、各々の型について複数の幅がサポートされ、幅を特定するためパラメータ<n>が使用される。したがって、例えば、.s16は16ビット符号付き整数を、.f32は32ビット浮動小数点数を表し、以下同様である。表600で示されるように、幾つかの変数型は或る一定の幅へ制限される。例えば、浮動小数点変数は少なくとも16ビットでなければならず、整数型は少なくとも8ビットでなければならない。仮想ISAの現実化は、特定された幅の全てをサポートすることを期待される。もしプロセッサのデータ経路及び/又はレジスタが、最も広い幅よりも狭ければ、当技術分野で知られるように、複数のレジスタ及びプロセッサ・サイクルを使用して、より広い型を取り扱うことができる。
[0098]注意すべきは、本明細書で使用されるデータの型及び幅は、例であって、本発明を限定しないことである。
[0099]図7は、例示的仮想ISAでサポートされる仮想状態空間を列挙する表である。9つの状態空間が定義され、これらは図3の仮想アーキテクチャにおける異なる共有レベル及び可能な記憶場所に対応する。
[0100]最初の3つの状態空間はスレッド・レベルで共有される。これは、各々のCTAスレッドが変数の別々のインスタンスを有し、CTAスレッドは他のCTAスレッドのインスタンスへのアクセスを有しないことを意味する。仮想レジスタ(.reg)状態空間は、有利には、オペランド、一時値(temporary value)、及び/又は各々のCTAスレッドによって行われる計算の結果を定義するために使用される。プログラムは任意の数の仮想レジスタを宣言できる。仮想レジスタは、計算されたアドレスではなく、静的なコンパイル時の名前によってのみアドレス可能である。この状態空間は、図3の仮想アーキテクチャ300におけるローカル仮想レジスタ314に対応する。
[0101]特殊レジスタ(.sreg)状態空間は、仮想アーキテクチャ300の特殊レジスタ311に記憶される図5の既定の特殊変数に対応する。幾つかの実施形態において、仮想ISAコードは.sreg空間内で他の変数を宣言しなくてもよいが、特殊変数を計算への入力として使用してもよい。全てのCTAスレッドは.sreg状態空間内の変数を読み出すことができる。%tid(又はそのコンポーネント)の場合、各々のCTAスレッドは自分の一意のスレッド識別子を読み出す。.sreg状態空間内の他の変数の場合、同じCTA内の全てのCTAスレッドは、同じ値を読み出す。
[0102]スレッド単位(per−thread)ローカルメモリ(.local)変数は、CTAスレッド単位で割り振り及びアドレスされるグローバルメモリ304の領域に対応する。言い換えれば、CTAスレッドが.local変数にアクセスするとき、CTAスレッドは変数の自分自身のインスタンスにアクセスし、1つのCTAスレッド内で行われた.local変数への変更は、他のCTAスレッドに影響しない。.reg及び.sreg状態空間とは異なり、スレッド単位ローカルメモリは、計算されたアドレスを使用してアドレスされることが可能である。
[0103]次の2つの状態空間はCTA単位(per−CTA)変数を定義する。これは、各々のCTAが変数の1つのインスタンスを有し、このインスタンスがCTA(仮想)スレッドの任意のものによってアクセスされてもよいことを意味する。共有(.shared)変数は、CTAスレッドの任意のものによって読み出し又は書き込まれてもよい。幾つかの実施形態において、この状態空間は仮想アーキテクチャ300(図3)の仮想共有メモリ316へマップする。仮想アーキテクチャ300の現実化において、.shared状態空間はオンチップ共有メモリ実現(例えば、共有レジスタ・ファイル又は共有キャッシュメモリ)の上にマップでき、他の現実化においては、.shared状態空間は、他のグローバルアクセス可能メモリとして割り振り及びアドレスされるオフチップ・メモリのCTA単位領域にマップできる。
[0104]パラメータ(.param)変数は読み出し専用であり、CTA内の任意の(仮想)スレッドによって読み出されることが可能である。この状態空間は、仮想アーキテクチャ300のパラメータ・メモリ318へマップし、例えば、オンチップ共有パラメータ・メモリ又はキャッシュメモリの中で現実化されるか、他のグローバルアクセス可能メモリと同じように割り振り及びアドレスされるグローバルアクセス可能オフチップ・メモリの領域の中で現実化されることが可能である。これらの変数は、仮想ドライバ320からのドライバ・コマンドに応答して初期化されることが期待される。
[0105]定数(.const)状態空間は、グリッド内のCTAの(仮想)スレッドによって読み出されることの可能な(しかし、修正されない)グリッドごとの定数を定義するために使用される。仮想アーキテクチャ300において、.const状態空間は、CTAスレッドが読み出し専用アクセスを有するグローバルメモリ内の領域へマップされてよもい。.const状態空間は、オンチップ共有パラメータ・メモリ又はキャッシュメモリの中、又は他のグローバルアクセス可能メモリと同じように割り振り及びアドレスされるグローバルアクセス可能オフチップ・メモリのグリッド単位(per−grid)領域の中で現実化されることが可能である。.param状態空間と同じように、.const状態空間内の変数は、仮想ドライバ320からのドライバ・コマンドに応答して初期化されることが期待される。
[0106]残りの3つの状態空間は「文脈」変数を定義する。これらの変数は、アプリケーションに関連したCTA内の(仮想)スレッドからアクセス可能である。これらの状態空間は、仮想アーキテクチャ300内のメモリ304にマップする。グローバル(.global)変数は汎用目的に使用されることが可能である。幾つかの実施形態において、共有テクスチャ(.tex)及びサーフェイス(.surf)の特有状態空間も定義されてもよい。これらの状態空間は、例えば、グラフィックス関連アプリケーションに有用可能であり、2D(又は幾つかの実施形態では3D)アレイの各画素に対応するデータ値を提供するグラフィックス・テクスチャ及び画素サーフェイス・データ構造へのアクセスを定義及び提供するために使用できる。
[0107]図4の仮想ISAコード410において、変数は、状態空間、型、及び名前を特定することによって宣言される。名前はプレースホルダであり、プログラマ又はコンパイラによって選択されてもよい。したがって、例えば、
.reg .b32 vrl;
は、vrlと名前を付けられた仮想レジスタ状態空間で型のない32ビット変数を宣言する。仮想ISAコードの後続の行は、例えば、演算のソース又はデスティネーションとして、vrlを参照することができる。
.reg .b32 vrl;
は、vrlと名前を付けられた仮想レジスタ状態空間で型のない32ビット変数を宣言する。仮想ISAコードの後続の行は、例えば、演算のソース又はデスティネーションとして、vrlを参照することができる。
[0110]例示的仮想ISAは、更に、仮想変数のアレイ及びベクトルをサポートする。例えば、
.global .f32 resultArray[1000][1000];
は、32ビット浮動小数点数の仮想グローバルアクセス可能1000×1000アレイを宣言する。仮想命令トランスレータ412は、割り当てられた状態空間に対応するアドレス可能メモリ領域へアレイをマップできる。
.global .f32 resultArray[1000][1000];
は、32ビット浮動小数点数の仮想グローバルアクセス可能1000×1000アレイを宣言する。仮想命令トランスレータ412は、割り当てられた状態空間に対応するアドレス可能メモリ領域へアレイをマップできる。
[0113]1つの実施形態におけるベクトルは、ベクトル接頭辞.v<m>を使用して定義されてもよい。ここでmはベクトルの成分の数である。例えば、
.reg .v3 .f32 vpos;
は、スレッド単位仮想レジスタ状態空間内の32ビット浮動小数点数の3成分ベクトルを宣言する。ベクトルが一度宣言されると、その成分は接尾辞を使用して識別されてもよい。例えば、vpos.x、vpos.y、vpos.zのようになる。1つの実施形態において、m=2,3,又は4が許され、成分を識別するため、接尾辞、例えば、(.x,.y,.z,.w)、(.0,.1,.2,.3)、又は(.r,.g,.b,.a)が使用される。
.reg .v3 .f32 vpos;
は、スレッド単位仮想レジスタ状態空間内の32ビット浮動小数点数の3成分ベクトルを宣言する。ベクトルが一度宣言されると、その成分は接尾辞を使用して識別されてもよい。例えば、vpos.x、vpos.y、vpos.zのようになる。1つの実施形態において、m=2,3,又は4が許され、成分を識別するため、接尾辞、例えば、(.x,.y,.z,.w)、(.0,.1,.2,.3)、又は(.r,.g,.b,.a)が使用される。
[0116]変数は仮想であるから、仮想ISAコード410は、任意の状態空間の中で任意の数の変数を定義又は参照してもよい(例外は.sregである。この場合、変数は予め定義されている)。仮想ISAコード410の中で特定の状態空間について定義される変数の数は、特定のハードウェア実現の中の対応する型の記憶量を超過することが可能である。仮想命令トランスレータ412は、有利には、適切な記憶管理命令(例えば、レジスタとオフチップ・メモリ間でデータを移動すること)を含んで、必要とされるとき変数を利用可能にするように構成される。仮想命令トランスレータ412は、更に、一時変数が最早必要とされない場合を検出し、一時変数に割り振られた空間を他の変数に再使用させることができる。レジスタを割り振る従来のコンパイラ手法が使用できる。
[0117]更に、例示的仮想ISAはベクトル変数の型を定義するが、目標プラットフォームがベクトル変数をサポートすることは要求されない。仮想命令トランスレータ412は、適切な数(例えば、2、3、又は4)のスカラーの集まりとして、ベクトル変数を実現できる。
5.3.仮想命令
[0118]図8A〜図8Hは、例示的仮想ISAで定義された仮想命令を列挙する表である。命令は、命令の効果によって定義される。例えば、1つ又は複数のオペランドを使用して特定の結果を計算し、この結果をデスティネーション・レジスタの中に置き、レジスタ値を設定する、などである。大部分の仮想命令は、入力及び/又は出力の形式を識別するような型に分けられ、命令実行の様相は型に依存する可能性がある。命令の一般形式は、次のとおりである。
name.<type> result,operands;
[0118]図8A〜図8Hは、例示的仮想ISAで定義された仮想命令を列挙する表である。命令は、命令の効果によって定義される。例えば、1つ又は複数のオペランドを使用して特定の結果を計算し、この結果をデスティネーション・レジスタの中に置き、レジスタ値を設定する、などである。大部分の仮想命令は、入力及び/又は出力の形式を識別するような型に分けられ、命令実行の様相は型に依存する可能性がある。命令の一般形式は、次のとおりである。
name.<type> result,operands;
[0120]ここで、nameは命名の名前であり、.<type>は図6で列挙された型のプレースホルダであり、resultは結果を記憶する変数であり、operandsは1つ又は複数の変数であって、これらの変数は入力として命令へ提供される。1つの実施形態において、仮想アーキテクチャ300は、レジスタ間(register−to−register)プロセッサである。メモリ・アクセス(図8F)以外の演算のresult及びoperandsは、仮想レジスタ状態空間.reg(又は幾つかのオペランドの場合には、特殊レジスタ状態空間.sreg)の中の変数であることを要求される。
[0121]目標プラットフォームは、仮想ISAの命令の各々を現実化することを期待される。命令は、特定された効果を作成する対応機械命令(本明細書では、「ハードウェア・サポート」と呼ばれる)、又は実行されたとき、特定された効果を作成する機械命令のシーケンスとして現実化できる(本明細書では、「ソフトウェア・サポート」と呼ばれる)。特定の目標プラットフォームに対する仮想命令トランスレータ412は、有利には、各々の仮想命令に対応する機械命令又は機械命令シーケンスを識別するように構成される。
[0122]下記のサブセクションは、図8A〜図8Hで列挙された命令の様々な種類を説明する。理解すべきは、本明細書で呈示された命令のリストは例であること、及び仮想ISAは本明細書で明示的に説明されない追加の命令を含んでもよく、本明細書で説明された命令の幾つか又は全部を排除できることである。
5.3.1 仮想命令−算術
[0123]図8Aは、例示的仮想ISAで定義された算術演算を列挙する表800である。この実施形態において、仮想アーキテクチャはレジスタ間算術のみをサポートし、全ての算術演算は1つ又は複数の仮想レジスタ・オペランド(図8Aのa,b,cで表される)を操作して、仮想レジスタへ書き込まれる結果(d)を作成する。こうして、算術演算のオペランド及びデスティネーションは、常に、仮想レジスタ状態空間.regの中にあるが、例外として、図5の(特殊レジスタ状態空間.sregの中にある)特殊レジスタをオペランドとして使用することができる。
[0123]図8Aは、例示的仮想ISAで定義された算術演算を列挙する表800である。この実施形態において、仮想アーキテクチャはレジスタ間算術のみをサポートし、全ての算術演算は1つ又は複数の仮想レジスタ・オペランド(図8Aのa,b,cで表される)を操作して、仮想レジスタへ書き込まれる結果(d)を作成する。こうして、算術演算のオペランド及びデスティネーションは、常に、仮想レジスタ状態空間.regの中にあるが、例外として、図5の(特殊レジスタ状態空間.sregの中にある)特殊レジスタをオペランドとして使用することができる。
[0124]表800の算術演算のリストは、4つの基本的算術演算を含む。即ち、加算(add)、減算(sub)、乗算(mul)、及び除算(div)である。これらの演算は、整数及び浮動小数点データの全ての型で実行可能であり、入力と同じ型の結果を作成する。幾つかの実施形態において、丸めモード修飾子を命令に付加し、結果をどのように丸め、整数オペランドの場合に飽和限度を課すべきかどうかをプログラマに特定させることができる。
[0125]オペランドa,b,cを有する3複合算術演算もサポートされる。即ち、乗算・加算(mad)、融合乗算・加算(fma)、及び絶対差分合計(sad)である。乗算・加算は積a*bを計算し(丸めは括弧で指示される)、結果にcを加える。融合乗算・加算はmadとは異なり、cを加える前に積a*bは丸められない。絶対差分合計は、絶対値|a−b|を計算した後、cを加える。
[0126]剰余(rem)演算は整数オペランド上でのみ実行され、オペランドaがオペランドbによって割られるとき、剰余(a mod b)を計算する。絶対値(abs)及び否定(neg)は単項演算であって、浮動小数点又は符号付き整数形式のオペランドaへ適用できる。整数又は浮動小数点オペランドへ適用できる最小(min)及び最大(max)演算は、より小さいオペランド、又はより大きいオペランドへ、デスティネーション・レジスタを設定する。1つ又は双方のオペランドが、(例えば、IEEE754標準に従って)非正規数である特殊ケースも、特定されてもよい。
[0127]表800内の残りの演算は、浮動小数点型についてのみ行われる。端数(frc)演算は、その入力の端数部分を返す。正弦(sin)、余弦(cos)、及び比の逆正接(atan2)は、三角関数に対応する便利な命令を提供する。底2の対数(lg2)及び累乗法(ex2)もサポートされる。逆数(rcp)、平方根(sqrt)、及び逆平方根(rsqrt)もサポートされる。
[0128]算術演算のこのリストは例であって、本発明を限定しないことに注意すべきである。他の演算又は演算の組み合わせがサポートされてもよく、それらの演算の中には、十分な頻度の呼び出しを期待される演算が含まれる。
[0129]幾つかの実施形態において、仮想ISAは、更に、ベクトル演算を定義する。図8Bは、例示的仮想ISAによってサポートされるベクトル演算を列挙する表810である。ベクトル演算は、オペランド・ベクトルa及びbのスカラー点乗積dを計算する点乗積(dot)演算、オペランド・ベクトルa及びbのベクトル交差積dを計算する交差積(cross)演算、及びオペランド・ベクトルaのスカラー長dを計算する量(mag)演算を含む。ベクトル簡約(vred)演算は、ベクトル・オペランドaの要素を横切って特定演算<op>を反復して行うことによってスカラー結果dを計算する。1つの実施形態において、簡約演算add、mul、min、及びmaxのみが、浮動小数点ベクトルに対してサポートされる。整数ベクトルについては、追加の簡約演算(例えば、下記で説明するように、and、or、及びxor)もサポートされてもよい。
[0130]これらの演算に加えて、(図8Bには列挙されていない)他のベクトル演算、例えば、ベクトル加算、ベクトル・スケーリングなども、仮想ISAで定義されてもよい。
[0131]前述したように、仮想アーキテクチャ300の幾つかのハードウェア現実化は、ベクトル処理をサポートしないかも知れない。そのような現実化の仮想命令トランスレータ412は、有利には、スカラー機械命令の適切なシーケンスを生成して、これらの演算を行うように適応させられる。当業者は適切なシーケンスを決定することができよう。
5.3.2.仮想命令−選択及びレジスタ設定
[0132]図8Cは、例示的仮想ISA内で定義される選択及びレジスタ設定演算を列挙する表820である。任意の数値データ型の上で行うことのできるこれらの演算は、比較演算のアウトカムに基づいてデスティネーション・レジスタを設定する。基本的選択(sel)演算は、もしcが非ゼロであればオペランドaを選択し、もしcが0であれば、オペランドbを選択する。比較及び設定(set)は、オペランドa及びbの上で比較演算<cmp>を行って比較結果tを生成し、次に比較結果tが真(〜0)であるか偽(0)であるかに基づいて、デスティネーション・レジスタdをブールの真(〜0)又は偽(0)へ設定する。1つの実施形態において、許される比較演算<cmp>は、等しい(もしa=bであれば、tは真である)、より大きい(もしa>bであれば、tは真である)、より小さい(もしa<bであれば、tは真である)、より大きいか等しい(もしa≧bであれば、tは真である)、より小さいか等しい(もしa≦bであれば、tは真である)、及び他の比較を含む。他の比較には、例えば、a及び/又はbが数値であるか、又は無定義値であるかが含まれる。
[0132]図8Cは、例示的仮想ISA内で定義される選択及びレジスタ設定演算を列挙する表820である。任意の数値データ型の上で行うことのできるこれらの演算は、比較演算のアウトカムに基づいてデスティネーション・レジスタを設定する。基本的選択(sel)演算は、もしcが非ゼロであればオペランドaを選択し、もしcが0であれば、オペランドbを選択する。比較及び設定(set)は、オペランドa及びbの上で比較演算<cmp>を行って比較結果tを生成し、次に比較結果tが真(〜0)であるか偽(0)であるかに基づいて、デスティネーション・レジスタdをブールの真(〜0)又は偽(0)へ設定する。1つの実施形態において、許される比較演算<cmp>は、等しい(もしa=bであれば、tは真である)、より大きい(もしa>bであれば、tは真である)、より小さい(もしa<bであれば、tは真である)、より大きいか等しい(もしa≧bであれば、tは真である)、より小さいか等しい(もしa≦bであれば、tは真である)、及び他の比較を含む。他の比較には、例えば、a及び/又はbが数値であるか、又は無定義値であるかが含まれる。
[0133]setb演算は、比較及び設定の変形であって、比較演算<cmp>の結果tと第3のオペランドcとの間で更なるブール演算<bop>を行う。ブール演算t<bop>cの結果は、デスティネーション・レジスタdがブールの真又は偽へ設定されるかどうかを決定する。1つの実施形態において、許されるブール演算<bop>は、and、or、及びxorを含む(下記で説明する図8Cを参照)。setp演算はsetbと類似しているが、例外は、2つの1ビット「述語」デスティネーション・レジスタが設定されることである。即ち、デスティネーション・レジスタd1はt<bop>cの結果へ設定され、デスティネーション・レジスタd2は(!t)<bop>cの結果へ設定される。
5.3.3.仮想命令−論理及びビット操作
[0134]図8Dは、例示的仮想ISAの中で定義される論理及びビット操作演算を列挙する表830である。ビットごとのブール演算and、or、及びxorは、オペランドa及びbの各ビット上で特定演算を行い、レジスタd内の対応するビットを結果へ設定することによって行われる。ビットごとの否定(not)演算は、オペランドaの各ビットを反転し、論理否定(cnot)演算は、もしaが0(ブールの偽)であれば、デスティネーション・レジスタを1(ブールの真)へ設定し、そうでなければ0(ブールの偽)へ設定する。
[0134]図8Dは、例示的仮想ISAの中で定義される論理及びビット操作演算を列挙する表830である。ビットごとのブール演算and、or、及びxorは、オペランドa及びbの各ビット上で特定演算を行い、レジスタd内の対応するビットを結果へ設定することによって行われる。ビットごとの否定(not)演算は、オペランドaの各ビットを反転し、論理否定(cnot)演算は、もしaが0(ブールの偽)であれば、デスティネーション・レジスタを1(ブールの真)へ設定し、そうでなければ0(ブールの偽)へ設定する。
[0135]ビット・シフトは左シフト(shl)及び右シフト(shr)演算によってサポートされ、これらの演算はオペランドbによって特定されたビット数だけオペランドa内のビット・フィールドを左又は右へシフトする。符号付き形式の場合、右シフトは、有利には、符号ビットに基づいて先導ビットを充填し、符号が付いていない形式の場合、右シフトは先導ビットを0で充填する。
5.3.4.仮想命令−形式変換
[0136]図8Eは、例示的仮想ISAで定義される形式変換演算を列挙する表840である。形式変換(cvt)命令は、第1の型<atype>のオペランドaを目標型<dtype>の同値へ変換し、結果をデスティネーション・レジスタdの中に記憶する。1つの実施形態において、有効な型は図6に列挙される。型のない値(.b<n>)は整数又は浮動小数点型との間で変換することはできない。形式変換命令の変形によって、プログラマは丸めモード<mode>を指定することができる。目標型として表現されたときに飽和する数の取り扱いも、指定されてもよい。
[0136]図8Eは、例示的仮想ISAで定義される形式変換演算を列挙する表840である。形式変換(cvt)命令は、第1の型<atype>のオペランドaを目標型<dtype>の同値へ変換し、結果をデスティネーション・レジスタdの中に記憶する。1つの実施形態において、有効な型は図6に列挙される。型のない値(.b<n>)は整数又は浮動小数点型との間で変換することはできない。形式変換命令の変形によって、プログラマは丸めモード<mode>を指定することができる。目標型として表現されたときに飽和する数の取り扱いも、指定されてもよい。
5.3.5.仮想命令−データ移動及びデータ共有
[0137]図8Fは、例示的仮想ISAで定義されるデータ移動及びデータ共有命令を列挙する表850である。移動(mov)演算は、デスティネーション・レジスタdを即値オペランドaの値へ設定するか、もしオペランドaがレジスタであれば、レジスタaの内容へ設定する。移動演算は、仮想レジスタ・タイプの状態空間、例えば、図7の.reg及び.sregへ制限することができる。
[0137]図8Fは、例示的仮想ISAで定義されるデータ移動及びデータ共有命令を列挙する表850である。移動(mov)演算は、デスティネーション・レジスタdを即値オペランドaの値へ設定するか、もしオペランドaがレジスタであれば、レジスタaの内容へ設定する。移動演算は、仮想レジスタ・タイプの状態空間、例えば、図7の.reg及び.sregへ制限することができる。
[0138]ロード(ld)命令は、メモリ内のソース場所からデスティネーション・レジスタdへ値をロードする。デスティネーション・レジスタdは、1つの実施形態において、仮想レジスタ(.reg)状態空間の中に存在しなければならない。.<space>修飾子は、ソース場所の状態空間を指定し、図7のアドレス可能状態空間、例えば、.reg及び.sreg以外の空間へ限定されることができる(.reg及び.sreg空間では、代わりに移動演算を使用することができる)。この実施形態における仮想アーキテクチャ300はレジスタ間プロセッサであるから、ロード命令は、有利には、アドレス可能状態空間から仮想レジスタ.reg状態空間へ変数を転送するために使用され、変数はオペランドとして使用できるようになる。
[0139]特有のソース場所は、様々な方法で定義され、異なるアドレス・モードをサポートするソース・パラメータ<src>を使用して識別される。例えば、幾つかの実施形態において、ソース・パラメータ<src>は、dの中に記憶される値を有する名前付きアドレス可能変数、ソース・アドレスを保持するレジスタへの参照、(即値オペランドとして供給される)オフセット値へ加えられるアドレスを保持するレジスタへの参照、又は即値絶対アドレスの任意の1つであってもよい。
[0140]同様に、記憶(st)演算は、ソース・レジスタa内の値を、デスティネーション・パラメータ<dst>によって識別されたメモリ場所へ記憶する。1つの実施形態におけるソース・レジスタaは、.reg状態空間の中に存在しなければならない。デスティネーションは書き込み可能及びアドレス可能状態空間(例えば、図7の.local、.global、又は.shared)の中に存在しなければならない。デスティネーション・パラメータ<dst>は、ロード命令のソース・パラメータ<src>と同じように、様々な方法で定義されて異なるアドレス・モードをサポートすることができる。記憶命令は、例えば、レジスタからアドレス可能状態空間へ演算結果を転送するために使用できる。
[0141]テクスチャ及びサーフェイス状態空間が提供される実施形態において、追加の仮想命令は、テクスチャ・メモリ状態空間(tex)から読み出し、サーフェイスメモリ状態空間から読み出し(suld)及び書き込む(sust)ために使用されてもよい。テクスチャ読み出しのオペランド(t,x,y)は、テクスチャ識別(t)及び座標(x,y)を指定する。同様に、サーフェイス読み出し又は書き込みのオペランド(s,x,y)は、サーフェイス識別子(s)及び座標(x,y)を指定する。
[0142]CTAスレッドは、他のCTAスレッドとデータを共有することによって、他のCTAスレッドと協力してもよい。例えば、CTA内のデータを共有するため、CTAスレッドはロード及び記憶仮想命令(並びに下記で説明するアトミック更新命令atom)を使用して、CTA単位仮想状態空間との間でデータを読み書きすることができる。こうして、1つのCTAスレッドは、適切に定義されたデスティネーション・アドレスを有するst.shared命令を使用して.shared状態空間へデータを書き込むことができ、同じCTA内の他のCTAスレッドは、続いてld.shared命令内の同じアドレスを使用することによってデータを読み出すことができる。この後で説明する同期命令(例えば、bar及びmembar)は、CTAスレッドの全体でデータ共有演算の適切なシーケンス、例えば、データ作成CTAスレッドがデータを書き込み、その後でデータ消費CTAスレッドがデータを読み出すシーケンスを確保するために使用できる。同様に、st.global及びld.global命令は、同じCTA内のCTAスレッド、同じグリッド内のCTA、及び/又は同じアプリケーション内の異なるグリッド間で、協力してデータを共有するために使用できる。
5.3.6.仮想命令−プログラム制御
[0143]図8Gは、例示的仮想ISAで提供されるプログラム制御演算を列挙する表860である。これらの制御演算は、当業者には良く知られていると思われるが、ログラマがプログラム実行をリダイレクトすることを可能にする。分岐(bra)はプログラムの流れを目標場所<target>へリダイレクトする。幾つかの実施形態において、分岐目標は、仮想ISAコード内で英数字ラベルを目標命令の前に置き、このラベルを分岐命令の目標識別子<target>として使用することによって定義される。例えば、1つの実施形態において、
label: add.int32 d,vrl,vr2;
は、ラベルlabelを有する分岐目標としてadd命令を識別する。次の命令、
bra label;
は、コード内のどこかにあって、ラベル付けされた命令へ実行をリダイレクトする。
[0143]図8Gは、例示的仮想ISAで提供されるプログラム制御演算を列挙する表860である。これらの制御演算は、当業者には良く知られていると思われるが、ログラマがプログラム実行をリダイレクトすることを可能にする。分岐(bra)はプログラムの流れを目標場所<target>へリダイレクトする。幾つかの実施形態において、分岐目標は、仮想ISAコード内で英数字ラベルを目標命令の前に置き、このラベルを分岐命令の目標識別子<target>として使用することによって定義される。例えば、1つの実施形態において、
label: add.int32 d,vrl,vr2;
は、ラベルlabelを有する分岐目標としてadd命令を識別する。次の命令、
bra label;
は、コード内のどこかにあって、ラベル付けされた命令へ実行をリダイレクトする。
[0148]call及び返却(ret)命令は、関数及びサブルーチン呼び出しをサポートする。fnameは関数又はサブルーチンを識別する。(1つの実施形態において、「サブルーチン」とは、単に返却値が無視される関数である。)関数fnameは.funcディレクティブを使用して宣言でき、関数を定義する仮想ISAコードも提供されてもよい。中括弧{}又は他のグループ記号は、関数又はサブルーチンを定義するコードを、他の仮想ISAコードから分離するために使用できる。
[0149]関数の場合、パラメータ・リスト<rv>は、返却値をどこに記憶するかを識別するために指定できる。関数及びサブルーチンの両者において、入力引数は、引数リスト<args>で指定される。callが実行されるとき、次の命令のアドレスが記憶され、retが実行されるとき、記憶されたアドレスへの分岐が取られる。
[0150]exit命令は、この命令に遭遇したCTAスレッドを終了させる。トラップ命令は、プロセッサ定義又はユーザ定義のトラップ・ルーチンを起動する。中断点(brkpt)命令は実行を一時停止し、例えば、デバッグ目的に有用である。無操作命令(nop)は、実行されたとき効果を有しない命令である。無操作命令は、例えば、どれほど遅れて次の演算が実行できるかを制御するために使用されてもよい。
5.3.7.仮想命令−並列スレッド
[0151]図8Hは、本発明の実施形態に従った例示的仮想ISAで提供される明示並列仮想命令を列挙する表870である。これらの命令は、CTA実行に望まれる協調的スレッド振る舞い、例えば、CTAスレッド間のデータ交換をサポートする。
[0151]図8Hは、本発明の実施形態に従った例示的仮想ISAで提供される明示並列仮想命令を列挙する表870である。これらの命令は、CTA実行に望まれる協調的スレッド振る舞い、例えば、CTAスレッド間のデータ交換をサポートする。
[0152]障壁(bar)命令は、この命令に達するCTAスレッドが、更なる命令を実行する前に待機すべきことを指示する。待機する時間は、(同じCTA内の)他の全CTAスレッドが同じ障壁命令に達するまでである。任意の数の障壁命令がCTAプログラム内で使用されてもよい。1つの実施形態において、障壁命令は(どれほど多くの障壁が使用されようとも)パラメータを必要としない。なぜなら、どのスレッドも(n+1)番目の障壁へ進む前に、全てのCTAスレッドがn番目の障壁に到達しなければならないからである。
[0153]他の実施形態において、障壁命令は例えば、特定の障壁で待機すべきCTAスレッドの数(又は特定のCTAスレッドの識別子)を指定することによって、パラメータ化されても。
[0154]更に他の実施形態は、「待機」及び「非待機」障壁の両者を提供する。待機障壁命令において、CTAスレッドは、他の関連性のあるCTAスレッドも障壁に到達するまで待機する。非待機命令において、CTAスレッドは、このスレッドが到着したが、他のCTAスレッドが到着する前に継続できることを指示する。所与の障壁において、幾つかのCTAスレッドは、他のCTAスレッドの非待機中に待機していてもよい。
[0155]幾つかの実施形態において、bar仮想命令は、共有メモリ状態空間を使用して、協力又はデータを共有しているCTAスレッドを同期させるために使用されてもよい。例えば、(CTAスレッドの幾つか又は全部を含んでもよい)CTAスレッドのセットの各スレッドが、スレッド単位変数(例えば、.fp32仮想レジスタ変数myData)の中で或るデータを作成し、セット内の他のCTAスレッドによって、作成されたデータを読み出すと仮定する。命令シーケンスは、次のとおりである。
st.shared.fp32 myWriteAddress,myData;
bar;
Id.shared.fp32 myData,myReadAddress;
ここで、myWriteAddress及びmyReadAddressは、.shared状態空間内のアドレスに対応するスレッド単位変数であり、上記の命令シーケンスは所望の振る舞いを提供する。各々のCTAスレッドが、作成されたデータを共有メモリへ書き込んだ後、各々のCTAスレッドは、全てのCTAスレッドがデータを記憶してしまうまで待機し、次に共有メモリから(異なるCTAスレッドによって書き込まれた可能性のある)データを読み出すように進行する。
st.shared.fp32 myWriteAddress,myData;
bar;
Id.shared.fp32 myData,myReadAddress;
ここで、myWriteAddress及びmyReadAddressは、.shared状態空間内のアドレスに対応するスレッド単位変数であり、上記の命令シーケンスは所望の振る舞いを提供する。各々のCTAスレッドが、作成されたデータを共有メモリへ書き込んだ後、各々のCTAスレッドは、全てのCTAスレッドがデータを記憶してしまうまで待機し、次に共有メモリから(異なるCTAスレッドによって書き込まれた可能性のある)データを読み出すように進行する。
[0156]メモリ障壁(membar)命令は、各々のCTAスレッドの既に要求されたメモリ演算(又は少なくとも全ての書き込み演算)が完了されるまで各CTAスレッドが待機すべきことを指示する。この命令は、membar命令の後で起こるメモリ・アクセスが、メモリ・アクセス前の書き込み演算の結果を見ることを保証する。1つの実施形態におけるmembar命令は、任意的な状態空間名<space>を使用して、特定された状態空間を目標とするメモリ演算へmembar命令の範囲を制限する。指定された状態空間は、メモリ状態空間でなければならない(例えば、.reg又は.sreg状態空間であってはならない)。もし状態空間名が特定されなければ、CTAスレッドは、全てのメモリ状態空間を目標とする全ての係属演算が完了するまで待機する。
[0157]アトミック更新(atom)命令は、参照<ref>によって識別される共有変数aへのアトミック更新(読み出し−修正−書き込み)を生じる。共有変数aは任意の共有状態空間の中に存在可能であり、他のメモリ参照と同じように、様々なアドレス・モードが使用できる。例えば、<ref>は、名前付きアドレス可能変数a、変数aのアドレスを保持するレジスタへの参照、変数aを突き止めるため(即値オペランドとして供給される)オフセット値へ加えられるアドレスを保持するレジスタへの参照、又は変数aの即値絶対アドレスの任意のものであってもよい。CTAスレッドは、共有状態空間場所からデスティネーション・レジスタdへ変数aをロードし、オペランドa及び(演算に依存して)第2及び第3のオペランドb及びcの上で行われる特定演算<op>を使用して変数aを更新し、更新結果は、<ref>によって識別される場所へ戻される。デスティネーション・レジスタdは、aの最初のロード値を保持する。ロード、更新、及び記憶演算はアトミックに行われ、第1のCTAスレッドがアトミック更新を実行している間、他のCTAスレッドが変数aにアクセスしないように保証する。1つの実施形態において、変数aは.global又は.shared状態空間へ限定され、前述したロード及び記憶演算と同じやり方で指定されてもよい。
[0158]幾つかの実施形態において、或る一部の演算のみがアトミック更新として行われもよい。例えば、1つの実施形態において、もしaが浮動小数点型であれば、次の演算<op>のみが指定されてもよい。即ち、aをbへ加算する演算、a及びbの最小値又は最大値でaを置換する演算、及び、もしaがbに等しければaをcで置換し、そうでなければaを不変のままに残す三項比較・交換演算である。整数aの場合、追加の演算がサポートされてもよく、例えば、オペランドa及びbの間のビットごとのand、or、及びxor、並びにオペランドaの増分又は減分である。他のアトミック演算、又は演算の組み合わせもサポートできる。
[0159]vote命令は、CTAスレッドの既定のグループにわたってブール(例えば、型.bl)オペランドaに簡約演算,<op>を行う。1つの実施形態において、仮想アーキテクチャは、CTAスレッドがSIMDグループの中で実行され、既定のグループがSIMDグループに対応することを指定する。他の実施形態において、CTAスレッドの他のグループは、仮想アーキテクチャ又はプログラマによって定義されてもよい。簡約演算<op>は、グループ内のCTAスレッドにわたるオペランドaの簡約、及び.<op>修飾子によって特定された簡約演算に基づいて、結果の値dをブールの真又は偽の状態への設定を引き起こす。1つの実施形態において、許される簡約演算は次のとおりである。即ち、(1).all。この場合、もしグループ内の全てのCTAスレッドについてaが真であれば、dは真であり、そうでなければ偽である。(2).any。この場合、もしグループ内の任意のCTAスレッドについてaが真であれば、dは真である。(3).uni。この場合、グループ内の全ての活動CTAスレッドについてaが同値(真又は偽)を有するならば、dは真である。
5.3.8.仮想命令−断定実行(predicated execution)
[0160]幾つかの実施形態において、仮想ISAは命令の断定実行をサポートする。断定実行において、ブールの「保護述語」値が命令に関連づけられ、実行時に保護述語が真であると評価される場合にのみ、命令が実行する。
[0160]幾つかの実施形態において、仮想ISAは命令の断定実行をサポートする。断定実行において、ブールの「保護述語」値が命令に関連づけられ、実行時に保護述語が真であると評価される場合にのみ、命令が実行する。
[0161]例示的仮想ISAにおいて、保護述語は任意の1ビット・ブール仮想レジスタ変数(本明細書ではPで表記される)であってもよい。断定実行は、命令の操作符号の前に述語保護@P又は非述語保護@!Pを置くことによって示される。値は、例えば、表820(図8C)のsetp命令のように、ブール結果を作成する命令のデスティネーション・レジスタとしてPを識別することによって、述語レジスタの中に確立される。@P又は@!P保護述語に出会うと、仮想プロセッサはPレジスタを読み出す。@P保護の場合、もしPが真であれば、命令が実行され、真でなければ、命令は省略される。@!P保護の場合、もしPが偽であれば、命令は実行され、そうでなければ、省略される。述語Pは、断定命令に出会った各々のCTAスレッドについて実行時に評価される。こうして、幾つかのCTAスレッドは断定命令を実行してもよく、その間、他のCTAスレッドは実行しない。
[0162]幾つかの実施形態において、述語は、命令が実行するときに設定されてもよい。例えば、表800〜870(図8A〜図8H)の仮想命令の一部は、述語レジスタを出力として指定するパラメータを受け入れてもよい。そのような命令は、命令結果の或る特性に基づいて特定述語レジスタを更新する。例えば、述語レジスタは、算術演算の結果が特殊数値(例えば、ゼロ、無限大、又はIEEE754浮動小数点演算の非数値)であるかどうか、などを指示するために使用されてもよい。
6.仮想命令トランスレータ
[0163]図4を参照したとき言及したように、仮想命令トランスレータ412は特定のプラットフォーム・アーキテクチャを目標にする。仮想命令トランスレータ412は、図1のCPU102のようなプロセッサ上で実行するソフトウェア・プログラムとして実現可能であり、仮想ISAコード410を受け取り、それを目標ISAコード414へ翻訳する。目標ISAコード414は、仮想命令トランスレータ412が目標とする特定のプラットフォーム・アーキテクチャの上で(例えば、図1のPPU122によって)実行できる。仮想命令トランスレータ412は、仮想ISAコード410内で宣言された仮想変数を、プロセッサ・レジスタ、オンチップ・メモリ、オフチップ・メモリなどを含む利用可能な記憶装置上の場所にマップする。幾つかの実施形態において、仮想命令トランスレータ412は、仮想状態空間の各々を特定タイプの記憶装置の上にマップする。例えば、.reg状態空間はスレッド特有データ・レジスタの上、.shared状態空間はプロセッサの共有可能メモリの上に、.global状態空間は、アプリケーション・プログラムに割り振られた仮想メモリの領域の上に、マップできる。以下同様である。他のマッピングも可能である。
[0163]図4を参照したとき言及したように、仮想命令トランスレータ412は特定のプラットフォーム・アーキテクチャを目標にする。仮想命令トランスレータ412は、図1のCPU102のようなプロセッサ上で実行するソフトウェア・プログラムとして実現可能であり、仮想ISAコード410を受け取り、それを目標ISAコード414へ翻訳する。目標ISAコード414は、仮想命令トランスレータ412が目標とする特定のプラットフォーム・アーキテクチャの上で(例えば、図1のPPU122によって)実行できる。仮想命令トランスレータ412は、仮想ISAコード410内で宣言された仮想変数を、プロセッサ・レジスタ、オンチップ・メモリ、オフチップ・メモリなどを含む利用可能な記憶装置上の場所にマップする。幾つかの実施形態において、仮想命令トランスレータ412は、仮想状態空間の各々を特定タイプの記憶装置の上にマップする。例えば、.reg状態空間はスレッド特有データ・レジスタの上、.shared状態空間はプロセッサの共有可能メモリの上に、.global状態空間は、アプリケーション・プログラムに割り振られた仮想メモリの領域の上に、マップできる。以下同様である。他のマッピングも可能である。
[0164]仮想ISAコード410内の仮想命令は機械命令へ翻訳される。1つの実施形態において、仮想命令トランスレータ412は、対応する機械命令が、CTAスレッドを実行するプロセッサの命令セットの中に存在するかどうかに依存して、各々の仮想ISA命令を、対応する機械命令又は機械命令シーケンスへマップするように構成される。
[0165]仮想命令トランスレータ412は、更に、CTAスレッドを目標プラットフォーム・アーキテクチャ内の「物理」スレッド又はプロセスの上にマップする。例えば、もし目標プラットフォーム・アーキテクチャが少なくともn0個の並行スレッドをサポートするならば、各々のCTAスレッドは1つの物理スレッドの上にマップ可能であり、仮想命令トランスレータ412は、n0個の一意の識別子を有するn0個のスレッドのために目標プラットフォーム440がコードを実行することを期待して、単一のCTAスレッドのために仮想命令コードを生成することができる。もし目標プラットフォーム・アーキテクチャが、n0個よりも少ないスレッドをサポートするならば、仮想命令トランスレータ412は、このコードがCTAごとに一回実行されることを期待して、複数のCTAスレッドに対応する命令を組み込んだ仮想ISAコード410を生成し、複数のCTAスレッドを単一の物理スレッド又はプロセスへマップすることができる。
[0166]具体的には、データ共有(例えば、.shared又は.global状態空間にアクセスするロード、記憶、及びアトミック更新命令)及び/又は協調的スレッド振る舞い(例えば、障壁、アトミック更新、及び図8Hの他の命令)に関する仮想命令は、機械命令又は機械命令シーケンスへ翻訳される。CTA実行に最適化される目標プラットフォーム・アーキテクチャは、有利には、ハードウェアでサポートされた障壁命令を含み、例えば、命令装置内にカウンタ及び/又はレジスタを用いて、障壁命令に到着したスレッドの数をカウントしてフラグを設定する。このフラグは、スレッドが障壁で待機している間、更なる命令がスレッドのために出されないように防止する。他の目標アーキテクチャは、スレッドの同期について直接のハードウェア・サポートを提供しないこともある。その場合、所望の振る舞いを作り出すために他のスレッド間通信手法(例えば、セマフォ、メモリ内の状況アレイなど)が使用可能である。
[0167]断定命令も機械命令へ翻訳される。幾つかの場合、目標ハードウェアは断定実行を直接サポートする。他の場合、述語は、例えば、プロセッサ・レジスタに記憶されてもよい。レジスタに質問し断定命令の周りを条件的に分岐することによって所望の実行時振る舞いを作り出すため、条件付き分岐命令などが使用されてもよい。
[0168]図9は、本発明の実施形態に従って仮想命令トランスレータを使用するプロセス900の流れ図である。ステップ902では、プログラマがCTAプログラム・コードを高級言語で書く。1つの実施形態において、CTAプログラム・コードは、単一のCTAスレッドの所望される振る舞いを定義し、(CTA ID及び/又はグリッドIDを含む)スレッドIDをパラメータとして使用して、CTAスレッドの振る舞いの様相を定義又は制御してもよい。例えば、読み出されるか書き込まれる共有メモリの場所は、スレッドIDの関数として決定されてもよく、同じCTA内の異なるCTAスレッドは、共有メモリ内の異なるメモリ場所との間で読み書きする。1つの実施形態において、CTAプログラム・コードはアプリケーション・プログラム・コード(例えば、図4のプログラム・コード402)の一部分として含まれる。CTAスレッド振る舞いを定義することに加えて、アプリケーション・プログラム・コードは、更に、CTA及び/又はグリッドを定義し、入力及び出力データセットなどをセットアップすることができる。
[0169]ステップ904では、コンパイラ(例えば、図4のコンパイラ408)が、高級言語コードから単一(仮想)CTAスレッドの振る舞いを定義する仮想ISAコードを生成する。もしコードがCTAプログラム・コード及び他のコードの両者を含むならば、コンパイラ408はCTAプログラム・コードを残りのコードから分離してもよく、これによりCTAプログラム・コードのみを使用して仮想ISAコードを生成する。1つの言語で書かれたプログラム・コードを他の(仮想)言語へコンパイルする従来の手法が使用されてもよい。注意すべきは、生成されたコードは仮想言語であるから、コンパイラを特定のハードウェアへ束縛又は最適化する必要はないことである。コンパイラは、入力コードの特定のシーケンスから生成された仮想ISAコードを最適化できる(例えば、仮想ISA命令の短いシーケンスを選好する)。仮想ISA内のプログラム・コードは、ディスク上のメモリに記憶され、及び/又は多様なプラットフォーム・アーキテクチャへ配布されてもよい。そのようなアーキテクチャの中には、図3の仮想アーキテクチャ300と物理的に類似しないアーキテクチャが含まれる。仮想ISA内のコードは機械から独立しており、仮想命令トランスレータを利用できる任意の目標プラットフォーム上で実行できる。代替の実施形態において、プログラマがCTAプログラム・コードを仮想ISAの中に直接書き込んでもよく、仮想ISAコードがプログラムによって自動的に生成されてもよい。もしプログラム・コードが最初に仮想ISAコードとして作られるならば、コンパイル・ステップ904は省略されてもよい。
[0170]ステップ906では、仮想命令トランスレータ(例えば、図4のトランスレータ412)が仮想ISAコードを読み出し、目標プラットフォーム上で実行可能な目標ISAコードを生成する。コンパイラとは異なり、仮想命令トランスレータは特定の(現実)プラットフォーム・アーキテクチャを目標とし、有利には、このアーキテクチャ上で最良パフォーマンスを達成するように目標ISAコードを適応及び最適化するように構成される。目標アーキテクチャが少なくともn0個のスレッドをサポートする1つの実施形態では、CTAを実現化するためにn0個のスレッドの各々によって並行して実行可能な仮想命令トランスレータが目標スレッド・プログラムを生成する、他の実施形態において、仮想命令トランスレータは、ソフトウェア手法(例えば、命令シーケンス)を使用してn0個の並行スレッドをエミュレートする目標プログラムを生成する。各々の並行スレッドは、仮想ISAコードに対応する命令を実行する。トランスレータは、プログラムの導入時、プログラムの初期化中、又はプログラム実行のジャストインタイム時間に動作してもよい。
[0171]ステップ908では、目標プラットフォーム内のプロセッサ(例えば、図1のPPU122)が、目標ISAコードを実行してデータを処理する。幾つかの実施形態において、ステップ908は、更に下記で説明するように、プロセッサの振る舞いを制御するために、コマンド及び状態パラメータをプロセッサへ供給することを含んでもよい。
[0172]プロセス900は例示であり、変形及び修正が可能であることが理解される。順次に起こるものとして説明されたステップは、並列に実行されてもよく、ステップの順序が変更されてもよく、ステップが修正又は組み合わせられてもよい。例えば、幾つかの実施形態において、プログラマは仮想ISAコードを生成するコンパイラの必要性を除いて仮想ISAを直接使用してCTAプログラム・コードを書いてもよい。他の実施形態において、CTAプログラム・コードは大きなアプリケーション・プログラムの一部分として書かれる。この大きなアプリケーション・プログラムも、例えば、特定の問題を解くために実行されるCTA及び/又はCTAグリッドの次元を定義するコードを含む。1つの実施形態において、CTAプログラムを記述するコード部分のみが仮想ISAコードへコンパイルされ、他の部分は他の(現実又は仮想の)命令セットへコンパイルされてもよい。
[0173]他の実施形態では、1つの仮想命令トランスレータが、異なる目標プラットフォームに適応された目標コードの複数のバージョンを生成するように構成されてもよい。例えば、トランスレータは、高級言語(例えば、C)のプログラム・コード、PPUの機械コード、及び/又はソフトウェア手法を使用してPPUの振る舞いをエミュレートするシングルコア又はマルチコアCPUの機械コードを作成できる。
7.仮想実行ドライバ
[0174]幾つかの実施形態において、仮想ISAコード410及び仮想命令トランスレータ412は、CTAの各スレッドのために実行されるCTAプログラム・コードを生成するために使用される。図2A〜図2Bのプログラミング・モデルに関して、CTAプログラムの指定は、各々のCTAスレッド204のための処理タスクを定義する。モデルを完成するため、CTA202の次元、グリッド内のCTAの数、処理される入力データセット等を定義することも必要である。そのような情報は、本明細書では「CTA制御情報」と呼ばれる。
[0174]幾つかの実施形態において、仮想ISAコード410及び仮想命令トランスレータ412は、CTAの各スレッドのために実行されるCTAプログラム・コードを生成するために使用される。図2A〜図2Bのプログラミング・モデルに関して、CTAプログラムの指定は、各々のCTAスレッド204のための処理タスクを定義する。モデルを完成するため、CTA202の次元、グリッド内のCTAの数、処理される入力データセット等を定義することも必要である。そのような情報は、本明細書では「CTA制御情報」と呼ばれる。
[0175]図4で示されるように、幾つかの実施形態において、アプリケーション・プログラム402は、仮想ライブラリ404内の関数への呼び出しを使用することによって、CTA制御情報を指定する。1つの実施形態において、仮想ライブラリ404は様々な関数呼び出しを含み、プログラマは関数呼び出しを介してCTA又はCTAグリッドを定義し、いつ実行を開始するかを指示することができる。
[0176]図10は、例示的仮想ライブラリ404で利用可能な関数を列挙する表1000である。関数の第1のグループはCTAの定義に関する。具体的には、initCTA関数は新しいCTAを作るために呼び出される最初の関数である。この関数によって、プログラマはCTAの次元(ntid.x,ntid.y,ntid.z)を定義し、新しいCTAへ識別子cnameを割り当てることができる。setCTAProgram関数は、CTA cnameの各スレッドによって実行されるCTAプログラムを指定する。パラメータpnameは、所望のCTAプログラム(例えば、仮想ISAコードのプログラム)に対応する論理プログラム識別子である。setCTAInputArray関数は、プログラマが、CTAcnameが入力データを読み出すグローバルメモリ内のソース場所(開始アドレス及びサイズ)を指定できるようにし、setCTAOutputArray関数は、プログラマが、CTA cnameが出力データを書き込むグローバルメモリ内の目標場所(開始アドレス及びサイズ)を指定できるようにする。setCTAParams関数は、CTA cnameのためにランタイム定数パラメータを設定するために使用される。プログラマは、パラメートのリストを、例えば、(名前、値)の対として、関数へ提供する。
[0177]1つの実施形態において、setCTAParams関数は、更に、仮想ISAコード410を生成するとき、コンパイラ408によって使用されてもよい。setCTAParams関数はCTAのためにランタイムパラメータを定義するので、コンパイラ408は、この関数が各々のパラメータを.param状態空間内の仮想変数として定義するものと解釈することができる。
[0178]表1000は、更に、CTAグリッドの定義に関する関数を列挙する。initGrid関数は、新しいグリッドを作るために呼び出される最初の関数である。この関数によって、プログラマは、グリッドの次元(nctaid.x,nctaid.y,nctaid.z)を定義し、グリッド上で実行されるCTA cnameを認識し、新しく定義されたグリッドへ識別子gnameを割り当てることができる。setGridInputArray及びsetGridOutputArray関数は、CTAレベル関数と類似しており、単一の入力及び/又は出力アレイがグリッド内の全CTAの全スレッドのために定義されることを可能にする。setGridParams関数は、グリッドgname内の全CTAのためにランタイム定数パラメータを設定するために使用される。コンパイラ408は、この関数が.const状態空間内の仮想変数として各パラメータを定義するものと解釈することができる。
[0179]launchCTA及びlaunchGrid関数は、特定されたCTA cname又はグリッドgnameの実行を開始すべきことを知らせる。
[0180]仮想APIは、更に、他の関数を含んでもよい。例えば、幾つかの実施形態は、複数CTAの実行を調整するために使用可能な同期関数を提供する。例えば、もし第1のCTA(又はグリッド)の出力が第2のCTA(又はグリッド)の入力として使用されるのであれば、APIは関数(又はlaunch関数のパラメータ)を含んでもよく、この関数を介して、仮想実行ドライバに、第1のCTA(又はグリッド)の実行が完了するまで、第2のCTA(又はグリッド)が起動されるべきでないことを指令できる。
[0181]本発明の実施形態によれば、表1000の関数呼び出しのいずれか又は全部が、アプリケーション・プログラムの中に含められてもよい。このアプリケーション・プログラムも、実行されるCTAプログラム(又は、もしアプリケーション内に複数のCTAが存在するならば、複数のプログラム)を定義する。コンパイル時に、関数呼び出しはアプリケーション・プログラム・インタフェース(API)ライブラリ404への呼び出しとして処置され、仮想APIコード406を生成する。
[0182]仮想APIコードは、仮想ライブラリ内の各関数を実現する仮想実行ドライバ418を使用して現実化される。1つの実施形態において、仮想実行ドライバ418は、CTAスレッドを現実化するPPU122を制御する図1のCPU102上で実行するドライバ・プログラムである。図10の表1000内の様々な関数呼び出しは、ドライバがプッシュバッファを介してコマンドをPPU122へ提供するように実現される。他の実施形態において、CPUは1つ又は複数のプログラムを実行してCTAを現実化し、仮想実行ドライバ418はパラメータをセットアップして、CPUによるそのようなプログラムの実行を制御する。
[0183]本明細書で説明される仮想APIは例であって、変形及び修正が可能であることが理解される。他の関数、又は関数の組み合わせがサポートされることは可能である。当技術分野で公知の仮想API手法が、本発明の目的に適応可能である。
更なる実施形態
[0184]本発明は特有の実施形態に関して説明されたが、当業者は多数の変更が可能であることを認識する。例えば、本明細書で説明される特有の仮想アーキテクチャ、仮想命令、及び仮想API関数は必要でなく、並行の協調的スレッドをサポートする他の仮想のアーキテクチャ、命令、及び/又は関数が置き換えられることが可能である。更に、これまで説明された実施形態は、全てのブロックが同数の要素を有し、全てのCTAが同数のスレッドを有して同じCTAプログラムを実行する、などの場合を基準としてもよい。幾つかのアプリケーションにおいて、例えば、複数の従属グリッドが使用される場合、プログラムを実行する異なるグリッド内のCTAを、又は異なる数及び/又はサイズのグリッドを有することが望ましいかも知れない。
[0184]本発明は特有の実施形態に関して説明されたが、当業者は多数の変更が可能であることを認識する。例えば、本明細書で説明される特有の仮想アーキテクチャ、仮想命令、及び仮想API関数は必要でなく、並行の協調的スレッドをサポートする他の仮想のアーキテクチャ、命令、及び/又は関数が置き換えられることが可能である。更に、これまで説明された実施形態は、全てのブロックが同数の要素を有し、全てのCTAが同数のスレッドを有して同じCTAプログラムを実行する、などの場合を基準としてもよい。幾つかのアプリケーションにおいて、例えば、複数の従属グリッドが使用される場合、プログラムを実行する異なるグリッド内のCTAを、又は異なる数及び/又はサイズのグリッドを有することが望ましいかも知れない。
[0185]本明細書では「協調的スレッド・アレイ」を基準としたが、幾つかの実施形態は、並行スレッド間のデータ共有がサポートされないスレッド・アレイを使用できることを理解すべきである。そのようなデータ共有がサポートされる他の実施形態において、所与のアプリケーションについて定義されたスレッドは、データを実際に共有しても共有しなくてもよい。
[0186]更に、これまで説明された実施形態は、複数のスレッドを有するスレッド・アレイを基準にしたが、「退化した」場合、スレッド・アレイは1つだけのスレッドを有してもよいことを理解すべきである。こうして、本発明は、1つ又は複数のシングルスレッド・コア又はマルチスレッド・コアを有するCPU上で実行されるプログラムの拡張性を提供するように応用可能である。本明細書で説明された手法を使用して、スレッドが任意の数の利用可能なCPUコアに分散可能な仮想ISAコードの変更又は再コンパイルを必要仕方で、(例えば、オペレーティング・システムの機能を使用して)仮想ISAコードの変更又は再コンパイルを必要としないプログラムを書くことができる。
[0187]「仮想」及び「現実」の用語は、本明細書において、問題の解法を記述するためプログラマによって使用される概念的プログラミング・モデルの、プログラムが究極的に実行できる実際のコンピュータ・システムからの切り離しを反映するために使用される。「仮想」プログラミング・モデル及びその関連アーキテクチャによって、プログラマは並列処理タスクの高レベル観察を得ることができ、コンポーネントが本明細書で説明された仮想アーキテクチャ・コンポーネントと1対1でマップする実際のコンピューティング・システム又はデバイスは、存在しても存在しなくてもよいことを理解すべきである。仮想ISAコード及び仮想APIコードを含む仮想コードは、有利には、実際の処理デバイスの命令セットと1対1に対応しても対応しなくてもよい言語のコードとして現実化される。全てのプログラム・コードと同じように、本明細書で参照される仮想コードは、有形メディア(例えば、メモリ又はディスク)に記憶され、ネットワークで送信され、その他同様の取り扱いを受けることができる。
[0188]仮想ISA及び/又は仮想APIコード、仮想命令トランスレータ、仮想ドライバ、コンパイラ、仮想関数ライブラリなどを含むが、これらに限定されない、本発明の様々な特徴を組み込んだコンピュータ・プログラムは、記憶及び/又は伝送のために様々なコンピュータ読み取り可能メディアの上に符号化されてもよい。適切なメディアは、磁気ディスク又はテープ、光記憶メディア、例えば、コンパクト・ディスク(CD)、又はDVD(ディジタル万能ディスク)、フラッシュ・メモリなどを含む。そのようなプログラムは、更に、インターネットを含む多様なプロトコルに合致する有線、光、及び/又は無線ネットワークを介して伝送されるように適応した搬送波信号を使用して、符号化及び伝送されてもよい。プログラム・コードを符号化されたコンピュータ読み取り可能記憶メディアは、互換性デバイスと共にパッケージされてもよく、また、プログラム・コードは他のデバイスとは別個に(例えば、インターネットのダウンロードを介して)提供されてもよい。
[0189]更に、或る一部の行動は、本明細書では「プログラマ」によって取られると説明されてもよい。プログラマは、人間であるか、人間の介入を殆ど又は全く受けないでプログラム・コードを生成する自動プロセスであるか、プログラム・コードを生成するために、自動又は部分的自動プロセスと、交信する人間との組み合わせであることが考えられる。
[0190]更に、本明細書で説明される実施形態は、特定の目標プラットフォームの特徴を基準にするが、本発明はこのようなプラットフォームに限定されない。実際に、仮想アーキテクチャはハードウェア及び/又はソフトウェア・コンポーネントの任意の組み合わせで現実化可能である。同じ仮想アーキテクチャの異なる現実化は、効率及び/又はスループットで異なることが想定できることを、当業者は理解する。しかし、そのような差異は本発明と無関係である。
[0191]したがって、本発明は特定の実施形態に関して説明されたが、本発明は、続く請求の範囲内で全ての変更及び均等物を対象とするように意図されることが理解される。
100…コンピュータ・システム、102…中央処理装置(CPU)、104…システム・メモリ、105…メモリ・ブリッジ、106…通信経路、107…入力/出力ブリッジ、108…ユーザ入力デバイス、110…ディスプレイ、112…並列処理サブシステム、113…他の通信経路、114…システム・ディスク、116…スイッチ、118…ネットワーク・アダプタ、120…アドイン・カード、121…アドイン・カード、122…並列処理装置(PPU)、124…並列処理(PP)メモリ、200…グリッド、200(0)…グリッド、200(1)…グリッド、200(2)…グリッド、202…スレッド・アレイ、202(0,0)…協調的スレッド・アレイ、202(0,1)…協調的スレッド・アレイ、202(0,2)…協調的スレッド・アレイ、202(1,0)…協調的スレッド・アレイ、202(1,1)…協調的スレッド・アレイ、202(2,0)…協調的スレッド・アレイ、204…スレッド、204(0,0)…スレッド、204(0,1)…スレッド、204(0,2)…スレッド、204(1,0)…スレッド、204(1,1)…スレッド、204(2,0)…スレッド、300…仮想アーキテクチャ、302…仮想プロセッサ、304…グローバルメモリ、306…フロントエンド、308…仮想コア、310…仮想処理エンジン、311…特殊レジスタ、312…仮想命令装置、314…ローカルレジスタ、316…共有メモリ、318…パラメータ・メモリ、320…仮想ドライバ、322…メモリ・インタフェース、400…概念的モデル、402…アプリケーション・プログラム、404…ライブラリ、406…仮想APIコード、408…コンパイラ、410…仮想ISAコード、412…仮想命令トランスレータ、414…目標ISAコード、416…仮想実行ドライバ、418…目標APIコマンド、430…命令装置、432…並列処理装置ドライバ、434…フロントエンド、440…目標プロセッサ又は目標プラットフォーム、500…特殊変数の表、600…変数型の表、700…仮想状態空間の表、800…算術演算の表、810…ベクトル演算の表、820…選択及びレジスタ設定演算の表、830…論理及びビット操作演算の表、840…形式変換演算の表、850…データ移動及びデータ共有命令の表、860…プログラム制御演算の表、870…明示並列仮想命令の表、900…仮想命令トランスレータの使用プロセス、1000…仮想ライブラリで利用できる関数の表、API…アプリケーション・プログラム・インタフェース(application program interface)、CTA…協調的スレッド・アレイ(cooperative thread array)、ISA…命令セット・アーキテクチャ(instruction set architecture)、PP…並列処理(parallel processing)、PPU…並列処理装置(parallel processing unit)
Claims (21)
- 目標プラットフォーム・アーキテクチャに合致するコンピュータ・システム上で実行される並列処理動作を定義する方法であって、
協調的仮想スレッド・アレイを実行するための仮想並列アーキテクチャを定義するステップと、
前記協調的仮想スレッド・アレイにおける複数の仮想スレッドの各々について行われる動作シーケンスを定義する第1のプログラム・コードを提供するステップと、
前記複数の仮想スレッドのうちの一つの仮想スレッドについて実行されるスレッド単位命令シーケンスを定義する仮想スレッド・プログラムへ前記第1のプログラム・コードをコンパイルするステップであって、前記スレッド単位命令シーケンスが、前記一つの仮想スレッドと前記複数の仮想スレッドの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含むステップと、
前記仮想スレッド・プログラムを記憶するステップと
を備える方法。 - 前記記憶された仮想スレッド・プログラムを、前記目標プラットフォーム・アーキテクチャに合致する命令シーケンスへ翻訳するステップ
を備える、請求項1に記載の方法。 - 入力データセットを処理して出力データセットを生成するように適応させられた協調的仮想スレッドのアレイを定義する第2のプログラム・コードを提供するステップであって、前記アレイ内の各仮想スレッドが前記仮想スレッド・プログラムを並行して実行するスレッドであるステップと、
前記第2のプログラム・コードを前記仮想並列アーキテクチャにおける所定の関数ライブラリ内の関数呼び出しシーケンスへ変換するステップであって、前記ライブラリが前記協調的仮想スレッドのアレイを初期化及び実行させる関数を含むステップと、
前記関数呼び出しシーケンスを記憶するステップと
を備える、請求項1に記載の方法。 - 前記記憶された仮想スレッド・プログラム及び前記関数呼び出しシーケンスを、前記目標プラットフォーム・アーキテクチャ上で実行可能なプログラム・コードへ翻訳するステップであって、前記実行可能プログラム・コードが、前記協調的仮想スレッドのアレイを実行する1つ又は複数のプラットフォーム・スレッドを定義するステップ
を備える、請求項3に記載の方法。 - 前記目標プラットフォーム・アーキテクチャに合致するコンピュータ・システム上で前記実行可能プログラム・コードを実行して前記出力データセットを生成するステップと、
前記出力データセットを記憶メディア上に記憶するステップと
を備える、請求項4に記載の方法。 - 前記スレッド単位命令シーケンスが、前記1つ又は複数の他の仮想スレッドが前記シーケンス内の特定点に達する時間まで、前記一つの仮想スレッドの動作の実行を前記特定点で一時停止する命令を含む、請求項1に記載の方法。
- 前記スレッド単位命令シーケンスが、前記1つ又は複数の他の仮想スレッドがアクセスを有する共有メモリ内に前記一つの仮想スレッドがデータを記憶する命令を含む、請求項1に記載の方法。
- 前記スレッド単位命令シーケンスが、前記1つ又は複数の他の仮想スレッドがアクセスを有する共有メモリ内に記憶されたデータを前記一つの仮想スレッドがアトミックに読み出し及び更新する命令を含む、請求項1に記載の方法。
- 前記仮想スレッド・プログラムが、複数の仮想状態空間の1つで変数を定義する変数定義ステートメントを含み、前記複数の仮想状態空間の異なるものが前記仮想スレッド間のデータ共有の異なるモードに対応する、請求項1に記載の方法。
- 前記データ共有のモードが、スレッド単位非共有モード及びグローバル共有モードを含む、請求項9に記載の方法。
- 前記データ共有のモードが、スレッド単位非共有モード、1つの仮想スレッド・アレイ内の共有モード、及びグローバル共有モードを含む、請求項9に記載の方法。
- 前記データ共有のモードが、スレッド単位非共有モード、1つの仮想スレッド・アレイ内の共有モード、複数の仮想スレッド・アレイ間の共有モード、及びグローバル共有モードを含む、請求項9に記載の方法。
- 目標プラットフォーム・アーキテクチャに合致するコンピュータ・システムの目標プロセッサを動作させる方法であって、
入力データセットを処理して出力データセットを生成するように適応させられた協調的仮想スレッド・アレイを実行するための仮想並列アーキテクチャを定義するステップと、
前記協調的仮想スレッド・アレイにおける複数の仮想スレッドの各々について行われる動作シーケンスを定義する第1の部分を含む入力プログラム・コードを提供するステップであって、前記入力プログラム・コードが、更に、前記協調的仮想スレッド・アレイの次元を定義する第2の部分を含むステップと、
前記複数の仮想スレッドのうちの一つの仮想スレッドについて実行されるスレッド単位命令シーケンスを定義する仮想スレッド・プログラムへ前記入力プログラム・コードの前記第1の部分をコンパイルするステップであって、前記スレッド単位命令シーケンスが、前記一つの仮想スレッドと前記複数の仮想スレッドの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含むステップと、
前記入力プログラム・コードの前記第2の部分を、前記仮想並列アーキテクチャにおける所定の関数ライブラリへの関数呼び出しシーケンスへ変換するステップであって、前記ライブラリが、前記協調的仮想スレッド・アレイを初期化及び実行させる関数を含むステップと、
前記仮想スレッド・プログラム及び前記関数呼び出しシーケンスを前記目標プラットフォーム・アーキテクチャ上で実行可能なプログラム・コードへ翻訳するステップであって、前記実行可能なプログラム・コードが、前記協調的仮想スレッド・アレイを実行する1つ又は複数の現実スレッドを定義するステップと、
前記目標プラットフォーム・アーキテクチャに合致する前記コンピュータ・システム上で前記実行可能プログラム・コードを実行して前記出力データセットを生成するステップと、
前記出力データセットを記憶メディアに記憶するステップと
を備える方法。 - 前記入力プログラム・コードの前記第2の部分が、前記協調的仮想スレッド・アレイのために複数の次元を定義するプログラム・コードを含む、請求項13に記載の方法。
- 前記入力プログラム・コードの前記第2の部分が、更に、
仮想スレッド・アレイより成るグリッドの1つ又は複数の次元を定義する関数呼び出しを含み、前記グリッド内の各アレイが実行される、
請求項14に記載の方法。 - 前記目標プラットフォーム・アーキテクチャがマスタ・プロセッサ及びコプロセッサを含み、前記翻訳する行為が、
前記コプロセッサ上で定義された複数のスレッドによって並列に実行可能なプログラム・コードへ前記仮想スレッド・プログラムを翻訳するステップと、
前記コプロセッサ用のドライバ・プログラムへの呼び出しシーケンスへ前記関数呼び出しシーケンスを翻訳するステップであって、前記ドライバ・プログラムが前記マスタ・プロセッサ上で実行されるプログラムであるステップと
を含む、請求項13に記載の方法。 - 前記目標プラットフォーム・アーキテクチャが中央処理装置(CPU)を含み、前記翻訳する行為が、
前記仮想スレッド・プログラム及び前記関数呼び出しシーケンスの少なくとも一部分を、前記仮想スレッドの数よりも少ない数のCPUスレッドを使用して前記協調的仮想スレッド・アレイを実行する目標プログラム・コードへ翻訳するステップ
を含む、請求項13に記載の方法。 - 目標プラットフォーム・アーキテクチャに合致するコンピュータ・システムの目標プロセッサを動作させる方法であって、
入力データセットを処理して出力データセットを生成するように適応させられた協調的仮想スレッド・アレイを実行するための仮想並列アーキテクチャを定義するステップと、
前記協調的仮想スレッド・アレイにおける複数の仮想スレッドのうちの一つの仮想スレッドについて実行されるスレッド単位命令シーケンスを定義する仮想スレッド・プログラムを取得するステップであって、前記スレッド単位命令シーケンスが、前記一つの仮想スレッドと前記複数の仮想スレッドの1つ又は複数の他の仮想スレッドとの間の協調的振る舞いを定義する少なくとも1つの命令を含むステップと、
前記協調的仮想スレッド・アレイの次元を定義する追加のプログラム・コードを取得するステップと、
前記仮想スレッド・プログラム及び前記追加のプログラム・コードを、前記目標プラットフォーム・アーキテクチャ上で実行可能なプログラム・コードへ翻訳するステップであって、前記実行可能なプログラム・コードが、前記協調的仮想スレッド・アレイを実行する1つ又は複数のプラットフォーム・スレッドを定義するステップと、
前記目標プラットフォーム・アーキテクチャに合致する前記コンピュータ・システム上で前記実行可能なプログラム・コードを実行して前記出力データセットを生成し、前記出力データセットをメモリに記憶するステップと
を備える方法。 - 前記仮想スレッド・プログラムを取得する前記行為が、
高級プログラミング言語で書かれたソース・プログラム・コードを取得するステップと、
前記ソース・プログラム・コードをコンパイルして前記仮想スレッド・プログラムを生成するステップと
を含む、請求項18に記載の方法。 - 前記仮想スレッド・プログラムを取得する前記行為が、
前記仮想スレッド・プログラムを記憶メディアから読み出すステップ
を含む、請求項18に記載の方法。 - 前記仮想スレッド・プログラムを取得する前記行為が、
ネットワークを介して前記仮想スレッド・プログラムをリモート・コンピュータ・システムから受け取るステップ
を含む、請求項18に記載の方法。
Applications Claiming Priority (2)
Application Number | Priority Date | Filing Date | Title |
---|---|---|---|
US11/627,892 | 2007-01-26 | ||
US11/627,892 US8321849B2 (en) | 2007-01-26 | 2007-01-26 | Virtual architecture and instruction set for parallel thread computing |
Publications (3)
Publication Number | Publication Date |
---|---|
JP2008276740A JP2008276740A (ja) | 2008-11-13 |
JP2008276740A5 true JP2008276740A5 (ja) | 2011-05-19 |
JP4999183B2 JP4999183B2 (ja) | 2012-08-15 |
Family
ID=39587517
Family Applications (1)
Application Number | Title | Priority Date | Filing Date |
---|---|---|---|
JP2008015281A Active JP4999183B2 (ja) | 2007-01-26 | 2008-01-25 | 並列スレッド・コンピューティングの仮想のアーキテクチャ及び命令セット |
Country Status (7)
Country | Link |
---|---|
US (1) | US8321849B2 (ja) |
JP (1) | JP4999183B2 (ja) |
KR (1) | KR101026689B1 (ja) |
CN (1) | CN101231585B (ja) |
DE (2) | DE202008017916U1 (ja) |
SG (1) | SG144869A1 (ja) |
TW (1) | TWI363294B (ja) |
Families Citing this family (142)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
WO2005072307A2 (en) * | 2004-01-22 | 2005-08-11 | University Of Washington | Wavescalar architecture having a wave order memory |
US7490218B2 (en) * | 2004-01-22 | 2009-02-10 | University Of Washington | Building a wavecache |
RU2312388C2 (ru) * | 2005-09-22 | 2007-12-10 | Андрей Игоревич Ефимов | Способ организации многопроцессорной эвм |
US7861060B1 (en) * | 2005-12-15 | 2010-12-28 | Nvidia Corporation | Parallel data processing systems and methods using cooperative thread arrays and thread identifier values to determine processing behavior |
EP2011018B1 (en) | 2006-04-12 | 2016-07-13 | Soft Machines, Inc. | Apparatus and method for processing an instruction matrix specifying parallel and dependent operations |
US8082289B2 (en) | 2006-06-13 | 2011-12-20 | Advanced Cluster Systems, Inc. | Cluster computing support for application programs |
US7836116B1 (en) * | 2006-06-15 | 2010-11-16 | Nvidia Corporation | Fast fourier transforms and related transforms using cooperative thread arrays |
US7640284B1 (en) | 2006-06-15 | 2009-12-29 | Nvidia Corporation | Bit reversal methods for a parallel processor |
EP2122461A4 (en) | 2006-11-14 | 2010-03-24 | Soft Machines Inc | DEVICE AND METHOD FOR PROCESSING COMMUNICATIONS IN A MULTITHREAD ARCHITECTURE WITH CONTEXT CHANGES |
US8549500B2 (en) * | 2007-02-14 | 2013-10-01 | The Mathworks, Inc. | Saving and loading graphical processing unit (GPU) arrays providing high computational capabilities in a computing environment |
US8533697B2 (en) * | 2007-02-14 | 2013-09-10 | The Mathworks, Inc. | Graphical processing unit (GPU) arrays providing high computational capabilities in a computing environment |
US8286196B2 (en) | 2007-05-03 | 2012-10-09 | Apple Inc. | Parallel runtime execution on multiple processors |
US8276164B2 (en) | 2007-05-03 | 2012-09-25 | Apple Inc. | Data parallel computing on multiple processors |
EP3413198A1 (en) | 2007-04-11 | 2018-12-12 | Apple Inc. | Data parallel computing on multiple processors |
US11836506B2 (en) | 2007-04-11 | 2023-12-05 | Apple Inc. | Parallel runtime execution on multiple processors |
US8341611B2 (en) | 2007-04-11 | 2012-12-25 | Apple Inc. | Application interface on multiple processors |
US7725518B1 (en) | 2007-08-08 | 2010-05-25 | Nvidia Corporation | Work-efficient parallel prefix sum algorithm for graphics processing units |
US7877573B1 (en) * | 2007-08-08 | 2011-01-25 | Nvidia Corporation | Work-efficient parallel prefix sum algorithm for graphics processing units |
US8225295B2 (en) * | 2007-09-21 | 2012-07-17 | Jens Palsberg | Register allocation by puzzle solving |
US20090204173A1 (en) | 2007-11-05 | 2009-08-13 | Zi-Ping Fang | Multi-Frequency Neural Treatments and Associated Systems and Methods |
US8200947B1 (en) * | 2008-03-24 | 2012-06-12 | Nvidia Corporation | Systems and methods for voting among parallel threads |
US8286198B2 (en) * | 2008-06-06 | 2012-10-09 | Apple Inc. | Application programming interfaces for data parallel computing on multiple processors |
US8225325B2 (en) | 2008-06-06 | 2012-07-17 | Apple Inc. | Multi-dimensional thread grouping for multiple processors |
US8434093B2 (en) | 2008-08-07 | 2013-04-30 | Code Systems Corporation | Method and system for virtualization of software applications |
US8776038B2 (en) | 2008-08-07 | 2014-07-08 | Code Systems Corporation | Method and system for configuration of virtualized software applications |
US8436862B2 (en) * | 2008-12-11 | 2013-05-07 | Nvidia Corporation | Method and system for enabling managed code-based application program to access graphics processing unit |
US8570333B2 (en) * | 2008-12-11 | 2013-10-29 | Nvidia Corporation | Method and system for enabling managed code-based application program to access graphics processing unit |
US8307350B2 (en) * | 2009-01-14 | 2012-11-06 | Microsoft Corporation | Multi level virtual function tables |
JP5149840B2 (ja) * | 2009-03-03 | 2013-02-20 | 株式会社日立製作所 | ストリームデータ処理方法、ストリームデータ処理プログラム、および、ストリームデータ処理装置 |
KR101572879B1 (ko) | 2009-04-29 | 2015-12-01 | 삼성전자주식회사 | 병렬 응용 프로그램을 동적으로 병렬처리 하는 시스템 및 방법 |
US8542247B1 (en) | 2009-07-17 | 2013-09-24 | Nvidia Corporation | Cull before vertex attribute fetch and vertex lighting |
US8564616B1 (en) | 2009-07-17 | 2013-10-22 | Nvidia Corporation | Cull before vertex attribute fetch and vertex lighting |
CN102023844B (zh) * | 2009-09-18 | 2014-04-09 | 深圳中微电科技有限公司 | 并行处理器及其线程处理方法 |
US8266383B1 (en) | 2009-09-28 | 2012-09-11 | Nvidia Corporation | Cache miss processing using a defer/replay mechanism |
US10360039B2 (en) * | 2009-09-28 | 2019-07-23 | Nvidia Corporation | Predicted instruction execution in parallel processors with reduced per-thread state information including choosing a minimum or maximum of two operands based on a predicate value |
US9665920B1 (en) * | 2009-10-05 | 2017-05-30 | Nvidia Corporation | Simultaneous execution of compute and graphics applications |
US8976195B1 (en) | 2009-10-14 | 2015-03-10 | Nvidia Corporation | Generating clip state for a batch of vertices |
US8384736B1 (en) | 2009-10-14 | 2013-02-26 | Nvidia Corporation | Generating clip state for a batch of vertices |
EP2519876A1 (en) * | 2009-12-28 | 2012-11-07 | Hyperion Core, Inc. | Optimisation of loops and data flow sections |
KR101613971B1 (ko) * | 2009-12-30 | 2016-04-21 | 삼성전자주식회사 | 프로그램 코드의 변환 방법 |
US8954958B2 (en) | 2010-01-11 | 2015-02-10 | Code Systems Corporation | Method of configuring a virtual application |
US9104517B2 (en) | 2010-01-27 | 2015-08-11 | Code Systems Corporation | System for downloading and executing a virtual application |
US8959183B2 (en) | 2010-01-27 | 2015-02-17 | Code Systems Corporation | System for downloading and executing a virtual application |
US9229748B2 (en) | 2010-01-29 | 2016-01-05 | Code Systems Corporation | Method and system for improving startup performance and interoperability of a virtual application |
US8763009B2 (en) | 2010-04-17 | 2014-06-24 | Code Systems Corporation | Method of hosting a first application in a second application |
US8375373B2 (en) * | 2010-04-19 | 2013-02-12 | Microsoft Corporation | Intermediate language support for change resilience |
US8527866B2 (en) | 2010-04-30 | 2013-09-03 | Microsoft Corporation | Multi-threaded sort of data items in spreadsheet tables |
US20110276868A1 (en) * | 2010-05-05 | 2011-11-10 | Microsoft Corporation | Multi-Threaded Adjustment of Column Widths or Row Heights |
US9851969B2 (en) | 2010-06-24 | 2017-12-26 | International Business Machines Corporation | Function virtualization facility for function query of a processor |
US10521231B2 (en) * | 2010-06-24 | 2019-12-31 | International Business Machines Corporation | Function virtualization facility for blocking instruction function of a multi-function instruction of a virtual processor |
US8782106B2 (en) | 2010-07-02 | 2014-07-15 | Code Systems Corporation | Method and system for managing execution of virtual applications |
CN103250131B (zh) | 2010-09-17 | 2015-12-16 | 索夫特机械公司 | 包括用于早期远分支预测的影子缓存的单周期多分支预测 |
US9529574B2 (en) * | 2010-09-23 | 2016-12-27 | Apple Inc. | Auto multi-threading in macroscalar compilers |
KR20120031756A (ko) * | 2010-09-27 | 2012-04-04 | 삼성전자주식회사 | Cpu와 gpu를 사용하는 이종 시스템에서 가상화를 이용한 어플리케이션 컴파일 및 실행 방법 및 장치 |
KR101649925B1 (ko) * | 2010-10-13 | 2016-08-31 | 삼성전자주식회사 | 멀티 트레드 프로그램에서 변수의 단독 메모리 접근여부를 분석하는 방법 |
US20120096292A1 (en) * | 2010-10-15 | 2012-04-19 | Mosaid Technologies Incorporated | Method, system and apparatus for multi-level processing |
US9021015B2 (en) | 2010-10-18 | 2015-04-28 | Code Systems Corporation | Method and system for publishing virtual applications to a web server |
US9209976B2 (en) | 2010-10-29 | 2015-12-08 | Code Systems Corporation | Method and system for restricting execution of virtual applications to a managed process environment |
JP5490253B2 (ja) * | 2010-11-02 | 2014-05-14 | インターナショナル・ビジネス・マシーンズ・コーポレーション | 数値集約計算における文字列集約方法 |
US8819700B2 (en) * | 2010-12-22 | 2014-08-26 | Lsi Corporation | System and method for synchronous inter-thread communication |
KR101157596B1 (ko) * | 2010-12-24 | 2012-06-19 | 서울대학교산학협력단 | 개방형 범용 병렬 컴퓨팅 프레임워크(OpenCL)에서의 메모리 접근영역 분석장치 및 그 방법 |
CN108108188B (zh) | 2011-03-25 | 2022-06-28 | 英特尔公司 | 用于通过使用由可分区引擎实例化的虚拟核来支持代码块执行的存储器片段 |
US9842005B2 (en) | 2011-03-25 | 2017-12-12 | Intel Corporation | Register file segments for supporting code block execution by using virtual cores instantiated by partitionable engines |
KR101638225B1 (ko) | 2011-03-25 | 2016-07-08 | 소프트 머신즈, 인크. | 분할가능한 엔진에 의해 인스턴스화된 가상 코어를 이용한 명령어 시퀀스 코드 블록의 실행 |
WO2012157786A1 (ja) | 2011-05-19 | 2012-11-22 | 日本電気株式会社 | 並列処理装置、並列処理方法、最適化装置、最適化方法、および、コンピュータ・プログラム |
KR101639853B1 (ko) * | 2011-05-20 | 2016-07-14 | 소프트 머신즈, 인크. | 복수의 엔진에 의해 명령어 시퀀스들의 실행을 지원하기 위한 자원들 및 상호접속 구조들의 비집중 할당 |
CN103649931B (zh) | 2011-05-20 | 2016-10-12 | 索夫特机械公司 | 用于支持由多个引擎执行指令序列的互连结构 |
US8990830B2 (en) * | 2011-07-19 | 2015-03-24 | International Business Machines Corporation | Thread management in parallel processes |
KR101894752B1 (ko) | 2011-10-27 | 2018-09-05 | 삼성전자주식회사 | 가상 아키텍쳐 생성 장치, 런타임 시스템, 멀티 코어 시스템 및 그 동작 방법 |
US9009686B2 (en) * | 2011-11-07 | 2015-04-14 | Nvidia Corporation | Algorithm for 64-bit address mode optimization |
US9507638B2 (en) * | 2011-11-08 | 2016-11-29 | Nvidia Corporation | Compute work distribution reference counters |
US10191746B2 (en) | 2011-11-22 | 2019-01-29 | Intel Corporation | Accelerated code optimizer for a multiengine microprocessor |
WO2013077876A1 (en) | 2011-11-22 | 2013-05-30 | Soft Machines, Inc. | A microprocessor accelerated code optimizer |
US10255228B2 (en) * | 2011-12-06 | 2019-04-09 | Nvidia Corporation | System and method for performing shaped memory access operations |
US9507593B2 (en) * | 2011-12-23 | 2016-11-29 | Intel Corporation | Instruction for element offset calculation in a multi-dimensional array |
CN104137055B (zh) * | 2011-12-29 | 2018-06-05 | 英特尔公司 | 点积处理器、方法、系统和指令 |
KR101885211B1 (ko) * | 2012-01-27 | 2018-08-29 | 삼성전자 주식회사 | Gpu의 자원 할당을 위한 방법 및 장치 |
US9104421B2 (en) * | 2012-07-30 | 2015-08-11 | Nvidia Corporation | Training, power-gating, and dynamic frequency changing of a memory controller |
WO2014022980A1 (en) * | 2012-08-08 | 2014-02-13 | Intel Corporation | Isa bridging including support for call to overidding virtual functions |
US9274772B2 (en) | 2012-08-13 | 2016-03-01 | Microsoft Technology Licensing, Llc. | Compact type layouts |
CN102902576B (zh) * | 2012-09-26 | 2014-12-24 | 北京奇虎科技有限公司 | 一种渲染网页的方法、服务器和系统 |
US8982124B2 (en) | 2012-09-29 | 2015-03-17 | Intel Corporation | Load balancing and merging of tessellation thread workloads |
US9123167B2 (en) | 2012-09-29 | 2015-09-01 | Intel Corporation | Shader serialization and instance unrolling |
US9904625B2 (en) | 2013-03-15 | 2018-02-27 | Intel Corporation | Methods, systems and apparatus for predicting the way of a set associative cache |
WO2014150806A1 (en) | 2013-03-15 | 2014-09-25 | Soft Machines, Inc. | A method for populating register view data structure by using register template snapshots |
US10275255B2 (en) | 2013-03-15 | 2019-04-30 | Intel Corporation | Method for dependency broadcasting through a source organized source view data structure |
US10140138B2 (en) | 2013-03-15 | 2018-11-27 | Intel Corporation | Methods, systems and apparatus for supporting wide and efficient front-end operation with guest-architecture emulation |
WO2014150991A1 (en) | 2013-03-15 | 2014-09-25 | Soft Machines, Inc. | A method for implementing a reduced size register view data structure in a microprocessor |
EP2972845B1 (en) | 2013-03-15 | 2021-07-07 | Intel Corporation | A method for executing multithreaded instructions grouped onto blocks |
US9891924B2 (en) | 2013-03-15 | 2018-02-13 | Intel Corporation | Method for implementing a reduced size register view data structure in a microprocessor |
WO2014150971A1 (en) | 2013-03-15 | 2014-09-25 | Soft Machines, Inc. | A method for dependency broadcasting through a block organized source view data structure |
KR102083390B1 (ko) | 2013-03-15 | 2020-03-02 | 인텔 코포레이션 | 네이티브 분산된 플래그 아키텍처를 이용하여 게스트 중앙 플래그 아키텍처를 에뮬레이션하는 방법 |
US9811342B2 (en) | 2013-03-15 | 2017-11-07 | Intel Corporation | Method for performing dual dispatch of blocks and half blocks |
US9886279B2 (en) | 2013-03-15 | 2018-02-06 | Intel Corporation | Method for populating and instruction view data structure by using register template snapshots |
US9569216B2 (en) | 2013-03-15 | 2017-02-14 | Soft Machines, Inc. | Method for populating a source view data structure by using register template snapshots |
US9535686B2 (en) | 2013-03-15 | 2017-01-03 | International Business Machines Corporation | Dynamic library replacement |
US9772864B2 (en) * | 2013-04-16 | 2017-09-26 | Arm Limited | Methods of and apparatus for multidimensional indexing in microprocessor systems |
US9600852B2 (en) * | 2013-05-10 | 2017-03-21 | Nvidia Corporation | Hierarchical hash tables for SIMT processing and a method of establishing hierarchical hash tables |
US9507594B2 (en) * | 2013-07-02 | 2016-11-29 | Intel Corporation | Method and system of compiling program code into predicated instructions for execution on a processor without a program counter |
US10628156B2 (en) * | 2013-07-09 | 2020-04-21 | Texas Instruments Incorporated | Vector SIMD VLIW data path architecture |
GB2524063B (en) | 2014-03-13 | 2020-07-01 | Advanced Risc Mach Ltd | Data processing apparatus for executing an access instruction for N threads |
US9471283B2 (en) * | 2014-06-11 | 2016-10-18 | Ca, Inc. | Generating virtualized application programming interface (API) implementation from narrative API documentation |
CN104077233B (zh) * | 2014-06-18 | 2017-04-05 | 百度在线网络技术(北京)有限公司 | 多通道卷积层处理方法和装置 |
CN106406810B (zh) * | 2014-07-02 | 2019-08-06 | 上海兆芯集成电路有限公司 | 微处理器及其方法 |
US10067768B2 (en) * | 2014-07-18 | 2018-09-04 | Nvidia Corporation | Execution of divergent threads using a convergence barrier |
US20160026486A1 (en) * | 2014-07-25 | 2016-01-28 | Soft Machines, Inc. | An allocation and issue stage for reordering a microinstruction sequence into an optimized microinstruction sequence to implement an instruction set agnostic runtime architecture |
US9398019B2 (en) | 2014-08-07 | 2016-07-19 | Vmware, Inc. | Verifying caller authorization using secret data embedded in code |
US9411979B2 (en) * | 2014-08-07 | 2016-08-09 | Vmware, Inc. | Embedding secret data in code |
US10922402B2 (en) | 2014-09-29 | 2021-02-16 | Vmware, Inc. | Securing secret data embedded in code against compromised interrupt and exception handlers |
US9749548B2 (en) | 2015-01-22 | 2017-08-29 | Google Inc. | Virtual linebuffers for image signal processors |
CN107408035B (zh) * | 2015-03-27 | 2021-11-09 | 英特尔公司 | 用于缕程间通信的装置和方法 |
US9772852B2 (en) | 2015-04-23 | 2017-09-26 | Google Inc. | Energy efficient processor core architecture for image processor |
US9965824B2 (en) | 2015-04-23 | 2018-05-08 | Google Llc | Architecture for high performance, power efficient, programmable image processing |
US10095479B2 (en) * | 2015-04-23 | 2018-10-09 | Google Llc | Virtual image processor instruction set architecture (ISA) and memory model and exemplary target hardware having a two-dimensional shift array structure |
US9756268B2 (en) | 2015-04-23 | 2017-09-05 | Google Inc. | Line buffer unit for image processor |
US9785423B2 (en) | 2015-04-23 | 2017-10-10 | Google Inc. | Compiler for translating between a virtual image processor instruction set architecture (ISA) and target hardware having a two-dimensional shift array structure |
US9769356B2 (en) | 2015-04-23 | 2017-09-19 | Google Inc. | Two dimensional shift array for image processor |
US10291813B2 (en) | 2015-04-23 | 2019-05-14 | Google Llc | Sheet generator for image processor |
US10313641B2 (en) | 2015-12-04 | 2019-06-04 | Google Llc | Shift register with reduced wiring complexity |
US9830150B2 (en) | 2015-12-04 | 2017-11-28 | Google Llc | Multi-functional execution lane for image processor |
US10115175B2 (en) * | 2016-02-19 | 2018-10-30 | Qualcomm Incorporated | Uniform predicates in shaders for graphics processing units |
US10204396B2 (en) | 2016-02-26 | 2019-02-12 | Google Llc | Compiler managed memory for image processor |
US10387988B2 (en) | 2016-02-26 | 2019-08-20 | Google Llc | Compiler techniques for mapping program code to a high performance, power efficient, programmable image processing hardware platform |
US10380969B2 (en) | 2016-02-28 | 2019-08-13 | Google Llc | Macro I/O unit for image processor |
US20180007302A1 (en) | 2016-07-01 | 2018-01-04 | Google Inc. | Block Operations For An Image Processor Having A Two-Dimensional Execution Lane Array and A Two-Dimensional Shift Register |
US10546211B2 (en) | 2016-07-01 | 2020-01-28 | Google Llc | Convolutional neural network on programmable two dimensional image processor |
US20180005346A1 (en) | 2016-07-01 | 2018-01-04 | Google Inc. | Core Processes For Block Operations On An Image Processor Having A Two-Dimensional Execution Lane Array and A Two-Dimensional Shift Register |
US20180005059A1 (en) | 2016-07-01 | 2018-01-04 | Google Inc. | Statistics Operations On Two Dimensional Image Processor |
WO2018094087A1 (en) * | 2016-11-17 | 2018-05-24 | The Mathworks, Inc. | Systems and methods for generating code for parallel processing units |
US10204394B2 (en) | 2017-04-10 | 2019-02-12 | Intel Corporation | Multi-frame renderer |
US10437593B2 (en) * | 2017-04-27 | 2019-10-08 | Nvidia Corporation | Techniques for comprehensively synchronizing execution threads |
US10656964B2 (en) * | 2017-05-16 | 2020-05-19 | Oracle International Corporation | Dynamic parallelization of a calculation process |
US10754746B2 (en) | 2017-11-15 | 2020-08-25 | General Electric Company | Virtual processor enabling unobtrusive observation of legacy systems for analytics in SoC |
CN110825514B (zh) * | 2018-08-10 | 2023-05-23 | 昆仑芯(北京)科技有限公司 | 人工智能芯片以及用于人工智能芯片的指令执行方法 |
GB2580327B (en) * | 2018-12-31 | 2021-04-28 | Graphcore Ltd | Register files in a multi-threaded processor |
US20200264921A1 (en) * | 2019-02-20 | 2020-08-20 | Nanjing Iluvatar CoreX Technology Co., Ltd. (DBA "Iluvatar CoreX Inc. Nanjing") | Crypto engine and scheduling method for vector unit |
CN110321193B (zh) * | 2019-05-05 | 2022-03-18 | 四川盛趣时代网络科技有限公司 | 一种基于Direct3D共享纹理的交互方法及系统 |
US11216281B2 (en) | 2019-05-14 | 2022-01-04 | International Business Machines Corporation | Facilitating data processing using SIMD reduction operations across SIMD lanes |
CN110209509B (zh) * | 2019-05-28 | 2021-08-17 | 北京星网锐捷网络技术有限公司 | 多核处理器间的数据同步方法及装置 |
CN110765082B (zh) * | 2019-09-06 | 2023-11-24 | 深圳平安通信科技有限公司 | Hadoop文件处理方法、装置、存储介质及服务器 |
CN114816529A (zh) * | 2020-10-21 | 2022-07-29 | 上海壁仞智能科技有限公司 | 配置向量运算系统中的协作线程束的装置和方法 |
US20230090604A1 (en) * | 2021-09-16 | 2023-03-23 | T-Head (Shanghai) Semiconductor Co., Ltd. | Parallel processing unit virtualization |
Family Cites Families (19)
Publication number | Priority date | Publication date | Assignee | Title |
---|---|---|---|---|
DE69416152T2 (de) * | 1993-10-06 | 1999-07-01 | Honeywell Inc | Virtueller graphikprozessor und verfahren für eingebettete echtzeitanzeigesysteme |
US5692193A (en) * | 1994-03-31 | 1997-11-25 | Nec Research Institute, Inc. | Software architecture for control of highly parallel computer systems |
US5828880A (en) * | 1995-07-06 | 1998-10-27 | Sun Microsystems, Inc. | Pipeline system and method for multiprocessor applications in which each of a plurality of threads execute all steps of a process characterized by normal and parallel steps on a respective datum |
JPH10228382A (ja) * | 1997-02-14 | 1998-08-25 | Nec Corp | コンパイル方式 |
US5991794A (en) * | 1997-07-15 | 1999-11-23 | Microsoft Corporation | Component integration system for an application program |
GB2336919A (en) | 1998-04-30 | 1999-11-03 | Ibm | Pre-emptive threading in a virtual machine |
KR20010072477A (ko) | 1998-08-13 | 2001-07-31 | 썬 마이크로시스템즈, 인코포레이티드 | 가상 머신 환경에서 네이티브 코드를 변환하고 실행하는방법 및 장치 |
US7234139B1 (en) * | 2000-11-24 | 2007-06-19 | Catharon Productions, Inc. | Computer multi-tasking via virtual threading using an interpreter |
US6779049B2 (en) * | 2000-12-14 | 2004-08-17 | International Business Machines Corporation | Symmetric multi-processing system with attached processing units being able to access a shared memory without being structurally configured with an address translation mechanism |
US6839828B2 (en) * | 2001-08-14 | 2005-01-04 | International Business Machines Corporation | SIMD datapath coupled to scalar/vector/address/conditional data register file with selective subpath scalar processing mode |
US7275249B1 (en) * | 2002-07-30 | 2007-09-25 | Unisys Corporation | Dynamically generating masks for thread scheduling in a multiprocessor system |
JP4487479B2 (ja) * | 2002-11-12 | 2010-06-23 | 日本電気株式会社 | Simd命令シーケンス生成方法および装置ならびにsimd命令シーケンス生成用プログラム |
US7000233B2 (en) * | 2003-04-21 | 2006-02-14 | International Business Machines Corporation | Simultaneous multithread processor with result data delay path to adjust pipeline length for input to respective thread |
DE602004017879D1 (de) * | 2003-08-28 | 2009-01-02 | Mips Tech Inc | Integrierter mechanismus zum suspendieren und endznem prozessor |
US6897871B1 (en) * | 2003-11-20 | 2005-05-24 | Ati Technologies Inc. | Graphics processing architecture employing a unified shader |
US20060150165A1 (en) * | 2004-12-30 | 2006-07-06 | Intel Corporation | Virtual microengine systems and methods |
US7861060B1 (en) * | 2005-12-15 | 2010-12-28 | Nvidia Corporation | Parallel data processing systems and methods using cooperative thread arrays and thread identifier values to determine processing behavior |
US8914618B2 (en) * | 2005-12-29 | 2014-12-16 | Intel Corporation | Instruction set architecture-based inter-sequencer communications with a heterogeneous resource |
US8010953B2 (en) * | 2006-04-04 | 2011-08-30 | International Business Machines Corporation | Method for compiling scalar code for a single instruction multiple data (SIMD) execution engine |
-
2007
- 2007-01-26 US US11/627,892 patent/US8321849B2/en active Active
-
2008
- 2008-01-22 DE DE202008017916U patent/DE202008017916U1/de not_active Expired - Lifetime
- 2008-01-22 DE DE102008005515A patent/DE102008005515A1/de not_active Ceased
- 2008-01-24 SG SG200800670-2A patent/SG144869A1/en unknown
- 2008-01-25 TW TW097102940A patent/TWI363294B/zh active
- 2008-01-25 CN CN2008100070263A patent/CN101231585B/zh active Active
- 2008-01-25 JP JP2008015281A patent/JP4999183B2/ja active Active
- 2008-01-28 KR KR1020080008740A patent/KR101026689B1/ko active IP Right Grant
Similar Documents
Publication | Publication Date | Title |
---|---|---|
JP4999183B2 (ja) | 並列スレッド・コンピューティングの仮想のアーキテクチャ及び命令セット | |
JP2008276740A5 (ja) | ||
US9898292B2 (en) | Hardware instruction generation unit for specialized processors | |
US20080109795A1 (en) | C/c++ language extensions for general-purpose graphics processing unit | |
US7941791B2 (en) | Programming environment for heterogeneous processor resource integration | |
TWI806550B (zh) | 處理器操作方法、相關電腦系統、及非暫時性電腦可存取儲存媒體 | |
Hormati et al. | Macross: Macro-simdization of streaming applications | |
Mikushin et al. | KernelGen--The Design and Implementation of a Next Generation Compiler Platform for Accelerating Numerical Models on GPUs | |
US20220365783A1 (en) | Matrix multiplication and accumulation operations on compressed matrices | |
JP2023518833A (ja) | ハードウェアアクセラレーションリソースを有効にするためのコンパイラ主導のタイル置換 | |
Compute | PTX: Parallel thread execution ISA version 2.3 | |
Varbanescu et al. | Towards an effective unified programming model for many-cores | |
US20230305845A1 (en) | Techniques to selectively store data | |
Manilov et al. | Free rider: A tool for retargeting platform-specific intrinsic functions | |
US20220342710A1 (en) | Application programming interface to indicate memory information | |
US20230086989A1 (en) | Parallel processing of thread groups | |
Döbrich et al. | Exploring online synthesis for CGRAs with specialized operator sets | |
Varbanescu et al. | Programming Models for Multicore and Many‐Core Computing Systems | |
Kawakami et al. | A Binary Translator to Accelerate Development of Deep Learning Processing Library for AArch64 CPU | |
US20220365750A1 (en) | Datatype conversion technique | |
US20230005097A1 (en) | Memory deallocation using graphs | |
US20220413945A1 (en) | Synchronization barrier | |
US20230005096A1 (en) | Memory allocation using graphs | |
Manilov et al. | Free rider: A source-level transformation tool for retargeting platform-specific intrinsic functions | |
Clarkson | Compiler and Runtime Support for Heterogeneous Programming |