本開示は、マルチプルプロセッサ計算プラットフォーム内で用いることができる通信技法について説明する。それらの技法は、幾つかの例では、コマンド待ち行列を用いてタスクを開始させるマルチプルプロセッサ計算プラットフォーム内でのメッセージ渡しをサポートするために用いることができるソフトウェアインタフェースを提供することができる。それらの技法は、追加の例では、マルチプルプロセッサ計算プラットフォーム内での共有メモリプロセッサ間通信のために用いることができるソフトウェアインタフェースを提供することができる。さらなる例では、それらの技法は、グラフィックス処理装置(GPU)とホストCPUとの間でのメッセージ渡し及び/又は共有メモリ通信をサポートするためのハードウェアを含むGPUを提供することができる。
近年においては、リアルタイム3Dグラフィックスの処理のために元来は設計されたプロセッサ、例えば、グラフィックス処理装置(GPU)、が、汎用計算タスク(GPGPU)を実行するために一般化された。GPGPUの価値は、部分的ではあるが、産業全体にわたる規格、例えば、Open Computing Language(OpenCL(登録商標))、の採用によって実証されている。OpenCLは、マルチプルプロセッサ計算プラットフォームにおいてタスクレベルでの並列性及びデータレベルでの並列性を有するプログラムを実行するために用いることができる、プラットフォーム横断、売り主横断型の異種計算プラットフォーム、並列プログラミングAPIの一例である。APIは、グラフィックス専用でない形でGPUのデータフロー及び制御経路を一般化することによってGPU上でのより一般化されたデータ処理を可能にするように特に設計されている。このアプローチ法の1つの限界は、ホストCPUと計算デバイス、例えば、GPU、との間でのデータ通信の粗い粒度(granularity)である。
例えば、OpenCL APIは、ホストデバイスと1つ以上の計算デバイスとの間でのタスクレベルの粒度の通信をサポートするコマンド待ち行列インタフェースを提供する。各コマンド待ち行列は、特定の計算デバイスによって実行されることになるコマンドを概して保有する。ホストデバイスで実行中のホストプロセスは、メモリ転送を行うようにホストデバイスに命令するコマンドをコマンド待ち行列内に入れることによってホストメモリスペースとデバイスメモリスペースとの間でデータを転送することができる。同様に、ホストプロセスは、計算デバイスでタスクを実行するようにホストデバイスに命令するコマンドをコマンド待ち行列に入れることによって計算デバイスでの実行をタスクに開始させることができる。
コマンド待ち行列インタフェースは、コマンドのインオーダー(in−order)実行又はコマンドのアウトオブオーダー(out−of−order)実行のいずれかを提供するように構成することができる。コマンド待ち行列インタフェースがコマンドのインオーダー実行を提供するように構成されるときには、コマンド待ち行列インタフェースは、コマンドがコマンド待ち行列内に入れられた順序で実行されること及び後続するコマンドの実行は先行コマンドが実行を完了するまで開始しないことを保証する。従って、ホストプロセスがタスクを実行するためのコマンドをコマンド待ち行列に入れたときには、コマンド待ち行列は、コマンド待ち行列に後続して入れられたあらゆる追加のコマンドを実行する前にそのタスクが実行を完了するのを待つ。
ホストCPUとGPU及びインオーダーコマンドが関わる単純な設定では、ホストCPUとGPUとの間の通信方式は、次の動作、すなわち、(1)ホストCPUがデータを準備してGPUがアクセス可能なメモリ内にそれを入れる、(2)ホストCPUが、タスクを実行するようにGPUに指令する、(3)ホストCPUが、GPUがタスクの実行を終了させるのを待つ、及び(4)ホストCPUが、GPUによってアクセス可能なメモリからホストメモリにデータをコピーする、を含むことができる。該構成においては、GPUでのタスクの実行の開始前にタスクの実行のために必要な全データがGPUによってアクセス可能なメモリに転送され、GPUで実行中のタスクによって生成されたデータは、GPUで実行中のタスクが実行を完了するまでホストCPUは利用可能できない。ホストCPUとGPUとの間でのデータシェアリングのこの粗さ(coarseness)は、並列に基づくアプリケーションのための数多くの有用な動作、例えば、ホストデバイスで実行中のプロセスとGPUで実行中のタスクとの間でプロセス間メッセージを渡すこと、の有効な実装を妨げることがある。該メッセージは、例えば、GPUで実行中のタスクがホストCPUでリモートプロシージャコール(Remote Procedure Call(RPC)(遠隔手順呼び出し)を実行する能力を有することを可能にする上で役立つことができる。
コマンド待ち行列インタフェースがコマンドのアウトオブオーダー実行を提供するように構成されるときには、ホストプロセスは、特定のタスクの実行中には、特定のコマンドの実行がいつ生じるか制御することができない。従って、コマンド待ち行列のためのアウトオブオーダー実行モードは、ホストデバイスで実行中のプロセスとGPUで実行中のタスクとの間でのプロセス間メッセージ渡しの実装を有効な形で可能にしない。
OpenCL内で用いられるメモリモデルに関して、APIは、ホストCPUとGPUとの間でデータを共有するために又は複数のOpenCL計算デバイス間でデータを共有するために用いることができるいわゆるグローバルなCLバッファ及びグローバルなCL画像を定義する。しかしながら、CPU及びGPUは、同時にバッファから読み取ること又はバッファに書き込むことができない。典型的には、CPUは、ソースデータが入った1つ以上のバッファを準備し、それらのバッファを処理のためにGPUに渡す。GPUは、これらのバッファを変更するか又はCPUデータ変更を受信することを目的としてCPUで実行中のソフトウェアによって推測的にも割り当てられたその他のバッファ内に結果を入れる。
OpenCL内のメモリオブジェクトは、現在は、ホストメモリスペースの領域を計算デバイスによって用いられるバッファデータの格納のために使用するのを可能にしているが、本明細書は、計算デバイスがタスクのより効率的な実行のためにこのデータをキャッシングするのを可能にする。ホストデバイスは、概して、バッファデータをキャッシングするために用いられる計算デバイスキャッシュを直ちに無効にすることができない。従って、ホストデバイスが、ホストメモリスペース内に格納される一定のメモリバッファデータをオーバーライトしようとしても、計算デバイスが変更されたデータに直ちにアクセスすることを可能にするために計算デバイス内のキャッシュが更新されることが保証されない。さらに、計算デバイスによって行われた計算の結果は、計算デバイスのキャッシュに格納することができるため、ホストデバイスで実行中のホストプロセスは、バッファからのデータは計算デバイスキャッシュに格納されているより新しいデータに起因して無効になっていることがあるためバッファからはどのような一部の結果も読み取ることができない。従って、OpenCLにおけるメモリ管理モデルは、共有されるメモリを介してのインフライト(in−flight)データシェアリングを容易に可能にすることができない。
本開示において説明される技法は、幾つかの例では、OpenCL APIの上記の限界のうちの1つ以上を克服するために使用することができる。例えば、本開示の技法は、タスクレベルの粒度のコマンド待ち行列を用いてタスクを開始させるマルチプルプロセッサ計算プラットフォーム内でのプロセス間メッセージ渡しをサポートするために用いることができるソフトウェアインタフェースを提供することができる。他の例として、本開示の技法は、マルチプルプロセッサ計算プラットフォーム内の共有メモリを介してのインフライトデータシェアリングをサポートするために使用することができるソフトウェアインタフェースを提供することができる。
幾つかの例では、本開示の技法は、ソフトウェアレベルでのメッセージ渡しを容易にするGPUハードウェアアーキテクチャを提供することができる。例えば、本開示の技法は、ソフトウェアレベルメッセージ渡し命令の実行をサポートするように構成されるGPUハードウェアアーキテクチャを提供することができる。さらなる例では、本開示の技法は、GPUとホストCPUとの間の共有メモリ通信を容易にするGPUハードウェアアーキテクチャを提供することができる。例えば、本開示の技法は、共有メモリスペースのためのキャッシングサービスを選択的にイネーブル及びディスエーブルにするように及び/又は共有メモリスペースのためにキャッシュコヒーレンシーメカニズムを選択的にイネーブル及びディスエーブルにするように構成されるGPUハードウェアアーキテクチャを提供することができる。
本開示の第1の態様により、計算デバイスによるタスクの実行中におけるホストデバイスと1つ以上の計算デバイスとの間でのメッセージ渡し命令の実行を容易にするメッセージ渡しインタフェースが提供される。メッセージ渡しは、通信中のプロセスが各々メッセージを成功裏に渡すために補完的な組の動作を行う、プロセス間の、及び潜在的にデバイス間の、通信の1つの形態を意味することができる。例えば、メッセージ渡しプロトコルにより通信するプロセスの各々は、送信動作及び受信動作を実装することができる。本開示におけるメッセージ渡し技法は、CPU及び計算デバイス、例えば、GPU、が計算デバイスでのタスクの実行中に互いにメッセージを渡すことを可能にすることができる。この方法により、タスクレベルの粒度のコマンド待ち行列通信方式を実装するマルチプロセッサ計算プラットフォームは、プロセス間及び/又はデバイス間通信を容易にすることができる。
幾つかの例では、本開示において説明されるメッセージ渡し技法は、ホストデバイスと計算デバイス、例えば、GPU、との間での通信のためにOpenCLにおいて典型的に用いられるコマンド待ち行列インタフェース以外のインタフェースを用いることができるため、これらの技法は、“アウトオブバンドシグナリング”(out−of−band signaling)技法と呼ぶことができる。換言すると、本開示の技法は、OpenCL内に含められているインバンド(in band)コマンド待ち行列インタフェースと論理的に別個の新しいアウトオブバンド通信インタフェースを含むことができる。アウトオブバンド通信インタフェースは、コマンド待ち行列インタフェースの対象となるのと同じタスクレベルの粒度にならず、それにより、コマンド待ち行列のタスクレベルの粒度に関して上述される1つ以上の限界の解決策を提供することができる。
本開示の技法によりCPUとGPUとの間で転送されるメッセージは、あらゆるタイプのメッセージであることができる。異なるタイプのメッセージの例は、信号と、メモリ割り当て要求と、メモリ割り当て解除要求と、通知メッセージと、同期化メッセージと、遠隔手順呼び出しメッセージ(例えば、リモートプロシージャコール(RPC)の一部であるメッセージ)と、データパケットと、報告メッセージと、アサーションメカニズムメッセージと、ロギングメッセージと、を含む。
現在のOpenCL規範では、ホストCPUからGPUへの全要求がOpenCLコマンド待ち行列内に入れられ、次にGPUに送信される。特に、アプリケーションは、非常に多数のカーネル実行及びバッファ動作をコマンド待ち行列に入れることができる。他方、最初に待ち行列に入れられたタスク、例えば、カーネル実行、が、例えば、CPUに追加のメモリ割り当てを要求する必要がある場合は、複数の課題が発生する。第1に、GPUは、メモリ割り当てが必要であることを実行中のカーネル内部からどのようにしてCPUに通知するかということである。第2に、CPUは、メモリ割り当ての完了及び新たに割り当てられたメモリブロックのアドレスをどのようにしてGPUに通知するかということである。しかしながら、本開示のメッセージ渡しインタフェース技法は、上記の通知及び情報が入った1つ以上のメッセージをCPUとGPUとの間で渡すことができるようにすることによってこれらの課題を解決することができる。
本開示のアウトオブバンドシグナリング技法は、幾つかの例では、ホストCPUと1つ以上の計算デバイス、例えば、OpenCL計算デバイス、との間でのシグナリングを実装するために用いることができる。アウトオブバンドシグナリングは、例えば、プッシュ又はプルメカニズムを用いて、高速のアウトオブバンド通知を提供することができる。幾つかの例では、アウトオブバンドシグナリング技法は、相対的に少量のデータを搬送することができる。
本開示の第2の態様により、GPU以外のプロセッサにおいて実行中のプロセスにメッセージを送信すること及び実行中のプロセスからメッセージを受信することが可能なGPUが提供される。例えば、GPUは、メッセージを送信及び受信するための1つ以上の動作を実装するように構成されるハードウェアを含むことができる。幾つかの例では、本開示により設計されたGPUは、メッセージ渡しプロトコルと関連付けられた状態情報およびデータ情報を格納するように構成された1つ以上のホストがアクセス可能なレジスタを含むことができる。1つ以上のレジスタは、GPUで実行中のタスクとGPU以外のデバイスで実行中のプロセスとの間でのメッセージ渡しを容易にするように構成することができる。さらなる例では、ホストがアクセス可能なレジスタを介してメッセージを送信及び受信するためにGPUのALU処理ブロック(例えば、プログラマブルシェーダユニット)をホストがアクセス可能なレジスタに通信可能な形で結合することができる。GPUは、同期的及び/又は非同期的メッセージ渡し技法を実装するための様々なポーリング及び/又は割り込みメカニズムを含むように設計することもできる。
本開示の第3の態様により、即時(immediate)メモリオブジェクトを生成するのを可能にするメモリバッファインタフェースが提供される。即時メモリオブジェクトは、計算デバイスでタスクが実行している間にホストデバイスで実行中のプロセスと計算デバイスで実行中のタスクとの間でデータを共有するためにキャッシング不能な共有メモリスペース及び/又はキャッシュコヒーレントな共有メモリスペースを実装するために用いることができる。共有メモリスペースは、計算デバイス、例えば、GPU、によるタスクの実行中にホストデバイス及び計算デバイスの両方によってアクセス可能なメモリスペースであることができる。ここにおいて用いられる場合におけるキャッシング不能な共有メモリスペースとは、ホストデバイス及び計算デバイスのうちの1つ又は両方における1つ以上の対応するキャッシュがメモリスペースに関してディスエーブルにされる共有メモリスペースを意味することができる。ここにおいて用いられる場合におけるキャッシュコヒーレントな共有メモリスペースとは、ホストデバイス及び計算デバイスのうちの1つ又は両方における1つ以上の対応するキャッシュ内でキャッシュコヒーレンシーを維持するために共有メモリキャッシュコヒーレンシー技法が用いられる共有メモリスペースを意味することができる。キャッシング不能な共有メモリスペース及びキャッシュコヒーレントな共有メモリスペースは、何時でもデータシェアリングを可能にすることができる。即時メモリオブジェクトは、幾つかの例では、ホストデバイス及び計算デバイスに関するキャッシング不能な揮発性の共有メモリとして及び/又はキャッシュコヒーレントな揮発性の共有メモリとして実装することができる。
幾つかの例では、本開示の即時メモリオブジェクトは、メモリオブジェクトメモリ管理方式を含む、プラットフォーム横断、売り主横断型の異種計算プラットフォーム、並列プログラミングAPI内に組み入れることができる。例えば、即時メモリオブジェクトは、OpenCLメモリオブジェクトの追加属性、例えば、OpenCLバッファオブジェクト又はOpenCL画像オブジェクト、としてOpenCL内に組み入れることができる。該例では、即時メモリオブジェクトは、メモリオブジェクト生成関数を変更して、関数呼び出しによって生成されて結果的に得られたメモリオブジェクトが標準モードメモリオブジェクト又は即時モードメモリオブジェクトのいずれであるべきかを指定するパラメータ又はフラグを含めることによって、生成することができる。この方法により、本開示の技法は、メモリオブジェクトメモリ管理方式、例えば、OpenCL、を含むAPIを実装するマルチプルプロセッサ計算システムがキャッシュコヒーレンシー問題の影響を受けない共有メモリスペースを介してのインフライトデータシェアリングを実装するのを可能にすることができる。
さらなる例では、本開示の即時メモリオブジェクトは、ホストCPUとOpenCL計算デバイスとの間での又は異なるOpenCL計算デバイス間でのインフライトなデータシェアリングのために用いることができる。追加の例では、即時メモリオブジェクトは、内部同期化マーカを含むことができる。さらなる例では、即時メモリオブジェクトは、同期化のためにアウトオブバンド信号とともに用いることができる。
本開示の第4の態様により、キャッシング不能な共有メモリスペースを提供するために特定のメモリアドレス空間に関して選択的にディスエーブルにすることができる共有メモリスペースに対応するキャッシュを含むGPUが提供される。例えば、GPUは、共有メモリスペースに関して読み取り動作及び/又は書き込み動作を実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報を受信したことに応答して共有メモリスペースと関連付けられたキャッシュによって提供されるキャッシングサービスを選択的にイネーブル及びディスエーブルにすることができる。幾つかの例では、共有メモリスペースに関して読み取り動作及び/又は書き込み動作を実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報は、特定の命令を実行するためにキャッシュドモード(cached mode)又は即時モード(immdeiate mode)のいずれが使用されるべきかを指定するキャッシュドモード命令又は即時モード命令であることができる。さらなる例では、共有メモリスペースに関して読み取り動作及び/又は書き込み動作を実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報は、メモリオブジェクトのために即時モードがイネーブルにされるかどうかを指定する即時モードメモリオブジェクト属性であることができる。
さらなる例では、本開示の技法は、キャッシュコヒーレントな共有メモリスペースを提供するために選択的にイネーブルにすることができるキャッシュコヒーレンシーモードを含むGPUを提供することができる。幾つかの例では、GPUは、ホストデバイスから受信された1つ以上の命令に基づいて共有メモリスペースに対応するキャッシュの部分のためにキャッシュコヒーレンシーモードを選択的にイネーブルにすることができる。ホストデバイスは、ホストプロセスによって指定された即時モードパラメータに基づいてホストデバイスによる共有メモリスペースの割り当て時に共有メモリスペースに対応するキャッシュの部分のために共有メモリスペースキャッシュコヒーレンシーモードを選択的にイネーブルにするために1つ以上の命令をGPUに出すことができる。
本開示のアウトオブバンドシグナリング技法及び即時バッファリング技法は、OpenCLコマンド待ち行列インタフェースのみを用いて入手可能になるタスク結合と比較してホストCPUとGPUとの間において又は2つのOpenCL計算デバイス間でより微細な粒子のそれを提供することができる。本開示の技法は、並列の及び/又はマルチスレッドのプログラムの効率的な実行を援助するためにマルチプルプロセッサ計算プラットフォームが様々な動作を行うことを可能にすることができる。例えば、本開示の技法は、GPUで実行中のタスクがRPCを起動させることを可能にすることができる。他の例として、本開示の技法は、GPUで実行中のタスクが、CPUを介して、他のGPUタスクを起動させることを可能にすることができる。さらなる例として、本開示の技法は、GPUで実行中のタスクがCPU及び/又はCPUで実行中のドライバに対してリソース管理要求、例えば、メモリ割り当て要求及び/又はメモリ割り当て解除要求、を出すことを可能にすることができる。さらに他の例として、本開示の技法は、GPUで実行中のタスクが状態検査及びCPUへの一般的なメッセージ渡し、例えば、アサーションメカニズムの実装、進行状況報告、及び/又は診断ロギング、を行うことを可能にすることができる。
図1は、本開示による計算システム例10を示したブロック図である。計算システム10は、複数の処理デバイスで1つ以上のソフトウェアアプリケーションを処理するように構成される。幾つかの例では、1つ以上のソフトウェアアプリケーションは、ホストプロセスを含むことができ、計算システム10は、ホストプロセスを実行するように及び計算システム10内のその他の計算デバイスでホストプロセスによって開始された1つ以上のタスクの実行を分散させるように構成することができる。さらなる例では、計算システム10によって実行されるホストプロセス及び/又はタスクは、並列プログラミングモデルによりプログラミングすることができる。例えば、アプリケーションは、基礎になるハードウェアシステムのタスクレベルでの並列性及び/又はデータレベルでの並列性を利用するように設計された命令を含むことができる。
計算システム10は、パソコン、デスクトップコンピュータ、ラップトップコンピュータ、コンピュータワークステーション、ビデオゲームプラットフォーム又はコンソール、移動電話、例えば、セルラー又は衛星電話、携帯電話、ランドライン電話、インターネット電話、ハンドヘルドデバイス、例えば、ポータブルビデオゲーム機、又はパーソナルデジタルアシスタント(PDA)、デジタルメディアプレーヤー、例えば、パーソナル音楽プレーヤー、ビデオプレーヤー、表示装置、テレビ、テレビセットトップボックス、サーバ、中間ネットワークデバイス、メインフレームコンピュータ又は情報を処理するその他のあらゆるタイプのデバイスであることができる。
計算デバイス10は、ホストデバイス12と、グラフィックス処理装置(GPU)14と、メモリ16と、相互接続ネットワーク18と、を含む。ホストデバイス12は、マルチプルプロセッサ計算プラットフォームAPIのためのホストプロセス及びランタイムモジュールの実行のためのプラットフォームを提供するように構成される。典型的には、ホストデバイス12は、汎用CPUであるが、ホストデバイス12は、プログラムを実行することが可能なあらゆるタイプのデバイスであることができる。ホストデバイス12は、相互接続ネットワーク18を介してGPU14及びメモリ16に通信可能な形で結合される。ホストデバイス12は、ホストプロセス20と、ランタイムモジュール22と、を含み、それらの各々は、1つ以上のプログラミング可能なプロセッサのあらゆる組み合わせにおいて実行することができる。
ホストプロセス20は、計算システム10の計算システムプラットフォームでの実行のためのソフトウェアプログラムを形成する命令の組を含む。ソフトウェアプログラムは、エンドユーザ端末のために1つ以上の特定のタスクを実行するように設計することができる。該タスクは、幾つかの例では、計算システム10によって提供される複数の処理デバイス及び並列アーキテクチャを利用することができる計算集約型のアルゴリズムを含むことができる。
ランタイムモジュール22は、ホストプロセス20に含まれる命令のうちの1つ以上にサービスを提供するように構成された1つ以上のインタフェースを実装するホストデバイス12で実行するソフトウェアモジュールであることができる。ランタイムモジュール22によって実装されたインタフェースは、コマンド待ち行列インタフェース24と、ホストメッセージ渡しインタフェース26と、を含む。幾つかの例では、ランタイムモジュール22は、本開示で説明されるインタフェースに加えて標準的なマルチプルプロセッサシステムAPI内に含まれる1つ以上のインタフェースを実装することができる。幾つかの例では、標準的なAPIは、異種計算プラットフォームAPI、プラットフォーム横断型API、売り主横断型API、並列プログラミングAPI、タスクレベル並列プログラミングAPI、及び/又はデータレベル並列プログラミングAPIであることができる。さらなる例では、標準的なAPIは、OpenCL APIであることができる。該例では、ランタイムモジュール22は、OpenCL仕様のうちの1つ以上に準拠するように設計することができる。追加の例では、ランタイムモジュール22は、ドライバプログラム、例えば、GPUドライバ、の一部として実装することができる。
コマンド待ち行列インタフェース24は、ホストプロセス20から1つ以上の待ち行列内追加命令を受信するように及び受信された命令によって指定された機能を実行するように構成される。幾つかの例では、コマンド待ち行列インタフェース24は、OpenCL仕様に準拠して設計することができる。例えば、コマンド待ち行列インタフェース24は、コマンド待ち行列と対話するためにOpenCL仕様において指定された待ち行列内追加命令のうちの1つ以上を実装することができる。
本開示により、ホストメッセージ渡しインタフェース26は、ホストプロセス20から1つ以上のメッセージ渡し命令を受信するように及び受信された命令によって指定された機能を実行するように構成される。幾つかの例では、ホストメッセージ渡しインタフェース26は、既存の標準API、例えば、OpenCL API、の拡張として実装することができる。追加の例では、ホストメッセージ渡しインタフェース26は、既存の標準API、例えば、OpenCL API、内に組み入れることができる。
GPU14は、ホストデバイス12から受信された命令に応答して1つ以上のタスクを実行するように構成される。GPU14は、1つ以上のプログラミング可能な処理素子を含むあらゆるタイプのGPUであることができる。例えば、GPU14は、タスクのための複数の実行インスタンスを並行して実行するように構成される1つ以上のプログラマブルシェーダユニットを含むことができる。プログラマブルシェーダユニットは、バーテックス(頂点)シェーダユニット、フラグメントシェーダユニット、ジオメトリシェーダユニット及び/又は統合シェーダユニットを含むことができる。GPU14は、相互接続ネットワーク18を介してホストデバイス12及びメモリ16に通信可能な形で結合される。GPU14は、タスク28と、デバイスメッセージ渡しインタフェース30と、を含む。タスク28及びデバイスメッセージ渡しインタフェース30は、1つ以上のプログラミング可能な処理素子のあらゆる組み合わせにおいて実行することができる。
タスク28は、計算システム10内の計算デバイスでの実行のためのタスクを形成する命令の組を備える。幾つかの例では、タスク28のための命令の組は、ホストプロセス20において定義し、幾つかの事例では、ホストデバイス12で実行中のホストプロセス20に含まれる命令によってコンパイルすることができる。さらなる例では、タスク28は、GPU14で並行して実行中の複数の実行インスタンスを有するカーネルプログラムであることができる。該例においては、ホストプロセス20は、カーネル実行インスタンスを実行するために各々の処理素子にカーネル実行インスタンスをマッピングするカーネル用のインデックススペースを定義することができ、GPU14は、そのカーネル用に定義されたインデックススペースによりタスク28のための複数のカーネル実行インスタンスを実行することができる。
本開示により、デバイスメッセージ渡しインタフェース30は、ホストプロセス20から1つ以上のメッセージ渡し命令を受信するように及び受信された命令によって指定された機能を実行するように構成される。幾つかの例では、デバイスメッセージ渡しインタフェース30は、既存の標準のAPIの拡張として実装することができる。例えば、標準APIは、標準的な計算デバイスAPI、例えば、OpenCL C API、であることができる。追加の例では、デバイスメッセージ渡しインタフェース30は、既存の標準のAPI、例えば、OpenCL C API、内に組み入れることができる。
メモリ16は、ホストデバイス12及びGPU14のうちの1つ又は両方による使用のためにデータを格納するように構成される。メモリ16は、1つ以上の揮発性又は非揮発性のメモリ又は記憶デバイス、例えば、ランダムアクセスメモリ(RAM)、スタティックRAM(SRAM)、ダイナミックRAM(DRAM)、読み取り専用メモリ(ROM)、消去可能プログラマブルROM(EPROM)、電気的に消去可能なプログラマブルROM(EEPROM)、フラッシュメモリ、磁気データ記憶媒体又は光学記憶媒体、のあらゆる組み合わせを含むことができる。メモリ16は、相互接続ネットワーク18を介してホストデバイス12及びGPU14に通信可能な形で結合される。メモリ16は、コマンド待ち行列32を含む。
コマンド待ち行列32は、コマンド待ち行列インタフェース24から受信されたコマンドを格納及び取り出すメモリ16内に実装されたデータ構造であることができる。幾つかの例では、コマンド待ち行列32は、特定の実行順序でコマンドを格納するバッファであることができる。
相互接続ネットワーク18は、ホストデバイス12、GPU14及びメモリ16の間での通信を容易にするように構成される。相互接続ネットワーク18は、当業において知られるあらゆるタイプの相互接続ネットワークであることができる。図1の計算システム例10では、相互接続ネットワーク18は、バスである。バスは、様々なバス構造、例えば、第3世代バス(例えば、HyperTransportバス又はInfiniBandバス)、第2世代バス(例えば、Advanced Graphics Portバス、Peripheral Component Interconnect Express(PCIe)バス、又はAdvanced eXentisible Interface(AXI)バス)、又はその他のタイプのバスのうちの1つ以上を含むことができる。相互接続ネットワーク18は、ホストデバイス12、GPU14及びメモリ16に結合される。
今度は、計算システム10内のコンポーネントの構造及び機能がさらに詳細に説明される。上述されるように、ホストプロセス20は、命令の組を含む。命令の組は、例えば、1つ以上の待ち行列内追加命令と、1つ以上のホストメッセージ渡し命令と、を含むことができる。追加の例では、命令の組は、GPU14で実行されるタスク又はカーネルを指定する命令と、コマンド待ち行列を生成してそれらのコマンド待ち行列を特定のデバイスと関連付ける命令と、プログラムをコンパイル及びバインドする命令と、カーネルパラメータを設定する命令と、インデックススペースを定義する命令と、デバイスコンテキストを定義する命令と、ホストプロセス20によって提供される機能をサポートするその他の命令と、を含むことができる。
ホストプロセス20は、コマンド待ち行列32内に1つ以上のコマンドを入れるようにコマンド待ち行列インタフェース24に命令する1つ以上の待ち行列内追加命令をコマンド待ち行列インタフェース24に出すことによってコマンド待ち行列インタフェース24と対話することができる。1つ以上の待ち行列内追加命令は、コマンド待ち行列32内にメモリ転送コマンドを追加するようにコマンド待ち行列インタフェース24に命令するメモリ転送待ち行列内追加命令を含むことができる。例えば、1つ以上の待ち行列内追加命令は、ホストデバイス12と関連付けられたメモリスペースとGPU14と関連付けられたメモリスペースとの間でデータを転送するようにホストデバイス12、例えば、ホストデバイス12で実行中のランタイムモジュール22、に命令するコマンドを待ち行列に入れる命令を含むことができる。
メモリスペースは、ホストデバイス12によるホストプロセス20の実行中にホストデバイス12によってアクセス可能である場合にホストデバイス12と関連付けることができる。同様に、メモリスペースは、GPU14によるタスク28の実行中にGPU14によってアクセス可能である場合にGPU14と関連付けることができる。ホストデバイス12と関連付けられたメモリスペースは、ここでは、ホストメモリスペースと呼ぶことができ、GPU14と関連付けられたメモリスペースは、ここでは、デバイスメモリスペースと呼ぶことができる。幾つかの例では、メモリ16は、ホストメモリスペース及びデバイスメモリスペースの両方の一部分を含むことができる。さらなる例では、ホストメモリスペース及びデバイスメモリスペースのうちの1つの又は両方の一部分は、図1の計算システム10において示されていない1つ以上のその他のメモリデバイスに配置することができる。
幾つかの例では、ホストデバイス12と関連付けられたメモリスペースとGPU14と関連付けられたメモリスペースとの間でデータを転送するようにホストデバイス12に命令するコマンドは、ホストメモリスペースの一部分に格納されたデータをデバイスメモリスペース内で割り当てられたバッファオブジェクトに転送するようにランタイムモジュール22に命令するコマンドであることができる。該コマンドを待ち行列に入れるためにホストプロセス20によって出された命令は、ここでは、書き込みバッファ待ち行列内追加命令と呼ぶことができる。幾つかの場合は、書き込みバッファ待ち行列内追加命令は、OpenCL API仕様によって指定されたclEnqueueWriteBuffer()関数の形態をとることができる。
追加の例では、ホストデバイス12と関連付けられたメモリスペースとGPU14と関連付けられたメモリスペースとの間でデータを転送するようにホストデバイス12に命令するコマンドは、デバイスメモリスペース内で割り当てられたバッファオブジェクトに格納されたデータをホストメモリスペースの一部分に転送するようにランタイムモジュール22に命令するコマンドであることができる。該コマンドを待ち行列に入れるためにホストプロセス20によって出される命令は、ここでは、読み取りバッファ待ち行列内追加命令と呼ぶことができる。幾つかの場合は、読み取りバッファ待ち行列内追加命令は、OpenCL API仕様によって指定されたclEnqueueReadBuffer()関数の形態をとることができる。
1つ以上の待ち行列内追加命令は、コマンド待ち行列32内にタスク実行コマンドを入れるようにコマンド待ち行列インタフェース24に命令するタスク実行待ち行列内追加命令を含むこともできる。例えば、1つ以上の待ち行列内追加命令は、GPU14でタスクを実行するようにホストデバイス12、例えば、ホストデバイス12で実行中のランタイムモジュール22、に命令するコマンドを待ち行列内に入れるための命令を含むことができる。幾つかの例では、タスクを実行するためのコマンドは、GPU14内の複数の処理素子でタスクの複数の実行インスタンスを並行して実行するためのコマンドであることができる。例えば、タスクは、カーネルであることができ、ホストプロセス20は、カーネル実行インスタンスを実行するためにGPU14内の各々の処理素子にカーネル実行インスタンスをマッピングするカーネル用のインデックススペースを定義することができる。該例では、タスクを実行するためのコマンドは、GPU14のために定義されたインデックススペースによりGPU14でカーネルを実行するためのコマンドであることができる。幾つかの事例では、タスク実行待ち行列内追加命令は、OpenCL APIによって指定されたclEnqueueNDRangeKernel()関数の形態を取ることができる。
本開示により、ホストプロセス20は、ホストメッセージ渡しインタフェース26に1つ以上のホストメッセージ渡し命令を出してホストデバイス12で実行中のホストプロセス20とGPU14で実行中のタスク28との間で1つ以上のメッセージを渡すようにホストメッセージ渡しインタフェース26に命令することによってホストメッセージ渡しインタフェース26と対話することもできる。ホストメッセージ渡し命令は、ホストデバイス12によって実行することができる。
ホストメッセージ渡し命令は、幾つかの例では、指定されたデバイスに指定されたデータを送信するようにホストデバイス12に命令する送信命令を含むことができる。例えば、送信命令は、ホストデバイス12で実行中のホストプロセス20からGPU14で実行中のタスク28にメッセージを送信するようにホストメッセージ渡しインタフェース26に命令することができる。幾つかの例では、送信命令は、メッセージが送信されるべき特定のデバイスを指定する第1の入力パラメータと、送信されるメッセージの内容を指定する第2の入力パラメータと、を含むことができる。
送信命令は、ブロッキング(blocking)送信命令又は非ブロッキング(non−blocking)送信命令のいずれかであることができる。送信命令は、幾つかの例では、送信命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する第3の入力パラメータを含むことができる。ブロッキング送信命令は、呼び出しを行っているプロセス、例えば、ホストデバイス12で実行中のホストプロセス20、に戻る前に送信動作が完了されるまで待つことができる。非ブロッキング送信命令は、送信動作が完了されるまで待たずに呼び出しを行っているプロセスに戻ることができる。例えば、非ブロッキング送信命令は、特定の送信動作が成功であったかどうかを決定するために呼び出しを行っているプロセスによって後続して問い合わせることができるハンドルをその送信動作に戻すことができる。非ブロッキング送信命令は、失敗することがあり、失敗した場合は、呼び出しを行っているプロセスは、送信動作を再試行するために送信命令を再度出すことが必要な場合がある。
幾つかの例では、送信命令のためのインタフェースは、次の形態をとることができる。
ここで、clSendOutOfBandDataは、命令識別子であり、cl_device *deviceIdは、メッセージが送信されるべき特定のOpenCLデバイスを指定する入力パラメータであり、int OOB_dataは、送信されるメッセージの内容を指定する入力パラメータであり、bool blockingは、命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する入力パラメータである。ブロッキング命令の場合は、命令は、送信動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。非ブロッキング送信命令の場合は、命令は、呼び出しを行っているプロセスによる後続する状態問い合わせのためのハンドルパラメータを戻すことができる。
ホストメッセージ渡し命令は、幾つかの例では、非同期的な方法で指定されたデバイスからデータを受信するためにコールバックをレジスタに入れるようにホストデバイス12に命令するレジスタコールバックルーチン命令を含むことができる。例えば、レジスタコールバックルーチン命令は、GPU14で実行中のタスクがホストプロセス20にメッセージを送信していることを示す信号をGPU14から受信したことに応答してコールバックルーチンを呼び出すようにホストメッセージ渡しインタフェース26に命令することができる。レジスタコールバックルーチン命令は、コールバックルーチンをレジスタに入れるべき対象となる特定のデバイスを指定する第1の入力パラメータと、コールバックルーチンのメモリ記憶場所を指定する第2の入力パラメータと、を含むことができる。
幾つかの例では、レジスタコールバックルーチン命令のためのインタフェースは、次の形態をとることができる。
ここで、clRegisterOutOfBandDataCallbackは、命令識別子であり、cl_device *deviceIdは、メッセージが送信されるべき特定のOpenCLデバイスを指定する入力パラメータであり、void(*)(int) callBackPtrは、コールバックルーチンのメモリ記憶場所を指定する入力パラメータである。レジスタコールバックルーチン命令は、コールバックルーチンレジストレーション動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。
ホストメッセージ渡し命令は、幾つかの例では、指定されたデバイスからデータを読み取るのを試行するようにホストデバイス12に命令するポーリング命令を含むことができる。例えば、ポーリング命令は、GPU14で実行中のタスク28がメッセージを送信しているかどうかを示すメッセージ状態情報に関してGPU14をポーリングするようにホストメッセージ渡しインタフェース26に命令することができる。ポーリング命令は、ポーリングされるべき特定のデバイスを指定する入力パラメータと、存在する場合に、ポーリングの結果得られたデータを指定する出力パラメータと、を含むことができる。
幾つかの例では、ポーリング命令のためのインタフェースは、次の形態をとることができる。
ここで、clTryReadOutOfBandDataは、命令識別子であり、cl_device *deviceIdは、ポーリングされるべき特定のOpenCLデバイスを指定する入力パラメータであり、int*OOB_dataは、存在する場合に、ポーリングの結果得られたデータを指定する出力パラメータである。ポーリング命令は、ポーリング動作からデータが成功裏に得られたかどうかを示すパラメータを戻すことができる。
ホストプロセス20と同様に、タスク28は、計算デバイスによって実行される1つ以上のデバイスメッセージ渡し命令を含むことができる。デバイスメッセージ渡し命令は、ホストデバイス12に指定されたデータを送信するように計算デバイスに命令する送信命令を含むことができる。例えば、送信命令は、GPU14で実行中のタスク28からホストデバイス12で実行中のホストプロセス20にメッセージを送信するようにGPU14に命令することができる。
送信命令は、ブロッキング送信命令又は非ブロッキング送信命令のいずれかであることができる。送信命令は、幾つかの例では、送信命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する第1の入力パラメータを含むことができる。ブロッキング送信命令は、呼び出しを行っているプロセス、例えば、GPU14で実行中のタスク28、を停止させ、呼び出しを行っているプロセスに戻る前に送信動作が完了されるのを待つことができる。非ブロッキング送信命令は、送信動作が完了されるまで待たずに呼び出しを行っているプロセスに戻ることができる。例えば、非ブロッキング送信命令は、特定の送信動作が成功であったかどうかを決定するために後続して呼び出しを行っているプロセスによって問い合わせることができるハンドルをその送信動作に戻すことができる。非ブロッキング送信動作は、失敗することがあり、失敗した場合は、呼び出しを行っているプロセスは、送信動作を再試行するために送信命令を再度出すことが必要な場合がある。送信命令は、ホストデバイスに送信されるべきメッセージの内容を指定する第2の入力パラメータを含むことができる。
幾つかの例では、送信命令のためのインタフェースは、次の形態をとることができる。
ここで、send_oobdataは、命令識別子であり、bool blockingは、命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する入力パラメータであり、int dataは、送信されるメッセージの内容を指定する入力パラメータである。ブロッキング命令の場合は、命令は、送信動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。非ブロッキング命令の場合は、命令は、呼び出しを行っているプロセスによる後続する状態問い合わせのためのハンドルパラメータを戻すことができる。
デバイスメッセージ渡し命令は、幾つかの例では、ホストデバイス12からデータを受信するように計算デバイスに命令する受信命令を含むことができる。例えば、受信命令は、入手可能な場合にホストデバイス12で実行中のホストプロセス20からタスク28に送信されたメッセージをGPU14で実行中のタスク28に提供するようにGPU14、例えば、デバイスメッセージ渡しインタフェース30、に命令することができる。該命令は、ポーリングメカニズムをサポートするために用いることができる。
受信命令は、ブロッキング受信命令又は非ブロッキング受信命令のいずれかであることができる。受信命令は、幾つかの例では、受信命令がブロッキング受信命令であるか又は非ブロッキング受信命令であるかを指定する入力パラメータを含むことができる。ブロッキング受信命令は、呼び出しを行っているプロセス、例えば、GPU14で実行中のタスク28、を停止させ、呼び出しを行っているプロセスに戻る前にメッセージが入手可能になるまで待つことができる。非ブロッキング受信命令は、メッセージが入手可能になるまで待たずに呼び出しを行っているプロセスに戻ることができる。例えば、メッセージが入手可能である場合は、非ブロッキング送信命令は、そのメッセージを戻すことができる。しかしながら、メッセージが入手可能でない場合は、非ブロッキング受信命令は、失敗することがある。失敗した場合は、呼び出しを行っているプロセスは、受信動作を再試行するために受信命令を再度出すことが必要な場合がある。受信命令は、存在する場合に、受信動作の結果得られたデータを指定する出力パラメータを含むことができる。
幾つかの例では、受信命令のためのインタフェースは、次の形態をとることができる。
ここで、receive_oobdataは、命令識別子であり、bool blockingは、命令がブロッキング受信命令であるか又は非ブロッキング受信命令であるかを指定する入力パラメータであり、int dataは、存在する場合に、受信動作の結果得られたデータを指定する出力パラメータである。命令は、受信動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。
コマンド待ち行列インタフェース24は、コマンド待ち行列32内にコマンドを入れるように構成される。例えば、コマンド待ち行列インタフェース24は、ホストプロセス20から1つ以上の待ち行列内追加命令を受信し、ホストプロセス20から1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列32内に1つ以上のコマンドを入れることができる。1つ以上の待ち行列内追加命令は、タスク実行コマンド及びデータ転送コマンドをそれぞれの待ち行列に入れるようにコマンド待ち行列インタフェース24に命令するタスク実行待ち行列内追加命令とデータ転送待ち行列内追加命令とを含むことができる。
コマンド待ち行列インタフェース24は、待ち行列32に格納されたコマンドを実行するようにも構成される。データ転送コマンドに関しては、コマンド待ち行列インタフェース24は、ホストメモリスペースとデバイスメモリスペースとの間でデータを転送することができる。例えば、書き込みバッファコマンドに関しては、コマンド待ち行列インタフェース24は、ホストメモリスペースの一部分に格納されたデータをデバイスメモリスペース内で割り当てられたバッファオブジェクトに転送することができる。他の例として、読み取りバッファコマンドに関しては、コマンド待ち行列インタフェース24は、デバイスメモリスペース内で割り当てられたバッファオブジェクトに格納されたデータをホストメモリスペースの一部分に転送することができる。デバイスメモリスペースは、コマンド待ち行列32が関連付けられているデバイスに対応することができる。
タスク実行コマンドに関しては、コマンド待ち行列インタフェース24は、タスクの実行をコマンド待ち行列と関連付けられたデバイスで開始させることができる。例えば、図1の例では、コマンド待ち行列32は、ランタイムモジュール22のコンテキスト内でGPU14と関連付けられる。従って、タスク実行コマンドを実行するときには、コマンド待ち行列インタフェース24は、GPU14での実行をタスクに開始させることができる。幾つかの例では、コマンド待ち行列インタフェース24は、GPU14内に入っているローカルコマンド待ち行列内に1つ以上のコマンドを入れることによってGPU14での実行をタスクに開始させることができる。その他の例では、コマンド待ち行列インタフェース24は、タスクの実行を開始するようにGPU14に命令する1つ以上の命令をGPU14に送信することによってGPU14での実行をタスクに開始させることができる。コマンド待ち行列インタフェース24は、GPU14、メモリ16、ホストメモリスペース及びデバイスメモリスペースと通信するために相互接続ネットワーク18を用いることができる。
幾つかの例では、コマンド待ち行列インタフェース24は、コマンドを順に実行することができる。該例では、第2のコマンドの前に第1のコマンドが待ち行列に入れられた場合は、第2のコマンドの実行は、第1のコマンドが実行を完了した後に開始する。さらなる例では、コマンド待ち行列インタフェース24は、順不同でコマンドを実行することができる。該例では、第2のコマンドの前に第1のコマンドが待ち行列に入れられた場合でも、第2のコマンドの実行は、必ずしも第1のコマンドが実行を完了した後に開始するわけではない。
ホストメッセージ渡しインタフェース26は、ホストプロセス20から受信された1つ以上のメッセージ渡し命令を実行するように構成される。例えば、ホストプロセス20から1つ以上のメッセージ渡し命令を受信したことに応答して、ホストメッセージ渡しインタフェース26は、GPU14でタスク28が実行している間にホストデバイス12で実行中のホストプロセス20とGPU14で実行中のタスク28との間で1つ以上のメッセージを渡すことができる。幾つかの例では、ホストメッセージ渡しインタフェース26は、コマンド待ち行列32内にコマンドを入れずに1つ以上のメッセージ渡し命令を実行することができる。
第1の例により、ホストプロセス20から送信命令を受信したことに応答して、ホストメッセージ渡しインタフェース26は、GPU14でタスク28が実行している間にホストプロセス20からタスク28にメッセージを送信することができる。例えば、ホストメッセージ渡しインタフェース26は、送信命令内に含まれたメッセージデータに基づいて発信メッセージを構成し、指定されたデバイスで実行中のタスク、例えば、タスク28、への引き渡しのために、相互接続ネットワーク18を介して、送信命令内で指定されたデバイス、例えば、GPU14、に発信メッセージを転送することができる。
第2の例により、ホストプロセス20からレジスタコールバックルーチン命令を受信したことに応答して、ホストメッセージ渡しインタフェース26は、命令内で指定されたコールバックルーチンを、指定されたデバイスで実行中のタスク、例えば、タスク28、がメッセージを送信していることを示す、命令内で指定されたデバイス、例えば、GPU14、からの信号と関連付けることができる。幾つかの例では、デバイスからの信号は、割り込み信号であることができる。割り込み信号は、幾つかの例では、専用の割り込み信号ラインを介して引き渡すことができる。デバイスで実行中のタスクがメッセージを送信していることを示す信号を指定されたデバイスから受信したことに応答して、ホストメッセージ渡しインタフェース26は、レジスタコールバックルーチン命令内で指定されたコールバックルーチンの実行を開始することができる。コールバックルーチンは、指定されたデバイス、例えば、GPU14、からタスク、例えば、タスク28、によって送信されたメッセージを入手し、さらなる処理のためにホストプロセス20にメッセージを戻すことができる。
第3の例により、ポーリング命令を受信したことに応答して、ホストメッセージ渡しインタフェース26は、メッセージ状態情報に関して、命令内で指定されたデバイス、例えば、GPU14、をポーリングすることができる。ホストメッセージ渡しインタフェース26は、デバイスをポーリングするために相互接続ネットワーク18又は他のハードウェアに基づく通信経路を用いることができる。指定されたデバイス、例えば、GPU14、で実行中のタスク、例えば、タスク28、がメッセージを送信していることをメッセージ状態情報が示す場合は、ホストメッセージ渡しインタフェース26は、指定されたデバイスからメッセージを入手し、さらなる処理のためにホストプロセス20にメッセージを戻すことができる。
デバイスメッセージ渡しインタフェース30は、タスク28から受信された1つ以上のデバイスメッセージ渡し命令を実行するように構成される。例えば、タスク28から1つ以上のデバイスメッセージ渡し命令を受信したことに応答して、デバイスメッセージ渡しインタフェース30は、GPU14でタスク28が実行している間にGPU14で実行中のタスク28とホストデバイス12で実行中のホストプロセス20との間で1つ以上のメッセージを渡すことができる。
第1の例により、送信命令を受信したことに応答して、デバイスメッセージ渡しインタフェース30は、GPU14で実行中のタスク28からホストデバイス12で実行中のホストプロセス20にメッセージを送信することができる。例えば、デバイスメッセージ渡しインタフェース30は、送信命令内に含まれるメッセージデータに基づいて発信メッセージを構成し、ホストプロセス20への引き渡しのために、相互接続ネットワーク18を介して、ホストデバイス12に発信メッセージを転送することができる。
第2の例により、タスク28から受信命令を受信したことに応答して、デバイスメッセージ渡しインタフェース30は、ホストプロセス20からのメッセージを入手可能であるどうかを決定することができる。幾つかの例では、デバイスメッセージ渡しインタフェース30は、1つ以上のホストがアクセス可能なレジスタを検査してメッセージが入手可能であるかどうかを決定することができる。ホストプロセス20からのメッセージを入手可能である場合は、デバイスメッセージ渡しインタフェース30は、タスク28にメッセージを提供することができる。
コマンド待ち行列インタフェース24及びホストメッセージ渡しインタフェース26は、図1のホストプロセス20とは別個のコンポーネントとして例示されるが、幾つかの例では、コマンド待ち行列インタフェース24及びホストメッセージ渡しインタフェース26のうちの1つ又は両方の機能は、部分的に及び/又は完全にコンパイルしてホストプロセス20に入れることができる。同様に、幾つかの例では、デバイスメッセージ渡しインタフェース30の機能は、部分的に及び/又は完全にコンパイルしてタスク28内に入れることができる。
例示を容易にするために、図1において示される計算システム例10は、GPU14を計算デバイスとして使用する本開示のメッセージ渡し技法について説明する。しかしながら、本開示の技法は、GPU14に加えての又はGPU14に代わるGPU以外の計算デバイスを有するマルチプルプロセッサ計算システムに対して適用可能であることが認識されるべきである。幾つかの例では、計算デバイスは、OpenCL計算デバイスであることができる。OpenCL計算デバイスは、1つ以上の計算ユニットを含む。計算ユニットの各々は、1つ以上の処理素子を含む。例えば、計算ユニットは、計算ユニット内の全処理素子によって用いることができるオンチップ共有メモリを有する処理素子、例えば、ALU、の集まりであることができる。ワークアイテム(work item)は、コマンド待ち行列内に入れられたコマンドによってOpenCL計算デバイスで呼び出されたカーネル又はタスクの複数の並行実行のうちの1つであることができる。各ワークアイテムは、その他の処理素子で実行中のその他のワークアイテムと並行して計算ユニット内の個々の処理素子で実行することができる。ワークグループは、単一のカーネル実行コマンドの一部として計算デバイス内の単一の計算ユニットにおいて処理される1つ以上のワークアイテムの集合であることができる。OpenCLホストは、OpenCLランタイムレイヤを実行するために用いられるプラットフォームの中央CPUであることができる。
OpenCL APIは、ホストデバイスと異なるタイプの計算デバイスとの間での対話のための共通のインタフェースの組を提供することができる。例えば、OpenCL APIは、ホストとGPU計算デバイス及びホストと非GPU計算デバイスとの間の対話のための共通のインタフェースを提供することができる。OpenCL APIは、様々な計算デバイスでタスク(例えば、
OpenCLカーネル)を実行するために共通のインタフェースをホストが使用するのを可能にする。幾つかの例では、タスクは、汎用計算タスクであることができ、OpenCL APIは、ホストがGPU計算デバイスで汎用計算タスクを実行させることを可能にすることができる。
図1に示される計算システム例10は、ホストデバイスと計算デバイスとの間でのメッセージ渡し及び/又はアウトオブバンドシグナリングを容易にするためのインフラストラクチャ及び技法を例示する。しかしながら、その他の計算システム例では、それらの技法は、2つ以上の計算デバイスを有する計算システム内の異なる計算デバイス(例えば、OpenCL計算デバイス)間でのインフライトのメッセージ渡しを提供するために容易に拡張することができる。該例では、異なる計算デバイス間で1本以上の割り込みラインを配線することができる。
図2は、本開示による図1の計算システム10において使用することができるGPU例40を示したブロック図である。幾つかの例では、図1において例示されたGPU14を実装するためにGPU40を用いることができる。GPU40は、GPU処理ブロック42と、ホストがアクセス可能なGPUレジスタ44と、バスコントローラ46と、を含む。GPU40は、相互接続ネットワーク18を介して1つ以上のその他のホストデバイス又は計算デバイスと通信可能な形で結合させることができる。
GPU処理ブロック42は、タスクを実行するように及びGPU処理ブロック42で実行中のタスクとその他のホスト又は計算デバイスで実行中のプロセスとの間でのメッセージ渡しを容易にするように構成される。GPU処理ブロック42は、例えば、1本以上の制御及び/又はデータラインを介して、ホストがアクセス可能なGPUレジスタ44に通信可能な形で結合される。幾つかの例では、GPU処理ブロック42は、算術論理装置(ALU)ブロックと呼ぶことができる。GPU処理ブロック42は、タスク48と、メッセージ渡しモジュール50と、着信データレジスタ52と、発信データレジスタ54と、を含む。
ホストがアクセス可能なGPUレジスタ44は、ホストデバイスに又はホストデバイスから通信することができるデータを格納するように構成される。ホストがアクセス可能なGPUレジスタ44は、メッセージ状態レジスタ56と、メッセージカウントレジスタ58と、着信メッセージレジスタ60と、発信メッセージレジスタ62と、割り込み状態レジスタ64と、割り込み確認応答レジスタ66と、を含む。ホストがアクセス可能なGPUレジスタ44の各々は、ホストデバイス、例えば、図1のホストデバイス12、によってアクセス可能であることができる。幾つかの例では、ホストがアクセス可能なGPUレジスタ44は、メモリマッピングされたレジスタ、すなわち、ホストデバイスのメモリスペースにマッピングされてアドレス指定が可能なレジスタ、であることができる。さらなる例では、ホストがアクセス可能なGPUレジスタ44は、入力/出力マッピングされた(I/Oマッピングされた)レジスタ、すなわち、ホストデバイスのI/Oスペースにマッピングされたレジスタ、であることができる。ホストがアクセス可能なGPUレジスタ44は、1本以上の制御及び/又はデータラインを介してGPU処理ブロック42に通信可能な形で結合される。ホストがアクセス可能なGPUレジスタ44は、相互接続ネットワーク18を介してバスコントローラ46にも通信可能な形で結合される。
タスク48は、1つ以上のプログラミング可能なプロセッサで実行することができる。幾つかの例では、GPU処理ブロック42は、タスク48の複数の実行インスタンスを実行するように構成された複数のプロセッサ又は処理素子を含むことができる。タスク48は、図1に関して上述されるタスク28と実質的に類似することができ、従って、さらに詳細には説明されない。
メッセージ渡しモジュール50は、GPU40によって行われるメッセージ渡し動作を制御するように構成される。メッセージ渡しモジュール50は、ハードウェア、ソフトウェア、ファームウェア又はそれらのあらゆる組み合わせ内に実装することができる。幾つかの例では、メッセージ渡しモジュール50の機能の一部又は全部がソフトウェア内に実装される場合は、該実装のためのソフトウェア命令は、タスク48のためのソフトウェア命令を含む実行可能ファイルと同じ実行可能ファイル内に含めることができる。メッセージ渡しモジュール50は、タスク48、メッセージ渡しモジュール50、着信データレジスタ52及び発信データレジスタ54に通信可能な形で結合される。
メッセージ渡しモジュール50は、1つ以上のプロセッサでタスク48が実行している間に及びタスク48から1つ以上のメッセージ渡し命令を受信したことに応答してそれらの1つ以上のプロセッサで実行中のタスク48とホストデバイスで実行中のプロセスとの間で、ホストがアクセス可能なGPUレジスタ44を介して、1つ以上のメッセージを渡すことができる。幾つかの例では、1つ以上のメッセージ渡し命令は、タスク48からホストデバイスで実行中のプロセスにメッセージを送信するようにメッセージ渡しモジュール50に命令する送信命令を含むことができる。該例では、メッセージ渡しモジュール50は、ホストがアクセス可能なGPUレジスタ44のうちの1つ内にメッセージと関連付けられたメッセージデータを格納することができる。さらなる例では、1つ以上のメッセージ渡し命令は、入手可能な場合にホストデバイスで実行中のプロセスからタスク48に送信されたメッセージをタスク48に提供するようにメッセージ渡しモジュール50に命令する受信命令を含むことができる。該例では、メッセージ渡しモジュール50は、ホストがアクセス可能なGPUレジスタ44のうちの1つ以上からメッセージと関連付けられたメッセージデータを入手することができる。
図2の例における着信データレジスタ52は、着信メッセージレジスタ60を介して外部のデバイスから受信された着信データを格納するハードウェアレジスタである。着信データレジスタ52は、着信データレジスタ52内のデータが消費されているかどうか及び/又は着信データレジスタ52内のデータが読み取りのために入手可能であるかどうかを示す状態ビットも格納することができる。着信データレジスタ52は、1本以上のデータラインを介して着信メッセージレジスタ60と通信可能な形で結合される。幾つかの例では、データライン数は、着信データレジスタ52内のビット数と等しいことができ、それらの両方とも、メッセージ内のビット数と等しいことができる。さらなる例では、ビット数は、32ビットであることができる。幾つかの例では、GPU処理ブロック42は、着信データレジスタ52から受信された複数の着信メッセージを格納するための内部の先入れ先出し(FIFO)バッファを実装することができる。
図2の例における発信データレジスタ54は、タスク48によって出された1つ以上のメッセージ渡し命令から受信された発信データを格納するハードウェアレジスタである。発信データレジスタ54は、1本以上のデータラインを介して発信メッセージレジスタ62と通信可能な形で結合される。幾つかの例では、データライン数は、発信データレジスタ54内のビット数と等しいことができ、それらの両方とも、メッセージ内のビット数と等しいことができる。幾つかの例では、発信データレジスタ54及び発信メッセージレジスタ62は、メッセージ渡しモジュール50が発信データレジスタ54にデータを書き込んだときに、発信メッセージレジスタ62が発信データレジスタ54に書き込まれたデータによって自動的に更新されるような形で構成することができる。幾つかの例では、GPU処理ブロック42は、発信データレジスタ54に書き込まれるべき複数の発信メッセージを格納するための内部の先入れ先出し(FIFO)バッファを実装することができる。
図2の例におけるメッセージ状態レジスタ56は、着信メッセージがGPU40によって受け入れられたかどうかを示すデータを格納するように構成される。メッセージ状態レジスタ56は、メッセージが成功裏に送信されたかどうかを決定するために、そして幾つかの例では、バックオフ(back−off)及び/又はオーバーフローメカニズムを実装するためにホストデバイスによって用いることができる。着信メモリを受け入れた後は、メッセージ渡しモジュール50は、着信メッセージが受け入れられたことを示す特定の値にメッセージ状態レジスタ56を設定することができる。
図2の例におけるメッセージカウントレジスタ58は、着信メッセージレジスタ60に着信メッセージが入っているかどうかを示すデータを格納するように構成される。幾つかの例では、メッセージカウントレジスタ58は、メッセージカウントレジスタ58がホストデバイスによって増分されたときにメッセージの到着を示すための信号をメッセージ渡しモジュール50に送信することができる。幾つかの事例では、信号は、1ビットパルスラインであることができる。さらなる例では、メッセージ渡しモジュール50は、着信データレジスタ52からメッセージを読み取った後にメッセージカウントレジスタ58を減分することができる。
図2の例における着信データレジスタ60は、着信メッセージデータを格納するように構成される。例えば、ホストデバイスは、タスク48にメッセージを送信するために着信メッセージデータを着信メッセージレジスタ60内に入れることができる。着信メッセージレジスタ60は、着信データレジスタ52と通信可能な形で結合される。
図2の例における発信メッセージレジスタ62は、発信データレジスタ54から受信された発信メッセージデータを格納するように構成される。発信メッセージレジスタ62は、発信データレジスタ54に新しいデータが書き込まれたときに発信データレジスタ54に対応するために発信メッセージレジスタ62内のデータを自動的に更新することができる。幾つかの例では、メッセージ渡しモジュール50は、発信メッセージレジスタ62に発信メッセージが書き込まれたことに応答して割り込み信号を生成することができる。割り込み信号は、ホストデバイスに送信し、メッセージ渡しモジュール50がメッセージを送信していることを示すことができる。
図2の例における割り込み状態レジスタ64は、発信メッセージレジスタ62に発信メッセージが書き込まれているかどうかを示す状態ビットを格納するように構成される。例えば、割り込み状態レジスタ64及び発信メッセージレジスタ62は、発信メッセージレジスタ62に発信メッセージが書き込まれたときに割り込み状態レジスタ64内の状態ビットが設定されるように構成することができる。状態ビットは、ホストデバイスで実行中のプロセスがメッセージを入手可能であるかどうかを確認するためにGPU40をポーリングするのを可能にすることができる。
図2の例における割り込み確認応答レジスタ66は、ホストデバイスが発信メッセージレジスタ62に格納された発信メッセージを読み取ったかどうかを示す確認応答ビットを格納するように構成される。例えば、発信メッセージレジスタ62及び割り込み確認応答レジスタ66は、発信メッセージレジスタ62に発信メッセージが書き込まれたときに割り込み確認応答レジスタ66内の確認応答ビットが設定されるような形で構成することができる。該例では、ホストデバイスが発信メッセージレジスタ62を読み取った後に、ホストデバイスは、確認応答ビットをクリアし、それにより、ホストデバイスが発信メッセージを読み取っており新しい発信メッセージを発信メッセージレジスタ62に書き込むことができることを示すことができる。確認応答ビットは、発信メッセージデータのための流れ制御方式を実装するために用いることができる。
図2の例におけるバスコントローラ46は、外部のデバイスが相互接続ネットワーク18を介してホストがアクセス可能なGPUレジスタ44にアクセスするのを可能にするように構成される。例えば、バスコントローラ46は、バス信号を多重化及び多重解除し、バス信号によって指定された様々な受信及び送信動作を行うことができる。バスコントローラ46は、1つ以上の公的な又は独占的なバス規格により動作することができる。
今度は、マルチプルプロセッサ計算システム内でのメッセージ渡しのための様々な技法が、本開示の幾つかの態様により説明される。幾つかの例では、図1の計算システム10は、図3乃至19に示される技法例を実装するために用いることができる。説明を容易にするために、それらの技法は、図1に示される計算システム例10のコンポーネントに関して説明されるが、それらの技法は、同じ又は異なる構成内の同じ又は異なるコンポーネントを有するその他のシステムで実行できることが理解されるべきである。追加の例では、図3乃至19に示される技法の一部は、図2のGPU40の特定のコンポーネントに関して説明することができる。繰り返すと、図2は、本開示の技法を実装することができるGPUの一例にすぎないこと、及び該技法は、同じ又は異なる構成内の同じ又は異なるコンポーネントを有するその他のGPUによって実行できることが理解されるべきである。
図3は、本開示によるマルチプルプロセッサプラットフォーム環境におけるメッセージ渡しのための技法例を示す。幾つかの例では、図3に示される技法例を実装するために図1の計算システム10を用いることができる。コマンド待ち行列インタフェース24は、コマンド待ち行列32内にメモリ転送コマンドを入れる(70)。コマンド待ち行列インタフェース24は、コマンド待ち行列32内にタスク実行コマンドを入れる(72)。コマンド待ち行列インタフェース24は、GPU14でのタスクの実行を開始するためにタスク実行コマンドを実行する(74)。ホストメッセージ渡しインタフェース26は、GPU14でタスク28が実行している間にホストデバイス12とGPU14との間で1つ以上のメッセージを渡す(76)。例えば、ホストメッセージ渡しインタフェース26は、ホストプロセス20によって出された1つ以上の送信命令から生じるメッセージをGPU14に渡すことができる。1つ以上の送信命令は、GPU14又はGPU14で実行中のタスクがメッセージの行先であることを指定することができる。
図4は、本開示によるホストデバイスで実行中のプロセスによって出された送信命令を実行するための技法例である。幾つかの例では、図4に示される技法例を実装するために図1の計算システム10を用いることができる。ホストメッセージ渡しインタフェース26は、ホストプロセス20から送信命令を受信する(78)。ホストメッセージ渡しインタフェース26は、送信命令とともに含まれるメッセージデータに基づいて発信メッセージを生成する(80)。幾つかの例では、発信メッセージは、送信命令に含まれるメッセージデータと同一であることができる。追加の例では、ホストメッセージ渡しインタフェース26は、発信メッセージを生成するために送信命令に含まれるメッセージデータに1つ以上のヘッダ情報及び/又はルーティング情報を添付することができる。さらなる例では、ホストメッセージ渡しインタフェース26は、発信メッセージを生成するために送信命令に含まれるメッセージデータに対して1つ以上のコーディング又は変換動作を行うことができる。ホストメッセージ渡しインタフェース26は、GPU14に発信メッセージを送信することができる(82)。
ホストメッセージ渡しインタフェース26は、送信命令がブロッキング命令であるか又は非ブロッキング命令であるかを決定することができる(84)。幾つかの例では、ホストメッセージ渡しインタフェース26は、送信命令において指定された入力パラメータに基づいて送信命令がブロッキング命令であるか又は非ブロッキング命令であるかの決定を行うことができる。その他の例では、2つの異なるタイプの送信命令を用いることができ、及び命令のタイプ、例えば、命令の演算コード(オプコード)に基づいて命令がブロッキング命令であるか又は非ブロッキング命令であるかの決定を行うことができる。送信命令が非ブロッキング命令であるとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、呼び出しを行っているプロセスにハンドルを戻すことができる(86)。ハンドルは、呼び出しを行っているプロセスが、メッセージが成功裏に送信されているかどうかをのちに決定するためにハンドルに問い合わせるのを可能にすることができる。送信が失敗であったことを後続する問い合わせが示した場合は、呼び出しを行っているプロセスは、送信動作を再試行するための後続する送信命令を出す必要がある場合がある。幾つかの例では、呼び出しを行っているプロセスは、失敗した送信動作に応答してバックオフルーチン又はオーバーフローメカニズムを実装することができる。
送信命令がブロッキング命令であるとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、発信メッセージがGPU14によって成功裏に受信されたかどうかを決定することができる(88)。発信メッセージが成功裏に受信されたとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、送信命令に含まれているメッセージが成功裏に送信されたことを示す値を呼び出しを行っているプロセスに戻すことができる(90)。そうでない場合、発信メッセージが成功裏に受信されなかったとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、プロセスブロック82に進み、GPU14に発信メッセージを再送信することができる。ブロッキング命令は、幾つかの例では、メッセージが成功裏に受信されたか又は不成功の引き渡しの試行のスレショルド数に達しているとホストメッセージ渡しインタフェース26が決定したときには、完了することができる。
図5は、本開示による図4のプロセスブロック82を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図5に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。ホストメッセージ渡しインタフェース26は、GPU40の着信メッセージレジスタ60に発信メッセージを入れる又は格納することができる(92)。ホストメッセージ渡しインタフェース26は、新しいメッセージが到着していることをGPU14内のメッセージ渡しモジュール50に示すためにGPU40のメッセージカウントレジスタ58を増分することができる。幾つかの例では、ホストメッセージ渡しインタフェース26は、プロセスブロック92及び94のうちの1つ以上を実行するために当業において知られるメモリマッピングされたレジスタハードウェア及び/又はI/Oマッピングされたレジスタハードウェアを用いることができる。
図6は、本開示による図4の判断ブロック88を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図6に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。ホストメッセージ渡しインタフェース26は、GPU40のメッセージ状態レジスタ56内の状態ビットを検査することができる(96)。ホストメッセージ渡しインタフェース26は、メッセージ状態レジスタ56内の状態ビットに基づいて送信されたメッセージがGPU14によって受け入れられたかどうかを決定することができる(98)。送信されたメッセージがGPU14によって受け入れられたことを状態ビットが示す場合は、ホストメッセージ渡しインタフェース26は、発信メッセージが成功裏に受信されたと決定することができる(100)。他方、送信されたメッセージがGPU14によって受け入れられなかったことを状態ビットが示す場合は、ホストメッセージ渡しインタフェース26は、発信メッセージが成功裏に受信されなかったと決定することができる(102)。
図7は、計算デバイス、例えば、GPU、において受信されたメッセージを処理するための技法例を示した流れ図である。幾つかの例では、図7に示される技法例を実装するために図2のGPU40を用いることができる。GPU40内のメッセージ渡しモジュール50がメッセージ到着信号を受信する(104)。例えば、メッセージカウントレジスタ58は、ホストデバイスがメッセージカウントレジスタ58を増分するごとにメッセージ到着パルスがメッセージ渡しモジュール50に送信されるような形で構成することができる。メッセージ渡しモジュール50は、着信メッセージレジスタ60に格納されたデータを着信データレジスタ52に転送させることができる(106)。例えば、メッセージ渡しモジュール50は、着信データレジスタ52に制御信号を出し、着信メッセージレジスタ60に格納されたデータを着信データレジスタ52に格納された現在のデータに上書きするのを着信データレジスタ52に行わせることができる。メッセージ渡しモジュール50は、着信データレジスタ52内でデータを入手可能である、例えば、消費されていない、ことを示すように着信データレジスタ52内の状態ビットを設定することができる(108)。メッセージ渡しモジュール50は、着信メッセージがGPU40によって受け入れられていることを示すためにメッセージ状態レジスタ56内の状態ビットを設定することができる(110)
図8は、本開示による計算デバイスで実行中のタスクによって出された受信命令を実行するための技法例を示した流れ図である。幾つかの例では、図8に示される技法例を実装するために図1の計算デバイス10を用いることができる。デバイスメッセージ渡しインタフェース30が、タスク28から受信命令を受信する(112)。デバイスメッセージ渡しインタフェース30は、ホストデバイスからメッセージを入手可能であるかどうかを決定する(114)。
メッセージが入手可能でないとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、受信命令がブロッキング受信命令であるか又は非ブロッキング受信命令であるかを決定することができる(116)。幾つかの例では、メッセージ渡しモジュール50は、受信命令内で指定された入力パラメータに基づいて受信命令がブロッキング命令であるか又は非ブロッキング命令であるかの決定を行うことができる。その他の例では、2つの異なるタイプの受信命令を用いることができ、メッセージ渡しモジュール50は、命令のタイプ、例えば、命令の演算コード(オプコード)、に基づいて受信命令がブロッキング命令であるか又は非ブロッキング命令であるかの決定を行うことができる。受信命令がブロッキング命令であるとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、判断ブロック114に戻って着信メッセージを入手可能であるかどうかを決定することができる。そうでない場合、命令が非ブロッキング命令であるとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、受信命令が失敗したことを示す値を呼び出しを行っているプロセスに戻すことができる(118)。
ホストデバイスからメッセージを入手可能であるとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、呼び出しを行っているプロセスにメッセージデータを戻すことができる(120)。メッセージ渡しモジュール50は、メッセージデータが消費されていると表示すべきであるかどうかを決定する(122)。メッセージ渡しモジュール50は、1つ以上の消費モードに基づいてデータを消費されていると表示すべきかどうかを決定することができる。幾つかの例では、消費モードは、GPU14内にハードワイヤすることができる。追加の例では、消費モードは、タスク28及び/又はホストプロセス20のいずれかによってプログラミングすることができる。例えば、タスク28又はホストプロセス20における送信及び/又は受信命令は、特定の消費モードを指定するパラメータを入れることができる。例えば、一消費モードは、タスクの少なくとも1つの実行インスタンスがデータを読み取ったときにメッセージデータを消費されたと表示すべきであると指定することができる。他の例として、一消費モードは、タスクの少なくともスレショルド数の実行インスタンスがデータを読み取ったときにメッセージデータを消費されたと表示すべきであると指定することができる。
メッセージデータが消費されたと表示すべきであるとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、メッセージデータをクリアすることができる(124)。例えば、メッセージ渡しモジュール50は、着信データレジスタ52内の状態ビットをクリアすることができる。他方、メッセージデータが消費されたと表示すべきでないとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、メッセージデータを保持することができる(126)。例えば、メッセージ渡しモジュール50は、着信データレジスタ52内の状態ビットをクリアすることができない。
図9は、本開示による図8の判断ブロック114を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図9に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。メッセージ渡しモジュール50は、GPU40の着信データレジスタ52内の状態ビットを読み取ることができる(128)。メッセージ渡しモジュール50は、状態ビットが設定されているかどうかを決定することができる(130)。着信データレジスタ52内の状態ビットが設定されている場合は、メッセージ渡しモジュール50は、着信メッセージが入手可能であると決定することができる(132)。他方、着信データレジスタ52内の状態ビットが設定されていない場合は、メッセージ渡しモジュール50は、着信メッセージを入手できないと決定することができる(134)。
図10は、本開示による図8のプロセスブロック120を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図10に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。メッセージ渡しモジュール50は、GPU40内の着信データレジスタ52から着信メッセージデータを取り出すことができる(136)。メッセージ渡しモジュール50は、着信データレジスタ52から取り出されたメッセージデータに基づいてタスク48のための戻りメッセージデータを生成することができる(138)。幾つかの例では、戻されるメッセージデータは、着信データレジスタ52に入ったメッセージデータと同一であることができる。追加の例では、メッセージ渡しモジュール50は、戻りメッセージデータを生成するために着信データレジスタ52に入っているメッセージデータから1つ以上のヘッダ情報及び/又はルーティング情報を取り除くことができる。さらなる例では、メッセージ渡しモジュール50は、戻りメッセージデータを生成するために着信データレジスタ52に入っているメッセージデータに対して1つ以上の復号動作又は変換動作を行うことができる。メッセージ渡しモジュール50は、タスク48にメッセージデータを提供する(140)。
図11は、本開示による計算デバイス、例えば、GPU14、で実行中のプロセスによって出された送信命令を実行するための技法例である。幾つかの例では、図11に示される技法例を実装するために図1の計算システム10を用いることができる。メッセージ渡しモジュール50は、タスク28から送信命令を受信する(142)。メッセージ渡しモジュール50は、送信命令とともに含まれているメッセージデータに基づいて発信メッセージを生成する(144)。幾つかの例では、発信メッセージは、送信命令に入っているメッセージデータと同一であることができる。追加の例では、メッセージ渡しモジュール50は、発信メッセージを生成するために送信命令に入っているメッセージデータに1つ以上のヘッダ情報及び/又はルーティング情報を添付することができる。さらなる例では、メッセージ渡しモジュール50は、発信メッセージを生成するために送信命令に入っているメッセージデータに対して1つ以上のコーディング動作又は変換動作を行うことができる。メッセージ渡しモジュール50は、ホストデバイス12に発信メッセージを送信することができる(146)。
メッセージ渡しモジュール50は、送信命令がブロッキング命令であるか又は非ブロッキング命令であるかを決定することができる(148)。幾つかの例では、メッセージ渡しモジュール50は、送信命令において指定された入力パラメータに基づいて送信命令がブロッキング命令であるか又は非ブロッキング命令であるかの決定を行うことができる。その他の例では、2つの異なるタイプの送信命令を用いることができ、メッセージ渡しモジュール50は、命令のタイプ、例えば、命令の演算コード(オプコード)、に基づいて送信命令がブロッキング命令であるか又は非ブロッキング命令であるかの決定を行うことができる。送信命令が非ブロッキング命令であるとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、呼び出しを行っているプロセス、例えば、タスク28、にハンドルを戻すことができる(150)。ハンドルは、呼び出しを行っているプロセスが、メッセージが成功裏に送信されているかどうかをのちに決定するためにハンドルに問い合わせるのを可能にすることができる。送信動作が失敗であったことを後続する問い合わせが示した場合は、呼び出しを行っているプロセスは、送信動作を再試行するための後続する送信命令を出す必要があることがある。
送信命令がブロッキング命令であるとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、発信メッセージがホストデバイス12によって成功裏に受信されたかどうかを決定することができる(152)。例えば、メッセージ渡しモジュール50は、メッセージが受け入れられたことを示すホストデバイス12に入った状態レジスタをポーリングすることができる。発信メッセージが成功裏に受信されたとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、送信命令に入ったメッセージが成功裏に受信されたことを示す値を呼び出しを行っているプロセスに戻すことができる(154)。そうでない場合、発信メッセージが成功裏に受信されなかったとメッセージ渡しモジュール50が決定した場合は、メッセージ渡しモジュール50は、プロセスブロック146に進み、ホストデバイス12に発信メッセージを再送信することができる。ブロッキング命令は、幾つかの例では、メッセージが成功裏に受信されたか又はスレショルド数の不成功引き渡し試行に達しているとメッセージ渡しモジュール50が決定したときに完了することができる。
図12は、本開示による図11のプロセスブロック146を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図12に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。メッセージ渡しモジュール50は、発信データレジスタ54に発信メッセージを入れる又は格納することができる(156)。発信メッセージレジスタ62は、新しいデータが発信データレジスタ54内に入れられたことに応答して発信データレジスタ54に対応するために発信メッセージレジスタ62内のデータを更新することができる(158)。メッセージ渡しモジュール50は、GPU40のタスク28からメッセージを入手可能であることを示す割り込み信号を生成してホストデバイス12に送信することができる(160)。
図13は、本開示による図11のプロセスブロック146を実装するために用いることができる他の技法例である。幾つかの例では、図13に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。メッセージ渡しモジュール50は、発信データレジスタ54に発信メッセージを入れる又は格納することができる(162)。発信メッセージレジスタ62は、新しいデータが発信データレジスタ54内に入れられたことに応答して発信データレジスタ54に対応するために発信メッセージレジスタ62内のデータを更新することができる(164)。メッセージ渡しモジュール50は、GPU40のタスク28からメッセージを入手可能であることを示すために割り込み状態レジスタ64内の状態ビットを設定することができる。状態ビットは、ホストデバイス12がメッセージを入手可能であるかどうかを決定するためにGPU40をポーリングするのを可能にするように設定することができる(166)。
図14は、本開示によるホストデバイスで実行中のプロセスによって出されたレジスタコールバックルーチン命令を実行するための技法例を示した流れ図である。幾つかの例では、図14に示される技法例を実装するために図1の計算システム10を用いることができる。ホストメッセージ渡しインタフェース26は、ホストプロセス20からレジスタコールバックルーチン命令を受信する(168)。ホストメッセージ渡しインタフェース26は、レジスタコールバックルーチン命令において指定されたコールバックルーチンを、命令内で指定されたデバイス、例えば、GPU14、からの割り込み信号と関連付ける(170)。幾つかの例では、割り込み信号は、指定されたデバイスで実行中のタスク、例えば、GPU14で実行中のタスク28、がメッセージを送信していることを示すことができる。割り込み信号は、幾つかの例では、ホストデバイス12とGPU14との間で結合された専用割り込み信号ラインを介して引き渡すことができる。さらなる例では、割り込み信号は、タスク28がメッセージを送信することに加えてのその他のイベントを示すことができる。該例では、ホストメッセージ渡しインタフェース26は、複数のイベントのうちのいずれが割り込み信号によって示されているかを決定するためにその割り込み信号を受信後に追加処理を行うことができる。
ホストメッセージ渡しインタフェース26は、コールバックルーチンが割り込み信号と成功裏に関連付けられたかどうかを決定する(172)。コールバックルーチンが割り込み信号と成功裏に関連付けられた場合は、ホストメッセージ渡しインタフェース26は、レジスタコールバックルーチン動作が成功裏に完了されたことを示す値を呼び出しを行っているプロセスに戻すことができる(174)。そうでない場合、コールバックルーチンが割り込み信号と成功裏に関連付けられなかった、例えば、エラーが発生した、場合は、ホストメッセージ渡しインタフェース26は、レジスタコールバックルーチン動作が失敗したことを示す値を呼び出しを行っているプロセスに戻すことができる(176)。
図15は、本開示による計算デバイスから受信された割り込みを処理するための技法例を示した流れ図である。幾つかの例では、図15に示される技法例を実装するために図1の計算システム10を用いることができる。ホストメッセージ渡しインタフェース26は、計算デバイス、例えば、GPU14、から割り込み信号を受信する(178)。ホストメッセージ渡しインタフェース26は、割り込み信号がメッセージ受信イベントに応答して送信されかどうかを決定する(180)。換言すると、ホストメッセージ渡しインタフェース26は、デバイスで実行中のタスク、例えば、GPU14で実行中のタスク28、がメッセージを送信していることを割り込み信号が示すかどうかを決定することができる。
幾つかの例では、割り込み信号は、メッセージ受信イベントをシグナリングしてその他のイベントはシグナリングしない専用割り込み信号であることができる。該例では、ホストメッセージ渡しインタフェース26は、割り込み信号自体を受信したことによってメッセージ受信イベントに応答して割り込み信号が送信され、その他の動作は必ずしも行う必要がないと決定することができる。割り込み信号が複数の潜在的イベントをシグナリングする例では、ホストメッセージ渡しインタフェース26は、いずれのイベントがシグナリングされたかを決定するために計算デバイスに問い合わせることが必要な場合がある。
割り込み信号がメッセージ受信イベントに応答して送信されたのではないとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、その他のタイプのイベントの存在を検査することができる(182)。そうでない場合、割り込み信号がメッセージ受信イベントに応答して送信されたとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、メッセージを送信したデバイスと関連付けられたコールバックルーチンを実行することができる(184)。
図16は、本開示による図15の判断ブロック180を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図16に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。ホストメッセージ渡しインタフェース26は、GPU40内の割り込み状態レジスタ64を読み取ることができる(186)ホストメッセージ渡しインタフェース26は、ホストデバイスのために新しいメッセージを入手可能であることを割り込み状態レジスタ64内の状態ビットが示すかどうかを決定することができる(188)。例えば、メッセージ渡しモジュール50は、メッセージが入手可能であるときには割り込み状態レジスタ64内の状態ビットを設定することができ、ホストメッセージ渡しインタフェース26は、ホストデバイスのために新しいメッセージを入手可能であるかどうかを決定するために状態ビットが設定されているかどうかを決定するために割り込み状態レジスタ64をポーリングすることができる。ホストデバイスのために新しいメッセージを入手可能であると状態ビットが示す場合は、ホストメッセージ渡しインタフェース26は、割り込み信号がメッセージ受信イベントに応答して送信されたと決定することができる(190)。他方、ホストデバイスのために新しいメッセージを入手可能でないと状態ビットが示す場合は、ホストメッセージ渡しインタフェース26は、割り込み信号がメッセージ受信イベントに応答して送信されたのではないと決定することができる(192)。
図17は、本開示による図15のプロセスブロック184を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図17に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。ホストメッセージ渡しインタフェース26は、GPU40内の発信メッセージレジスタ62からメッセージを取り出すことができる(194)。ホストメッセージ渡しインタフェース26は、割り込み確認応答レジスタ66内の確認応答ビットをクリアすることができる(196)。確認応答ビットをクリアすることは、GPU40の流れ制御において役立つことができる。例えば、GPU40は、発信メッセージレジスタ62に発信メッセージが書き込まれるときに割り込み確認応答レジスタ66内の確認応答ビットを設定し、発信メッセージレジスタ62に追加データを書き込む前に確認応答ビットがクリアされるまで待つことができる。
図18は、本開示によるホストデバイスで実行中のプロセスによって出された読み取り命令を実行するための技法例を示した流れ図である。幾つかの例では、図18に示される技法例を実装するために図1の計算システム10を用いることができる。ホストメッセージ渡しインタフェース26は、データが読み取られる特定のデバイスを指定する読み取り命令を受信する(198)。ホストメッセージ渡しインタフェース26は、読み取り命令で指定されたデバイスをポーリングする(200)。ホストメッセージ渡しインタフェース26は、ポーリング動作から受信されたポーリングデータに基づいて受信命令で指定されたデバイスからメッセージを入手可能であるかどうかを決定する(202)。受信命令で指定されたデバイスからメッセージを入手可能であるとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、読み取り命令で指定されたデバイスからメッセージを取り出すことができる(204)。幾つかの例では、ホストメッセージ渡しインタフェース26は、ホストデバイス12にアクセス可能なデバイス内のレジスタ、例えば、GPU40内の発信メッセージレジスタ62、からメッセージを取り出すことができる。ホストメッセージ渡しインタフェース26は、呼び出しを行っているプロセス、例えば、ホストプロセス20、にメッセージデータを戻すことができる(206)。受信命令で指定されたデバイスからメッセージを入手可能でないとホストメッセージ渡しインタフェース26が決定した場合は、ホストメッセージ渡しインタフェース26は、読み取り命令が失敗したことを示す値を戻すことができる(208)。呼び出しを行っているプロセスは、読み取り動作を再試行するために読み取り命令を再度出すことが必要になる場合がある。
図19は、本開示による図18の判断ブロック202を実装するために用いることができる技法例を示した流れ図である。幾つかの例では、図19に示される技法例を実装するために図1の計算システム10及び/又は図2のGPU40を用いることができる。ホストメッセージ渡しインタフェース26は、GPU40内の割り込み状態レジスタ64を読み取ることができる(210)。ホストメッセージ渡しインタフェース26は、ホストデバイスのために新しいメッセージを入手可能であることを割り込み状態レジスタ64内の状態ビットが示すかどうかを決定することができる(212)。例えば、メッセージ渡しモジュール50は、メッセージが入手可能なときは割り込み状態レジスタ64内の状態ビットを設定することができ、ホストメッセージ渡しインタフェース26は、ホストデバイスのために新しいメッセージが入手可能であるかどうかを決定するために状態ビットが設定されているかどうかを決定するために割り込み状態レジスタ64をポーリングすることができる。状態ビットが設定されている場合は、ホストメッセージ渡しインタフェース26は、メッセージを入手可能であると決定することができる(214)。他方、状態ビットが設定されていない場合は、ホストメッセージ渡しインタフェース26は、メッセージが入手可能でないと決定することができる(216)。
ホストメッセージ渡しインタフェース26及びデバイスメッセージ渡しインタフェース30によって実装されたメッセージ渡し技法は、ホストデバイス12とGPU14との間でのアウトオブバンドシグナリングを提供するものとして上述されているが、その他の例では、アウトオブバンドシグナリングを提供するためにその他の技法を用いることができる。例えば、幾つかの例では、特殊な高優先度待ち行列を定義することができ、アウトオブバンドメッセージを送信するためにそれを用いることができる。
図20は、本開示による即時メモリオブジェクトの使用を容易にすることができる計算システム例310を示したブロック図である。計算システム310は、複数の処理デバイスにおいて1つ以上のソフトウェアアプリケーションを処理するように構成される。幾つかの例では、1つ以上のアプリケーションは、ホストプロセスを含むことができ、計算システム310は、ホストプロセスを実行するように及び計算310内のその他の計算デバイス上でホストデバイスによって開始された1つ以上のタスクの実行を分散させるように構成することができる。さらなる例では、計算システム310によって実行されるホストプロセス及び/又はタスクは、並列プログラミングモデルによりプログラミングすることができる。例えば、アプリケーションは、基礎になるハードウェアシステムのタスクレベルでの並列性及び/又はデータレベルでの並列性を利用するように設計される命令を含むことができる。
計算システム310は、パソコン、デスクトップコンピュータ、ラップトップコンピュータ、コンピュータワークステーション、ビデオゲームプラットフォーム又はコンソール、移動電話、例えば、セルラー又は衛星電話、携帯電話、ランドライン電話、インターネット電話、ハンドヘルドデバイス、例えば、ポータブルビデオゲーム機又はパーソナルデジタルアシスタント(PDA)、デジタルメディアプレーヤー、例えば、パーソナル音楽プレーヤー、ビデオプレーヤー、表示装置、テレビ、テレビセットトップボックス、サーバ、中間ネットワークデバイス、メインフレームコンピュータ又は情報を処理するその他のあらゆるタイプのデバイスを備えることができる。
計算デバイス310は、ホストデバイス312と、GPU314と、メモリ316と、相互接続ネットワーク318と、を含む。ホストデバイス312は、マルチプルプロセッサ計算プラットフォームAPIのためのホストプロセス及びランタイムモジュールの実行のためのプラットフォームを提供するように構成される。典型的には、ホストデバイス312は、汎用CPUであるが、ホストデバイス12は、プログラムを実行することが可能なあらゆるタイプのデバイスであることができる。ホストデバイス12は、相互接続ネットワーク318を介してGPU314及びメモリ316に通信可能な形で結合される。ホストデバイス312は、ホストプロセス320と、ランタイムモジュール322と、ホストキャッシュ324と、ホストキャッシュ制御モジュール制御326と、を含む。ホストプロセス320及びランタイムモジュール322は、1つ以上のプログラミング可能なプロセッサのあらゆる組み合わせにおいて実行することができる。
ホストプロセス320は、計算システム310プラットフォームでの実行のためのソフトウェアプログラムを形成する命令の組を含む。ソフトウェアプログラムは、エンドユーザのために1つ以上の特定のタスクを実行するように設計することができる。該タスクは、幾つかの例では、計算システム310によって提供される複数の処理デバイス及び並列アーキテクチャを利用することができる計算集約型アルゴリズムを含むことができる。
ランタイムモジュール322は、ホストプロセス320に含まれる命令のうちの1つ以上にサービスを提供するように構成された1つ以上のインタフェースを実装するソフトウェアモジュールであることができる。ランタイムモジュール322によって実装されたインタフェースは、メモリバッファインタフェース328を含む。幾つかの例では、ランタイムモジュール322は、メモリバッファインタフェース328に加えて、図1に示されるコマンド待ち行列インタフェース24及び図1に示されるホストメッセージ渡しインタフェース26のうちの1つ以上を実装することができる。さらなる例では、ランタイムモジュール322は、本開示で説明されるインタフェースに加えて標準的なマルチプルプロセッサシステムAPI内に含まれる1つ以上のインタフェースを実装することができる。幾つかの例では、標準的なAPIは、異種計算プラットフォームAPI、プラットフォーム横断型API、売り主横断型API、並列プログラミングAPI、タスクレベル並列プログラミングAPI、及び/又はデータレベル並列プログラミングAPIであることができる。さらなる例では、標準的なAPIは、OpenCL APIであることができる。該例では、ランタイムモジュール322は、OpenCL仕様のうちの1つ以上に準拠するように設計することができる。追加の例では、ランタイムモジュール322は、ドライバプログラム、例えば、GPUドライバ、の一部として実装することができる。
メモリバッファインタフェース328は、ホストプロセス20から1つ以上のメモリオブジェクト生成命令を受信するように及び受信された命令によって指定された機能を実行するように構成される。幾つかの例では、メモリバッファインタフェース328は、既存の標準API、例えば、OpenCL API、の拡張として、実装することができる。追加の例では、コマンド待ち行列インタフェース24は、既存の標準API、例えば、OpenCL API、内に組み入れることができる。
ホストキャッシュ324は、ホストデバイス312内で実行中のプロセスによる使用のためのデータを格納するように構成される。幾つかの例では、ホストキャッシュ324に格納されたデータと関連付けられたメモリスペースは、メモリ316内のメモリスペースの一部分とオーバーラップすることができる。ホストキャッシュ324は、当業において知られるあらゆるタイプのキャッシュであることができる。例えば、ホストキャッシュ324は、キャッシュレベル(例えば、L1、L2、等)及び/又はマッピング方式(例えば、直接マッピング、完全関連付け、セット関連付け、等)のあらゆる組み合わせを含むことができる。ホストキャッシュ制御モジュール326は、ホストキャッシュ324の動作を制御するように構成される。
GPU314は、ホストデバイス312から受信された命令に応答して1つ以上のタスクを実行するように構成される。GPU314は、1つ以上のプログラミング可能なプロセッサ又は処理素子を含むあらゆるタイプのGPUであることができる。例えば、GPU314は、タスクのための複数の実行インスタンスを並行して実行するように構成される1つ以上のプログラマブルシェーダユニットを含むことができる。プログマブルシェーダユニットは、バーテックスシェーダユニット、フラグメントシェーダユニット、ジオメトリシェーダユニット及び/又は統合シェーダユニットを含むことができる。GPU314は、相互接続ネットワーク318を介してホストデバイス312及びメモリ316に通信可能な形で結合される。GPU314は、タスク330と、GPUキャッシュ332と、GPUキャッシュ制御モジュール334と、を含む。タスク330は、1つ以上のプログラミング可能な処理素子のあらゆる組み合わせにおいて実行することができる。
タスク330は、計算システム310内の計算デバイスでの実行のためのタスクを形成する命令の組を備える。幾つかの例では、タスク330のための命令の組は、ホストプロセス320において定義し、幾つかの事例では、ホストプロセス320に含まれる命令によってコンパイルすることができる。さらなる例では、タスク330は、GPU314で並行して実行中の複数の実行インスタンスを有するカーネルプログラムであることができる。該例では、ホストプロセス320は、カーネル実行インスタンスを実行するために各々の処理素子にカーネル実行インスタンスをマッピングするカーネル用のインデックススペースを定義することができ、GPU314は、そのカーネル用に定義されたインデックススペースによりタスク330のための複数のカーネル実行インスタンスを実行することができる。
GPUキャッシュ332は、GPU314内で実行中のタスクによる使用のためのデータを格納するように構成される。幾つかの例では、GPUキャッシュ332に格納されたデータと関連付けられたメモリスペースは、メモリ316内のメモリスペースの一部分とオーバーラップすることができる。GPUキャッシュ332は、当業において知られるあらゆるタイプのキャッシュであることができる。例えば、GPUキャッシュ332は、キャッシュレベル(例えば、L1、L2、等)及び/又はマッピング方式(例えば、直接マッピング、完全関連付け、セット関連付け、等)のあらゆる組み合わせを含むことができる。GPUキャッシュ制御モジュール334は、GPUキャッシュ332の動作を制御するように構成される。
メモリ316は、ホストデバイス312及びGPU314のうちの1つ又は両方による使用のためにデータを格納するように構成される。メモリ316は、1つ以上の揮発性又は非揮発性のメモリ又は記憶デバイス、例えば、ランダムアクセスメモリ(RAM)、スタティックRAM(SRAM)、ダイナミックRAM(DRAM)、読み取り専用メモリ(ROM)、消去可能プログラマブルROM(EPROM)、電気的に消去可能なプログラマブルROM(EEPROM)、フラッシュメモリ、磁気データ記憶媒体又は光学記憶媒体、のあらゆる組み合わせを含むことができる。メモリ316は、相互接続ネットワーク318を介してホストデバイス312及びGPU314に通信可能な形で結合される。メモリ316は、共有メモリスペース336を含む。共有メモリスペース336は、ホストデバイス312及びGPU314の両方によってアクセス可能なメモリスペースであることができる。
相互接続ネットワーク318は、ホストデバイス312、GPU314及びメモリ316の間での通信を容易にするように構成される。相互接続ネットワーク318は、当業において知られるあらゆるタイプの相互接続ネットワークであることができる。図20の計算システム例310では、相互接続ネットワーク318は、バスである。バスは、様々なバス構造、例えば、第3世代バス(例えば、HyperTransportバス又はInfiniBandバス)、第2世代バス(例えば、Advanced Graphics Portバス、Peripheral Component Interconnect Express(PCIe)バス、又はAdvanced eXentisible Interface(AXI)バス)、又はその他のタイプのバスのうちの1つ以上を含むことができる。相互接続ネットワーク318は、ホストデバイス312、GPU314及びメモリ316に結合される。
今度は、計算システム310内のコンポーネントの構造及び機能がさらに詳細に説明される。上述されるように、ホストプロセス320は、命令の組を含む。命令の組は、例えば、1つ以上のメモリオブジェクト生成命令を含むことができる。追加の例では、命令の組は、GPU14で実行されるタスク又はカーネルを指定する命令と、コマンド待ち行列を生成してそれらのコマンド待ち行列を特定のデバイスと関連付ける命令と、プログラムをコンパイル及びバインドする命令と、カーネルパラメータを設定する命令と、インデックススペースを定義する命令と、デバイスコンテキストを定義する命令と、待ち行列内追加命令と、メッセージ渡し命令と、ホストプロセス320によって提供される機能をサポートするその他の命令と、を含むことができる。
本開示により、ホストプロセス320は、メモリオブジェクトのために即時モードがイネーブルにされるかどうかを指定する命令に含まれる情報に基づいてメモリオブジェクトを生成するようにメモリバッファインタフェース328に命令する1つ以上のメモリオブジェクト生成命令をメモリバッファインタフェース328に出すことによってメモリバッファインタフェース328と対話することができる。ここにおいて用いられる場合において、メモリオブジェクトは、GPU314によってアクセス可能なメモリスペースの領域を表すソフトウェアオブジェクトを意味することができる。幾つかの例では、メモリスペースの領域は、ホストデバイス312によってもアクセス可能であることができる。メモリオブジェクトは、メモリオブジェクトと関連付けられたメモリスペース内に入ったデータを含むことができる。メモリオブジェクトは、メモリスペースと関連付けられた1つ以上の特徴をさらに含むことができる。幾つかの例では、メモリオブジェクトは、グローバルメモリ、例えば、メモリ316、の基準のカウントされた領域へのハンドルを含むことができる。
メモリオブジェクトは、バッファオブジェクトと画像オブジェクトと、を含むことができる。バッファオブジェクトは、一次元のバイトの集合を格納するメモリオブジェクトであることができる。一次元のバイトの集合は、メモリオブジェクトと関連付けられたデータであることができる。バッファオブジェクトは、情報、例えば、バイトが単位のバッファオブジェクトと関連付けられたメモリスペースのサイズ、バッファオブジェクトのための使用法情報、及びバッファオブジェクトのために割り当てられたメモリスペースの領域を含むこともできる。画像オブジェクトは、二次元又は三次元の配列のデータ、例えば、テクスチャ、フレームバッファ又は画像、を格納する。画像オブジェクトは、情報、例えば、画像の次元、画像内の各要素の記述、画像オブジェクトのための使用法情報、画像オブジェクトのために割り当てられたメモリスペースの領域、を含むこともできる。
本開示の幾つかの態様により、メモリオブジェクト生成命令は、生成されるべきメモリオブジェクトのために即時モードがイネーブルにされるべきかどうかを指定する入力パラメータを含むことができる。ここにおいてさらに詳細に説明されるように、即時モードがイネーブルにされるときには、メモリオブジェクトは、キャッシング不能な共有メモリとして及び/又はキャッシュコヒーレントな共有メモリとして実装することができる。即時モードがディスエーブルにされるときには、メモリオブジェクトは、必ずしもキャッシング不能な共有メモリとして及び/又はキャッシュコヒーレントな共有メモリとして実装されないことがある。
幾つかの例では、メモリオブジェクトは、メモリオブジェクトが即時モードメモリオブジェクトであるかどうかを示す即時モード属性を含むことができる。該例では、メモリバッファインタフェース328は、生成されるべきメモリオブジェクトのための即時モード属性を、メモリオブジェクトのために即時モードがイネーブルにされるべきかどうかを指定する情報に基づいてメモリオブジェクトのために即時モードがイネーブルにされるべきかどうかを示す値に設定するように構成することができる。メモリオブジェクトの即時モード属性は、メモリオブジェクトをキャッシング不能な共有メモリとして及び/又はキャッシュコヒーレントな共有メモリとして実装すべかどうかを決定するために計算システム310によって用いることができる。
メモリオブジェクト生成命令は、幾つかの例では、バッファオブジェクトのために即時モードがイネーブルにされるかどうかを指定する命令内の情報に基づいてバッファオブジェクトを生成するようにメモリバッファインタフェース328に命令するバッファオブジェクト生成命令を含むことができる。メモリオブジェクト生成命令は、さらなる例では、画像オブジェクトのために即時モードがイネーブルにされるかどうかを指定する命令内の情報に基づいて画像オブジェクトを生成するようにメモリバッファインタフェース328に命令する画像オブジェクト生成命令を含むことができる。
幾つかの例では、バッファオブジェクト生成命令のためのインタフェースは、次の形態をとることができる。
ここで、clCreateBufferは、命令識別子であり、cl_context contextは、バッファオブジェクトを生成するために用いられる有効なコンテキスト、例えば、OpenCLコンテキストであり、cl_mem_flags flagsは、バッファオブジェクトのための割り当て及び使用法情報を指定するために用いられるビットフィールドであり、size_t sizeは、割り当てられるべきバッファメモリオブジェクトのサイズをバイト単位で指定するパラメータであり、void *host_ptrは、アプリケーションによって割り当て済みであることができるバッファデータのポインタであり、cl_int *errcode_retは、1つ以上のエラーコードを戻す。その命令は、生成されたバッファオブジェクトをcl_memメモリオブジェクトとして戻すことができる。この例では、画像オブジェクトのために即時モードがイネーブルにされるべきであるかどうかを指定する入力パラメータは、例えば、cl_mem_flags flagsフィールドで指定されたCL_IMMEDIATEフラグであることができる。
さらなる例では、画像オブジェクト生成命令のためのインタフェースは、次の形態をとることができる。
ここで、clCreateImage2Dは、命令識別子であり、cl_context contextは、バッファオブジェクトを生成するために用いられる有効なコンテキスト、例えば、OpenCLコンテキスト、であり、cl_mem_flags flagsは、画像オブジェクトのための割り当て及び使用法情報を指定するために用いられるビットフィールドであり、const cl_image_format *image_formatは、割り当てられるべき画像のフォーマットプロパティを記述する構造のポインタであり、size_t image_widthは、ピクセルを単位とする画像の幅であり、size_t image_heightは、ピクセルを単位とする画像の高さであり、size_t image_row_pitchは、バイトを単位とするスキャンラインのピッチであり、void *host_ptrは、アプリケーションによって割り当て済みであることができる画像データのポインタであり、cl_int *errcode_retは、1つ以上のエラーコードを戻す。その命令は、生成された画像オブジェクトをcl_memメモリオブジェクトとして戻すことができる。この例では、画像オブジェクトのために即時モードがイネーブルにされるべきであるかどうかを指定する入力パラメータは、例えば、cl_mem_flags flagsフィールドで指定されたCL_IMMEDIATEフラグであることができる。
幾つかの例では、メモリオブジェクト生成インタフェースは、読み取り/書き込み属性に関してWRITE_ONLY属性又はREAD_ONLY属性のいずれかのみを許容するように構成することができる。換言すると、該例では、メモリバッファインタフェース328は、READ_WRITE属性は拒否することができる。即時でないCL画像は、OpenCL仕様によって提供された該特徴を既に有していることができる。READ_WRITE属性を拒否することは、キャッシュのコヒーレンシーを維持する上での複雑さを低減することができる。
本開示により、メモリバッファインタフェース328は、ホストデバイス312及びGPU14の両方によってアクセス可能である共有メモリスペース336のために即時モードがイネーブルにされるべきであるかどうかを指定する命令を受信するように、及び、共有メモリスペース336のために即時モードがイネーブルにされるべきであるかどうかを指定する受信された命令に基づいて共有メモリスペース336のために即時モードを選択的にイネーブルにように構成される。例えば、メモリバッファインタフェース328は、共有メモリスペース336のために即時モードがイネーブルにされるべきであることを命令が指定する場合は共有メモリスペース336のために即時モードをイネーブルにすることができ、共有メモリスペース336のために即時モードがディスエーブルにされるべきであることを命令が指定する場合は共有メモリスペース336のために即時モードをディスエーブルにすることができる。命令は、例えば、メモリオブジェクト生成命令、バッファオブジェクト生成命令又は画像オブジェクト生成命令のうちの1つであることができる。共有メモリスペース336は、例えば、メモリオブジェクト、バッファオブジェクト又は画像オブジェクトに対応することができる。
幾つかの例では、メモリバッファインタフェース328が共有メモリスペース336のために即時モードをイネーブルしたときには、メモリバッファインタフェース328は、共有メモリスペース336のためのキャッシングサービスをディスエーブルにさせることができる。同様に、メモリバッファインタフェース328が共有メモリスペース336のために即時モードをディスエーブルにしたときには、メモリバッファインタフェース328は、共有メモリスペース336のためのキャッシングサービスを共有メモリスペース336のためにイネーブルにさせることができる。キャッシングサービスは、ホストキャッシュ324及びGPUキャッシュ332のうちの1つ又は両方によって行うことができる。ここにおいて用いられる場合のキャッシングサービスとは、当業において知られるキャッシュによって典型的に行われるサービスを意味することができる。
さらなる例では、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられた即時モード属性を、共有メモリスペースのために即時モードがイネーブルにされるかどうかを示す値に設定することによって、共有メモリスペース336のために即時モードをイネーブル及びディスエーブルにすることができる。例えば、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられた即時モード属性を、共有メモリスペース336のために即時モードがイネーブルにされること、例えば、即時モード属性=真、を示す値に設定することによって、共有メモリスペース336のために即時モードをイネーブルにすることができる。同様に、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられた即時モード属性を、共有メモリスペース336のために即時モードがディスエーブルにされること、例えば、即時モード属性=偽、を示す値に設定することによって、共有メモリスペース336のために即時モードをディスエーブルにすることができる。即時モード属性は、幾つかの場合においては、GPU314で実行中のタスク330によってアクセス可能であるグローバル変数、例えば、ブール変数、であることができる。幾つかの例では、即時モード属性は、共有メモリスペース336内に格納することができる。その他の例では、即時モード属性は、共有メモリスペース336以外のGPU314で実行中のタスク330によってアクセス可能な記憶場所内に格納することができる。共有メモリスペース336がメモリオブジェクトの一部である事例では、即時モード属性は、メモリオブジェクトのその他の属性が格納されるメモリスペースの記憶場所内に格納することができる。
メモリバッファインタフェース328が共有メモリスペース336と関連付けられた即時モード属性を設定することによって共有メモリスペース336のために即時モードをイネーブル及びディスエーブルにする例では、タスク330のためのソースコードは、幾つかの事例では、共有メモリスペース336に関するメモリ読み取り又は書き込み動作を行う前に、タスク330が共有メモリスペース336と関連付けられた即時モード属性にアクセスし、共有メモリスペース336のための即時モード属性に基づいて共有メモリスペース336のために即時モードがイネーブルにされるかどうかを決定するような形でコンパイルすることができる。共有メモリスペース336のために即時モードがイネーブルにされる場合は、タスク330は、共有メモリスペース336からデータを読み取るための又は共有メモリスペース336にデータを書き込むための即時モード読み取り又は書き込み命令を実行するようにプログラミングすることができる。他方、共有メモリスペースのために即時モードがイネーブルにされない場合は、タスク330は、共有メモリスペース336からデータを読み取るための又は共有メモリスペース336にデータを書き込むためのキャッシュドモード読み取り又は書き込み命令、例えば、キャッシングされた読み取り又は書き込み命令、を実行するようにプログラミングすることができる
即時モード読み取り及び書き込み命令は、例えば、キャッシングサービスを使用せずに読み取り及び書き込み動作をそれぞれ行うことができる。例えば、即時モード読み取り命令は、読み取り動作を行う前にキャッシュを無効にさせることができ及び/又は読み取り動作を行うときにキャッシュをバイパスすることができる。即時モード書き込み命令は、例えば、書き込み動作を行うときに即時ライトバックをキャッシュに行わせることができ及び/又は書き込み動作を行うときにキャッシュをバイパスすることができる。キャッシングされた読み取り命令及び書き込み命令は、例えば、GPUキャッシュ332のうちの1つ又は両方のキャッシングサービスを用いて、読み取り及び書き込み動作をそれぞれ実行することができる。
追加の事例では、タスク330のためのコンパイラは、タスク330のためのソースコードをコンパイルするときに、共有メモリスペース336のために即時モードがイネーブルにされるかどうかを示す情報へのアクセスを有することができる。例えば、タスク330のためのソースコード、例えば、カーネルソースコード、は、タスク330によって用いられ及び共有メモリスペース336と関連付けられたメモリオブジェクトのために即時モードがイネーブルにされるかどうかを示すフラグを含むことができる。幾つかの例では、フラグは、OpenCL属性修飾子、例えば、_cl_immediate属性修飾子、の形態をとることができる。共有メモリスペース336と関連付けられたメモリオブジェクトのために即時モードがイネーブルにされる場合は、コンパイラは、タスク330のためのコンパイルされたコードが、共有メモリスペース336に関して生じる読み取り又は書き込み動作のための即時モード読み取り及び/又は書き込み命令を含むような形でタスク330をコンパイルすることができる。他方、共有メモリスペース336と関連付けられたメモリオブジェクトに関して即時モードがイネーブルにされない場合は、コンパイラは、タスク330のためのコンパイルされたコードが、共有メモリスペース336に関して生じる読み取り又は書き込み動作のための即時モード読み取り及び/又は書き込み命令を含まないような形でタスク330をコンパイルすることができる。例えば、コンパイラは、タスク330のためのコンパイルされたコードが、共有メモリスペース336に関して生じる読み取り又は書き込み動作のためのキャッシングされた読み取り及び/又は書き込み命令を含むような形でタスク330をコンパイルすることができる。
さらなる例では、メモリバッファインタフェース328は、ホストデバイス312内のホストキャッシュ324及びGPU314内のGPUキャッシュ332のうちの少なくとも1つによる共有メモリスペース336のためのキャッシングサービスの実施をイネーブル及びディスエーブルにすることによって共有メモリスペース336のために即時モードをイネーブル及びディスエーブルにすることができる。例えば、メモリバッファインタフェース328は、ホストデバイス312内のホストキャッシュ324及びGPU314内のGPUキャッシュ332のうちの少なくとも1つによる共有メモリスペース336のためのキャッシングサービスの実施をディスエーブルにすることによって共有メモリスペース336のために即時モードをイネーブルにすることができる。同様に、メモリバッファインタフェース328は、ホストデバイス312内のホストキャッシュ324及びGPU314内のGPUキャッシュ332のうちの少なくとも1つによる共有メモリスペース336のためのキャッシングサービスの実施をイネーブルにすることによって共有メモリスペース336のために即時モードをディスエーブルにすることができる。
該例では、メモリバッファインタフェース328は、共有メモリスペース336のためのキャッシングサービスを行うキャッシュと関連付けられたハードウェアに基づくキャッシュ制御モジュール及び/又はハードウェアに基づくメモリ管理ユニットを構成することによって共有メモリスペース336のためのキャッシングサービスの実施をイネーブル及びディスエーブルにすることができる。例えば、ホストキャッシュ324による共有メモリスペース336のためのキャッシングサービスの実施をイネーブルにするためには、メモリバッファインタフェース328は、共有メモリスペース336のためにホストキャッシュ324によってキャッシングサービスが提供されるような形でホストキャッシュ制御モジュール326を構成することができる。ホストキャッシュ324による共有メモリスペース336のためのキャッシングサービスの実施をディスエーブルにするためには、メモリバッファインタフェース328は、例えば、共有メモリスペース336のためにホストキャッシュ324によってキャッシングサービスが提供されないような形でホストキャッシュ制御モジュール326を構成することができる。同様に、GPUキャッシュ332による共有メモリスペース336のためのキャッシングサービスの実施をイネーブルにするためには、メモリバッファインタフェース328は、例えば、共有メモリスペース336のためにホストキャッシュ324によってキャッシングサービスが提供されるような形でGPUキャッシュ制御モジュール334を構成することができる。GPUキャッシュ332による共有メモリスペース336のためのキャッシングサービスの実施をディスエーブルにするためには、メモリバッファインタフェース328は、例えば、共有メモリスペース336のためにGPUキャッシュ332によってキャッシングサービスが提供されないような形でGPUキャッシュ制御モジュール334を構成することができる。
幾つかの例では、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられた1つ以上のハードウェアに基づく即時フラグを、共有メモリスペース336のためにキャッシングサービスが提供されるべきかどうかを示す値に設定することによって、ホストキャッシュ制御モジュール326及びGPUキャッシュ制御モジュール334のうちの1つ又は両方を構成することができる。1つ以上のハードウェアに基づく即時フラグは、幾つかの例では、1つ以上のレジスタであることができる。さらなる例では、ハードウェアに基づく即時フラグは、即時フラグのテーブルの一部であることができ、即時フラグのテーブル内の各即時フラグは、メモリ316内の特定のアドレス空間に対応する。いずれの場合も、共有メモリスペース336と関連付けられた1つ以上の即時フラグが、キャッシングサービスが提供されるべきであることを示す値に設定されたときには、ホストキャッシュ制御モジュール326及び/又はGPUキャッシュ制御モジュール334は、ホストキャッシュ324及び/又はGPUキャッシュ332を用いて共有メモリスペース336のためのキャッシングサービスを提供することができる。同様に、共有メモリスペース336と関連付けられた1つ以上の即時フラグが、キャッシングサービスが提供されるべきでないことを示す値に設定されたときには、ホストキャッシュ制御モジュール326及び/又はGPUキャッシュ制御モジュール334は、共有メモリスペース336のためのキャッシングサービスを提供することができない。
該例では、GPUキャッシュ制御モジュール334は、メモリ316のアドレス空間内のメモリアドレスのための読み取り命令及び/又は書き込み命令を処理するように構成することができる。読み取り及び書き込み命令は、例えば、GPU314で実行中のタスク330によってGPUキャッシュ制御モジュール334に出すことができる。メモリ316の所定のアドレス空間内のメモリ記憶場所からデータを読み取る又はメモリ記憶場所にデータを書き込むための読み取り又は書き込み命令を受信したことに応答して、GPUキャッシュ制御モジュール334は、アドレス空間と関連付けられたハードウェアに基づくフラグを識別し、ハードウェアに基づくフラグの値に基づいて読み取り又は書き込み命令を処理するときにGPUキャッシュ332のキャッシングサービスを使用するかどうかを決定することができる。GPUキャッシュ制御モジュール334がGPUキャッシュ332のキャッシングサービスを使用することを決定した場合は、GPUキャッシュ制御モジュール334は、例えば、データが有効である場合にGPUキャッシュ332からデータを読み取ること及び/又はGPUキャッシュ332にデータを書き込むことを試行することができる。GPUキャッシュ制御モジュール334がGPUキャッシュ332のキャッシングサービスを使用しないことを決定した場合は、GPUキャッシュ制御モジュール334は、幾つかの例では、GPUキャッシュ332をバイパスし、メモリ316から直接データを読み取ること又はメモリ316に直接データを書き込むことができる。追加の例では、GPUキャッシュ制御モジュール334がGPUキャッシュ332のキャッシングサービスを使用しないことを決定した場合は、GPUキャッシュ制御モジュール334は、読み取り命令を実行する前にアドレス空間と関連付けられたキャッシュ332の部分を無効にすること及び/又は書き込み命令を実行するときにキャッシュライトバック(write back)又はキャッシュライトスルー(write through)技法を実施することができる。ホストキャッシュ制御モジュール334は、ホストデバイス312で実行中のホストプロセス320から受信された読み取り及び書き込み命令に応答してホストキャッシュ324に関して同様の方法で動作することができる。
追加の例では、メモリバッファインタフェース328は、ホストデバイス312内のホストキャッシュ324及びGPU314内のGPUキャッシュ332のうちの少なくとも1つのために共有メモリキャッシュコヒーレンシーモードをイネーブル及びディスエーブルにすることによって共有メモリスペース336のために即時モードをイネーブル及びディスエーブルにすることができる。例えば、共有メモリスペース336のために即時モードをイネーブルにするめには、メモリバッファインタフェース328は、ホストデバイス312内のホストキャッシュ324及びGPU314内のGPUキャッシュ332のうちの少なくとも1つのために共有メモリキャッシュコヒーレンシーモードをイネーブルにすることができる。同様に、共有メモリスペース336のために即時モードをディスエーブルにするためには、メモリバッファインタフェース328は、ホストデバイス312内のホストキャッシュ324及びGPU314内のGPUキャッシュ332のうちの少なくとも1つのために共有メモリキャッシュコヒーレンシーモードをディスエーブルにすることができる。該例では、メモリバッファインタフェース328は、幾つかの事例では、共有メモリキャッシュコヒーレンシーモードをイネーブルにするようにホストキャッシュ制御モジュール326及びGPUキャッシュ制御モジュール334のうちの1つ又は両方を構成することによってホストキャッシュ324のために共有メモリキャッシュコヒーレンシーモードをイネーブルにすることができ、共有メモリキャッシュコヒーレンシーモードをディスエーブルにするようにホストキャッシュ制御モジュール326及びGPUキャッシュ制御モジュール334のうちの1つ又は両方を構成することによってホストキャッシュ324のために共有メモリキャッシュコヒーレンシーモードをディスエーブルにすることができる。
ホストキャッシュ324のための共有メモリキャッシュコヒーレンシーモードがイネーブルにされたときには、ホストキャッシュ制御モジュール326は、既知の方法により共有メモリスペース336に関して共有メモリキャッシュコヒーレンシー技法を実行することができる。ホストキャッシュ324のための共有メモリキャッシュコヒーレンスモードがディスエーブルにされたときには、
ホストキャッシュ324は、共有メモリスペース336に関して共有メモリキャッシュコヒーレンシー技法を実行することができない。同様に、GPUキャッシュ332のための共有メモリキャッシュコヒーレンシーモードがイネーブルにされた時には、GPUキャッシュ制御モジュール334は、既知の方法により共有メモリスペース336に関して共有メモリキャッシュコヒーレンシー技法を実行することができる。GPUキャッシュ332のための共有メモリキャッシュコヒーレンスモードがディスエーブルにされたときには、GPUキャッシュ制御モジュール334は、共有メモリスペース336に関して共有メモリキャッシュコヒーレンシー技法を実行することができない。
例示を容易にするため、図20において例示された計算システム例310は、GPU314を計算デバイスとして使用した本開示の即時バッファリング技法について説明する。本開示の技法は、GPU314に加えて又はGPU314の代わりにGPU以外の計算デバイスを有するマルチプルプロセッサ計算システムに適用できることが認識されるべきである。幾つかの例では、計算システム310内の計算デバイスは、OpenCL計算デバイスであることができる。さらに、図20に示される計算システム例310は、ホストデバイスと計算デバイスとの間でのインフライトデータシェアリングを容易にする即時メモリオブジェクトを実装するためのインフラストラクチャ及び技法を例示する。しかしながら、その他の計算システム例では、それらの技法は、2つ以上の計算デバイスを有する計算システム内で異なる計算デバイス(例えば、OpenCL計算デバイス)間でのインフライトデータシェアリングを提供するために容易に拡張することができる。該例では、異なる計算デバイス間で1本以上の割り込みラインを配線することができる。
図21は、本開示によるホストデバイスで実行中のプロセスによって出されたメモリオブジェクト生成命令を実行するための技法例を示した流れ図である。幾つかの例では、図21に示される技法例を実装するために図20の計算システム310を用いることができる。メモリオブジェクト生成命令は、バッファオブジェクト生成命令又は画像オブジェクト生成命令であることができる。メモリバッファインタフェース328は、メモリオブジェクト生成命令を受信する(340)。メモリバッファインタフェース328は、メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定するかどうかを決定する(342)。例えば、メモリバッファインタフェース328は、メモリオブジェクト生成命令用のパラメータリストに即時フラグパラメータが含まれるかどうかを決定することができる。
メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定しないとメモリバッファインタフェース328が決定した場合は、メモリバッファインタフェース328は、生成されるべきメモリオブジェクトのための共有メモリスペース336を割り当て(344)、ホストキャッシュ324及びGPUキャッシュ332のうちの1つ又は両方によるキャッシングサービスの実施を共有メモリスペース336のためにイネーブルにさせ(346)、生成されたメモリオブジェクトへの参照を戻すことができる(348)。メモリオブジェクト生成命令は、例えば、即時フラグパラメータを含めないことによって又は即時モードがイネーブルにされるべきでないことを他のパラメータ値で指定することによって即時モードがイネーブルにされるべきでないことを指定することができる。逆に、メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定するとメモリバッファインタフェース328が決定した場合は、メモリバッファインタフェース328は、生成されるべきメモリオブジェクトのための共有メモリスペース336を割り当て(350)、ホストキャッシュ324及びGPUキャッシュ332のうちの1つ又は両方によるキャッシングサービスの実施を共有メモリスペース336のためにディスエーブルにさせ(352)、生成されたメモリオブジェクトへの参照を戻すことができる(354)。メモリオブジェクト生成命令は、例えば、即時フラグパラメータを含めることによって又は即時モードがイネーブルにされるべきであることを他のパラメータ値で指定することによって即時モードがイネーブルにされるべきであることを指定することができる。
幾つかの例では、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられたメモリオブジェクトの即時モード属性を、共有メモリスペース336と関連付けられたメモリオブジェクトのためにキャッシングサービスが提供されるべきであることを示す値に設定することによって、共有メモリスペース336のためにキャッシングサービスの実施をイネーブルにさせることができる。同様に、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられたメモリオブジェクトの即時モード属性を、共有メモリスペース336と関連付けられたメモリオブジェクトのためにキャッシングサービスが提供されるべきでないことを示す値に設定することによって、共有メモリスペース336のためにキャッシングサービスの実施をディスエーブルにさせることができる。戻されたメモリオブジェクトは、即時モード属性を含むことができる。該例では、メモリオブジェクトのための即時モード属性は、ホストデバイス312で実行中のホストプロセス320及びGPU314で実行中のタスク330のうちの1つ又は両方によってアクセス可能であることができる。ホストプロセス320及び/又はタスク330は、共有メモリスペース336と関連付けられたメモリオブジェクトの即時モード属性に基づいて共有メモリスペース336に関して特定の読み取り及び書き込み命令を実行するときにキャッシングサービスを使用すべきかどうかを決定することができる。
さらなる例では、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられた1つ以上のハードウェアに基づく即時フラグを、共有メモリスペース336のためにキャッシングサービスが提供されるべきであることを示す値に設定することによって、共有メモリスペース336のためにキャッシングサービスの実施をイネーブルにさせることができる。同様に、メモリバッファインタフェース328は、共有メモリスペース336と関連付けられた1つ以上のハードウェアに基づく即時フラグを、共有メモリスペース336のためにキャッシングサービスが提供されるべきでないことを示す値に設定することによって、共有メモリスペース336のためにキャッシングサービスの実施をディスエーブルにさせることができる。1つ以上のハードウェアに基づく即時フラグは、ホストキャッシュ制御モジュール326及びGPUキャッシュ制御モジュール334のうちの1つ以上に又は他のローカルな又はグローバルなメモリ管理ユニット(示されていない)に配置することができる。
追加の例では、メモリバッファインタフェース328は、データを格納するためにメモリ316内の物理的メモリスペースを割り当てる前に、呼び出しを行っているプロセス、例えば、ホストプロセス320、にメモリオブジェクトを戻すことができる。該例では、メモリバッファインタフェース328は、戻されたメモリオブジェクト内に即時モード属性を含めることができる。メモリオブジェクトのためにのちにメモリ316が割り当てられるときには、メモリバッファインタフェース328又は他のモジュールは、メモリオブジェクトの即時モード属性に基づいて1つ以上のハードウェアに基づく即時フラグを設定することができる。
図22は、本開示によるホストデバイスで実行中のプロセスによって出されたメモリオブジェクト生成命令を実行するための他の技法例を示した流れ図である。幾つかの例では、図22に示される技法例を実装するために図20の計算システム310を用いることができる。メモリオブジェクト生成命令は、バッファオブジェクト生成命令又は画像オブジェクト生成命令であることができる。メモリバッファインタフェース328は、メモリオブジェクト生成命令を受信する(356)。メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定するかどうかを決定する(358)。例えば、メモリバッファインタフェース328は、メモリオブジェクト生成命令用のパラメータリストに即時フラグパラメータが含められるかどうかを決定することができる。
メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定しないとメモリバッファインタフェース328が決定した場合は、メモリバッファインタフェース328は、生成されるべきメモリオブジェクトのための共有メモリスペース336を割り当て(360)、共有メモリスペース336のために共有メモリキャッシュコヒーレンシーモードをディスエーブルにし(362)、生成されたメモリオブジェクトへの参照を戻す(364)ことができる。逆に、メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定するとメモリバッファインタフェース328が決定した場合は、メモリバッファインタフェース328は、生成されるべきメモリオブジェクトのための共有メモリスペース336を割り当て(366)、共有メモリスペース336のために共有メモリキャッシュコヒーレンシーモードをイネーブルにし(368)、生成されたメモリオブジェクトへの参照を戻す(370)ことができる。
幾つかの例では、メモリバッファインタフェース328は、データを格納するためにメモリ316内の物理的メモリスペースを割り当てる前に、呼び出しを行っているプロセス、例えば、ホストプロセス320、にメモリオブジェクトを戻すことができる。該例では、メモリバッファインタフェース328は、戻されたメモリオブジェクト内に即時モード属性を含めることができる。メモリオブジェクトのためにのちにメモリ316が割り当てられるときには、メモリバッファインタフェース328又は他のモジュールは、メモリオブジェクトの即時モード属性に基づいて共有メモリスペースキャッシュコヒーレンシーモードをイネーブル又はディスエーブルにすることができる。
図23乃至28は、GPUが本開示による即時モード及びキャッシュドモードのロード命令及び格納命令を処理するために用いることができる技法例を示す。上述されるように、タスク330のためのソースコードは、幾つかの例では、即時メモリオブジェクト及びキャッシングされたメモリオブジェクトの両方をサポートするためにコンパイルされたコードがキャッシュドモード命令及び即時モード命令の両方を含むことができるような形でコンパイルすることができる。キャッシュドモード命令は、基礎になるメモリと関連付けられたキャッシュのキャッシングサービスを用いてメモリに関して読み取り及び書き込み動作を実行することができ、即時モード命令は、基礎になるメモリと関連付けられたキャッシュのキャッシングサービスを用いずにメモリに関して読み取り及び書き込み動作を実行することができる。キャッシュドモード命令は、ここでは代替として非即時モード命令と呼ぶことができる。ロード及び格納命令は、ここでは代替として読み取り及び書き込み命令とそれぞれ呼ぶことができる。
幾つかの例では、ロード又は格納命令のキャッシュドモードバージョン及びロード又は格納命令の即時モードバージョンは、例えば、各々が異なる演算コード、すなわちオプコード、を有する異なる命令であることができる。さらなる例では、ロード又は格納命令のキャッシュドモードバージョン及びロード又は格納命令の即時モードバージョンは、例えば、両方が同じオプコードを有する同じ命令であることができる。該例では、命令とともに提供されるパラメータは、命令がキャッシュドモードであるか又は即時モードであるかを指定することができる。
図23は、本開示によるキャッシュドモード及び即時モード命令を処理するための技法例を示す流れ図である。幾つかの例では、図23に示される技法例を実装するために図20の計算システム310を用いることができる。図23の例では、即時モードは、バイパスキャッシュモードと呼ばれ、即時モード命令は、バイパスキャッシュモード命令に対応する。GPUキャッシュ制御モジュール334は、メモリ記憶場所及びバイパスキャッシュモードがイネーブルにされるかどうかを指定するロード命令を受信する(372)。GPUキャッシュ制御モジュール334は、バイパスキャッシュモードがイネーブルにされるようにロード命令が指定するかどうかを決定する(374)。幾つかの事例では、GPUキャッシュ制御モジュール334は、命令のタイプ、例えば、命令のオプコード、に基づいてバイパスキャッシュモードがイネーブルにされるようにロード命令が指定するかどうかを決定することができる。追加の事例では、GPUキャッシュ制御モジュール334は、ロード命令とともに含まれておりバイパスキャッシュモードがイネーブルにされるかどうかを示すパラメータに基づいてバイパスキャッシュモードがイネーブルにされるようにロード命令が指定するかどうかを決定することができる。バイパスキャッシュモードがイネーブルにされないとGPUキャッシュ制御モジュール334が決定した場合は、GPUキャッシュ制御モジュール334は、ロード命令において指定されたメモリ記憶場所と関連付けられたキャッシュ記憶場所において、キャッシュ、例えば、GPUキャッシュ332、からデータを取り出す(376)。他方、バイパスキャッシュモードがイネーブルにされるとGPUキャッシュ制御モジュール334が決定した場合は、GPUキャッシュ制御モジュール334は、ロード命令において指定されたメモリ記憶場所において、メモリ、例えば、共有メモリスペース336、からデータを取り出す(378)。
図24は、本開示によるキャッシュドモード命令及び即時モード命令を処理するための他の技法例を示した流れ図である。幾つかの例では、図24に示される技法例を実装するために図20の計算システム310を用いることができる。図24の例では、即時モードは、バイパスキャッシュモードと呼ばれ、即時モード命令は、バイパスキャッシュモード命令に対応する。GPUキャッシュ制御モジュール334は、メモリ記憶場所、格納すべきデータ及びバイパスキャッシュモードがイネーブルにされるかどうかを指定する格納命令を受信する(380)。GPUキャッシュ制御モジュール334は、バイパスキャッシュモードがイネーブルにされるように格納命令が指定するかどうかを決定する(382)。幾つかの事例では、GPUキャッシュ制御モジュール334は、命令のタイプ、例えば、命令のオプコード、に基づいてバイパスキャッシュモードがイネーブルにされるように格納命令が指定するかどうかを決定することができる。追加の事例では、GPUキャッシュ制御モジュール334は、ロード命令とともに含まれておりバイパスキャッシュモードがイネーブルにされるかどうかを示すパラメータに基づいてバイパスキャッシュモードがイネーブルにされるように格納命令が指定するかどうかを決定することができる。バイパスキャッシュモードがイネーブルにされないとGPUキャッシュ制御モジュール334が決定した場合は、GPUキャッシュ制御モジュール334は、格納命令において指定されたメモリ記憶場所と関連付けられたキャッシュ記憶場所において、キャッシュ、例えば、GPUキャッシュ332に、格納命令内で指定されたデータを格納する(384)。他方、バイパスキャッシュモードがイネーブルにされるとGPUキャッシュ制御モジュール334が決定した場合は、GPUキャッシュ制御モジュール334は、格納命令において指定されたメモリ記憶場所において、メモリ、例えば、共有メモリスペース336、に格納命令内で指定されたデータを格納する(386)。
図25は、本開示によるキャッシュドモード命令及び即時モード命令を処理するための他の技法例を示した流れ図である。幾つかの例では、図25に示される技法例を実装するために図20の計算システム310を用いることができる。GPUキャッシュ制御モジュール334は、メモリ記憶場所、格納すべきデータ及び即時モードがイネーブルにされるかどうかを指定する格納命令を受信する。GPUキャッシュ制御モジュール334は、格納命令内で指定されたメモリ記憶場所と関連付けられたキャッシュ記憶場所において、キャッシュ、例えば、GPUキャッシュ332、に格納命令内で指定されたデータを格納する(390)。GPUキャッシュ制御モジュール334は、即時モードがイネーブルにされるかどうかを指定する格納命令内の情報に基づいて即時モードがイネーブルにされるかどうかを決定する(392)。即時モードがイネーブルにされるかどうかを指定する情報は、幾つかの例では、命令のタイプ、例えば、命令のためのオプコード、及び/又は命令に関して即時モードがイネーブルにされるかどうかを指定する命令とともに含まれているパラメータであることができる。即時モードがイネーブルにされない場合は、GPUキャッシュ制御モジュール334は、即時のキャッシュライトバック動作を行わない(394)。他方、即時モードがイネーブルにされる場合は、GPUキャッシュ制御モジュール334は、即時のキャッシュライトバック動作を行う(396)。
図26は、本開示によるキャッシュドモード及び即時モード命令を処理するための他の技法例を示す流れ図である。幾つかの例では、図26に示される技法例を実装するために図20の計算システム310を用いることができる。GPUキャッシュ制御モジュール334は、メモリ記憶場所及び即時モードがイネーブルにされるかどうかを指定するロード命令を受信する(398)。GPUキャッシュ制御モジュール334は、即時モードがイネーブルにされるかどうかを指定するロード命令内の情報に基づいて即時モードがイネーブルにされるかどうかを決定する(400)。即時モードがイネーブルにされるかどうかを指定する情報は、幾つかの例では、命令のタイプ、例えば、命令のためのオプコード、及び/又は命令とともに含まれており、命令に関して即時モードがイネーブルにされるかどうかを指定するパラメータであることができる。即時モードがイネーブルにされない場合は、GPUキャッシュ制御モジュール334は、キャッシュをフラッシング(flushing)して無効にしない(402)。GPUキャッシュ制御モジュール334は、ロード命令内で指定されたデータを、キャッシュ、例えば、GPUキャッシュ332、内で入手可能な場合はキャッシュから、又は、キャッシュ内でデータを入手可能でない場合は、基礎となるメモリから、データを取り出す(404)。即時モードがイネーブルにされる場合は、GPUキャッシュ制御モジュール334は、キャッシュをフラッシングして無効化する(406)。GPUキャッシュ制御モジュール334は、ロード命令内で指定されたデータを基礎となるメモリから取り出す(408)。キャッシュは、フラッシングされて無効化されているためデータを戻さない。
図27は、本開示による図20の計算システム310において用いることができるGPU例420を示したブロック図である。幾つかの例では、図20に例示されるGPU314を実装するためにGPU420を用いることができる。GPU420は、GPU処理モジュール422と、GPUキャッシュ制御モジュール424と、GPUキャッシュ426と、キャッシュバス428と、バイパスバス430と、を含む。GPU処理モジュール422は、キャッシュバス428を介してGPUキャッシュ制御モジュール424に通信可能な形で結合される。GPU処理モジュール422は、バイパスバス430を介してメモリ316にも通信可能な形で結合される。GPUキャッシュ制御モジュール424及びGPUキャッシュ426は、図20のGPUキャッシュ制御モジュール334及びGPUキャッシュ332と実質的に類似しており、さらに詳細には説明されない。GPU処理モジュール422は、処理素子432と、バスコントローラ434と、を含む。処理素子432は、バスコントローラ434にロード及び格納命令を出すように構成される。
バスコントローラ434は、キャッシュバス428及びバイパスパス430を介して該当する記憶場所にロード及び格納命令を転送するように構成することができる。バスコントローラ434は、命令が即時モード命令であるか又はキャッシュドモード命令であるかを示すロード命令又は格納命令内の情報に基づいて即時モード又は非即時モードで動作するように構成することができる。バスコントローラ434が非即時モード、すなわち、キャッシュドモード、で動作するように構成されるときには、バスコントローラ434は、実行のためにGPUキャッシュ制御モジュール424にロード及び格納命令を転送するためにキャッシュバス428を用いることができる。他方、バスコントローラ434が即時モードで動作するように構成されるときには、バスコントローラ434は、実行のためにメモリ316にロード及び格納命令を転送するためにバイパスバス430を用いることができる。
図28は、本開示によるキャッシュドモード命令及び即時モード命令を処理するための技法例を示した流れ図である。幾つかの例では、図28に示される技法例を実装するために図27のGPU420を用いることができる。バスコントローラ434は、ロード又は格納命令を受信する(440)。バスコントローラ434は、即時モードがイネーブルにされるかどうかを指定するロード又は格納命令内の情報に基づいて即時モードがイネーブルにされるかどうかを決定する(442)。即時モードがイネーブルにされるかどうかを指定する情報は、幾つかの例では、命令のタイプ、例えば、命令のためのオプコード、及び/又は命令に関して即時モードがイネーブルにされるかどうかを指定する、命令とともに含まれているパラメータであることができる。即時モードがイネーブルにされないとバスコントローラ434が決定した場合は、バスコントローラ434は、GPUキャッシュ制御モジュール424に受信された命令を転送するためにキャッシュバス428を使用する(444)。そうでない場合、即時モードがイネーブルにされるとバスコントローラ434が決定した場合は、バスコントローラ434は、メモリ316に受信された命令を転送するためにバイパスバス430を使用する(446)。
図29は、本開示によるホストデバイスで実行中のプロセスによって出されたメモリオブジェクト生成命令を実行するための他の技法例を示した流れ図である。幾つかの例では、図29に示される技法例を実装するために図20の計算システム310を用いることができる。メモリオブジェクト生成命令は、バッファオブジェクト生成命令又は画像オブジェクト生成命令であることができる。メモリバッファインタフェース328がメモリオブジェクト生成命令を受信する(448)。メモリバッファインタフェース328は、メモリオブジェクトのために即時モードがイネーブルにされるべきかどうかをメモリオブジェクト生成命令が指定するかどうかを決定する(450)。例えば、メモリバッファインタフェース328は、メモリオブジェクト生成命令用のパラメータリストに即時フラグパラメータが含められているかどうかを決定することができる。
メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定しないとメモリバッファインタフェース328が決定した場合は、メモリバッファインタフェース328は、生成されたメモリオブジェクトのための即時モード属性を、即時モードがイネーブルにされない、例えば、“偽”、であることを示す値に設定する(452)。他方、メモリオブジェクトのために即時モードがイネーブルにされるべきであることをメモリオブジェクト生成命令が指定するとメモリバッファインタフェース328が決定した場合は、メモリバッファインタフェース328は、生成されたメモリオブジェクトのための即時モード属性を、即時モードがイネーブルにされる、例えば、“真”、であることを示す値に設定する(454)。メモリオブジェクトの即時モード属性は、幾つかの例では、特定のメモリオブジェクトに格納されたデータにアクセスするときにキャッシュドモード読み取り及び書き込み動作又は即時モード読み取り及び書き込み動作を実行すべきかどうかを決定するために、ホストデバイス312及び/又はGPU314によって用いることができる。
幾つかの例では、ホストプロセス320及び/又はタスク330は、幾つかのメモリオブジェクトが即時メモリオブジェクトであるように及びその他のオブジェクトがキャッシングされたメモリオブジェクト、すなわち、非即時メモリオブジェクト、であるようにプログラミングすることを希望することができる。本開示の技法は、幾つかの例では、コンパイルされたタスク330がキャッシングされたメモリオブジェクト及び即時メモリオブジェクトの両方に関する読み取り及び書き込み動作を行うのを可能にする専門化されたコンパイル技法を含むことができる。第1のコンパイル技法例は、所定の読み取り動作又は書き込み動作をコンパイルして命令のシーケンスにすることができる。命令のシーケンスは、読み取られる又は書き込まれるメモリオブジェクトのための即時モード属性の値を検査し、即時モード属性の値に基づいてキャッシュドモード命令又は即時モード命令のいずれを実行するかを決定することができる。第2のコンパイル技法例は、メモリオブジェクトにアクセスするためにコンパイルされたコード内で使用するキャッシュモード命令又は即時モード命令のいずれかを選択するためにメモリオブジェクトが即時モードオブジェクトであるかどうかを示すソースコード内の情報を用いることができる。
第1のコンパイル技法例により、コンパイラは、タスク330のためのコンパイルされたコードが次の擬似コード例による読み取りシーケンスを含むような形でタスク330のためのソースコードをコンパイルすることができる。
ここで、“isImmediate”は、データが読み取られるメモリオブジェクトのためのブール即時モード属性を表し、“immediate_read(...)”は、即時モード読み取り命令を表し、“cached_read(...)”は、キャッシュドモード読み取り命令を表す。
GPUキャッシュ制御モジュール334は、使用されている場合は、GPUキャッシュ332からデータを読み取る前に、例えば、GPUキャッシュ332を無効化することによってimmediate_read(...)命令を処理することができる。GPUキャッシュ制御モジュール334は、例えば、読み取りを行う前にGPUキャッシュ332を無効化することなしに、通常の方法でGPUキャッシュからデータを読み取ることによってcached_read(...)命令を処理することができる。
第1のコンパイル技法例により、コンパイラは、タスク330のためのコンパイルされたコードが次の擬似コード例による書き込みシーケンスを含むような形でタスク330のためのソースコードをコンパイルすることができる。
ここで、“isImmediate”は、データが書き込まれるメモリオブジェクトのためのブール即時モード属性を表し、“immediate_write(...)”は、即時モード書き込み命令を表し、“cached_write(...)”は、キャッシュドモード書き込み命令を表す。
GPUキャッシュ制御モジュール334は、幾つかの例では、キャッシュが用いられる場合は、GPUキャッシュ332のためにライトスルーモードを用いることによって、immediate_write(...)命令を処理することができる。さらなる例では、GPUキャッシュ制御モジュール334は、キャッシュが用いられる場合は、GPUキャッシュ332にデータを書き込み、GPUキャッシュ332にデータを書き込むことに応答してGPUキャッシュ332のためのキャッシュフラッシュを行うことによってimmediate_write(...)命令を処理することができる。GPUキャッシュ制御モジュール334は、例えば、ライトスルーモードを使用せずに及び/又は書き込み動作に応答してキャッシュをフラッシングすることなしに通常の方法でGPUキャッシュ332にデータを書き込むことによってcached_write(...)命令を処理することができる。
図30は、GPUが上記の第1のコンパイル技法によりコンパイルされた命令のシーケンスをどのようにして処理することができるかを例示した流れ図である。幾つかの例では、図30に例示された技法は、読み取り及び書き込みシーケンスのために上記において提供された擬似コード例を実装するために用いることができる。タスク330は、読み取りシーケンス又は書き込みシーケンスを開始する(456)。例えば、タスク330は、タスク330が、タスク330の実行の際に特定のメモリオブジェクトのための読み取り又は書き込み命令が発生すべきである時点に達したときに読み取りシーケンス又は書き込みシーケンスを開始することができる。タスク330は、データが読み取られる又はデータが書き込まれるメモリオブジェクトと関連付けられた即時モード属性にアクセスする(458)。タスク330は、メモリオブジェクトのための属性が、即時モードがイネーブルにされることを示す値、例えば、“真”、に設定されているかどうかを決定する。メモリオブジェクトのための属性が、即時モードがイネーブルにされないことを示す値に設定されているとタスク330が決定した場合は、タスク330は、キャッシングされた読み取り又は書き込み命令を用いてメモリオブジェクトのためのキャッシングされた読み取り動作又は書き込み動作を行う(462)。そうでない場合、メモリオブジェクトのための属性が、即時モードがイネーブルにされることを示す値に設定されているとタスク330が決定した場合は、タスク330は、即時の読み取り又は書き込み命令を用いてメモリオブジェクトのための即時の読み取り又は書き込み動作を行う(464)。
第2のコンパイル技法例により、ソースコードをコンパイルするときに、コンパイラは、タスク330によって読み取られる又は書き込まれる特定のメモリオブジェクトのために即時モードがイネーブルにされるかどうかを示す情報へのアクセスを有することができる。コンパイラは、タスク330が特定のメモリオブジェクトから読み取る又は特定のメモリオブジェクトに書き込むときにキャッシュドモード読み取り及び書き込み命令又は即時モード読み取り及び書き込み命令の間でタスク330が選択するためにソースコードをコンパイルするためにこの情報を用いることができる。
幾つかの例では、特定のメモリオブジェクトのために即時モードがイネーブルにされるかどうかを示す情報は、タスク330のためのソースコードによってアクセスされた1つ以上のメモリオブジェクトのために即時モードがイネーブルにされるかどうかを示すコンパイル時間属性であることができる。例えば、タスク330のためのソースコード、例えば、カーネルソースコード、は、タスク330によって用いられる1つ以上のメモリオブジェクトのために即時モードがイネーブルにされるかどうかを示すコンパイル時間属性を含むことができる。コンパイル時間属性は、幾つかの事例では、OpenCL属性修飾子、例えば、_cl_immediate、の形態をとることができる。属性修飾子は、1つ以上の特定のメモリオブジェクト及び/又は1つ以上のメモリオブジェクト内に格納される1つ以上の変数と関連付けることができる。属性修飾子が特定のメモリオブジェクトと関連付けられるときには、コンパイラは、メモリオブジェクトのために即時モードがイネーブルにされると決定することができる。同様に、属性修飾子が特定のメモリオブジェクトと関連付けられないときには、コンパイラは、メモリオブジェクトのために即時モードがイネーブルにされないと決定することができる。該属性を用いることは、コンパイラにとっての作業を軽減し、潜在的にカーネルのサイズを小さくすることができる。幾つかの例では、ソフトウェアアプリケーションは、即時バッファの使用を、該バッファが必要とされる事例に制限することができる。該例では、即時バッファを使用すべきかどうかの判断は、コンパイル時間判断であることができる。
共有メモリスペース336と関連付けられたメモリオブジェクトのために即時モードがイネーブルにされることをコンパイル時間属性が示す場合は、コンパイラは、タスク330のためのコンパイルされたコードが共有メモリスペース336に関して生じる読み取り又は書き込み動作のための即時モード読み取り及び/又は書き込み命令を含むような形でタスク330をコンパイルすることができる。そうでない場合、共有メモリスペース336と関連付けられたメモリオブジェクトのために即時モードがイネーブルにされない場合は、コンパイラは、タスク330のためのコンパイルされたコードが共有メモリスペース336に関して生じる読み取り又は書き込み動作のための即時モード読み取り及び/又は書き込み命令を含まないような形でタスク330をコンパイルことができる。例えば、コンパイラは、タスク330のためのコンパイルされたコードが共有メモリスペース336に関して生じる読み取り又は書き込み動作のためのキャッシングされた読み取り及び/又は書き込み命令を含むような形でタスク330をコンパイルことができる。
図31は、本開示によるタスクのためのソースコードをコンパイルするための技法例を示した流れ図である。図31の技法を用いてコンパイルされた結果得られたコードは、幾つかの例では、図20のタスク330に対応することができる。図31の技法例では、タスク330は、カーネルと呼ばれる。コンパイラは、メモリオブジェクトによって実装されるカーネル引数を処理する(466)。コンパイラは、メモリオブジェクトが即時モードメモリオブジェクトであるかどうかを決定する(468)。幾つかの例では、コンパイラは、カーネルのソースコードに含まれる情報、例えば、カーネル引数と関連付けられたコンパイル時間属性、に基づいてメモリオブジェクトが即時モードメモリオブジェクトであるかどうかを決定することができる。メモリオブジェクトが即時モードメモリオブジェクトでないとコンパイラが決定した場合は、コンパイラは、キャッシングされた読み取り及び書き込み命令を用いて特定のカーネル引数と関連付けられた読み取り動作及び書き込み動作をコンパイルする(470)。他方、メモリオブジェクトが即時モードメモリオブジェクトであるとコンパイラが決定した場合は、コンパイラは、即時モード読み取り及び書き込み命令を用いて特定のカーネル引数と関連付けられた読み取り動作及び書き込み動作をコンパイルする(472)。
図32は、本開示によるキャッシングサービスを選択的に使用するためにGPUによって用いることができる技法例を示した流れ図である。例えば、それらの技法は、メモリのメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにキャッシングサービスが使用されるべきであるかを指定する情報を受信したことに応答してメモリスペースに関して読み取り動作又は書き込み動作のうちの少なくとも1つを実行するためにメモリと関連付けられたGPUキャッシュをGPUが選択的に使用するのを可能にすることができる。幾つかの例では、図32において例示された技法を実装するために図20において例示されたGPU314及び/又は図27において例示されたGPU420を用いることができる。
GPU314が、処理するために読み取り命令又は書き込み命令のうちの少なくとも1つを受信する(474)。受信された命令は、メモリのメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するようにGPU314に命令することができる。GPU314は、メモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにキャッシングサービスが使用されるべきかどうかを指定するキャッシュモード情報を受信する(476)。幾つかの例では、キャッシュモード情報は、受信された命令内に含めることができる。さらなる例では、キャッシュモード情報は、メモリスペースと関連付けられたメモリオブジェクトの即時モード属性であることができる。GPU314は、キャッシュモード情報に基づいてキャッシングサービスを使用すべきかどうかを決定する(478)。受信された命令を実行するためにキャッシングサービスが使用されるべきであることを指定する情報を受信したことに応答して、GPU314は、受信された命令を実行するためにキャッシングサービスを使用することができる(480)。受信された命令を実行するためにキャッシングサービスが使用されるべきでないことを指定する情報を受信したことに応答して、GPU314は、受信された命令を実行するためにキャッシングサービスを使用することができない(482)。幾つかの例では、GPU314は、判断ボックス478及びプロセスボックス480及び482のうちの1つ以上を実装するために図23乃至28及び30に例示された技法のうちの1つ以上を用いることができる。幾つかの事例では、図32に示される技法を実装するためにGPUキャッシュ制御モジュール又はメモリ管理ユニット、例えば、図20において例示されたGPUキャッシュ制御モジュール334、を用いることができる。追加の事例では、図32に示される技法を実装するためにバスコントローラ、例えば、図27において例示されたバスコントローラ434、を用いることができる。
幾つかの例では、即時メモリオブジェクトを実装するために、GPU ALUはグローバルメモリキャッシュ及び/又はALU命令内で指定されたグローバルメモリキャッシュの特定の部分を無効にするALU命令を実行するように設計することができる。概して、ホストデバイス312は、即時メモリオブジェクトを実装するために既存のCPU能力を用いることができる。
今度は、本開示において説明されるアウトオブバンドシグナリング技法、例えば、ここにおいて説明されるメッセージ渡し技法、及び本開示において説明される即時メモリオブジェクトに関する様々な使用事例がさらに詳細に説明される。第1の使用事例により、アウトオブバンドシグナリングは、アウトオブバンドシグナリング技法に加えて即時メモリオブジェクトを必ずしも使用せずに独立型の機能として用いることができる。アウトオブバンドシグナリングは、
同期化のため及び相対的に少量のデータを素早く渡すために用いることができる。幾つかの例では、アウトオブバンドシグナリングは、即時メモリオブジェクトよりも低いレーテンシーを有することができるが、即時メモリオブジェクトよりも低い帯域幅を有することができる。
アウトオブバンドシグナリングは、メモリ割り当て動作のために第1の使用事例により用いることもできる。例えば、GPUは、ホストCPUが新しいバッファを割り当てるように要求するためにアウトオブバンドシグナリングを用いることができる。GPUは、要求されたバッファ長をホストCPUに指定するためにアウトオブバンドシグナリングを用いることもできる。他の例として、CPUは、バッファのためのメモリ記憶場所を指定するポインタをGPUに送信するためのバッファを割り当てた後にアウトオブバンドシグナリングを用いることができる。
アウトオブバンドシグナリングは、少量のデータがやり取りされることになるリモートプロシージャコールのために第1の使用事例により用いることもできる。例えば、計算デバイス内の計算ユニットで実行中のカーネルが同じ計算デバイス又は他の計算デバイス内の他の計算ユニットで他のカーネルを起動させるためにRPCを用いる事例においては、RPCのためのデータは、起動を行う計算ユニットのローカルメモリに格納することができる。本開示のアウトオブバンドシグナリング技法は、起動を行う計算ユニットのローカルメモリから新しく起動されたカーネルを実行する計算ユニットのローカルメモリにデータを転送するために用いることができる。
アウトオブバンドシグナリングは、進行状況の報告のために第1の事例により用いることもできる。例えば、GPUは、現在のタスクの完了率をホストCPUに報告するためにアウトオブバンドシグナリングを用いることができる。
アウトオブバンドシグナリングは、エラー報告のために第1の使用事例により用いることもできる。例えば、GPUは、ホストCPUにエラーコードを報告するためにアウトオブバンドシグナリングを用いることができる。
アウトオブバンドシグナリングは、コンテキストスイッチを援助するために第1の使用事例により用いることもできる。例えば、ホストCPUは、GPUがコンテキストスイッチのために準備するために状態を保存するように要求するためにアウトオブバンドシグナリングを用いることができる。
第2の使用事例により、即時メモリオブジェクトは、即時メモリオブジェクトに加えてアウトオブバンドシグナリングを必ずしも使用せずに独立型機能として用いることができる。例えば、即時バッファは、相対的に大量のデータのやり取りを達成させるために用いることができる。即時バッファは、データだけでなく、同期化マーカも入れることができる。この場合は、データ生成器は、最初にバッファにデータを書き込み、次に、データが準備完了であること及び/又は記憶場所をコンシューマ(consumer)に示す同期化マーカを書き込むことができる。コンシューマは、推測で決定された記憶場所をポーリングすることによって、このメモリ記憶場所、例えば、バッファの見出し部、内の同期化データを探す。同期化マーカが入手された時点で、コンシューマはデータを読み取る。同様の技法を即時画像オブジェクトに対して適用することができる。
これらの技法のために様々な同期化プロトコルを採用することができる。例えば、同期化マーカをデータバッファ内部に埋め込むことができ、又は、別個のバッファに配置することができる。該技法は、可変長符号化又はラン長符号化方式を用いて圧縮される圧縮されたデータの送信に適用することができる。
第3の事例により、例えば、相対的に大量のデータのやり取りを達成させるために即時メモリオブジェクトをアウトオブバンドシグナリングとともに用いることができる。この場合は、アウトオブバンドシグナリングは、即時メモリオブジェクトがデータを格納する間に同期化のために用いることができる。例えば、データ生成器は、即時バッファ内にデータを入れ、アウトオブバンドシグナリングを用いてデータの準備完了及び記憶場所及び/又はサイズをコンシューマに通知することができる。流れが制御されたシナリオにおいては、コンシューマは、データを読み取り、バッファを再使用可能であることを生成器に通知する。同通知は、アウトオブバンドシグナリングを用いて達成させることもできる。
該技法は、流れが制御されたデータパイプライニングを要求するアルゴリズムにおいて用いることができる。ホストCPU及びGPUに関しては、該技法は、例えば、診断ロギングのために用いることができる。複数のOpenCLデバイスに関しては、これらの技法は、非同期的な流れが制御されたデータパイプライン内に複数のデバイスを接続するために用いることができる。これは、各CPU又はGPUにより適するブロックにアプリケーションを分割すること、複数のデバイスで様々なパイプライン処理段を起動させること及び/又はほとんど、さらにはすべてのデータ同期化をホストCPUからオフロードすることを可能にすることができる。
幾つかの例では、本開示の技法は、コマンド待ち行列を用いてタスクを開始させるマルチプルプロセッサ計算プラットフォームに関してホストデバイスで実行中のプロセスと計算デバイスで実行中のタスクとの間でのメッセージの送信及び受信を容易にするメッセージ渡しインタフェースを提供することができる。計算デバイスは、幾つかの事例では、GPUであることができる。追加の事例では、計算デバイスは、プラットフォーム横断、売り主横断型の異種計算プラットフォームAPIによって定義されたあらゆるタイプの計算デバイスであることができる。
さらなる例では、本開示の技法は、ホストデバイスによってアクセス可能である1つ以上のレジスタを含むGPUを提供することができる。1つ以上のレジスタは、GPUで実行中のタスクとGPU以外のデバイスで実行中のプロセスとの間でのメッセージ渡しを容易にするように構成することができる。
追加の例では、本開示の技法は、即時メモリオブジェクトを生成するのを可能にするメモリバッファインタフェースを提供することができる。即時メモリオブジェクトは、計算デバイスでタスクが実行している間にホストデバイスで実行中のプロセスと計算デバイスで実行中のタスクとの間でデータを共有するためにキャッシング不能な共有メモリスペース及び/又はキャッシュコヒーレントな共有メモリスペースを実装するために用いることができる。計算デバイスは、幾つかの事例においては、グラフィックス処理装置(GPU)であることができる。追加の事例では、計算デバイスは、プラットフォーム横断、売り主横断型の異種計算プラットフォームAPIによって定義されたあらゆるタイプの計算デバイスであることができる。
さらなる例では、本開示の技法は、キャッシング不能な共有メモリスペースを提供するために選択的にディスエーブルにすることができる共有メモリスペースのためのキャッシュを含むGPUを提供することができる。追加の例では、本開示の技法は、キャッシュコヒーレントな共有メモリスペースを提供するために選択的にイネーブルにすることができるキャッシュコヒーレンシーモードを含むGPUを提供することができる。
本開示において説明される技法は、少なくとも部分的には、ハードウェア、ソフトウェア、ファームウェア又はそれらの組み合わせ内に実装することができる。例えば、説明される技法の様々な態様を、1つ以上のプロセッサ内に実装することができ、1つ以上のマイクロプロセッサ、デジタル信号プロセッサ(DSP)、特定用途向け集積回路(ASIC)、フィールドプログラマブルゲートアレイ(FPGA)、その他の同等の集積回路又はディスクリート論理回路、及び、該コンポーネントのあらゆる組み合わせを含む。用語“プロセッサ”又は“処理回路”は、概して、単独での又はその他の論理回路と組み合わせた上記の論理回路、又は、あらゆるその他の同等の回路、例えば、処理を行うディスクリートハードウェアのうちのいずれかを意味することができる。
該ハードウェア、ソフトウェア、及びファームウェアは、本開示において説明される様々な動作及び機能をサポートするために同じデバイス内に又は別々のデバイス内に実装することができる。さらに、説明されるユニット、モジュール又はコンポーネントのいずれも、個別のただし相互運用可能な論理デバイスとしてまとめて又は別々に実装することができる。異なる特徴をモジュール又はユニットとして描写することは、異なる機能上の態様を強調することが意図されており、該モジュール又はユニットを別個のハードウェアコンポーネント又はソフトウェアコンポーネントによって実現させなければならないということは必ずしも意味しない。むしろ、1つ以上のモジュール又はユニットと関連付けられた機能は、別個のハードウェア、ファームウェア、及び/又はソフトウェアコンポーネントによって実行すること、又は、共通の又は別個のハードウェアコンポーネント又はソフトウェアコンポーネント内に組み入れることができる。
本開示において説明される技法は、コンピュータによって読み取り可能な媒体、例えば、命令を格納するコンピュータによって読み取り可能な記憶媒体、において格納すること、具現化すること又は符号化することもできる。コンピュータによって読み取り可能な媒体において埋め込まれた又は符号化された命令は、例えば、1つ以上のプロセッサによって命令が実行されるときに、ここにおいて説明される技法を実行することを1つ以上のプロセッサに行わせることができる。コンピュータによって読み取り可能な記憶媒体は、ランダムアクセスメモリ(RAM)、読み取り専用メモリ(ROM)、プログラマブル読み取り専用メモリ(PROM)、消去可能プログラマブル読み取り専用メモリ(EPROM)、電子的に消去可能なプログラマブル読み取り専用メモリ(EEPROM)、フラッシュメモリ、ハードディスク、CD−ROM、フロッピー(登録商標)ディスク、カセット、磁気媒体、光学媒体、又は有形であるその他のコンピュータによって読み取り可能な記憶媒体を含むことができる。
コンピュータによって読み取り可能な媒体は、有形な記憶媒体、例えば、上記のそれら、に対応するコンピュータによって読み取り可能な記憶媒体を含むことができる。コンピュータによって読み取り可能な媒体は、例えば、通信プロトコルによる1カ所から他へのコンピュータプログラムの転送を容易にするあらゆる媒体を含む通信媒体も備えることができる。このように、句“コンピュータによって読み取り可能な媒体”は、概して、(1)非一時的である有形なコンピュータによって読み取り可能な記憶媒体、及び(2)非有形なコンピュータによって読み取り可能な通信媒体、例えば、一時的な信号又は搬送波、に対応することができる。