JP2013537993A - マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法 - Google Patents

マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法 Download PDF

Info

Publication number
JP2013537993A
JP2013537993A JP2013530215A JP2013530215A JP2013537993A JP 2013537993 A JP2013537993 A JP 2013537993A JP 2013530215 A JP2013530215 A JP 2013530215A JP 2013530215 A JP2013530215 A JP 2013530215A JP 2013537993 A JP2013537993 A JP 2013537993A
Authority
JP
Japan
Prior art keywords
gpu
message
task
host device
instructions
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
Application number
JP2013530215A
Other languages
English (en)
Other versions
JP5738998B2 (ja
Inventor
ボード、アレクセイ・ブイ.
シャープ、コリン・クリストファー
ガルシア・ガルシア、デイビッド・リゲル
ジャン、チホン
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Qualcomm Inc
Original Assignee
Qualcomm Inc
Priority date (The priority date 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 date listed.)
Filing date
Publication date
Application filed by Qualcomm Inc filed Critical Qualcomm Inc
Publication of JP2013537993A publication Critical patent/JP2013537993A/ja
Application granted granted Critical
Publication of JP5738998B2 publication Critical patent/JP5738998B2/ja
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46Multiprogramming arrangements
    • G06F9/54Interprogram communication
    • G06F9/544Buffers; Shared memory; Pipes
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46Multiprogramming arrangements
    • G06F9/54Interprogram communication
    • G06F9/546Message passing systems or structures, e.g. queues
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06FELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00Arrangements for program control, e.g. control units
    • G06F9/06Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46Multiprogramming arrangements
    • G06F9/54Interprogram communication
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06TIMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00General purpose image data processing
    • G06T1/20Processor architectures; Processor configuration, e.g. pipelining
    • GPHYSICS
    • G06COMPUTING; CALCULATING OR COUNTING
    • G06TIMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00General purpose image data processing
    • G06T1/60Memory management

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • Memory System Of A Hierarchy Structure (AREA)
  • Advance Control (AREA)
  • Multi Processors (AREA)
  • Image Generation (AREA)

Abstract

本開示は、マルチプルプロセッサ計算プラットフォーム内で用いることができる通信技法について説明する。それらの技法は、幾つかの例では、コマンド待ち行列を用いてタスクを開始させるマルチプルプロセッサ計算プラットフォーム内でのメッセージ渡しをサポートするために用いることができるソフトウェアインタフェースを提供することができる。それらの技法は、追加の例では、マルチプルプロセッサ計算プラットフォーム内における共有メモリプロセッサ間通信のために用いることができるソフトウェアインタフェースを提供することができる。さらなる例では、それらの技法は、グラフィックス処理装置(GPU)とホストCPUとの間でのメッセージ渡し及び/又は共有メモリ通信をサポートするためのハードウェアを含むGPUを提供することができる。

Description

本開示は、計算プラットフォームに関するものである。本開示は、より具体的には、複数のプロセッサを含む計算プラットフォームに関するものである。
高い計算集約度及び/又は高いデータスループットが要求されるアプリケーションの性能を向上させるために複数のプロセッサを含む計算プラットフォームが用いられる。マルチプルプロセッサ計算プラットフォームは、ホストデバイスとして働くことができる汎用中央処理装置(CPU)と、ホストCPUが計算集約型タスクを実施する負担を軽減してそれによってシステム全体の性能を向上させるために使用することができる1つ以上の計算デバイスと、を含むことができる。幾つかの場合においては、1つ以上の計算デバイスは、幾つかのタイプのタスクをホストCPUよりも効率的に処理してシステム全体のためのさらなる性能向上を提供することができるように特に設計することができる。例えば、1つ以上の計算デバイスは、並列アルゴリズムをホストCPUよりも効率的に実行するように特に設計することができる。
マルチプルプロセッサ計算システムにおいて用いることができる1つのタイプの計算デバイスは、グラフィックス処理装置(GPU)である。伝統的には、CPUは、表示デバイスへの三次元(3D)グラフィックスのリアルタイムレンダリング用に特に設計された固定された機能のハードウェアが含まれていたが、典型的にはプログラミングすることができない、すなわち、コンパイルされたプログラムをGPUにダウンロードしてGPU上で実行することができなかった。しかしながら、より最近においては、プログラマブルシェーダユニット(shader unit)の開発に伴い、GPUのアーキテクチャの多くは、数多くの並列処理素子を含むプログラミング可能なアーキテクチャに移行している。プログラミング可能なアーキテクチャは、GPUがグラフィックス動作だけでなく汎用の計算タスクを高度に並列な形で実行することを容易にすることができる。
GPUを用いて汎用の非グラフィックス専用の計算タスクを実行することは、ここでは、グラフィックス処理装置上での汎用計算(GPGPU)、又は代替としてGPU計算と呼ぶことができる。幾つかの事例においては、GPUは、グラフィック専用でないアプリケーションプログラミングインタフェース(API)を利用可能にし、それによって汎用計算タスクの実行のためのGPUのプログラミングを容易にすることができる。GPU計算タスクは、計算集約型のタスク、及び/又は高度な並列性、例えば、行列計算、信号処理計算、統計アルゴリズム、分子モデル作成アプリケーション、金融アプリケーション、医療画像撮影、暗号解読アプリケーション、等を含むタスクを含むことができる。
GPUは、マルチプルプロセッサ計算プラットフォームにおいて用いることができる1つのタイプの計算デバイスであるにすぎず、GPUに加えて又はGPUの代わりにその他のタイプの計算デバイスを使用することもできる。例えば、マルチプルプロセッサ計算プラットフォームにおいて用いることができるその他のタイプの計算デバイスは、例えば、追加のCPU、デジタル信号プロセッサ(DSP)、セルブロードバンドエンジン(Cell/BE)プロセッサ又はその他のタイプの処理ユニットを含む。
複数の計算デバイスを有するマルチプルプロセッサ計算プラットフォームは、同種プラットフォーム又は異種プラットフォームであることができる。同種プラットフォームでは、すべての計算デバイスが共通の命令セットアーキテクチャ(ISA)を共有する。対照的に、異種プラットフォームは、異なるISAを有する2つ以上の計算デバイスを含むことができる。概して、異なるタイプの計算デバイスは、異なるISAを有することができ、異なるブランドの同じタイプの計算デバイスも異なるISAを有することができる。
マルチプルプロセッサ計算プラットフォームの性能は、マルチコア計算デバイス及び/又は多コア計算デバイスを利用することによってさらに向上させることができる。マルチコア計算デバイスの一例は、複数の処理コアを有するプログラマブルシェーダユニットを内蔵する上記のGPUである。しかしながら、CPUも、複数の処理コアを含むように設計することができる。概して、複数の処理コアを含むチップ又はダイは、マルチコアプロセッサとみなすことができる。処理コアは、1つの特定のデータに関する命令を実行することが可能な処理ユニットを意味することができる。例えば、GPU内の単一の算術論理装置(ALU)ユニット又はベクトルプロセッサは、処理コアとみなすことができる。多コアプロセッサは、概して、相対的に数多くのコア、例えば、10個よりも多いコア、を有するマルチコアプロセッサを意味し、典型的には、それよりも少ない数のコアを有するマルチコアプロセッサを設計するために用いられる技法と異なるそれらを用いて設計される。マルチコアプロセッサは、ソフトウェアプログラムが単一チップ上の複数のコアにおいて並列で、例えば、同時並行して、実行するのを可能にすることによって性能向上を提供する。
並列プログラミングモデルは、プログラムを複数の処理コア上で同時並行して実行するのを可能にするように設計されるプログラミングモデルを意味する。プログラムは、マルチスレッドプログラムであることができ、その場合は、単一のスレッドが各処理コアにおいて動作することができる。幾つかの例では、単一の計算デバイスは、プログラムを実行するために用いられる全処理コアを含むことができる。その他の例では、プログラムを実行するために用いられる処理コアの一部を、同じタイプ又は異なるタイプの異なる計算デバイス上に配置することができる。
異なるISAを実装する異なる売り主によって製造される可能性がある異なるタイプの計算デバイスを含む異種のマルチコア計算プラットフォームの並列プログラミングのための共通言語仕様を提供するためにプラットフォーム横断、売り主横断型の、異種計算プラットフォーム、並列プログラミングモデルアプリケーションプログラミングインタフェース(API)を用いることができる。Open Computing Language(OpenCL(登録商標))は、プラットフォーム横断、売り主横断型の、異種計算プラットフォーム、並列プログラミングAPIの一例である。該APIは、GPUでのより一般化されたデータ処理を考慮して設計することができる。例えば、これらのAPIは、共通言語を介して拡張シェーダサブシステム能力を明らかにすることに加えて、非グラフィックス専用の方法でGPU内へのデータフロー及び制御経路を一般化することができる。しかしながら、現在では、該APIによって提供された命令セットは、GPUのハードウェアアーキテクチャに基づくものであり、従って、既存のGPUアーキテクチャと互換可能な機能に制限される。
本開示は、マルチプルプロセッサ計算プラットフォーム内で使用することができる通信技法について説明する。それらの技法は、幾つかの例では、コマンド待ち行列を用いてタスクを開始させるマルチプルプロセッサ計算プラットフォーム内でのメッセージ渡し(message passing)をサポートするために用いることができるソフトウェアインタフェースを提供することができる。それらの技法は、追加の例では、マルチプルプロセッサ計算プラットフォーム内での共有メモリプロセッサ間通信のために用いることができるソフトウェアインタフェースを提供する。さらなる例では、それらの技法は、GPUとホストCPUとの間でのメッセージ渡し及び/又は共有メモリ通信をサポートするためのハードウェアを含むグラフィックス処理装置(GPU)を提供することができる。
一例では、本開示は、1つ以上のプロセッサを含むホストデバイスについて説明する。デバイスは、1つ以上のプロセッサにおいて実行し、ホストデバイス上で実行中のプロセスから1つ以上の待ち行列内追加命令(enqueue instruction)を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れるように構成されたコマンド待ち行列インタフェース、をさらに含む。複数のコマンドは、ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するようにホストデバイスに命令する第1のコマンドを含む。複数のコマンドは、GPUでのタスクの実行を開始するようにホストデバイスに命令する第2のコマンドをさらに含む。デバイスは、1つ以上のプロセッサ上で実行し、GPU上でタスクが実行している間に及びホストデバイスで実行中のプロセスから1つ以上のメッセージ渡し命令を受信したことに応答してホストデバイスで実行中のプロセスとGPUで実行中のタスクとの間で1つ以上のメッセージを渡すように構成されたメッセージ渡しインタフェース、をさらに含む。
他の例では、本開示は、ホストデバイスの1つ以上のプロセッサで実行するコマンド待ち行列インタフェースを用いて、ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れることを含む方法について説明する。複数のコマンドは、ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するようにホストデバイスに命令する第1のコマンドを含む。複数のコマンドは、GPUでのタスクの実行を開始するようにホストデバイスに命令する第2のコマンドをさらに含む。方法は、ホストデバイスの1つ以上のプロセッサで実行するメッセージ渡しインタフェースを用いて、GPUでタスクが実行している間に及びホストデバイス上で実行中のプロセスから1つ以上のメッセージ渡し命令を受信したことに応答してホストデバイス上で実行中のプロセスとGPU上で実行中のタスクとの間で1つ以上のメッセージを渡すことをさらに含む。
他の例では、本開示は、ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れるための手段を含む装置について説明する。複数のコマンドは、ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するようにホストデバイスに命令する第1のコマンドを含む。複数のコマンドは、GPUでのタスクの実行を開始するようにホストデバイスに命令する第2のコマンドをさらに含む。装置は、GPU上でタスクが実行中に及びホストデバイスで実行中のプロセスから1つ以上のメッセージ渡し命令を受信したことに応答してホストデバイスで実行中のプロセスとGPUで実行中のタスクとの間で1つ以上のメッセージを渡すための手段をさらに含む。
他の例では、本開示は、ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れることを1つ以上のプロセッサに行わせる命令を含むコンピュータによって読み取り可能な記憶媒体について説明する。複数のコマンドは、ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するようにホストデバイスに命令する第1のコマンドを含む。複数のコマンドは、GPUでのタスクの実行を開始するようにホストデバイスに命令する第2のコマンドをさらに含む。コンピュータによって読み取り可能な記憶媒体は、GPUでタスクが実行している間に及びホストデバイスで実行中のプロセスから1つ以上のメッセージ渡し命令を受信したことに応答してホストデバイスで実行中のプロセスとGPUで実行中のタスクとの間で1つ以上のメッセージを渡すことを1つ以上のプロセッサに行わせる命令をさらに含む。
他の例では、本開示は、タスクを実行するように構成された1つ以上のプロセッサを含むグラフィックス処理装置(GPU)について説明する。GPUは、ホストデバイスによってアクセス可能な1つ以上のレジスタをさらに含む。GPUは、1つ以上のプロセッサでタスクが実行している間に及び1つ以上のプロセッサで実行中のタスクから1つ以上のメッセージ渡し命令を受信したことに応答して1つ以上のプロセッサで実行中のタスクとホストデバイスで実行中のプロセスとの間で、1つ以上のレジスタを介して、1つ以上のメッセージを渡すように構成されたメッセージ渡しモジュールをさらに含む。
他の例では、本開示は、グラフィックス処理装置(GPU)のメッセージ渡しモジュールを用いて、GPUで実行中のタスクから1つ以上のメッセージ渡し命令を受信することを含む方法について説明する。方法は、GPUでタスクが実行している間に及びGPUで実行中のタスクから1つ以上のメッセージ渡し命令を受信したことに応答してGPUで実行中のタスクとホストデバイスで実行中のプロセスとの間で、ホストデバイスによってアクセス可能であるGPU内の1つ以上のレジスタを介して、1つ以上のメッセージを渡すことをさらに含む。
他の例では、本開示は、グラフィックス処理装置(GPU)で実行中のタスクから1つ以上のメッセージ渡し命令を受信するための手段を含む装置について説明する。装置は、GPUでタスクが実行している間に及びGPUで実行中のタスクから1つ以上のメッセージ渡し命令を受信したことに応答してGPUで実行中のタスクとホストデバイスで実行中のプロセスとの間で、ホストデバイスによってアクセス可能であるGPU内の1つ以上のレジスタを介して、1つ以上のメッセージを渡すための手段をさらに含む。
他の例では、本開示は、グラフィックス処理装置(GPU)で実行中のタスクから1つ以上のメッセージ渡し命令を受信することを1つ以上のプロセッサに行わせる命令を備えるコンピュータによって読み取り可能な媒体について説明する。コンピュータによって読み取り可能な記憶媒体は、GPUでタスクが実行している間に及びGPUで実行中のタスクから1つ以上のメッセージ渡し命令を受信したことに応答してGPUで実行中のタスクとホストデバイスで実行中のプロセスとの間で、ホストデバイスによってアクセス可能であるGPU内の1つ以上のレジスタを介して、1つ以上のメッセージを渡すことを1つ以上のプロセッサに行わせる命令をさらに含む。
他の例では、本開示は、ホストデバイスの1つ以上のプロセッサで実行するメモリバッファインタフェースを用いて、ホストデバイスによって及びグラフィックス処理装置(GPU)によってアクセス可能である共有メモリスペースのために即時モード(immediate mode)がイネーブルにされるべきであるかどうかを指定する情報を含む命令を受信することを含む方法について説明する。方法は、メモリバッファインタフェースを用いて、即時モードがイネーブルにされるべきであるかどうかを指定する情報に基づいて共有メモリスペースのために即時モードを選択的にイネーブルにすることをさらに含む。
他の例では、本開示は、1つ以上のプロセッサを含むホストデバイスについて説明する。デバイスは、1つ以上のプロセッサで実行し、共有メモリスペースのために即時モードがイネーブルにされるべきかどうかを指定する情報を含む命令を受信するように、及び即時モードがイネーブルにされるべきかどうかを指定する情報に基づいて共有メモリスペースのために即時モードを選択的にイネーブルにするように構成されたメモリバッファインタフェース、をさらに含み、共有メモリスペースは、ホストデバイスによって及びグラフィックス処理装置(GPU)によってアクセス可能である。
他の例では、本開示は、ホストデバイスによって及びグラフィックス処理装置(GPU)によってアクセス可能である共有メモリスペースのために即時モードがイネーブルにされるべきかどうかを指定する情報を含む命令を受信するための手段を含む装置について説明する。装置は、即時モードがイネーブルにされるべきかどうかを指定する情報に基づいて共有メモリスペースのために即時モードを選択的にイネーブルにするための手段をさらに含む。
他の例では、本開示は、ホストデバイスによって及びグラフィックス処理装置(GPU)によってアクセス可能である共有メモリスペースのために即時モードがイネーブルにされるべきかどうかを指定する情報を含む命令を受信することを1つ以上のプロセッサに行わせる命令を備えるコンピュータによって読み取り可能な媒体について説明する。コンピュータによって読み取り可能な記憶媒体は、即時モードがイネーブルにされるべきかどうかを指定する情報に基づいて共有メモリスペースのために即時モードを選択的にイネーブルにすることを1つ以上のプロセッサに行わせる命令をさらに含む。
他の例では、本開示は、メモリと関連付けられたグラフィックス処理装置(GPU)キャッシュを含むGPUについて説明する。デバイスは、メモリのメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報を受信したことに応答してメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにGPUのキャッシュのキャッシングサービスを選択的に使用するように構成された1つ以上の処理モジュールをさらに含む。
他の例では、本開示は、メモリのメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報を受信したことに応答してメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにメモリと関連付けられたグラフィックス処理装置(GPU)キャッシュのキャッシングサービスを選択的に使用することを含む方法について説明する。
他の例では、本開示は、メモリと関連付けられたGPUキャッシュを含む装置について説明する。装置は、メモリのメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報を受信したことに応答してメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにGPUキャッシュのキャッシングサービスを選択的に使用するための手段をさらに含む。
他の例では、本開示は、メモリのメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにキャッシングサービスが使用されるべきかどうかを指定する情報を受信したことに応答してメモリスペースに関して読み取り動作及び書き込み動作のうちの少なくとも1つを実行するためにメモリと関連付けられたグラフィックス処理装置(GPU)キャッシュのキャッシングサービスを選択的に使用することを1つ以上のプロセッサに行わせる命令を備えるコンピュータによって読み取り可能な媒体について説明する。
本開示によるメッセージ渡し技法を実行するために用いることができる計算システム例を示したブロック図である。 本開示による図1の計算システムにおいて用いることができるGPU例を示したブロック図である。 本開示によるマルチプルプロセッサプラットフォーム環境におけるメッセージ渡し技法例を示した流れ図である。 本開示によるホストデバイスで実行中のプロセスによって出された送信命令を実行するための技法例を示した流れ図である。 本開示による図4において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による図4において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による、計算デバイス、例えば、GPU、で受信されたメッセージを処理するための技法例を示した流れ図である。 本開示による、計算デバイス、例えば、GPU、で実行中のタスクによって出された受信命令を実行するための技法例を示した流れ図である。 本開示による図8において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による図8において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による計算デバイス、例えば、GPU、で実行中のプロセスによって出された送信命令を実行するための技法例を示した流れ図である。 本開示による図11において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による図11において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示によるホストデバイスで実行中のプロセスによって出されたレジスタコールバックルーチン命令を実行するための技法例を示した流れ図である。 本開示による計算デバイスから受信された割り込みを処理するための技法例を示した流れ図である。 本開示による図15において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による図15において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示によるホストデバイスで実行中のプロセスによって出された読み取り命令を実行するための技法例を示した流れ図である。 本開示による図18において例示された技法の一部分を実装するために用いることができる技法例を示した流れ図である。 本開示による即時メモリオブジェクトの使用を容易にすることができる計算システム例を示したブロック図である。 本開示によるホストデバイスで実行中のプロセスによって出されたメモリオブジェクト生成命令を実行するための技法例を示した流れ図である。 本開示によるホストデバイスで実行中のプロセスによって出されたメモリオブジェクト生成命令を実行するための他の技法例を示した流れ図である。 本開示によるキャッシュドモード及び即時モード命令を処理するための技法例を示した流れ図である。 本開示によるキャッシュドモード及び即時モード命令を処理するための技法例を示した流れ図である。 本開示によるキャッシュドモード及び即時モード命令を処理するための技法例を示した流れ図である。 本開示によるキャッシュドモード及び即時モード命令を処理するための技法例を示した流れ図である。 本開示による図20の計算システムにおいて用いることができるGPU例を示したブロック図である。 本開示によるキャッシュドモード及び即時モード命令を処理するための技法例を示した流れ図である。 本開示によるホストデバイスで実行中のプロセスによって出されたメモリオブジェクト生成命令を実行するための他の技法例を示した流れ図である。 本開示による第1のコンパイル技法によってコンパイルされた命令のシーケンスをGPUがどのようにして処理することができるかを例示した流れ図である。 本開示によるタスクのためのソースコードをコンパイルするための技法例を示した流れ図である。 本開示によるキャッシングサービスを選択的に使用するためにGPUによって用いることができる技法例を示した流れ図である。
本開示は、マルチプルプロセッサ計算プラットフォーム内で用いることができる通信技法について説明する。それらの技法は、幾つかの例では、コマンド待ち行列を用いてタスクを開始させるマルチプルプロセッサ計算プラットフォーム内でのメッセージ渡しをサポートするために用いることができるソフトウェアインタフェースを提供することができる。それらの技法は、追加の例では、マルチプルプロセッサ計算プラットフォーム内での共有メモリプロセッサ間通信のために用いることができるソフトウェアインタフェースを提供することができる。さらなる例では、それらの技法は、グラフィックス処理装置(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、に戻る前に送信動作が完了されるまで待つことができる。非ブロッキング送信命令は、送信動作が完了されるまで待たずに呼び出しを行っているプロセスに戻ることができる。例えば、非ブロッキング送信命令は、特定の送信動作が成功であったかどうかを決定するために呼び出しを行っているプロセスによって後続して問い合わせることができるハンドルをその送信動作に戻すことができる。非ブロッキング送信命令は、失敗することがあり、失敗した場合は、呼び出しを行っているプロセスは、送信動作を再試行するために送信命令を再度出すことが必要な場合がある。
幾つかの例では、送信命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、clSendOutOfBandDataは、命令識別子であり、cl_device *deviceIdは、メッセージが送信されるべき特定のOpenCLデバイスを指定する入力パラメータであり、int OOB_dataは、送信されるメッセージの内容を指定する入力パラメータであり、bool blockingは、命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する入力パラメータである。ブロッキング命令の場合は、命令は、送信動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。非ブロッキング送信命令の場合は、命令は、呼び出しを行っているプロセスによる後続する状態問い合わせのためのハンドルパラメータを戻すことができる。
ホストメッセージ渡し命令は、幾つかの例では、非同期的な方法で指定されたデバイスからデータを受信するためにコールバックをレジスタに入れるようにホストデバイス12に命令するレジスタコールバックルーチン命令を含むことができる。例えば、レジスタコールバックルーチン命令は、GPU14で実行中のタスクがホストプロセス20にメッセージを送信していることを示す信号をGPU14から受信したことに応答してコールバックルーチンを呼び出すようにホストメッセージ渡しインタフェース26に命令することができる。レジスタコールバックルーチン命令は、コールバックルーチンをレジスタに入れるべき対象となる特定のデバイスを指定する第1の入力パラメータと、コールバックルーチンのメモリ記憶場所を指定する第2の入力パラメータと、を含むことができる。
幾つかの例では、レジスタコールバックルーチン命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、clRegisterOutOfBandDataCallbackは、命令識別子であり、cl_device *deviceIdは、メッセージが送信されるべき特定のOpenCLデバイスを指定する入力パラメータであり、void(*)(int) callBackPtrは、コールバックルーチンのメモリ記憶場所を指定する入力パラメータである。レジスタコールバックルーチン命令は、コールバックルーチンレジストレーション動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。
ホストメッセージ渡し命令は、幾つかの例では、指定されたデバイスからデータを読み取るのを試行するようにホストデバイス12に命令するポーリング命令を含むことができる。例えば、ポーリング命令は、GPU14で実行中のタスク28がメッセージを送信しているかどうかを示すメッセージ状態情報に関してGPU14をポーリングするようにホストメッセージ渡しインタフェース26に命令することができる。ポーリング命令は、ポーリングされるべき特定のデバイスを指定する入力パラメータと、存在する場合に、ポーリングの結果得られたデータを指定する出力パラメータと、を含むことができる。
幾つかの例では、ポーリング命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、clTryReadOutOfBandDataは、命令識別子であり、cl_device *deviceIdは、ポーリングされるべき特定のOpenCLデバイスを指定する入力パラメータであり、int*OOB_dataは、存在する場合に、ポーリングの結果得られたデータを指定する出力パラメータである。ポーリング命令は、ポーリング動作からデータが成功裏に得られたかどうかを示すパラメータを戻すことができる。
ホストプロセス20と同様に、タスク28は、計算デバイスによって実行される1つ以上のデバイスメッセージ渡し命令を含むことができる。デバイスメッセージ渡し命令は、ホストデバイス12に指定されたデータを送信するように計算デバイスに命令する送信命令を含むことができる。例えば、送信命令は、GPU14で実行中のタスク28からホストデバイス12で実行中のホストプロセス20にメッセージを送信するようにGPU14に命令することができる。
送信命令は、ブロッキング送信命令又は非ブロッキング送信命令のいずれかであることができる。送信命令は、幾つかの例では、送信命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する第1の入力パラメータを含むことができる。ブロッキング送信命令は、呼び出しを行っているプロセス、例えば、GPU14で実行中のタスク28、を停止させ、呼び出しを行っているプロセスに戻る前に送信動作が完了されるのを待つことができる。非ブロッキング送信命令は、送信動作が完了されるまで待たずに呼び出しを行っているプロセスに戻ることができる。例えば、非ブロッキング送信命令は、特定の送信動作が成功であったかどうかを決定するために後続して呼び出しを行っているプロセスによって問い合わせることができるハンドルをその送信動作に戻すことができる。非ブロッキング送信動作は、失敗することがあり、失敗した場合は、呼び出しを行っているプロセスは、送信動作を再試行するために送信命令を再度出すことが必要な場合がある。送信命令は、ホストデバイスに送信されるべきメッセージの内容を指定する第2の入力パラメータを含むことができる。
幾つかの例では、送信命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、send_oobdataは、命令識別子であり、bool blockingは、命令がブロッキング送信命令であるか又は非ブロッキング送信命令であるかを指定する入力パラメータであり、int dataは、送信されるメッセージの内容を指定する入力パラメータである。ブロッキング命令の場合は、命令は、送信動作が成功裏に完了されたかどうかを示すパラメータを戻すことができる。非ブロッキング命令の場合は、命令は、呼び出しを行っているプロセスによる後続する状態問い合わせのためのハンドルパラメータを戻すことができる。
デバイスメッセージ渡し命令は、幾つかの例では、ホストデバイス12からデータを受信するように計算デバイスに命令する受信命令を含むことができる。例えば、受信命令は、入手可能な場合にホストデバイス12で実行中のホストプロセス20からタスク28に送信されたメッセージをGPU14で実行中のタスク28に提供するようにGPU14、例えば、デバイスメッセージ渡しインタフェース30、に命令することができる。該命令は、ポーリングメカニズムをサポートするために用いることができる。
受信命令は、ブロッキング受信命令又は非ブロッキング受信命令のいずれかであることができる。受信命令は、幾つかの例では、受信命令がブロッキング受信命令であるか又は非ブロッキング受信命令であるかを指定する入力パラメータを含むことができる。ブロッキング受信命令は、呼び出しを行っているプロセス、例えば、GPU14で実行中のタスク28、を停止させ、呼び出しを行っているプロセスに戻る前にメッセージが入手可能になるまで待つことができる。非ブロッキング受信命令は、メッセージが入手可能になるまで待たずに呼び出しを行っているプロセスに戻ることができる。例えば、メッセージが入手可能である場合は、非ブロッキング送信命令は、そのメッセージを戻すことができる。しかしながら、メッセージが入手可能でない場合は、非ブロッキング受信命令は、失敗することがある。失敗した場合は、呼び出しを行っているプロセスは、受信動作を再試行するために受信命令を再度出すことが必要な場合がある。受信命令は、存在する場合に、受信動作の結果得られたデータを指定する出力パラメータを含むことができる。
幾つかの例では、受信命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、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に命令する画像オブジェクト生成命令を含むことができる。
幾つかの例では、バッファオブジェクト生成命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、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フラグであることができる。
さらなる例では、画像オブジェクト生成命令のためのインタフェースは、次の形態をとることができる。
Figure 2013537993
ここで、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のためのソースコードをコンパイルすることができる。
Figure 2013537993
ここで、“isImmediate”は、データが読み取られるメモリオブジェクトのためのブール即時モード属性を表し、“immediate_read(...)”は、即時モード読み取り命令を表し、“cached_read(...)”は、キャッシュドモード読み取り命令を表す。
GPUキャッシュ制御モジュール334は、使用されている場合は、GPUキャッシュ332からデータを読み取る前に、例えば、GPUキャッシュ332を無効化することによってimmediate_read(...)命令を処理することができる。GPUキャッシュ制御モジュール334は、例えば、読み取りを行う前にGPUキャッシュ332を無効化することなしに、通常の方法でGPUキャッシュからデータを読み取ることによってcached_read(...)命令を処理することができる。
第1のコンパイル技法例により、コンパイラは、タスク330のためのコンパイルされたコードが次の擬似コード例による書き込みシーケンスを含むような形でタスク330のためのソースコードをコンパイルすることができる。
Figure 2013537993
ここで、“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)非有形なコンピュータによって読み取り可能な通信媒体、例えば、一時的な信号又は搬送波、に対応することができる。

Claims (34)

  1. ホストデバイスであって、
    1つ以上のプロセッサと、
    前記1つ以上のプロセッサで実行し、前記ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れるように構成されたコマンド待ち行列インタフェースであって、前記複数のコマンドは、前記ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するように前記ホストデバイスに命令する第1のコマンドを含み、前記複数のコマンドは、前記GPUでのタスクの実行を開始するように前記ホストデバイスに命令する第2のコマンドをさらに含む、コマンド待ち行列インタフェースと、
    前記1つ以上のプロセッサで実行し、前記GPUでタスクが実行している間に及び前記ホストデバイスで実行中の前記プロセスから1つ以上のメッセージ渡し命令を受信したことに応答して前記ホストデバイスで実行中の前記プロセスと前記GPUで実行中の前記タスクとの間で1つ以上のメッセージを渡すように構成されたメッセージ渡しインタフェースと、を備える、ホストデバイス。
  2. 前記1つ以上のメッセージ渡し命令は、前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクにメッセージを送信するように前記メッセージ渡しインタフェースに命令する送信命令を備え、
    前記メッセージ渡しインタフェースは、前記送信命令を受信したことに応答して、前記GPUで前記タスクが実行している間に前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクに前記メッセージを送信するようにさらに構成される請求項1に記載のデバイス。
  3. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す信号を前記GPUから受信したことに応答してコールバックルーチンを呼び出すように前記メッセージ渡しインタフェースに命令するレジスタコールバックルーチン命令を備え、
    前記メッセージ渡しインタフェースは、前記GPUで実行中の前記タスクがメッセージを送信していることを示す前記信号を前記GPUから受信したことに応答して前記レジスタコールバックルーチン命令内で指定された前記コールバックルーチンの実行を開始するようにさらに構成される請求項1に記載のデバイス。
  4. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中のタスクがメッセージを送信しているかどうかを示すメッセージ状態情報に関して前記GPUをポーリングするように前記メッセージ渡しインタフェースに命令するポーリング命令を備え、
    前記メッセージ渡しインタフェースは、前記ポーリング命令を受信したことに応答して前記メッセージ状態情報に関して前記GPUをポーリングするようにさらに構成され、前記GPUで実行中の前記タスクがメッセージを送信していることを前記メッセージ状態情報が示す場合は、前記GPUから前記メッセージを入手するようにさらに構成される請求項1に記載のデバイス。
  5. 前記GPUで実行中の前記タスクは、前記GPUで実行中の前記タスクから前記ホストデバイスで実行中の前記プロセスにメッセージを送信するように前記GPUに命令する命令を含む請求項1に記載のデバイス。
  6. 前記GPUで実行中の前記タスクは、入手可能な場合に前記ホストデバイスで実行中の前記プロセスから前記タスクに送信されたメッセージを前記タスクに提供するように前記GPUに命令する命令を含む請求項1に記載のデバイス。
  7. 前記メッセージ渡しインタフェースは、前記コマンド待ち行列内にコマンドを入れずに前記1つ以上のメッセージ渡し命令を実行するようにさらに構成される請求項1に記載のデバイス。
  8. ホストデバイスの1つ以上のプロセッサで実行するコマンド待ち行列インタフェースを用いて、前記ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れることであって、前記複数のコマンドは、前記ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するように前記ホストデバイスに命令する第1のコマンドを含み、前記複数のコマンドは、前記GPUでのタスクの実行を開始するように前記ホストデバイスに命令する第2のコマンドをさらに含むことと、
    前記ホストデバイスの前記1つ以上のプロセッサで実行するメッセージ渡しインタフェースを用いて、前記GPUでタスクが実行している間に及び前記ホストデバイスで実行中の前記プロセスから1つ以上のメッセージ渡し命令を受信したことに応答して前記ホストデバイスで実行中の前記プロセスと前記GPUで実行中の前記タスクとの間で1つ以上のメッセージを渡すことと、を備える、方法。
  9. 前記1つ以上のメッセージ渡し命令は、前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクにメッセージを送信するように前記メッセージ渡しインタフェースに命令する送信命令を備え、
    前記方法は、前記メッセージ渡しインタフェースを用いて、前記GPUで前記タスクが実行している間に及び前記送信命令を受信したことに応答して前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクに前記メッセージを送信することをさらに備える請求項8に記載の方法。
  10. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す信号を前記GPUから受信したことに応答してコールバックルーチンを呼び出すように前記メッセージ渡しインタフェースに命令するレジスタコールバックルーチン命令を備え、
    前記方法は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す前記信号を前記GPUから受信したことに応答して前記レジスタコールバックルーチン命令内で指定された前記コールバックルーチンの実行を開始することをさらに構成される、請求項8に記載の方法。
  11. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信しているかどうかを示すメッセージ状態情報に関して前記GPUをポーリングするように前記メッセージ渡しインタフェースに命令するポーリング命令を備え、
    前記方法は、
    前記メッセージ渡しインタフェースを用いて、前記ポーリング命令を受信したことに応答して前記メッセージ状態情報に関して前記GPUをポーリングすることと、
    前記GPUで実行中の前記タスクがメッセージを送信していることを前記メッセージ状態情報が示す場合は、前記GPUから前記メッセージを入手することと、をさらに備える請求項8に記載の方法。
  12. 前記GPUで実行中の前記タスクは、前記GPUで実行中の前記タスクから前記ホストデバイスで実行中の前記プロセスにメッセージを送信するように前記GPUに命令する命令を含む請求項8に記載の方法。
  13. 前記GPUで実行中の前記タスクは、入手可能な場合に前記ホストデバイスで実行中の前記プロセスから前記タスクに送信されたメッセージを前記タスクに提供するように前記GPUに命令する命令を含む請求項8に記載の方法。
  14. 前記メッセージ渡しインタフェースを用いて、前記コマンド待ち行列内にコマンドを入れずに前記1つ以上のメッセージ渡し命令を実行することをさらに備える請求項8に記載の方法。
  15. ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れるための手段であって、前記複数のコマンドは、前記ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するように前記ホストデバイスに命令する第1のコマンドを含み、前記複数のコマンドは、前記GPUでのタスクの実行を開始するように前記ホストデバイスに命令する第2のコマンドをさらに含む手段と、
    前記GPUでタスクが実行している間に及び前記ホストデバイスで実行中の前記プロセスから1つ以上のメッセージ渡し命令を受信したことに応答して前記ホストデバイスで実行中の前記プロセスと前記GPUで実行中の前記タスクとの間で1つ以上のメッセージを渡すための手段と、を備える、装置。
  16. 前記1つ以上のメッセージ渡し命令は、前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクにメッセージを送信するように前記1つ以上のメッセージを渡すための前記手段に命令する送信命令を備え、
    前記装置は、前記送信命令を受信したことに応答して、前記GPUで前記タスクが実行している間に前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクに前記メッセージを送信するための手段をさらに備える請求項15に記載の装置。
  17. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す信号を前記GPUから受信したことに応答してコールバックルーチンを呼び出すように前記1つ以上のメッセージを渡すための前記手段に命令するレジスタコールバックルーチン命令を備え、
    前記装置は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す前記信号を前記GPUから受信したことに応答して前記レジスタコールバックルーチン命令内で指定された前記コールバックルーチンの実行を開始するための手段をさらに備える請求項15に記載の装置。
  18. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信しているかどうかを示すメッセージ状態情報に関して前記GPUをポーリングするように前記1つ以上のメッセージを渡すための前記手段に命令するポーリング命令を備え、
    前記装置は、
    前記ポーリング命令を受信したことに応答して前記メッセージ状態情報に関して前記GPUをポーリングするための手段と、
    前記GPUで実行中の前記タスクがメッセージを送信していることを前記メッセージ状態情報が示す場合は、前記GPUから前記メッセージを入手するための手段と、をさらに備える請求項15に記載の装置。
  19. ホストデバイスで実行中のプロセスから1つ以上の待ち行列内追加命令を受信したことに応答してコマンド待ち行列内に複数のコマンドを入れること、及び
    前記GPUでタスクが実行している間に及び前記ホストデバイスで実行中の前記プロセスから1つ以上のメッセージ渡し命令を受信したことに応答して前記ホストデバイスで実行中の前記プロセスと前記GPUで実行中の前記タスクとの間で1つ以上のメッセージを渡すことを1つ以上のプロセッサに行わせる命令を備えるコンピュータによって読み取り可能な媒体であって、前記複数のコマンドは、前記ホストデバイスと関連付けられた第1のメモリスペースとグラフィックス処理装置(GPU)と関連付けられた第2のメモリスペースとの間でデータを転送するように前記ホストデバイスに命令する第1のコマンドを含み、前記複数のコマンドは、前記GPUでのタスクの実行を開始するように前記ホストデバイスに命令する第2のコマンドをさらに含む、コンピュータによって読み取り可能な媒体。
  20. 前記1つ以上のメッセージ渡し命令は、前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクにメッセージを送信するように前記1つ以上のプロセッサに命令する送信命令を備え、
    前記コンピュータによって読み取り可能な媒体は、前記送信命令を受信したことに応答して、前記GPUで前記タスクが実行している間に前記ホストデバイスで実行中の前記プロセスから前記GPUで実行中の前記タスクに前記メッセージを送信することを前記1つ以上のプロセッサに行わせる命令をさらに備える請求項19に記載のコンピュータによって読み取り可能な媒体。
  21. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す信号を前記GPUから受信したことに応答してコールバックルーチンを呼び出すように前記1つ以上のプロセッサに命令するレジスタコールバックルーチン命令を備え、
    前記コンピュータによって読み取り可能な媒体は、前記GPUで実行中の前記タスクがメッセージを送信していることを示す前記信号を前記GPUから受信したことに応答して前記レジスタコールバックルーチン命令内で指定された前記コールバックルーチンの実行を開始することを前記1つ以上のプロセッサに行わせる命令をさらに備える請求項19に記載のコンピュータによって読み取り可能な媒体。
  22. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクがメッセージを送信しているかどうかを示すメッセージ状態情報に関して前記GPUをポーリングするように前記1つ以上のプロセッサに命令するポーリング命令を備え、
    前記コンピュータによって読み取り可能な媒体は、
    前記ポーリング命令を受信したことに応答して前記メッセージ状態情報に関して前記GPUをポーリングすること、及び
    前記GPUで実行中の前記タスクがメッセージを送信していることを前記メッセージ状態情報が示す場合は、前記GPUから前記メッセージを入手することを前記1つ以上のプロセッサに行わせる命令をさらに備える請求項19に記載のコンピュータによって読み取り可能な媒体。
  23. タスクを実行するように構成された1つ以上のプロセッサと、
    ホストデバイスによってアクセス可能な1つ以上のレジスタと、
    前記1つ以上のプロセッサで前記タスクが実行している間に及び前記1つ以上のプロセッサで実行中の前記タスクから1つ以上のメッセージ渡し命令を受信したことに応答して、前記1つ以上のプロセッサで実行中の前記タスクと前記ホストデバイスで実行中のプロセスとの間で、前記1つ以上のレジスタを介して、1つ以上のメッセージを渡すように構成されたメッセージ渡しモジュールと、を備える、グラフィックス処理装置(GPU)。
  24. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクから前記ホストデバイスで実行中の前記プロセスにメッセージを送信するように前記メッセージ渡しモジュールに命令する送信命令を備え、
    前記メッセージ渡しモジュールは、前記1つ以上のレジスタに前記メッセージと関連付けられたメッセージデータを格納するようにさらに構成される請求項23に記載のGPU。
  25. 前記1つ以上のメッセージ渡し命令は、入手可能な場合に前記ホストデバイスで実行中の前記プロセスから前記タスクに送信されたメッセージを前記タスクに提供するように前記メッセージ渡しモジュールに命令する受信命令を備え、
    前記メッセージ渡しモジュールは、前記1つ以上のレジスタから前記メッセージと関連付けられたメッセージデータを入手するようにさらに構成される請求項23に記載のGPU。
  26. グラフィックス処理装置(GPU)のメッセージ渡しモジュールを用いて、前記GPUで実行中のタスクから1つ以上のメッセージ渡し命令を受信することと、
    前記GPUで前記タスクが実行している間に及び前記GPUで実行中の前記タスクから前記1つ以上のメッセージ渡し命令を受信したことに応答して前記GPUで実行中の前記タスクと前記ホストデバイスで実行中のプロセスとの間で、ホストデバイスによってアクセス可能な前記GPU内の1つ以上のレジスタを介して、1つ以上のメッセージを渡すことと、を備える、方法。
  27. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクから前記ホストデバイスで実行中の前記プロセスにメッセージを送信するように前記メッセージ渡しモジュールに命令する送信命令を備え、
    前記方法は、前記1つ以上のレジスタに前記メッセージと関連付けられたメッセージデータを格納することをさらに備える請求項26に記載の方法。
  28. 前記1つ以上のメッセージ渡し命令は、入手可能な場合に前記ホストデバイスで実行中の前記プロセスから前記タスクに送信されたメッセージを前記タスクに提供するように前記メッセージ渡しモジュールに命令する受信命令を備え、
    前記方法は、前記1つ以上のレジスタから前記メッセージと関連付けられたメッセージデータを入手することをさらに備える請求項26に記載の方法。
  29. グラフィックス処理装置(GPU)で実行中のタスクから1つ以上のメッセージ渡し命令を受信するための手段と、
    前記GPUで前記タスクが実行している間に及び前記GPUで実行中の前記タスクから前記1つ以上のメッセージ渡し命令を受信したことに応答して前記GPUで実行中の前記タスクと前記ホストデバイスで実行中のプロセスとの間で、ホストデバイスによってアクセス可能な前記GPU内の1つ以上のレジスタを介して、1つ以上のメッセージを渡すための手段と、を備える、装置。
  30. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクから前記ホストデバイスで実行中の前記プロセスにメッセージを送信するように渡すための前記手段に命令する送信命令を備え、
    前記装置は、前記1つ以上のレジスタに前記メッセージと関連付けられたメッセージデータを格納するための手段をさらに備える請求項29に記載の装置。
  31. 前記1つ以上のメッセージ渡し命令は、入手可能な場合に前記ホストデバイスで実行中の前記プロセスから前記タスクに送信されたメッセージを前記タスクに提供するように渡すための前記手段に命令する受信命令を備え、
    前記装置は、前記1つ以上のレジスタから前記メッセージと関連付けられたメッセージデータを入手するための手段をさらに備える請求項29に記載の装置。
  32. グラフィックス処理装置(GPU)で実行中のタスクから1つ以上のメッセージ渡し命令を受信し、及び
    前記GPUで前記タスクが実行している間に及び前記GPUで実行中の前記タスクから前記1つ以上のメッセージ渡し命令を受信したことに応答して前記GPUで実行中の前記タスクと前記ホストデバイスで実行中のプロセスとの間で、ホストデバイスによってアクセス可能な前記GPU内の1つ以上のレジスタを介して、1つ以上のメッセージを渡すことを1つ以上のプロセッサに行わせる命令を備える、コンピュータによって読み取り可能な媒体。
  33. 前記1つ以上のメッセージ渡し命令は、前記GPUで実行中の前記タスクから前記ホストデバイスで実行中の前記プロセスにメッセージを送信するように前記1つ以上のプロセッサに命令する送信命令を備え、
    前記コンピュータによって読み取り可能な媒体は、前記1つ以上のレジスタに前記メッセージと関連付けられたメッセージデータを格納することを前記1つ以上のプロセッサに行わせる命令をさらに備える請求項32に記載のコンピュータによって読み取り可能な媒体。
  34. 前記1つ以上のメッセージ渡し命令は、入手可能な場合に前記ホストデバイスで実行中の前記プロセスから前記タスクに送信されたメッセージを前記タスクに提供するように1つ以上のプロセッサに命令する受信命令を備え、
    前記コンピュータによって読み取り可能な媒体は、前記1つ以上のレジスタから前記メッセージと関連付けられたメッセージデータを入手することを前記1つ以上のプロセッサに行わせる命令をさらに備える請求項32に記載のコンピュータによって読み取り可能な媒体。
JP2013530215A 2010-09-20 2011-09-19 マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法 Active JP5738998B2 (ja)

Applications Claiming Priority (9)

Application Number Priority Date Filing Date Title
US38457110P 2010-09-20 2010-09-20
US61/384,571 2010-09-20
US201161515182P 2011-08-04 2011-08-04
US61/515,182 2011-08-04
US13/235,236 US9645866B2 (en) 2010-09-20 2011-09-16 Inter-processor communication techniques in a multiple-processor computing platform
US13/235,266 2011-09-16
US13/235,266 US8937622B2 (en) 2010-09-20 2011-09-16 Inter-processor communication techniques in a multiple-processor computing platform
US13/235,236 2011-09-16
PCT/US2011/052196 WO2012040121A1 (en) 2010-09-20 2011-09-19 Inter-processor communication techniques in a multiple-processor computing platform

Publications (2)

Publication Number Publication Date
JP2013537993A true JP2013537993A (ja) 2013-10-07
JP5738998B2 JP5738998B2 (ja) 2015-06-24

Family

ID=45817338

Family Applications (2)

Application Number Title Priority Date Filing Date
JP2013530215A Active JP5738998B2 (ja) 2010-09-20 2011-09-19 マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法
JP2013530216A Expired - Fee Related JP5815712B2 (ja) 2010-09-20 2011-09-19 マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法

Family Applications After (1)

Application Number Title Priority Date Filing Date
JP2013530216A Expired - Fee Related JP5815712B2 (ja) 2010-09-20 2011-09-19 マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法

Country Status (10)

Country Link
US (3) US8937622B2 (ja)
EP (2) EP2619965A1 (ja)
JP (2) JP5738998B2 (ja)
KR (2) KR101564816B1 (ja)
CN (2) CN103119912B (ja)
BR (1) BR112013006488A2 (ja)
ES (1) ES2617303T3 (ja)
HU (1) HUE033041T2 (ja)
IN (1) IN2013MN00405A (ja)
WO (2) WO2012040122A1 (ja)

Cited By (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2017507398A (ja) * 2014-01-06 2017-03-16 ジョンソン コントロールズ テクノロジー カンパニーJohnson Controls Technology Company 複数のユーザインターフェース動作ドメインを有する車両
JP2019532427A (ja) * 2016-10-18 2019-11-07 アドバンスト・マイクロ・ディバイシズ・インコーポレイテッドAdvanced Micro Devices Incorporated トリガ動作を用いたgpuリモート通信

Families Citing this family (82)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
KR101332840B1 (ko) * 2012-01-05 2013-11-27 서울대학교산학협력단 병렬 컴퓨팅 프레임워크 기반의 클러스터 시스템, 호스트 노드, 계산 노드 및 어플리케이션 실행 방법
US8937622B2 (en) 2010-09-20 2015-01-20 Qualcomm Incorporated Inter-processor communication techniques in a multiple-processor computing platform
US9239793B2 (en) * 2011-12-13 2016-01-19 Ati Technologies Ulc Mechanism for using a GPU controller for preloading caches
US9170820B2 (en) * 2011-12-15 2015-10-27 Advanced Micro Devices, Inc. Syscall mechanism for processor to processor calls
JP5238876B2 (ja) * 2011-12-27 2013-07-17 株式会社東芝 情報処理装置及び情報処理方法
US8943516B2 (en) * 2012-05-31 2015-01-27 International Business Machines Corporation Mechanism for optimized intra-die inter-nodelet messaging communication
WO2013185015A2 (en) * 2012-06-08 2013-12-12 Advanced Micro Devices, Inc. System and method for providing low latency to applications using heterogeneous processors
US8869176B2 (en) 2012-11-09 2014-10-21 Qualcomm Incorporated Exposing host operating system services to an auxillary processor
US20140149528A1 (en) * 2012-11-29 2014-05-29 Nvidia Corporation Mpi communication of gpu buffers
KR102030733B1 (ko) 2013-01-02 2019-10-10 삼성전자주식회사 메모리 시스템 및 이의 구동 방법
US20140208134A1 (en) * 2013-01-21 2014-07-24 Texas Instruments Incorporated Host controller interface for universal serial bus (usb) power delivery
US9086813B2 (en) * 2013-03-15 2015-07-21 Qualcomm Incorporated Method and apparatus to save and restore system memory management unit (MMU) contexts
US9563561B2 (en) * 2013-06-25 2017-02-07 Intel Corporation Initiation of cache flushes and invalidations on graphics processors
CN103353851A (zh) * 2013-07-01 2013-10-16 华为技术有限公司 一种管理任务的方法和设备
US20150149745A1 (en) * 2013-11-25 2015-05-28 Markus Eble Parallelization with controlled data sharing
EP3084622A4 (en) * 2013-12-20 2018-02-28 Intel Corporation Execution offloading
US9632761B2 (en) * 2014-01-13 2017-04-25 Red Hat, Inc. Distribute workload of an application to a graphics processing unit
KR102100161B1 (ko) * 2014-02-04 2020-04-14 삼성전자주식회사 Gpu 데이터 캐싱 방법 및 그에 따른 데이터 프로세싱 시스템
US10055342B2 (en) * 2014-03-19 2018-08-21 Qualcomm Incorporated Hardware-based atomic operations for supporting inter-task communication
WO2015180668A1 (en) * 2014-05-28 2015-12-03 Mediatek Inc. Memory pool management method for sharing memory pool among different computing units and related machine readable medium and memory pool management apparatus
US10417052B2 (en) 2014-10-31 2019-09-17 Hewlett Packard Enterprise Development Lp Integrated heterogeneous processing units
GB2533768B (en) * 2014-12-19 2021-07-21 Advanced Risc Mach Ltd Cleaning a write-back cache
JP6338152B2 (ja) * 2015-04-16 2018-06-06 日本電信電話株式会社 通信装置、通信方法、及びプログラム
JP2016208105A (ja) * 2015-04-16 2016-12-08 日本電信電話株式会社 通信装置、通信方法、及びプログラム
US9996487B2 (en) * 2015-06-26 2018-06-12 Intel Corporation Coherent fabric interconnect for use in multiple topologies
US10664751B2 (en) 2016-12-01 2020-05-26 Via Alliance Semiconductor Co., Ltd. Processor with memory array operable as either cache memory or neural network unit memory
US11029949B2 (en) 2015-10-08 2021-06-08 Shanghai Zhaoxin Semiconductor Co., Ltd. Neural network unit
US10776690B2 (en) 2015-10-08 2020-09-15 Via Alliance Semiconductor Co., Ltd. Neural network unit with plurality of selectable output functions
US10228911B2 (en) 2015-10-08 2019-03-12 Via Alliance Semiconductor Co., Ltd. Apparatus employing user-specified binary point fixed point arithmetic
US11216720B2 (en) 2015-10-08 2022-01-04 Shanghai Zhaoxin Semiconductor Co., Ltd. Neural network unit that manages power consumption based on memory accesses per period
US11221872B2 (en) 2015-10-08 2022-01-11 Shanghai Zhaoxin Semiconductor Co., Ltd. Neural network unit that interrupts processing core upon condition
US10725934B2 (en) 2015-10-08 2020-07-28 Shanghai Zhaoxin Semiconductor Co., Ltd. Processor with selective data storage (of accelerator) operable as either victim cache data storage or accelerator memory and having victim cache tags in lower level cache wherein evicted cache line is stored in said data storage when said data storage is in a first mode and said cache line is stored in system memory rather then said data store when said data storage is in a second mode
US11226840B2 (en) 2015-10-08 2022-01-18 Shanghai Zhaoxin Semiconductor Co., Ltd. Neural network unit that interrupts processing core upon condition
US10380481B2 (en) 2015-10-08 2019-08-13 Via Alliance Semiconductor Co., Ltd. Neural network unit that performs concurrent LSTM cell calculations
US10366050B2 (en) 2015-10-08 2019-07-30 Via Alliance Semiconductor Co., Ltd. Multi-operation neural network unit
US9965417B1 (en) * 2016-01-13 2018-05-08 Xilinx, Inc. Use of interrupt memory for communication via PCIe communication fabric
KR101842764B1 (ko) * 2016-03-18 2018-03-28 연세대학교 산학협력단 하드웨어 가속기와 호스트 시스템 사이의 데이터 일관성 유지 장치 및 방법
JP6146508B1 (ja) 2016-03-31 2017-06-14 日本電気株式会社 同期処理ユニット、デバイス、システムおよび方法
KR102589298B1 (ko) * 2016-05-11 2023-10-13 삼성전자주식회사 그래픽스 프로세싱 장치 및, 그래픽스 프로세싱 장치에서 캐시 바이패스를 제어하는 방법
CN106127673B (zh) * 2016-07-19 2019-02-12 腾讯科技(深圳)有限公司 一种视频处理方法、装置及计算机设备
US10152243B2 (en) 2016-09-15 2018-12-11 Qualcomm Incorporated Managing data flow in heterogeneous computing
US10248565B2 (en) * 2016-09-19 2019-04-02 Qualcomm Incorporated Hybrid input/output coherent write
US10423876B2 (en) 2016-12-01 2019-09-24 Via Alliance Semiconductor Co., Ltd. Processor with memory array operable as either victim cache or neural network unit memory
US10438115B2 (en) 2016-12-01 2019-10-08 Via Alliance Semiconductor Co., Ltd. Neural network unit with memory layout to perform efficient 3-dimensional convolutions
US10395165B2 (en) 2016-12-01 2019-08-27 Via Alliance Semiconductor Co., Ltd Neural network unit with neural memory and array of neural processing units that collectively perform multi-word distance rotates of row of data received from neural memory
US10417560B2 (en) 2016-12-01 2019-09-17 Via Alliance Semiconductor Co., Ltd. Neural network unit that performs efficient 3-dimensional convolutions
US10430706B2 (en) 2016-12-01 2019-10-01 Via Alliance Semiconductor Co., Ltd. Processor with memory array operable as either last level cache slice or neural network unit memory
US10515302B2 (en) 2016-12-08 2019-12-24 Via Alliance Semiconductor Co., Ltd. Neural network unit with mixed data and weight size computation capability
KR102576707B1 (ko) 2016-12-26 2023-09-08 삼성전자주식회사 전자 시스템 및 그 동작 방법
US10565494B2 (en) 2016-12-31 2020-02-18 Via Alliance Semiconductor Co., Ltd. Neural network unit with segmentable array width rotator
US10565492B2 (en) 2016-12-31 2020-02-18 Via Alliance Semiconductor Co., Ltd. Neural network unit with segmentable array width rotator
US10140574B2 (en) 2016-12-31 2018-11-27 Via Alliance Semiconductor Co., Ltd Neural network unit with segmentable array width rotator and re-shapeable weight memory to match segment width to provide common weights to multiple rotator segments
US10586148B2 (en) 2016-12-31 2020-03-10 Via Alliance Semiconductor Co., Ltd. Neural network unit with re-shapeable memory
US10331532B2 (en) * 2017-01-19 2019-06-25 Qualcomm Incorporated Periodic non-intrusive diagnosis of lockstep systems
JP2018165913A (ja) * 2017-03-28 2018-10-25 富士通株式会社 演算処理装置、情報処理装置、及び演算処理装置の制御方法
US10503652B2 (en) 2017-04-01 2019-12-10 Intel Corporation Sector cache for compression
US10373285B2 (en) 2017-04-09 2019-08-06 Intel Corporation Coarse grain coherency
US10325341B2 (en) 2017-04-21 2019-06-18 Intel Corporation Handling pipeline submissions across many compute units
JP7032631B2 (ja) * 2017-07-04 2022-03-09 富士通株式会社 送受信システム、送受信システムの制御方法、及び送信装置
JP6907787B2 (ja) * 2017-07-28 2021-07-21 富士通株式会社 情報処理装置および情報処理方法
KR102403379B1 (ko) * 2017-09-12 2022-06-02 주식회사 코코링크 다중 gpu간 데이터 공유 방법
KR102384759B1 (ko) * 2017-11-13 2022-04-11 삼성전자주식회사 호스트 메모리 버퍼를 사용하기 위해 호스트 장치와 속성 정보를 공유하는 스토리지 장치 및 그것을 포함하는 전자 장치
US10303384B1 (en) * 2017-11-28 2019-05-28 Western Digital Technologies, Inc. Task readiness for queued storage tasks
KR102442921B1 (ko) 2017-12-11 2022-09-13 삼성전자주식회사 디지털 시그널 프로세서(dsp)의 태스크 관리 효율을 높일 수 있는 전자 장치
KR102533241B1 (ko) 2018-01-25 2023-05-16 삼성전자주식회사 적응적으로 캐시 일관성을 제어하도록 구성된 이종 컴퓨팅 시스템
US10671460B2 (en) 2018-02-05 2020-06-02 Micron Technology, Inc. Memory access communications through message passing interface implemented in memory systems
US10776281B2 (en) * 2018-10-04 2020-09-15 International Business Machines Corporation Snoop invalidate filter for distributed memory management unit to reduce snoop invalidate latency
US10761822B1 (en) * 2018-12-12 2020-09-01 Amazon Technologies, Inc. Synchronization of computation engines with non-blocking instructions
US10628342B1 (en) * 2018-12-20 2020-04-21 Dell Products, L.P. System and method for accelerating performance of non-volatile memory RAID stacks
CN111381958B (zh) * 2018-12-29 2022-12-09 上海寒武纪信息科技有限公司 通信装置、神经网络处理芯片、组合装置和电子设备
KR20200083048A (ko) * 2018-12-31 2020-07-08 삼성전자주식회사 폴링 시간을 예측하는 뉴럴 네트워크 시스템 및 이를 이용한 뉴럴 네트워크 모델 처리 방법
US20200342109A1 (en) * 2019-04-29 2020-10-29 Hewlett Packard Enterprise Development Lp Baseboard management controller to convey data
US11256423B2 (en) * 2019-10-14 2022-02-22 Western Digital Technologies, Inc. Efficiently identifying command readiness based on system state and data spread in multi queue depth environment
CN112764668A (zh) * 2019-11-01 2021-05-07 伊姆西Ip控股有限责任公司 扩展gpu存储器的方法、电子设备和计算机程序产品
JP2021149549A (ja) * 2020-03-19 2021-09-27 キオクシア株式会社 ストレージ装置およびアドレス変換テーブルのキャッシュ制御方法
US20210311897A1 (en) 2020-04-06 2021-10-07 Samsung Electronics Co., Ltd. Memory with cache-coherent interconnect
US20210373951A1 (en) * 2020-05-28 2021-12-02 Samsung Electronics Co., Ltd. Systems and methods for composable coherent devices
CN111897653A (zh) * 2020-07-30 2020-11-06 云知声智能科技股份有限公司 一种协同计算方法、装置、系统及介质
CN112100169B (zh) * 2020-08-05 2021-09-21 中科驭数(北京)科技有限公司 数据库交互数据编码方法及装置
CN112416851B (zh) * 2020-11-30 2023-07-18 中国人民解放军国防科技大学 一种可扩展的多核片上共享存储器
US11861758B2 (en) * 2021-03-12 2024-01-02 Nvidia Corporation Packet processing acceleration using parallel processing
CN115934385B (zh) * 2023-02-08 2023-05-23 苏州浪潮智能科技有限公司 一种多核核间通信方法、系统、设备及存储介质

Citations (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2005182825A (ja) * 2003-12-18 2005-07-07 Nvidia Corp マルチスレッド式マイクロプロセッサのスレッドにまたがるアウト・オブ・オーダー命令ディスパッチ
JP2010020755A (ja) * 2008-04-08 2010-01-28 Avid Technology Inc 複数のハードウェア・ドメイン、データ・タイプ、およびフォーマットの処理を統合し抽象化するフレームワーク

Family Cites Families (39)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JPS6242247A (ja) 1985-08-20 1987-02-24 Fujitsu Ltd キヤツシユメモリ制御方式
JP2736352B2 (ja) 1988-11-21 1998-04-02 日本電信電話株式会社 マルチプロセッサシステムにおけるキャッシュメモリ制御方法
JPH03127146A (ja) 1989-10-12 1991-05-30 Koufu Nippon Denki Kk 情報処理装置
JP3056606B2 (ja) 1993-02-17 2000-06-26 住友重機械工業株式会社 液状放射性薬剤小分け装置
US5604909A (en) 1993-12-15 1997-02-18 Silicon Graphics Computer Systems, Inc. Apparatus for processing instructions in a computing system
JPH0950399A (ja) 1995-08-10 1997-02-18 Fujitsu Ltd 多次元空間に配列されたデータの処理に適したキャッシュメモリシステム
US5854637A (en) 1995-08-17 1998-12-29 Intel Corporation Method and apparatus for managing access to a computer system memory shared by a graphics controller and a memory controller
US6088740A (en) 1997-08-05 2000-07-11 Adaptec, Inc. Command queuing system for a hardware accelerated command interpreter engine
US6212667B1 (en) * 1998-07-30 2001-04-03 International Business Machines Corporation Integrated circuit test coverage evaluation and adjustment mechanism and method
US6618759B1 (en) 2000-01-31 2003-09-09 Hewlett-Packard Development Company, L.P. Immediate mode computer graphics command caching
US6801208B2 (en) 2000-12-27 2004-10-05 Intel Corporation System and method for cache sharing
KR100960413B1 (ko) 2001-12-14 2010-05-28 엔엑스피 비 브이 데이터 처리 시스템, 통신 수단 및 데이터 처리 방법
US6891543B2 (en) * 2002-05-08 2005-05-10 Intel Corporation Method and system for optimally sharing memory between a host processor and graphics processor
US7958144B2 (en) 2002-08-30 2011-06-07 Boss Logic, Llc System and method for secure reciprocal exchange of data
CN100422974C (zh) * 2003-05-07 2008-10-01 皇家飞利浦电子股份有限公司 处理系统和用于传输数据的方法
US7015915B1 (en) 2003-08-12 2006-03-21 Nvidia Corporation Programming multiple chips from a command buffer
GB0319697D0 (en) 2003-08-21 2003-09-24 Falanx Microsystems As Method of and apparatus for differential encoding and decoding
TW200517825A (en) 2003-11-28 2005-06-01 First Int Computer Inc Over-clocking method used in VGA card and its application method for computer system
GB2409303B (en) 2003-12-18 2006-10-18 Advanced Risc Mach Ltd Inter-processor communication mechanism
US7023445B1 (en) * 2004-04-12 2006-04-04 Advanced Micro Devices, Inc. CPU and graphics unit with shared cache
US7305524B2 (en) * 2004-10-08 2007-12-04 International Business Machines Corporation Snoop filter directory mechanism in coherency shared memory system
US7302528B2 (en) * 2004-11-19 2007-11-27 Intel Corporation Caching bypass
US7583268B2 (en) 2005-11-10 2009-09-01 Via Technologies, Inc. Graphics pipeline precise interrupt method and apparatus
US8212832B2 (en) * 2005-12-08 2012-07-03 Ati Technologies Ulc Method and apparatus with dynamic graphics surface memory allocation
US7526634B1 (en) 2005-12-19 2009-04-28 Nvidia Corporation Counter-based delay of dependent thread group execution
US7353991B2 (en) 2006-02-21 2008-04-08 David Benjamin Esplin System and method for managing wireless point-of-sale transactions
US7814486B2 (en) * 2006-06-20 2010-10-12 Google Inc. Multi-thread runtime system
JP2008140078A (ja) 2006-11-30 2008-06-19 Toshiba Corp バスブリッジ装置、情報処理装置、およびデータ転送制御方法
JP5101128B2 (ja) 2007-02-21 2012-12-19 株式会社東芝 メモリ管理システム
US8031194B2 (en) * 2007-11-09 2011-10-04 Vivante Corporation Intelligent configurable graphics bandwidth modulator
US9035959B2 (en) 2008-03-28 2015-05-19 Intel Corporation Technique to share information among different cache coherency domains
GB2462860B (en) 2008-08-22 2012-05-16 Advanced Risc Mach Ltd Apparatus and method for communicating between a central processing unit and a graphics processing unit
US8675000B2 (en) 2008-11-07 2014-03-18 Google, Inc. Command buffers for web-based graphics rendering
US20100123717A1 (en) 2008-11-20 2010-05-20 Via Technologies, Inc. Dynamic Scheduling in a Graphics Processor
US8589629B2 (en) * 2009-03-27 2013-11-19 Advanced Micro Devices, Inc. Method for way allocation and way locking in a cache
US9354944B2 (en) * 2009-07-27 2016-05-31 Advanced Micro Devices, Inc. Mapping processing logic having data-parallel threads across processors
US8400458B2 (en) 2009-09-09 2013-03-19 Hewlett-Packard Development Company, L.P. Method and system for blocking data on a GPU
US8859133B2 (en) 2010-08-17 2014-10-14 GM Global Technology Operations LLC Repeating frame battery with compression joining of cell tabs to welded connection terminals
US8937622B2 (en) 2010-09-20 2015-01-20 Qualcomm Incorporated Inter-processor communication techniques in a multiple-processor computing platform

Patent Citations (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2005182825A (ja) * 2003-12-18 2005-07-07 Nvidia Corp マルチスレッド式マイクロプロセッサのスレッドにまたがるアウト・オブ・オーダー命令ディスパッチ
JP2010020755A (ja) * 2008-04-08 2010-01-28 Avid Technology Inc 複数のハードウェア・ドメイン、データ・タイプ、およびフォーマットの処理を統合し抽象化するフレームワーク

Non-Patent Citations (4)

* Cited by examiner, † Cited by third party
Title
CSNB201100917001; 奥薗 隆司: 「OpenCL入門」 第1版, 20100520, p.15-32, 株式会社 秀和システム *
CSND199800708010; 川村 竹弥,小野澤 隆: '「研究・DSP用リアルタイムOS SPOXのシステム・インターフェース」' インターフェース 第20巻 第9号(通巻第208号), 19940901, p.139-154, CQ出版株式会社 *
JPN6014017559; 奥薗 隆司: 「OpenCL入門」 第1版, 20100520, p.15-32, 株式会社 秀和システム *
JPN6014048834; 川村 竹弥,小野澤 隆: '「研究・DSP用リアルタイムOS SPOXのシステム・インターフェース」' インターフェース 第20巻 第9号(通巻第208号), 19940901, p.139-154, CQ出版株式会社 *

Cited By (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JP2017507398A (ja) * 2014-01-06 2017-03-16 ジョンソン コントロールズ テクノロジー カンパニーJohnson Controls Technology Company 複数のユーザインターフェース動作ドメインを有する車両
JP2019532427A (ja) * 2016-10-18 2019-11-07 アドバンスト・マイクロ・ディバイシズ・インコーポレイテッドAdvanced Micro Devices Incorporated トリガ動作を用いたgpuリモート通信
US10936533B2 (en) 2016-10-18 2021-03-02 Advanced Micro Devices, Inc. GPU remote communication with triggered operations

Also Published As

Publication number Publication date
US9645866B2 (en) 2017-05-09
JP5815712B2 (ja) 2015-11-17
CN103109274B (zh) 2016-08-17
JP2013546035A (ja) 2013-12-26
EP2619666B1 (en) 2016-11-23
US9626234B2 (en) 2017-04-18
US8937622B2 (en) 2015-01-20
WO2012040121A1 (en) 2012-03-29
US20150097849A1 (en) 2015-04-09
EP2619965A1 (en) 2013-07-31
KR101564816B1 (ko) 2015-10-30
US20120069029A1 (en) 2012-03-22
WO2012040122A1 (en) 2012-03-29
HUE033041T2 (hu) 2017-11-28
US20120069035A1 (en) 2012-03-22
CN103119912B (zh) 2016-06-01
CN103109274A (zh) 2013-05-15
ES2617303T3 (es) 2017-06-16
IN2013MN00405A (ja) 2015-05-29
KR20130094322A (ko) 2013-08-23
JP5738998B2 (ja) 2015-06-24
BR112013006488A2 (pt) 2016-07-26
EP2619666A1 (en) 2013-07-31
CN103119912A (zh) 2013-05-22
KR101564815B1 (ko) 2015-10-30
KR20130060337A (ko) 2013-06-07

Similar Documents

Publication Publication Date Title
JP5738998B2 (ja) マルチプルプロセッサ計算プラットフォームにおけるプロセッサ間通信技法
US11354251B2 (en) Apparatus and methods implementing dispatch mechanisms for offloading executable functions
JP6649267B2 (ja) タスク間通信をサポートするためのハードウェアベースのアトミック動作
US11119944B2 (en) Memory pools in a memory model for a unified computing system
US20120162234A1 (en) Device Discovery and Topology Reporting in a Combined CPU/GPU Architecture System
US9299121B2 (en) Preemptive context switching
KR101900436B1 (ko) 결합된 cpu/gpu 아키텍처 시스템에서의 디바이스의 발견 및 토폴로지 보고
US20130135327A1 (en) Saving and Restoring Non-Shader State Using a Command Processor
US20130141446A1 (en) Method and Apparatus for Servicing Page Fault Exceptions
WO2016067496A1 (ja) 情報処理装置
EP4020226A1 (en) Inter-node messaging controller
US9378139B2 (en) System, method, and computer program product for low latency scheduling and launch of memory defined tasks

Legal Events

Date Code Title Description
A621 Written request for application examination

Free format text: JAPANESE INTERMEDIATE CODE: A621

Effective date: 20130521

A977 Report on retrieval

Free format text: JAPANESE INTERMEDIATE CODE: A971007

Effective date: 20140415

A131 Notification of reasons for refusal

Free format text: JAPANESE INTERMEDIATE CODE: A131

Effective date: 20140507

A601 Written request for extension of time

Free format text: JAPANESE INTERMEDIATE CODE: A601

Effective date: 20140730

A602 Written permission of extension of time

Free format text: JAPANESE INTERMEDIATE CODE: A602

Effective date: 20140806

A131 Notification of reasons for refusal

Free format text: JAPANESE INTERMEDIATE CODE: A131

Effective date: 20141118

TRDD Decision of grant or rejection written
A01 Written decision to grant a patent or to grant a registration (utility model)

Free format text: JAPANESE INTERMEDIATE CODE: A01

Effective date: 20150324

A61 First payment of annual fees (during grant procedure)

Free format text: JAPANESE INTERMEDIATE CODE: A61

Effective date: 20150422

R150 Certificate of patent or registration of utility model

Ref document number: 5738998

Country of ref document: JP

Free format text: JAPANESE INTERMEDIATE CODE: R150

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250

R250 Receipt of annual fees

Free format text: JAPANESE INTERMEDIATE CODE: R250