JP5711853B2 - 異種コアの自動カーネル移行 - Google Patents

異種コアの自動カーネル移行 Download PDF

Info

Publication number
JP5711853B2
JP5711853B2 JP2014511476A JP2014511476A JP5711853B2 JP 5711853 B2 JP5711853 B2 JP 5711853B2 JP 2014511476 A JP2014511476 A JP 2014511476A JP 2014511476 A JP2014511476 A JP 2014511476A JP 5711853 B2 JP5711853 B2 JP 5711853B2
Authority
JP
Japan
Prior art keywords
kernel
processor core
code
microarchitecture
computational
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.)
Active
Application number
JP2014511476A
Other languages
English (en)
Other versions
JP2014513853A5 (ja
JP2014513853A (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.)
Advanced Micro Devices Inc
Original Assignee
Advanced Micro Devices 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 Advanced Micro Devices Inc filed Critical Advanced Micro Devices Inc
Publication of JP2014513853A publication Critical patent/JP2014513853A/ja
Publication of JP2014513853A5 publication Critical patent/JP2014513853A5/ja
Application granted granted Critical
Publication of JP5711853B2 publication Critical patent/JP5711853B2/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/48Program initiating; Program switching, e.g. by interrupt
    • G06F9/4806Task transfer initiation or dispatching
    • G06F9/4843Task transfer initiation or dispatching by program, e.g. task dispatcher, supervisor, operating system
    • G06F9/485Task life-cycle, e.g. stopping, restarting, resuming execution
    • G06F9/4856Task life-cycle, e.g. stopping, restarting, resuming execution resumption being on a different machine, e.g. task migration, virtual machine migration
    • 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/50Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5061Partitioning or combining of resources
    • G06F9/5066Algorithms for mapping a plurality of inter-dependent sub-tasks onto a plurality of physical CPUs

Landscapes

  • Engineering & Computer Science (AREA)
  • Software Systems (AREA)
  • Theoretical Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • General Physics & Mathematics (AREA)
  • Devices For Executing Special Programs (AREA)
  • Advance Control (AREA)

Description

本発明は、コンピューティングシステム、より具体的には、複数の異種コア間で作業単位の実行を自動的に移行させることに関する。
コンピュータシステムのスループットを増加させるために、タスクの並列化が用いられている。この目的を達成するために、コンパイラは、プログラムコードから並列化タスクを抽出し、システムハードウェア上で並行して実行し得る。単一コアアーキテクチャでは、単一のコアは、マルチスレッディングを行うように構成されたディープパイプラインを含んでもよい。ハードウェア上での並列実行をさらに増加させるために、マルチコアアーキテクチャは、複数の汎用コアを含んでもよい。この種類のアーキテクチャは、同種マルチコアアーキテクチャと呼ばれてもよく、単一コアアーキテクチャよりも高い命令スループットを提供し得る。
いくつかのソフトウェアアプリケーションは、度々、並列タスクに分割されない場合がある。加えて、特定のタスクは、汎用コア上で効率的に実行しない場合がある。計算集約的タスクの特定の命令は、共有リソースの不均衡な配分を引き起こし得、共有リソースの割り当て解除を遅延する。そのような特定のタスクの例には、暗号作成、ビデオグラフィックスレンダリング、およびガーベジコレクションが挙げられ得る。
従来の汎用コアの性能限界を打開するために、コンピュータシステムは、特定のタスクを特殊用途のハードウェアにオフロードし得る。このハードウェアは、単一命令複数データ(SIMD)並列アーキテクチャ、フィールドプログラマブルゲートアレイ(FPGA)および他の特殊のコアを含み得る。異なる種類のコアを有する一種のアーキテクチャは、異種マルチコアアーキテクチャと称され得る。この種のアーキテクチャは、タスクのスケジューリングに応じて、同種マルチコアアーキテクチャよりも高い命令スループットを提供し得る。
多くの場合、特定のソフトウェアアプリケーションは、それぞれの作業項目または並列関数呼び出しの実行がその内部でデータ依存性である、データ並列性を有する。例えば、第1の作業項目は、第2の作業項目から独立したデータであってよく、第1の作業項目および第2の作業項目の各々は、SIMDマイクロアーキテクチャを有するコア内の別々のパス上にスケジュールされる。しかしながら、第1の作業項目および第2の作業項目の各々で実行される命令の量は、データ依存性があり得る。分岐命令として実装される条件付きテストは、第1の作業項目を通過し得るが、それぞれの作業項目のデータに依存した第2の作業項目に対して失敗し得る。
第2の作業項目が実行を停止して待機し、第1の作業項目が進行中の実行を続けるため、並列実行の効率性は低減するおそれがある。わずかな作業項目のみがテスト通過のために実行を継続する一方で、大半の作業項目がテスト失敗のために休止状態にあるときに、非効率性が増大する。異種マルチコアアーキテクチャにおけるOSスケジューラによる作業項目の効率的な機能性一致割り当て後に、システム性能は、特定のソフトウェアアプリケーションのデータ依存性挙動のために、さらに低減するおそれがある。
複数の異種コア間で作業単位の実行を自動的に移行させるためのシステムおよび方法が企図される。
一実施形態において、コンピューティングシステムは、第1のマイクロアーキテクチャを有する第1のプロセッサコアと、第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを有する第2のプロセッサコアとを含む。一実施形態において、第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャであって、第2のマイクロアーキテクチャは、汎用マイクロアーキテクチャである。コンピューティングシステムは、第1および第2のプロセッサコアの各々に接続されたメモリを含む。メモリは、1つ以上の計算カーネルまたは関数呼び出しを含むコンピュータプログラムを記憶する。コンパイラが所与の関数呼び出しの命令をトラバース(traverse)するときに、コンパイラは、関数呼び出しの実行が、所与の位置で異なるプロセッサコアに移行することを予測するように構成されている。コンパイラは、所与の位置での関数呼び出しの実行と関連付けられたライブ値の移動を支援するデータ構造を作成する。そのようなライブ値は、「コンテキスト」と称され得る。
オペレーティングシステム(OS)内のスケジューラは、プログラム順における所与の位置前に、少なくともコードを第1のプロセッサコアにスケジュールする。移行条件が満たされているという指標の受信に応じて、OSスケジューラは、第2のプロセッサコアによるアクセスのために、データ構造によって示された位置にライブ値を移動させ、且つ、プログラム順における所与の位置後に、コードを第2のプロセッサにスケジュールする。移行条件が満たされているか否かを判定するために、第1および第2のプロセッサコアの各々は、出口点に到達した関数呼び出しの反復の並列実行数が所与の閾値を超えるか否かを判定するように構成されている。
これらおよび他の実施形態は、以下の説明および図面を参照してさらに理解される。
異種マルチコアアーキテクチャを備えた例示的な処理ノードの一実施形態の一般化されたブロック図である。 計算カーネルを定義するソースコードの一実施形態の一般化されたブロック図である。 条件文を有する計算カーネルを定義するソースコードの一実施形態の一般化されたブロック図である。 ハードウェアリソースと計算カーネルとの間でスケジュールされた割り当ての一実施形態の一般化されたブロック図である。 2種類のプロセッサコアのマイクロアーキテクチャの論理レイアウトの一実施形態の一般化されたブロック図である。 汎用パイプライン実行フローの一実施形態の一般化されたブロック図である。 SIMDパイプライン実行フローの一実施形態の一般化されたブロック図である。 SIMDパイプライン実行フローの一実施形態の別の一般化されたブロック図である。 移行タグ付け分岐を有するプログラムコードの一実施形態の一般化されたブロック図である。 計算カーネル移行のためにコードを計測する方法の一実施形態を説明する一般化されたフロー図である。 プログラム実行中に計算カーネルを移行させる方法の一実施形態を説明する一般化されたフロー図である。
本発明の様々な修正および代替形態が可能であるが、特定の実施形態は、図面において一例として示されており、本明細書において詳細に説明される。しかしながら、図面および図面の詳細説明が、本発明を開示された特定の形態に限定するように意図されず、それとは逆に、本発明は、添付の特許請求の範囲によって定義される本発明の精神および範囲内に収まるすべての修正、等価物および代替物を包含すると理解されるべきである。
以下の説明では、本発明の完全なる理解を提供するために、多数の具体的詳細が記載されている。しかしながら、当業者であれば、これらの具体的詳細を伴わずに、本発明が実践され得ることを認識するはずである。場合によっては、本発明を曖昧にすることを回避するために、周知の回路、構造および技法が詳細に示されていない。
図1を参照すると、異種マルチコアアーキテクチャを備えた例示的な処理ノード110の一実施形態が示されている。処理ノード110は、1つ以上のプロセッサコア112と、関連キャッシュメモリサブシステム114とを含み得る1つ以上の処理ユニット115を含んでもよい。一実施形態では、プロセッサコア112は、汎用マイクロアーキテクチャを利用する。
また、処理ノード110は、1つ以上のプロセッサコア172と、データ記憶バッファ174とを備え得る1つ以上の処理ユニット170を含んでもよい。プロセッサコア172は、プロセッサコア112のミラーリングされたシリコンイメージでなくてもよい。プロセッサコア172は、プロセッサコア112によって用いられるマイクロアーキテクチャとは異なるマイクロアーキテクチャを有し得る。一実施形態では、プロセッサコア172は、プロセッサコア112と同一のプロセッサ群の異なる世代のものであってもよい。別の実施形態では、プロセッサコア172は、プロセッサコア112の電圧および/または周波数を拡大/縮小したものであってもよい。言い換えれば、プロセッサコア172は、同一の機能性および命令セットアーキテクチャ(ISA)、同一のクロック周波数、同一のキャッシュサイズ、同一のメモリモデル等を備えたプロセッサコア112のシリコンコピーではない。
プロセッサコア172のマイクロアーキテクチャの参照を続けると、さらに別の実施形態では、プロセッサコア172は、計算集約的タスクのための高い命令スループットを提供するマイクロアーキテクチャを備えてもよい。プロセッサコア172は、並列アーキテクチャを有し得る。例えば、プロセッサコア172は、単一命令複数データ(SIMD)コアであってもよい。SIMDコアの例は、グラフィックス処理ユニット(GPU)、デジタル信号処理(DSP)コア等を含む。一実施形態では、処理ノード110は、単一命令セットアーキテクチャ(ISA)を備える。典型的には、当技術分野において周知であるように、単一ISAマルチコアアーキテクチャは、チップマルチプロセッサ(CMP)のためにより高い出力およびスループット性能を提供することが示されている。
処理ノード110上の高い命令スループットは、ソフトウェアアプリケーションのスレッドが効率的にスケジューリングされたときに、所与の電力限界内の測定された電力消費とともに達成されてもよい。スレッドは、プロセッサコア112および172のランタイムのハードウェアリソースの少なくとも一部に基づき、各スレッドが最高命令スループットを有する方式で、プロセッサコア112および172のうち1つでスケジューリングされてもよい。
処理ノード110内の構成要素の参照を続けると、処理ノード110は、メモリコントローラ120と、インターフェース論理140とを含んでもよい。一実施形態では、処理ノード110の説明される機能性は、単一の集積回路に組み込まれる。一実施形態では、プロセッサコア112は、所定の汎用命令セットに従って命令を実行するための回路を含む。例えば、SPARC(登録商標)命令セットアーキテクチャ(ISA)が選択されてもよい。代替として、x86、x86−64(登録商標)、Alpha(登録商標)、PowerPC(登録商標)、MIPS(登録商標)、PA−RISC(登録商標)または任意の他の命令セットアーキテクチャが選択されてもよい。概して、プロセッサコア112は、データおよび命令のそれぞれのために、キャッシュメモリサブシステム114にアクセスする。要求されたブロックが、キャッシュメモリサブシステム114内または共有キャッシュメモリサブシステム118内で検出されない場合には、読取要求が生成され、欠落したブロックがマップされたノード内のメモリコントローラに当該読取要求が伝送されてもよい。
一実施形態では、処理ユニット170は、グラフィックス処理ユニット(GPU)である。現代のGPUは、コンピュータグラフィックスを操作および表示することに非常に有能である。GPUの高度並列構造は、様々な複雑なアルゴリズムに対して、処理ユニット115等の汎用中央処理ユニット(CPU)よりも効果的である。典型的には、GPUは、グラフィックスおよびビデオに使用される計算を実行し、CPUは、グラフィックス単独よりもより多くのシステムプロセスのための計算を実行する。従来のGPUは、画像レンダリングアプリケーションにおいて高いスループットを達成するために、非常に幅広い単一命令複数データ(SIMD)アーキテクチャを利用している。そのようなアプリケーションは、概して、多数のオブジェクト(頂点またはピクセル)上で、頂点シェーダまたはピクセルシェーダ等の同一のプログラムを実行することを必要とする。各オブジェクトは、他のオブジェクトとは独立して処理されるが、同一の一連の演算が使用されるため、SIMDアーキテクチャは、相当の性能強化を提供する。また、GPUは、非図式計算のためにも考慮されている。
一実施形態では、GPU170は、ビデオカード上に配置され得る。別の実施形態では、GPU170は、マザーボード上に統合されてもよい。さらに別の実施形態では、処理ノード110の説明される機能性は、単一の集積回路上に組み込まれてもよい。そのような実施形態では、CPU115およびGPU170は、異なる設計センターからの専有コアであってもよい。また、GPU170は、現在、インターフェース140を介してオフチップでメモリアクセスを行うよりも、処理ノード110からメモリコントローラ120を介して、ローカルメモリ114および118と、メインメモリとの両方に直接アクセスすることが可能であってもよい。この実施形態は、GPU170のメモリアクセスの待ち時間をより低下させてもよく、それが、より高い性能につながり得る。
図1の処理ノード110の構成要素の参照を続けると、キャッシュサブシステム114および118は、データのブロックを記憶するように構成された高速キャッシュメモリを備えてもよい。キャッシュメモリサブシステム114は、個々のプロセッサコア112内に統合されてもよい。代替として、キャッシュメモリサブシステム114は、所望に応じて、バックサイドキャッシュ構成またはインライン構成で、プロセッサコア114に連結されてもよい。さらにまた、キャッシュメモリサブシステム114は、キャッシュの階層として実装されてもよい。(階層内で)プロセッサコア112に最も近く位置するキャッシュは、所望であれば、プロセッサコア112内に統合されてもよい。一実施形態では、キャッシュメモリサブシステム114の各々はL2キャッシュ構造を表し、共有キャッシュサブシステム118はL3キャッシュ構造を表す。キャッシュメモリサブシステム114および共有キャッシュメモリサブシステム118の両方は、対応するキャッシュコントローラに連結されたキャッシュメモリを含んでもよい。
概して、パケット処理論理116は、処理ノード110が連結されたリンク上で受信した制御パケットに応答するように構成されており、プロセッサコア112および/またはキャッシュメモリサブシステム114に応じて制御パケットを生成するように構成されており、サービスのためのメモリコントローラ120によって選択されたトランザクションに応じてプローブコマンドおよび応答パケットを生成するように構成されており、ノード110が中間ノードであるパケットを、インターフェース論理140を通して他のノードへ送るように構成されている。インターフェース論理140は、パケットを受信し、パケットを、パケット処理論理116によって使用される内部クロックに同期させる論理を含んでもよい。
ここで図2を参照すると、計算カーネルを利用するソースコードの一実施形態が示されている。OpenCL(登録商標)(Open Computing Language)は、異種計算のためのアプリケーションプログラミングインターフェース(API)の一例である。OpenCL(登録商標)は、それぞれOpenCL(登録商標)デバイスと関連付けられた実行待ち行列を定義するC風の言語を含む。OpenCL(登録商標)デバイスは、CPU、GPU、または異種マルチコアアーキテクチャ内の少なくとも1つのプロセッサコアを備えた他のユニットであってもよい。関数呼び出しは、OpenCL(登録商標)カーネル、または単に「計算カーネル」と呼ばれてもよい。OpenCL(登録商標)フレームワークは、ゲーム、娯楽、科学および医療分野で使用される多種多様のデータ並列アプリケーションの計算性能を向上させてもよい。異種アーキテクチャについては、コンピュータプログラムは、典型的に、計算カーネルおよび内部関数の集合を備える。ソフトウェアプログラマは、計算カーネルを定義し得る一方で、内部関数は、所与のライブラリ内で定義されてもよい。
データ並列ソフトウェアアプリケーションについては、N次元計算ドメインは、「実行ドメイン」の組織化を定義し得る。また、N次元計算ドメインは、N次元グリッドまたはN次元範囲(「NDRange」)と呼ばれてもよい。NDRangeは、1、2または3次元空間であってもよい。いくつかの実施形態が、3次元を超えるデータを可能にし得ることに留意する。この次元空間は、インデックス空間と呼ばれてもよい。例えば、ソフトウェアアプリケーションは、画像ファイル等のデータの2次元(2D)アレイのデータにデータ処理を行ってもよい。ソフトウェアアプリケーションは、2次元画像のピクセルごとに、または2次元マトリックスの要素ごとに、ソフトウェアプログラマによって開発されたアルゴリズムを行い得る。所与の計算カーネルは、インデックス空間(NDRange)にわたって呼び出され得る。他の実施形態において、ソフトウェアアプリケーションは、3次元格子上での静電電位マッピングおよび高分子モデリングに使用される直接クーロン総和のために、データ並列プログラミングを利用するアルゴリズムを含み得る。
典型的には、コンパイル後、各計算カーネルの引数およびパラメータが設定される。加えて、関連メモリオブジェクトおよびバッファが作成される。計算カーネルの所与のインスタンスは、独自のソフトウェアスレッドとして実行されてもよい。しかしながら、計算カーネルは、フォークを作成する制御フロー転送命令を含み得るが、コンピュータプログラム内のフォークは、典型的に、共通の定義によって、ソフトウェアスレッドを作成する。インデックス空間内の所与の点における計算カーネルの所与のインスタンスは、「作業項目」と呼ばれてもよい。また、作業項目は、作業単位と呼ばれてもよい。作業単位は、2次元画像の所与のピクセル(所与のインデックス)に対応するデータのレコード上で、計算カーネル内の1つ以上の命令とともに動作し得る。典型的には、作業単位は、関連する一意の識別子(ID)を有する。別の実施例では、「Hello World」という文字列を処理する入門コンピュータプログラムは、文字列内の各文字を計算するための1つの作業単位を有し得る。
NDRangeは、十分なハードウェアサポートがある場合に並行して実行する作業単位の総数を定義し得る。例えば、NDRangeは、280の作業単位を定義し得るが、GPUは、任意の所与の時に64個の作業単位の同時実行をサポートし得る。作業単位の総数は、全体的な作業サイズを定義し得る。当業者に周知であるように、作業単位は、作業グループにさらにグループ化されてもよい。各作業グループは、一意の識別子(ID)を有し得る。所与の作業グループ内の作業単位は、相互に通信し、実行を同期させ、メモリアクセスを協調させることが可能であってもよい。いくつかの作業単位が、SIMD方式で、GPU上の同時実行のためのウェーブフロントに分けられ(クラスタ化され)てもよい。合計280個の作業単位についての上記の実施例に関して、ウェーブフロントは、64個の作業単位を含んでもよい。
OpenCL(登録商標)フレームワークは、種々の計算デバイス、またはOpenCL(登録商標)デバイス用のオープンプログラミング標準である。ソフトウェアプログラマは、ベンダ特有のコードを書くことを回避してもよく、それが向上したコード移植性をもたらし得る。他のフレームワークが利用可能であり、当該他のフレームワークは、よりベンダ特有のコーディングを異種アーキテクチャに提供し得る。例えば、NVIDIAは、Compute Unified Device Architecture(CUDA(登録商標))を提供し、AMDは、ATI Stream(登録商標)を提供する。CUDA(登録商標)フレームワークを用いると、計算カーネルは、典型的には、コンピュータプログラムがコンパイルされたときに、静的にコンパイルされる。OpenCL(登録商標)フレームワークを用いると、計算カーネルは、典型的には、ジャストインタイム(JIT)方法でコンパイルされる。JIT方法は、システム構成を取得した後に、適切なバイナリコードを生成し得る。JITコンパイル方法を用いると、コンパイル時間が総実行時間とともに含まれる。したがって、コンパイラ最適化は、実行時間を増加させ得る。加えて、ランタイムに、OpenCLコンパイラは、計算カーネルの複数のバージョンを生成し得る。計算カーネルの1つのバージョンは、汎用CPU、SIMD GPU等の各種のOpenCL(登録商標)デバイス種類のために生成されてもよい。
OpenCL(登録商標)およびCUDA(登録商標)といった2つのフレームワークは、それぞれの実行モデル間で用語の違いを有する。例えば、OpenCL(登録商標)における作業単位、作業グループ、ウェーブフロントおよびNDRangeは、スレッド、スレッドブロック、ワープおよびグリップ等のCUDA(登録商標)における対応する用語を有する。残りの説明の全体を通して、OpenCL(登録商標)に対応する用語が使用される。しかしながら、説明されるシステムおよび方法は、CUDA(登録商標)、ATI Streamおよび他のフレームワークに適用されてもよい。
図2に示されるように、コード210は、一般に「doWorkA」および「doWorkB」という、2つの関数呼び出しを定義する。各関数呼び出しは、「計算カーネル」と呼ばれてもよい。計算カーネルは、1つ以上の計算の作業単位を生じさせるために、1つ以上のデータのレコードと合致させられてもよい。したがって、2つ以上の作業単位は、単一の関数呼び出しの同一の命令を利用し得るが、異なるデータのレコードに作用する。例えば、コード220内の関数呼び出し「Power2」は、アレイ「INPUT」内の各データ値に1つずつ、10個の作業単位を実行するために使用されてもよい。ここで、レコードは、単一のデータ値を備える。他の実施例では、レコードは、それぞれデータ値を含む2つ以上のフィールドを備えてもよい。SIMDマイクロアーキテクチャは、カーネル「Power2」の命令を効率的に実行し、INPUTアレイ内の値について2の累乗を計算し、RESULTアレイに出力を書いてもよい。
OpenCL(登録商標)フレームワークは、並行して複数回、計算カーネルのインスタンスを呼び出し得る。計算カーネルの呼び出しは、get_global_id(0)と名付けられた内部関数を呼び出すことによって取り出され得る、1つの関連する一意のID(作業単位ID)を有する。コード220における上記の例に関して、計算カーネル「Power2」は、INPUTアレイ内の各データ値に対して1回呼び出される。この場合、計算カーネル「Power2」は、10回呼び出される。したがって、10個の一意の作業単位IDが取り出される。JITコンパイル方法を用いて、これらのインスタンスはランタイムに呼び出される。OpenCL(登録商標)フレームワークは、一意の作業単位IDを利用することによって、これらの異なるインスタンスを区別し得る。また、INPUTアレイ内の特定のデータ値等が作用するデータ(レコード)が特定されてもよい。したがって、ランタイムに、関連付けられた計算カーネルがスケジュールされると、作業単位は、デフォルト設定で同一のOpenCL(登録商標)デバイスへスケジュールされてもよい。
ここで図3を参照すると、条件文を有する計算カーネルを定義するソースコードの一実施形態が示されている。コード210と同様に、図3に示されるコード230は、一般に「doWorkA」および「doWorkB」という、2つの関数呼び出しを定義する。この場合においても、各関数呼び出しは、「計算カーネル」と呼ばれてよい。ここで、2つの計算カーネルのうち1つのみが、ランタイム中に実行される。いずれの計算カーネルが実行されるかという選択は、関数呼び出し「EvaluateFunction」によって提供される条件付きテストに基づいて実行される。所与の命令の結果、または所与の命令が実行されるか否かは、以前の命令の実行と、関連付けられたレコードに対応するデータとにデータ依存している。条件付きテストの結果が作業単位のウェーブフロントの間で一致していない場合、SIMDマイクロアーキテクチャの利益が低減する場合がある。例えば、所与のSIMDコアは、64個の作業単位の同時実行に利用可能な64個の並列計算ユニットを有し得る。しかしながら、64個の作業単位の半分が条件付きテストを通過する一方で、他の半分が条件付きテストに失敗する場合には、並列計算ユニットの半分のみが、所与の処理段階中に利用される。
ここで図4を参照すると、ハードウェアリソースと計算カーネルとの間でスケジュールされた割り当て400の一実施形態を図示する一般化されたブロック図が示されている。ここで、ハードウェアおよびソフトウェアリソースの分割と、1つ以上のソフトウェアアプリケーション430の実行中のハードウェアおよびソフトウェアリソースの相互関係および割り当てが示されている。一実施形態では、オペレーティングシステム420は、計算カーネル440a〜440jおよび440k〜440qのためのメモリの領域を割り付ける。アプリケーション430またはコンピュータプログラムが実行される場合、各アプリケーションは、複数の計算カーネルを備えてもよい。例えば、第1の実行アプリケーションは、計算カーネル440a〜440jを備えてもよく、第2の実行アプリケーションは、計算カーネル440k〜440qを備えてもよい。カーネル440a〜440qを1つずつ用いて、1つ以上のデータレコード(図示されず)と組み合わせることによって、1つ以上の作業単位を生成することができる。例えば、計算カーネル440aは、作業単位442a〜442dをもたらし、計算カーネル440jは、作業単位442e〜442hをもたらし、計算カーネル440kは、442j〜442mをもたらし、計算カーネル440qは、作業単位442n〜442qをもたらし得る。作業単位は、他の作業単位とは独立して実行され、且つ、他の作業単位と同時に実行され得る。
図4に示される計算カーネルのそれぞれは、メモリのイメージ等の独自のリソース、またはアプリケーション実行前の命令およびデータのインスタンスを所有し得る。また、計算カーネルのそれぞれは、例えばコード、データ、ならびに可能性としてヒープおよびスタックをアドレスするアドレス空間等のプロセス特有の情報と、例えばスタックポインタ、汎用および浮動小数点レジスタ、プログラムカウンタおよびその他等のデータおよび制御レジスタ内の変数と、例えばstdin、stdoutおよびその他等のオペレーティングシステム記述子と、例えば一式の許可等のセキュリティ属性とを備えてもよい。
一実施形態では、ハードウェアコンピュータシステム410は、それぞれ1つ以上の作業単位を処理するように構成された汎用プロセッサコア112およびSIMDプロセッサコア172を組み込む。別の実施形態では、システム410は、2つの他の異種プロセッサコアを含む。一般に、所与のアプリケーションについては、オペレーティングシステム420は、当該アプリケーション用のアドレス空間を設定し、当該アプリケーションのコードをメモリにロードし、プログラム用のスタックを設定し、アプリケーション内の所与の位置に分岐し、当該アプリケーションの実行を開始する。典型的には、オペレーティングシステム420のうち、そのようなアクティビティを管理する部分は、オペレーティングシステム(OS)カーネル422である。OSカーネル422は、計算カーネルまたは関数呼び出しと混同しないために、「OSカーネル」と呼ばれる。さらに、OSカーネル422は、アプリケーションの実行に利用可能なメモリが不十分であるときに、一連の行動を決定し得る。前述のように、アプリケーションは、2つ以上の計算カーネルに分割されてもよく、システム410は、2つ以上のアプリケーションを実行していてもよい。したがって、並行して起動するいくつかの計算カーネルが存在し得る。OSカーネル422は、任意の時点で、どの同時実行計算カーネルがプロセッサコア112および172に割り付けられるかを判定し得る。OSカーネル422は、タイムスライスと呼ばれる所与の時間にわたって、1つ以上のコアを有し得るプロセッサのコア上で、プロセスが起動することを可能にし得る。オペレーティングシステム420内のOSスケジューラ424は、計算カーネルをコアに割り当てるための決定論理を備えてもよい。
一実施形態では、1つだけの計算カーネルが、任意の時点で、ハードウェア計算ユニット412a〜412gおよび412h〜412rのうち何れか1つで実行することができる。これらのハードウェア計算ユニットは、関連付けられたデータを用いて所与の作業単位の所与の命令の実行を取り扱うことが可能なハードウェアを備える。このハードウェアは、加算、乗算、ゼロ検出、ビット単位シフト、除算、ビデオグラフィックスおよびマルチメディア命令、またはプロセッサ設計の当業者に公知である他の演算を行うように構成された算術論理演算ユニットを含んでもよい。これらのハードウェア計算ユニットは、マルチスレッドプロセッサ内のハードウェアスレッド、SIMDマイクロアーキテクチャ内の並列ハードウェアカラム等を含んでもよい。
図4の破線は、割り当てを表し、必ずしも直接物理接続を表すとは限らない。したがって、例えば、ハードウェア計算ユニット412aは、作業単位442dを実行するように割り当てられてもよい。しかしながら、以降で(例えば、コンテキスト切り替え後に)、ハードウェア計算ユニット412aは、作業単位442hを実行するように割り当てられてもよい。一実施形態では、OSスケジューラ424は、ラウンドロビン方式を用いて、ハードウェア計算ユニット412a〜412rへ作業単位442a〜442qをスケジューリングしてもよい。代替として、OSスケジューラ424は、ラウンドロビン方式を用いて、コア112および172へ作業単位442a〜442qをスケジューリングしてもよい。所与のハードウェア計算ユニットへの所与の作業単位の割り当ては、関連付けられたプロセッサコアによって行われてもよい。別の実施形態では、OSスケジューラ424は、プロセッサコア112および172の可用性に基づいてスケジューリングを行ってもよい。さらに別の実施形態では、OSスケジューラ424は、OpenCL(登録商標)APIまたは別の類似APIを利用して、プログラマによって作成された割り当てに従ってスケジューリングを行ってもよい。これらのスケジューリング方式は、作業単位割り当てとハードウェアリソースとの間に不一致があるときに、移植性および性能を制限する場合がある。
図5を参照すると、2種類のプロセッサコア用のマイクロアーキテクチャの論理レイアウトの一実施形態を示す一般化ブロック図が示されている。汎用コア510および単一命令複数データ(SIMD)コア560のそれぞれが示されているが、他の種類の異種コアも可能であり、考慮される。コア510および560のそれぞれは、データおよび命令の記憶のために、ダイナミックランダムアクセスメモリ(DRAM)550aおよび550bを有する。一実施形態では、コア510および560は、同一のDRAMを共有する。別の実施形態では、DRAMに加えて、所与のレベルのキャッシュメモリサブシステム(図示せず)が共有される。例えば、再び図1を参照すると、キャッシュメモリサブシステム118は、コア112および172によって共有される。
コア510および560のそれぞれは、キャッシュメモリサブシステム530を含んでもよい。示されるように、汎用コア510は、制御論理520および算術論理演算ユニット(ALU)540とは別にキャッシュメモリサブシステム530を論理的に有する。コア510内のデータフローは、パイプライン型であり得るが、パイプラインレジスタ等の記憶要素は、説明を簡略化するために示されていない。所与のパイプライン段階で、この段階での命令がある種類のALUを利用しない場合、または別の作業単位(あるいは汎用コアのための別のスレッド)がこの段階中にALUを消費する場合には、ALUは未使用であってもよい。
示されるように、SIMDコア560は、計算ユニット542の各行ごとの制御論理520とグループ化されたキャッシュメモリサブシステム530を有する。コア560内のデータフローは、パイプライン型であり得るが、パイプラインレジスタ等の記憶要素は、説明を簡略化するために示されていない。所与のパイプライン段階で、この段階での関連付けられた命令が、例えば分岐しない等のように事前に失敗したテストに基づいて実行されない場合には、計算ユニットは未使用であってもよい。
ここで図6を参照すると、汎用パイプライン実行フロー600の一実施形態を示す一般化されたブロック図が示されている。命令602〜608が取り出され、汎用パイプラインに入力し得る。命令606は、計算集約的命令であってもよい。パイプライン実行フローの特定の段階中に、命令602〜608のうち1つ以上の命令が、例えばデコーダ論理、命令スケジューラ入力、リオーダバッファ入力、ALU、レジスタファイル入力、分岐予測ユニット等の汎用プロセッサコア112内のリソースを消費する。
均衡のとれたスキームにおいて、命令602〜608のそれぞれは、各段階で等量のリソースを消費する。しかしながら、典型的には、汎用コアは、半導体不動産費用、電力消費および他の設計配慮により、各命令用のリソースを複製しない。したがって、作業負荷が不均衡になり得る。例えば、命令606は、その計算集約的挙動により、1つ以上のパイプ段階のためにより多くのリソースを消費し得る。示されるように、この命令によって消費されるリソース630は、他の命令によって消費されるリソースよりもはるかに大きくなり得る。実際に、計算集約的命令は、他の命令によるハードウェアリソースの使用を阻止し得る。
いくつかの計算集約的タスクは、図1に示された汎用コア112内の共有リソースに圧力を加え得る。したがって、スループット損失が、計算集約的プロセスと、共有リソースを待つ他のプロセスとの両方に起こる。加えて、いくつかの命令は、共有リソース上で行われている計算をサポートするように、共有リソースおよび他のリソースを占有する。そのような長い待ち時間の命令は、長い待ち時間の間に、他のプロセスがいくつかのリソースを使用することを同時に阻止し得る。
ここで図7Aを参照すると、SIMDパイプライン実行フロー700の一実施形態を示す一般化されたブロック図が示されている。命令702〜708が取り出され、関連付けられたデータを有するSIMDパイプラインに入力し得る。命令704は、条件付き分岐等の制御フロー転送命令であってもよい。命令706は、条件が真であるときに実行されるパスにおける第1の命令であってもよい。命令708は、条件が偽であるときに実行されるパスにおける第1の命令であってもよい。例えば、分岐命令704は、高レベル言語プログラム内のIF文と関連付けられてもよい。命令706は、高レベル言語プログラム内のTHEN文と関連付けられてもよい。命令708は、高レベル言語プログラム内のELSE文と関連付けられてもよい。
所与の行内の計算ユニットのそれぞれは、同一の計算ユニットであってもよい。これらの計算ユニットのそれぞれは、同一の命令に動作し得るが、異なる作業単位と関連付けられた異なるデータ上では動作しなくてもよい。示されるように、いくつかの作業単位は、条件付き分岐命令704によって提供されるテストを通過し、他の作業単位はテストに失敗する。SIMDコア172は、利用可能なパスのそれぞれを実行し、現在のパスを選択しなかった作業項目に対応する計算ユニット等の実行単位を選択的に無効にし得る。例えば、If−Then−Else構築文の実行中に、SIMDアーキテクチャの各列内に「Then」(パスA)および「Else」(パスB)パスを実行するように構成された実行単位がある。第1および第2の作業単位が実行を停止して待機し、第3の作業単位が継続中の実行を続ける際に、並列実行の効率性が低減し得る。したがって、計算ユニットの全てが、分岐命令704の実行後に、所与の行内で動作中の計算ユニット710であるわけではない。示されるように、1つ以上の計算ユニットは、実行が無効にされた、動作していない計算ユニット711である。多数の計算ユニットが所与のパイプ段階中で動作していない場合、SIMDコアの効率性およびスループットが低減する。
一実施形態において、「Else」パスは、計算カーネルへの戻りである。計算カーネルの実行が終了し、対応する作業単位が休止状態になる。しかしながら、SIMDコア内の隣接する作業単位は、実行を継続し得る。ここで図7Bを参照すると、SIMDパイプライン実行フロー720の別の実施形態を説明する一般化されたブロック図が示されている。実行フロー700と同様に、命令702〜706は、1つ以上の計算ユニットを、SIMDコアの特定の列内で無効にし得る。ここで、各「Else」パスは、計算カーネルへの戻りであり得る。したがって、所与の作業単位について、分岐しない方向に決定する分岐は、所与の作業単位に計算カーネルのさらなる実行を中止させ得る。実行フロー720では、説明の簡略化のために、1つのみの命令が、第1の分岐命令704と第2の分岐命令712との間に示されている。しかしながら、分岐命令704と712との間には、複数の命令が存在し得る。分岐704と712との間の命令数にかかわらず、分岐しない方向に第1の分岐704を決定する作業単位は、実行を完了し得る。分岐712も同様に、分岐しない方向に第2の分岐を決定する作業単位は、実行を完了し得る。SIMDコアの後の段階の計算ユニットは、これらの作業単位のために無効にされ得る。多数の計算ユニットが所与のパイプ段階中に動作しない場合、SIMDコアの効率性およびスループットは低減する。
隣接する作業単位が継続し得るが、複数の作業単位にテストを失敗させ、かつ実行を停止させ得るアプリケーションの一例は、顔検出である。当業者に既知であるように、OpenCv(オープンコンピュータビジョンライブラリ)において実装される顔検出は、Viola−Jonesオブジェクト検出アルゴリズムの1つのアプリケーションである。Viola−Jonesアルゴリズムは、データ依存性の実行パターンを示す。検索計算カーネルは、1つ以上のピクセルを含み得るデータレコードに適用される。検索計算カーネルは、2次元または3次元画像のサブウィンドウ内で顔を検索する。計算カーネル内で、分岐命令等の制御フロー転送命令として実装される一連のテストが存在し得る。1つの典型的な例において、一連のテストは、22段階、または22個のテストを含む。この一連のテストは、入力ウィンドウが顔を含むか否かを判定し得る。
Viola−Jonesアルゴリズムにおける一連のテストは、見込みのないパスを素早く取り除くように設計され得る。したがって、ほとんどの作業単位は、顔の不在を決定し、終了し得る。作業単位の実行は、顔を含有する可能性の高い残りのピクセル上で続けられる。ほんのわずかのピクセル(すなわち、作業単位実行)は、22段階にわたって続けられ得るが、ほとんどのピクセルは、最初の数段階のテスト後に顔を含有しないことが判明する。かなりのタスク並列性を伴っても、ウェーブフロント上での数個の継続中の作業単位の存在は、SIMDコア利用の低下を引き起こし得る。以下に記載の1つの方法は、さらなる処理のためにSIMDコアを解除しながら、別々の異種コアを利用する。この方法は、少しのSIMD並列性が存在することが検出されるときに、全体の計算性能を増大させ得る。
ここで図8を参照すると、移行点を定義するタグ付けされた分岐を含むコード800の一実施形態が示される。コード800は、一般に「foo」と題される計算カーネルを含む。実行中、コード800の一部は、別々の異種コアに移動し得る。示される例において、外側ループはデータ依存性である。一実施形態において、コンパイラは、「while」ループテストに対応する分岐命令においてタグビットを用いることによって、データ依存性のSIMDコアに情報提供する。実行中、移行条件が検出されるとき、例えば、測定されたSIMD利用が所与の閾値を下回るとき、中間局所値は、別々の異種コアによってアクセスされるように、メモリ内のデータ構造に移動し得る。例えば、汎用コアは、タグ付けされた分岐点から移行点まで計算カーネルの実行を継続し得る。例えば、while文における黙示的な条件付き分岐は、「secondary_entry」というラベルでタグ付けされる。別々の異種コアは、コンパイラによって生成されたデータ構造を使用し得る。別の実施形態において、このデータは、キャッシュされてもよく、移行コストを軽減する。一例において、ライブデータは、「tmp」アレイのローカルスライスと、local_temp変数の現在の値との両方を含み得る。移動中、このデータは、計算カーネルの継続実行を、「secondary_entry」というラベルによって表示される二次エントリポイントに指向するランタイム環境に伝達され得る。
ここで図9を参照すると、プリランタイムデータ情報を利用してプロセッサ内の複数の作業単位の並列実行を最適化するための方法900の一実施形態が示されている。処理ノード110内に統合されるコンポーネントおよび上述の図4に示されたハードウェアリソース割り当ては、概して、方法900に従って動作し得る。考察目的のために、この実施形態および後に説明される方法の後続の実施形態におけるステップは、起こった順に示される。しかしながら、他の実施形態において、いくつかのステップは、示された順序とは異なる順序で起こってもよく、いくつかのステップは、同時に行われてもよく、いくつかのステップは、他のステップと組み合わせられてもよく、いくつかのステップは、行われなくてもよい。
ブロック902において、ソフトウェアプログラムまたはサブルーチンが特定され、分析され得る。このソフトウェアプログラムは、異種マルチコアアーキテクチャ上でのコンパイルおよび実行のために書き込まれ得る。プログラムコードは、ソフトウェアアプリケーション、サブルーチン、動的なリンクされたライブラリまたはその他の任意の部分を参照し得る。パス名は、ユーザによるコマンドプロンプトにおいて入力されてもよく、ソースコードのコンパイルを開始するために、所与のディレクトリ位置またはその他から読み出され得る。プログラムコードは、C等の高水準言語、OpenCL(登録商標)等のC風の言語等で設計者によって書き込まれ得る。一実施形態において、ソースコードは、静的にコンパイルされる。そのような実施形態において、静的なフロントエンドコンパイル中に、ソースコードは、中間表現(IR)に変換され得る。バックエンドコンパイルステップは、IRを機械コードに変換し得る。静的バックエンドコンパイルは、より多くの変換および最適化を行い得る。別の実施形態において、ソースコードは、ジャストインタイム(JIT)方法でコンパイルされる。JIT方法は、システム構成を取得した後に、適切なバイナリコードを生成し得る。いずれかの方法を用いて、コンパイラは、プログラムコード内の計算カーネルを識別し得る。一実施形態において、OpenCL(登録商標)コンパイラ等のコンパイラは、複数のバージョンの計算カーネルを生成し得る。計算カーネルの1つのバージョンが、汎用CPU、SIMD GPU等の各種のOpenCL(登録商標)デバイス種類のために生成され得る。
ブロック904では、コンパイラは、計算カーネルの1つ以上の命令を読み取り、それらを分析し得る。条件文は、分岐等の制御フロー転送命令であってもよい。異なる種類の制御フロー転送命令は、フォワード/バックワード分岐、直接/間接分岐、ジャンプ等を含み得る。コンパイラまたは他のツールが分岐方向および/または分岐先を静的に決定することが可能であり得る。しかしながら、一実施形態において、関連付けられたデータ上でランタイム中に通常行われたいくつかの処理は、コンパイル中に行われ得る。例えば、分岐の方向(分岐する、分岐しない)を決定する単純なテストが行われて得る。コンパイルは「静的コンパイル」と称され得るが、1つ以上の小さい動的演算が行われ得る。このコンパイルは、「プリランタイムコンパイル」とも称され得る。この時点で行われる動的ステップの別の例は、If−Then−ElseIf−Else構築文のTHEN、ELSEIFおよびELSEブロックのそれぞれにおいて実行される次の命令を識別することである。例えば、条件付き分岐が失敗した場合、return文が実行され得る。したがって、コンパイラは、実行中、このコンピュータカーネルの対応する作業単位が、分岐テストが失敗するときに休止状態になり得ることを認識する。
ブロック906において、計算カーネルにおける特定のコードラインは、移行点を作成するために選択される。移行点は、インフライトの実行が異なる異種コアに転送するコンピュータカーネル内の位置であり得る。一実施形態において、この計算サブカーネル移行は、プロセス移行に類似した機構によって達成されてよく、実行状態は、第1の異種コアから第1のコアとは異なる可能性のあるマイクロアーキテクチャを有する第2の異種コアに移動する。別の実施形態において、この計算サブカーネル移行は、後に送信される複数の計算サブカーネルを作成することによって達成され得る。
一実施形態において、コンパイラは、移行点を自動的に識別し得る。本明細書で使用されるとき、移行点は、切替点とも称され得る。コンパイラは、制御フロー分析を使用し得る。移行点の識別は、計算カーネルの出口または戻りにつながるデータ依存性ループを見つけ出すために、静的制御フロー分析の利用を含み得る。出口または戻りを含むパスでそれぞれの分岐を識別するのではなく、コンパイラは、いくつかの移行点を減少させるためにカウントを使用し得る。例えば、計算カーネルにおいて見つけられた第1の5つの分岐は、移行点としてタグ付けするための候補ではない場合がある。第1の5つの分岐後のあらゆる第3の分岐は、移行点としてタグ付けするための候補であり得る。カウントに基づく他のフィルタリングアルゴリズムが可能であり、企図される。
加えて、コンパイラは、以前の実行からのプロファイル入力を用いて、移行点を識別し得る。例えば、所与の分岐と関連付けられた条件付きテストは、いくつかのデータレコードが所与の閾値を上回るため失敗し得る。したがって、この分岐は、移行点として識別され得る。さらに、移行点を示すプログラムマ注釈は、「pragmas」として、またはOpenCL(登録商標)フレームワークの拡張子として追加され得る。
ブロック908において、コンパイラは、コンパイルされたコードのバージョンごとに、コード内の選択された点をタグ付けし得る。それぞれのバージョンは、各OpenCL(登録商標)デバイスに対する送り先計算カーネルと称され得る。この場合もやはり、コンパイラは、識別された計算カーネルをコンパイルして、コンパイルされたコードの2つ以上のバージョンをもたらしてよく、それぞれは、OpenCLデバイスのうち1つの上で起動することができる。再び図9のコード800を参照すると、「secondary_entry」というラベルによって表示された二次エントリポイントは、分岐に対する移行タグの例である。コンパイラ内のコードジェネレータは、タグを挿入し、他のコードを挿入して、移行中にライブ値を呼び出し得る。ライブ値の呼び出しは、ライブ値の送り先OpenCL(登録商標)デバイスへの移動、およびその値の送り先OpenCL(登録商標)デバイス上での初期化を含み得る。コード生成および挿入プロセスは、デビュー点で挿入されるデバッガコードおよび動的挙動を測定するための計測に類似し得る。
一実施形態において、計算カーネルは、上述のように、タグ付けされ、移行点を識別し得る。別の実施形態において、計算カーネルは、独立してスケジュールおよび送信される複数の計算サブカーネルに分割され得る。ランタイムプロファイル情報またはコンパイラ静的推定を用いて、分岐命令によって実装される条件付きテストの通過/失敗統計を決定することができる。「ホット」実行パスは、複数のデータレコードの条件付きテストの所与の閾値を上回る多数のパスを含み得る。「コールド」実行パスは、複数のデータレコードの条件付きテストの第2の所与の閾値を下回る少数のパスを含み得る。計算カーネルは、「ホット」および「コールド」実行パスに基づいて計算サブカーネルに分割され得る。
対応する計算サブカーネルの生成は、汎用コア上で実行を継続する「コールド」実行パス等のそれらの計算サブカーネルに対応する実行範囲(NDRange)の生成に加えて、同様のランタイムコード生成機構を利用し得る。これは、汎用コア上で実行されるように、OpenCL(登録商標)指定を利用し得る計算サブカーネル識別子(ID)を含有する潜在的に希薄なアレイを作成することによって行われ得る。所与の計算カーネルは、適切な計算サブカーネルおよび後の作業単位を識別するために、このアレイへの間接アクセスを利用し得る。あるいは、コンパイラは、これらのIDのリストを生成してもよく、対応する計算サブカーネルは、実行作業単位のそれぞれに対して呼び出しおよびマッピングされる。
プロファイル起動または静的推定後、「ホット」実行パスに対応する計算サブカーネルは、SIMDコアにコンパイルされ得る。「コールド」実行パスに対応する計算サブカーネルは、汎用コアにコンパイルされ得る。一連のテストの初期段階は、通過する可能性が高くあり得る。したがって、これらの実行パスは、SIMDコア上で実行される「ホット」計算サブカーネルにおいて実装され得る。これらの特定の「ホット」計算サブカーネルの実行後、関連付けられた生成データは、メモリ内を移動し得る。このデータ移動は、グローバルデータに対してライブであるローカルデータを促進する。「ホット」計算サブカーネルに対応する作業単位は、関連付けられた「コールド」計算サブカーネルがその後汎用コア上で実行し続けるか否かを示すために、その作業単位IDに基づいてビットアレイに書き込み得る。
ブロック910において、コンパイラは、識別された移行点で一組のライブ値を識別する。ライブ値は、中間計算値およびローカルアレイを含み得る。再び図8のコード800を参照すると、ライブデータは、コード内の「tmp」アレイのローカルスライスと、local_temp変数の現在の値との両方を含み得る。移行が後の関連付けられた作業単位の実行中に生じた場合、ライブ値は転送され、送り先OpenCL(登録商標)デバイス上で初期化され得る。上述のように、コンパイラ内のコードジェネレータは、タグを挿入し、他のコードを挿入して、移動中にライブ値を呼び出し得る。送り先OpenCL(登録商標)デバイスにおいて、移行エントリポイントのコードジェネレータは、ライブ値を含有するデータ構造を初期化し、カーネル実行を始める。あるいは、コンパイラは、上述のように、計算サブカーネルを作成して、実行を始め得る。ブロック912において、コンパイラは、少なくとも2つの異種プロセッサコアに対して計算カーネルのコンパイルを終了する。他のデバッグおよび計測コードが挿入され得る。
一実施形態において、コンパイラは、複数のデータ構造を生成する。2つ以上のデータ構造は、汎用コアおよびSIMDコア等の所与のターゲットOpenCLデバイス上の計算サブカーネルごとに、実行可能なオブジェクトコードを含む。別のデータ構造は、移行時に転送およびアクセスされるライブデータを含む。計算カーネルにおける潜在的な移行点を表すラベルを考慮して、コンパイラは、データフロー分析を利用して、転送され得るライブ値を決定する。実行時に定義されないライブ値、例えば、レジスタでキャッシュされるライブ値は、ランタイム環境にアクセス可能な位置に配置される。これらの位置の例には、保存されるコンテンツを保持する、関連付けられた最初のメモリ位置およびレジスタが挙げられる。一実施形態において、ヒューリスティックチェックを利用して、データ転送のサイズが異種コア間の有益な変化実行を可能にするか否かを判定することができる。
さらに、コンパイラは、ランタイム環境によって解釈される別のデータ構造を生成して、ライブデータを、関連付けられた送り先OpenCL(登録商標)デバイスに転送し得る。このデータ構造は、転送されるライブデータの位置およびサイズ、ならびにソースおよび送り先OpenCL(登録商標)デバイスの両方のアドレス空間におけるそれらの位置を提供し得る。また、コンパイラは、送り先デバイスに対応するバージョンのカーネルを生成する。OpenCLデバイスのそれぞれに対する各コンパイルされたコードは、指定された位置でライブデータにアクセスし、移行点で実行を開始する。
ここで図10を参照すると、プリランタイムデータ情報を利用することによって、プロセッサにおける複数の作業単位の並列実行を最適化するための方法1000の一実施形態が示されている。処理ノード110内に統合されるコンポーネントおよび上述の図4に示されるハードウェアリソース割り当ては、概して、方法1000に従って動作し得る。考察目的のために、この実施形態および後に説明される方法の後続の実施形態におけるステップは、起こった順に示される。しかしながら、いくつかのステップは、示された順序とは異なる順序で起こってもよく、いくつかのステップは、同時に行われてもよく、いくつかのステップは、他のステップと組み合わせられてもよく、いくつかのステップは、別の実施形態では不在であり得る。
ブロック1002において、関連付けられたデータレコードは、所与の計算カーネルのそれぞれの作業単位に割り当てられる。ブロック1004において、OSスケジューラ424は、作業単位を異種コアにスケジュールする。ブロック1006において、異種プロセッサコアは、対応するスケジュールされた作業単位を実行する。
ブロック1008において、所与のタグ付けされた移行点に到達する。一実施形態において、現在使用されるOpenCL(登録商標)デバイスの利用の測定が行われ得る。測定結果が、利用または性能が所与の閾値を下回ることを示す場合、関連付けられた計算カーネルまたは計算サブカーネルは、異なるマイクロアーキテクチャを有する異種コア等の別のOpenCL(登録商標)デバイスに移動し得る。一実施形態において、この測定結果は、関連付けられた計算カーネルまたは計算サブカーネル内で出口または戻りに到達したSIMDコア上のいくつかの現在実行している作業単位のカウントである。あるいは、ウェーブフロントにおけるいくつかの無効にされた計算ユニットのカウントは、同一の数を提供し得る。このカウントが所与の閾値を超える場合、まだ出口点に到達していない作業単位は、汎用コア等の別の異種コアに移行し得る。その後、SIMDコア上のウェーブフロントは解放されてよく、他のスケジュールされた作業単位に使用可能である。
他の実施形態において、上述の技法は、SIMDコア上のウェーブフロントにおいて作業単位を実行する並列の大部分が休止状態であることが決定され、且つ、残りの作業単位が実質的な実行を継続することが見込まれる任意の状況において、移行を開始するまでに拡大され得る。例えば、生成されたデータ構造は、共有メモリおよび1つ以上のキャッシュ内に存在し得る。仮想メモリ支援を有するシステムにおいて、作業単位のサブセットは、キャッシュをヒットし得るが、残りの作業単位は、長い待ち時間イベントである仮想メモリミスを経験する。この場合、全体の計算性能は、さらなる実行が、現在の実行によって有効にされるプリフェッチ技法の恩恵を受け得るため、汎用コア上での継続的な実行を伴ってより改善し得る。
実行効率性が所与の閾値を下回ると判定されない場合(条件付きブロック1010)、方法1000の制御フローは、ブロック1006に戻り、実行が続く。実行効率性が所与の閾値を下回ると判定された場合(条件付きブロック1010)、ブロック1012において、1つ以上の作業単位が識別されて、第1のプロセッサコアのマイクロアーキテクチャとは異なるマイクロアーキテクチャを有する第2のプロセッサコアに移行する。識別された作業単位は、上述の測定結果を、所与の閾値未満にする可能性がある。ブロック1014において、第1のプロセッサコアによって生成される関連付けられたローカルデータは、グローバルデータに昇格する。ブロック1016において、コンパイルされたバージョンの移行した作業単位は、移行タグ付け点で開始する第2のプロセッサコア上で実行されるようにスケジュールされる。
上述の実施形態がソフトウェアを備え得ることに留意する。そのような実施形態において、方法および/または機構を実装するプログラム命令は、コンピュータ可読媒体上で伝達または記憶され得る。プログラム命令を記憶するように構成される多くの種類の媒体が利用可能であり、ハードディスク、フロッピー(登録商標)ディスク、CD−ROM、DVD、フラッシュメモリ、プログラム可能なROM(PROM)、ランダムアクセスメモリ(RAM)および様々な他の形態の揮発性または不揮発性記憶装置を含む。一般的に言えば、コンピュータアクセス可能な記憶媒体は、命令および/またはデータをコンピュータに提供するために、使用するときにコンピュータによってアクセス可能な任意の記憶媒体を含み得る。例えば、コンピュータアクセス可能な記憶媒体は、磁気または光学媒体等の記憶媒体、例えば、ディスク(固定もしくは取り外し可能)、テープ、CD−ROMもしくはDVD−ROM、CD−R、CD−RW、DVD−R、DVD−RW、またはブルーレイを含み得る。記憶媒体は、ユニバーサルシリアルバス(USB)インターフェース等の周辺インターフェースを介してアクセス可能なRAM(例えば、同期ダイナミックRAM(SDRAM)、ダブルデータレート(DDR、DDR2、DDR3等)SDRAM、低出力DDR(LPDDR2等)SDRAM、Rambus DRAM(RDRAM)、スタティックRAM(SRAM)等)、ROM、フラッシュメモリ、不揮発性メモリ(例えば、フラッシュメモリ)等の揮発性もしくは不揮発性メモリ媒体をさらに含み得る。記憶媒体は、微小電気機械システム(MEMS)、ならびにネットワークおよび/または無線リンク等の通信媒体を介してアクセス可能な記憶媒体を含み得る。
さらに、プログラム命令は、C等の高レベルプログラミング言語、またはVerilog、VHDL若しくはGDS IIストリームフォーマット(GDSII)等のデータベースフォーマット等の設計言語(HDL)で、ハードウェア機能性の挙動レベル記述またはレジスタ転送レベル(RTL)記述を含み得る。いくつかの場合において、記述は、合成ライブラリからのゲートのリストを備えるネットリストを生成するように記述を合成し得る合成ツールによって読み取られてもよい。ネットリストは、システムを含むハードウェアの機能性も表す一組のゲートを含む。次いで、ネットリストは、マスクに適用される幾何学形状を描写するデータセットを生成するように配置され、ルートされ得る。次いで、マスクは、システムに対応する1つまたは複数の半導体回路を生成するために、様々な半導体製造ステップで使用され得る。あるいは、コンピュータアクセス可能記憶媒体上の命令は、所望に応じて、ネットリスト(合成ライブラリを有する、または有しない)またはデータセットであり得る。さらに、命令は、Cadence(登録商標)、EVE(登録商標)およびMentor Graphics(登録商標)等のベンダからのハードウェアベース型のエミュレータによって、エミュレーションの目的で利用され得る。
上述の実施形態がかなり詳細に説明されているが、上述の開示が完全に理解されるときに、多数の変形および修正が当業者に明らかになるであろう。以下の特許請求の範囲が、すべてのそのような変形および修正を包含すると解釈されることが意図される。

Claims (18)

  1. 複数の命令を含む計算カーネル内の位置であって、前記計算カーネルの実行中に前記計算カーネルの実行が移行可能な位置を、前記計算カーネルのコンパイル中に特定するステップと、
    前記計算カーネルのコンテキストを維持及び移行するためのデータ構造を作成するステップと、
    第1のマイクロアーキテクチャを有する第1のプロセッサコア上で実行するために、前記位置前の前記計算カーネル内のコードをスケジュールするステップと、
    移行条件を満たしているという指標を受信したことに応じて、
    前記コンテキストを、前記第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを有する第2のプロセッサコアがアクセス可能な位置に移動させるステップと、
    前記位置後の前記計算カーネル内のコードを、前記第2のプロセッサコアにスケジュールするステップと、
    を含み、
    前記移行条件を満たしていることを判定するために、出口点に到達した前記計算カーネルの並列実行反復回数が所与の閾値を超えているか判定するステップをさらに含む、
    方法。
  2. 前記第1のプロセッサコアに対応する前記計算カーネルの第1のバージョンのコードを生成するステップと、
    前記第2のプロセッサコアに対応する前記計算カーネルの第2のバージョンのコードを生成するステップと、をさらに含む、請求項1に記載の方法。
  3. 前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャであり、前記第2のマイクロアーキテクチャは、汎用マイクロアーキテクチャである、請求項2に記載の方法。
  4. プロファイルランタイム情報及び静的情報の少なくとも1つに基づいて前記特定を行うステップをさらに含む、請求項2に記載の方法。
  5. 前記移行条件を満たしているか否かを判定する命令を用いて、前記第1のプロセッサコアの第1のバージョンのコードを計測するステップと、
    前記データ構造によって示された位置でライブ値を検出するとともに実行を開始する命令を用いて、前記第2のプロセッサコアの第2のバージョンのコードを計測するステップと、
    をさらに含む、請求項2に記載の方法。
  6. 前記位置は、条件付き分岐命令の直前である、請求項1に記載の方法。
  7. 異種マルチコアアーキテクチャを含むコンピューティングシステムであって、
    第1のマイクロアーキテクチャを有する第1のプロセッサコアと、
    前記第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを有する第2のプロセッサコアと、
    複数の命令を含む計算カーネルであって、前記計算カーネルの実行中に前記計算カーネルの実行が移行可能な位置を含む計算カーネルと、
    前記計算カーネルのコンテキストの維持及び移行に使用可能なデータ構造と、
    スケジューラを含むオペレーティングシステムと、を備え、
    前記スケジューラは、
    第1のマイクロアーキテクチャを有する第1のプロセッサコア上で実行するために、前記位置前の前記計算カーネル内のコードをスケジュールし、
    移行条件を満たしたという指標を受信したことに応じて、
    前記コンテキストを、前記第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを有する第2のプロセッサコアがアクセス可能な位置に移動させ、
    前記位置後の前記計算カーネル内のコードを、前記第2のプロセッサコアにスケジュールするように構成されており、
    前記第1のプロセッサコア及び前記第2のプロセッサコアの各々は、前記移行条件を満たしていることを判定するために、出口点に到達した前記計算カーネルの並列実行反復回数が所与の閾値を超えているか判定するように構成されている、
    コンピューティングシステム。
  8. 前記第1のプロセッサコアに対応する前記計算カーネルの第1のバージョンのコードを生成し、前記第2のプロセッサコアに対応する前記計算カーネルの第2のバージョンのコードを生成するように構成されたコンパイラをさらに含む、
    請求項7に記載のコンピューティングシステム。
  9. 前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャであり、前記第2のマイクロアーキテクチャは、汎用マイクロアーキテクチャである、請求項8に記載のコンピューティングシステム。
  10. 前記コンパイラは、プロファイルランタイム情報及び静的情報の少なくとも1つに基づいて前記移行可能な位置の特定を行うように構成されている、請求項8に記載のコンピューティングシステム。
  11. 前記コンパイラは、
    前記移行条件を満たしているか否かを判定する命令を用いて、前記第1のプロセッサコアの第1のバージョンのコードを計測し、
    前記データ構造によって示された位置でライブ値を検出するとともに実行を開始する命令を用いて、前記第2のプロセッサコアの第2のバージョンのコードを計測するように構成されている、
    請求項8に記載のコンピューティングシステム。
  12. 前記コンパイラは、
    前記計算カーネルの後の並列実行反復回数が前記移条件を満たすという予測に応じて、前記計算カーネルを前記位置で2つの計算サブカーネルに分割し、
    前記位置前のコードを含む第1の計算サブカーネルを、前記第1のプロセッサコアにスケジュールし、
    前記位置後のコードを含む第2の計算サブカーネルを、前記第2のプロセッサコアにスケジュールするように構成されている、
    請求項に記載のコンピューティングシステム。
  13. 前記位置は、条件付き分岐命令の直前である、請求項7に記載のコンピューティングシステム。
  14. プログラム命令を記憶するコンピュータ可読記憶媒体であって、
    前記プログラム命令は、
    複数の命令を含む計算カーネル内の位置であって、前記計算カーネルの実行中に前記計算カーネルの実行が移行可能な位置を、前記計算カーネルのコンパイル中に特定し、
    前記計算カーネルのコンテキストを維持及び移行するためのデータ構造を作成し、
    第1のマイクロアーキテクチャを有する第1のプロセッサコア上で実行するために、前記位置前の前記計算カーネル内のコードをスケジュールし、
    移行条件を満たしているという指標を受信したことに応じて、
    前記コンテキストを、前記第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを有する第2のプロセッサコアがアクセス可能な位置に移動させ、
    前記位置後の前記計算カーネル内のコードを、前記第2のプロセッサコアにスケジュールし、
    前記移行条件を満たしていることを判定するために、出口点に到達した前記計算カーネルの並列実行反復回数が所与の閾値を超えているか判定するように実行可能である、
    コンピュータ可読記憶媒体。
  15. 前記プログラム命令は、前記第1のプロセッサコアに対応する前記計算カーネルの第1のバージョンのコードを生成し、前記第2のプロセッサコアに対応する前記計算カーネルの第2のバージョンのコードを生成するように実行可能である、請求項14に記載のコンピュータ可読記憶媒体。
  16. 前記プログラム命令は、
    前記移行条件を満たしているか否かを判定する命令を用いて、前記第1のプロセッサコアの第1のバージョンのコードを前記位置で計測し、
    前記データ構造によって示された位置でライブ値を検出するとともに実行を開始する命令を用いて、前記第2のプロセッサコアの第2のバージョンのコードを前記位置で計測するように実行可能である、
    請求項14に記載のコンピュータ可読記憶媒体。
  17. 複数の命令を含む計算カーネル内の位置であって、前記計算カーネルの実行中に前記計算カーネルの実行が移行可能な位置を、前記計算カーネルのコンパイル中に特定するステップと、
    前記計算カーネルのコンテキストを維持及び移行するためのデータ構造を作成するステップと、
    前記計算カーネルの後の並列実行反復回数が移動条件を満たすという予測に応じて、前記計算カーネルを前記位置で2つの計算サブカーネルに分割するステップと、
    前記位置前のコードを含む第1の計算サブカーネルを、第1のプロセッサコアにスケジュールするステップと、
    前記位置後のコードを含む第2の計算サブカーネルを、第2のプロセッサコアにスケジュールするステップと、
    を含む、方法。
  18. 異種マルチコアアーキテクチャを含むコンピューティングシステムであって、
    第1のマイクロアーキテクチャを有する第1のプロセッサコアと、
    前記第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを有する第2のプロセッサコアと、
    複数の命令を含む計算カーネルであって、前記計算カーネルの実行中に前記計算カーネルの実行が移行可能な位置を含む計算カーネルと、
    前記計算カーネルのコンテキストの維持及び移行に使用可能なデータ構造と、
    前記計算カーネルの後の並列実行反復回数が移動条件を満たすという予測に応じて、前記計算カーネルを前記位置で2つの計算サブカーネルに分割するように構成されたコンパイラと、
    スケジューラを含むオペレーティングシステムと、を備え、
    前記スケジューラは、
    前記位置前のコードを含む第1の計算サブカーネルを、第1のプロセッサコアにスケジュールし、
    前記位置後のコードを含む第2の計算サブカーネルを、第2のプロセッサコアにスケジュールする、ように構成されている、
    コンピューティングシステム。
JP2014511476A 2011-05-16 2012-05-16 異種コアの自動カーネル移行 Active JP5711853B2 (ja)

Applications Claiming Priority (3)

Application Number Priority Date Filing Date Title
US13/108,438 US8683468B2 (en) 2011-05-16 2011-05-16 Automatic kernel migration for heterogeneous cores
US13/108,438 2011-05-16
PCT/US2012/038057 WO2012158753A1 (en) 2011-05-16 2012-05-16 Automatic kernel migration for heterogeneous cores

Publications (3)

Publication Number Publication Date
JP2014513853A JP2014513853A (ja) 2014-06-05
JP2014513853A5 JP2014513853A5 (ja) 2014-11-06
JP5711853B2 true JP5711853B2 (ja) 2015-05-07

Family

ID=46147108

Family Applications (1)

Application Number Title Priority Date Filing Date
JP2014511476A Active JP5711853B2 (ja) 2011-05-16 2012-05-16 異種コアの自動カーネル移行

Country Status (6)

Country Link
US (1) US8683468B2 (ja)
EP (1) EP2710467B1 (ja)
JP (1) JP5711853B2 (ja)
KR (1) KR101559090B1 (ja)
CN (1) CN103534686B (ja)
WO (1) WO2012158753A1 (ja)

Families Citing this family (67)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US8789063B2 (en) * 2007-03-30 2014-07-22 Microsoft Corporation Master and subordinate operating system kernels for heterogeneous multiprocessor systems
US9092267B2 (en) * 2011-06-20 2015-07-28 Qualcomm Incorporated Memory sharing in graphics processing unit
US9195501B2 (en) * 2011-07-12 2015-11-24 Qualcomm Incorporated Instruction culling in graphics processing unit
US9378120B2 (en) * 2011-11-09 2016-06-28 Tata Consultancy Services Limited Automated test execution plan derivation system and method
US9430807B2 (en) 2012-02-27 2016-08-30 Qualcomm Incorporated Execution model for heterogeneous computing
US20120222043A1 (en) * 2012-05-01 2012-08-30 Concurix Corporation Process Scheduling Using Scheduling Graph to Minimize Managed Elements
US8595743B2 (en) 2012-05-01 2013-11-26 Concurix Corporation Network aware process scheduling
US8650538B2 (en) 2012-05-01 2014-02-11 Concurix Corporation Meta garbage collection for functional code
US8726255B2 (en) 2012-05-01 2014-05-13 Concurix Corporation Recompiling with generic to specific replacement
US9417935B2 (en) 2012-05-01 2016-08-16 Microsoft Technology Licensing, Llc Many-core process scheduling to maximize cache usage
EP2742425A1 (en) * 2012-05-29 2014-06-18 Qatar Foundation Graphics processing unit controller, host system, and methods
US8700838B2 (en) 2012-06-19 2014-04-15 Concurix Corporation Allocating heaps in NUMA systems
US9047196B2 (en) 2012-06-19 2015-06-02 Concurix Corporation Usage aware NUMA process scheduling
US8707326B2 (en) 2012-07-17 2014-04-22 Concurix Corporation Pattern matching process scheduler in message passing environment
US9575813B2 (en) 2012-07-17 2017-02-21 Microsoft Technology Licensing, Llc Pattern matching process scheduler with upstream optimization
US8793669B2 (en) 2012-07-17 2014-07-29 Concurix Corporation Pattern extraction from executable code in message passing environments
US9043788B2 (en) 2012-08-10 2015-05-26 Concurix Corporation Experiment manager for manycore systems
US10187452B2 (en) * 2012-08-23 2019-01-22 TidalScale, Inc. Hierarchical dynamic scheduling
US8656134B2 (en) 2012-11-08 2014-02-18 Concurix Corporation Optimized memory configuration deployed on executing code
US8607018B2 (en) 2012-11-08 2013-12-10 Concurix Corporation Memory usage configuration based on observations
US8656135B2 (en) 2012-11-08 2014-02-18 Concurix Corporation Optimized memory configuration deployed prior to execution
US10585801B2 (en) 2012-11-26 2020-03-10 Advanced Micro Devices, Inc. Prefetch kernels on a graphics processing unit
US9823927B2 (en) * 2012-11-30 2017-11-21 Intel Corporation Range selection for data parallel programming environments
GB2508433A (en) * 2012-12-03 2014-06-04 Ibm Migration of processes in heterogeneous computing environments using emulating and compiling source code on target system
US20130219372A1 (en) 2013-03-15 2013-08-22 Concurix Corporation Runtime Settings Derived from Relationships Identified in Tracer Data
US9298511B2 (en) * 2013-03-15 2016-03-29 International Business Machines Corporation Resolving deployment conflicts in heterogeneous environments
US9292349B2 (en) * 2013-03-15 2016-03-22 International Business Machines Corporation Detecting deployment conflicts in heterogenous environments
US9479449B2 (en) 2013-06-03 2016-10-25 Advanced Micro Devices, Inc. Workload partitioning among heterogeneous processing nodes
US9170854B2 (en) 2013-06-04 2015-10-27 Advanced Micro Devices, Inc. Thread assignment for power and performance efficiency using multiple power states
CN105531668B (zh) * 2013-08-08 2019-04-23 英派尔科技开发有限公司 用于执行进程的迁移的方法、迁移器及计算机可读介质
CN104254020B (zh) * 2013-09-25 2015-12-02 腾讯科技(深圳)有限公司 媒体数据的播放方法、装置及终端
US20150220340A1 (en) * 2013-10-04 2015-08-06 Rajkishore Barik Techniques for heterogeneous core assignment
JP5946068B2 (ja) * 2013-12-17 2016-07-05 インターナショナル・ビジネス・マシーンズ・コーポレーションInternational Business Machines Corporation 演算コア上で複数の演算処理単位が稼働可能なコンピュータ・システムにおける応答性能を評価する計算方法、計算装置、コンピュータ・システムおよびプログラム
CN104793924B (zh) * 2014-01-21 2019-03-15 中兴通讯股份有限公司 计算任务的处理方法及装置
CN105210059B (zh) * 2014-04-04 2018-12-07 华为技术有限公司 一种数据处理方法及系统
US10133572B2 (en) * 2014-05-02 2018-11-20 Qualcomm Incorporated Techniques for serialized execution in a SIMD processing system
EP3152859A1 (en) * 2014-06-04 2017-04-12 Giesecke & Devrient GmbH Method for enhanced security of computational device with multiple cores
US9880918B2 (en) * 2014-06-16 2018-01-30 Amazon Technologies, Inc. Mobile and remote runtime integration
US9959142B2 (en) * 2014-06-17 2018-05-01 Mediatek Inc. Dynamic task scheduling method for dispatching sub-tasks to computing devices of heterogeneous computing system and related computer readable medium
US10769175B1 (en) 2014-06-20 2020-09-08 Amazon Technologies, Inc. Real-time hosted system analytics
US10776397B2 (en) * 2014-06-20 2020-09-15 Amazon Technologies, Inc. Data interest estimation for n-dimensional cube computations
US11868372B1 (en) 2014-06-20 2024-01-09 Amazon Technologies, Inc. Automated hierarchy detection for cloud-based analytics
EP3158478B1 (en) 2014-06-20 2023-06-07 Amazon Technologies, Inc. Embeddable cloud analytics
US9898348B2 (en) * 2014-10-22 2018-02-20 International Business Machines Corporation Resource mapping in multi-threaded central processor units
US9286196B1 (en) * 2015-01-08 2016-03-15 Arm Limited Program execution optimization using uniform variable identification
US9400685B1 (en) * 2015-01-30 2016-07-26 Huawei Technologies Co., Ltd. Dividing, scheduling, and parallel processing compiled sub-tasks on an asynchronous multi-core processor
US9529950B1 (en) 2015-03-18 2016-12-27 Altera Corporation Systems and methods for performing profile-based circuit optimization using high-level system modeling
CN104899385B (zh) * 2015-06-16 2018-01-26 北京思朗科技有限责任公司 异构多核的SoC设计评估系统
US9501304B1 (en) 2015-06-16 2016-11-22 Architecture Technology Corporation Lightweight application virtualization architecture
US9983857B2 (en) * 2015-06-16 2018-05-29 Architecture Technology Corporation Dynamic computational acceleration using a heterogeneous hardware infrastructure
US20170052799A1 (en) * 2015-08-21 2017-02-23 Microchip Technology Incorporated Integrated Circuit Device With Selectable Processor Core
WO2017074377A1 (en) * 2015-10-29 2017-05-04 Intel Corporation Boosting local memory performance in processor graphics
US11513805B2 (en) * 2016-08-19 2022-11-29 Wisconsin Alumni Research Foundation Computer architecture with synergistic heterogeneous processors
US10353736B2 (en) 2016-08-29 2019-07-16 TidalScale, Inc. Associating working sets and threads
US10460513B2 (en) * 2016-09-22 2019-10-29 Advanced Micro Devices, Inc. Combined world-space pipeline shader stages
CN107291535B (zh) * 2017-05-18 2020-08-14 深圳先进技术研究院 多核系统的资源管理方法、资源管理设备及电子设备
US10585703B2 (en) * 2017-06-03 2020-03-10 Apple Inc. Dynamic operation allocation for neural networks
US10228972B2 (en) * 2017-06-22 2019-03-12 Banuba Limited Computer systems and computer-implemented methods for dynamically adaptive distribution of workload between central processing unit(s) and graphics processing unit(s)
US11023135B2 (en) 2017-06-27 2021-06-01 TidalScale, Inc. Handling frequently accessed pages
KR20240010541A (ko) * 2017-07-04 2024-01-23 삼성전자주식회사 다중 코어 변환에 의한 비디오 복호화 방법 및 장치, 다중 코어 변환에 의한 비디오 부호화 방법 및 장치
US10628223B2 (en) * 2017-08-22 2020-04-21 Amrita Vishwa Vidyapeetham Optimized allocation of tasks in heterogeneous computing systems
US11119835B2 (en) 2017-08-30 2021-09-14 Intel Corporation Technologies for providing efficient reprovisioning in an accelerator device
US10817347B2 (en) 2017-08-31 2020-10-27 TidalScale, Inc. Entanglement of pages and guest threads
US11334469B2 (en) * 2018-04-13 2022-05-17 Microsoft Technology Licensing, Llc Compound conditional reordering for faster short-circuiting
CN110716750A (zh) * 2018-07-11 2020-01-21 超威半导体公司 用于部分波前合并的方法和系统
CN110968320A (zh) * 2018-09-30 2020-04-07 上海登临科技有限公司 针对异构硬件架构的联合编译方法和编译系统
WO2021117186A1 (ja) * 2019-12-12 2021-06-17 三菱電機株式会社 データ処理実行装置、データ処理実行方法及びデータ処理実行プログラム

Family Cites Families (21)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
JPH09269903A (ja) 1996-04-02 1997-10-14 Hitachi Ltd プロセス管理方式
US6345041B1 (en) 1996-10-24 2002-02-05 Hewlett-Packard Company Method and apparatus for automatic load-balancing on multisegment devices
US6480930B1 (en) 1999-09-15 2002-11-12 Emc Corporation Mailbox for controlling storage subsystem reconfigurations
US6560717B1 (en) 1999-12-10 2003-05-06 Art Technology Group, Inc. Method and system for load balancing and management
US7437536B2 (en) * 2004-05-03 2008-10-14 Sony Computer Entertainment Inc. Systems and methods for task migration
US7437581B2 (en) * 2004-09-28 2008-10-14 Intel Corporation Method and apparatus for varying energy per instruction according to the amount of available parallelism
JP4936517B2 (ja) * 2006-06-06 2012-05-23 学校法人早稲田大学 ヘテロジニアス・マルチプロセッサシステムの制御方法及びマルチグレイン並列化コンパイラ
US20080005591A1 (en) * 2006-06-28 2008-01-03 Trautman Mark A Method, system, and apparatus for dynamic thermal management
KR101366075B1 (ko) * 2007-12-20 2014-02-21 삼성전자주식회사 멀티코어 플랫폼에서의 태스크 이동 방법 및 장치
US8615647B2 (en) * 2008-02-29 2013-12-24 Intel Corporation Migrating execution of thread between cores of different instruction set architecture in multi-core processor and transitioning each core to respective on / off power state
US9223677B2 (en) * 2008-06-11 2015-12-29 Arm Limited Generation of trace data in a multi-processor system
JP5300005B2 (ja) 2008-11-28 2013-09-25 インターナショナル・ビジネス・マシーンズ・コーポレーション スレッド実行制御方法、およびシステム
WO2010067377A2 (en) * 2008-12-08 2010-06-17 Kpit Cummins Infosystems Ltd. Method for reorganizing tasks for optimization of resources
US8528001B2 (en) * 2008-12-15 2013-09-03 Oracle America, Inc. Controlling and dynamically varying automatic parallelization
US8881157B2 (en) * 2009-09-11 2014-11-04 Empire Technology Development Llc Allocating threads to cores based on threads falling behind thread completion target deadline
US9354944B2 (en) * 2009-07-27 2016-05-31 Advanced Micro Devices, Inc. Mapping processing logic having data-parallel threads across processors
US8418187B2 (en) * 2010-03-01 2013-04-09 Arm Limited Virtualization software migrating workload between processing circuitries while making architectural states available transparent to operating system
US9542231B2 (en) * 2010-04-13 2017-01-10 Et International, Inc. Efficient execution of parallel computer programs
CN101923491A (zh) * 2010-08-11 2010-12-22 上海交通大学 多核环境下线程组地址空间调度和切换线程的方法
US20120079501A1 (en) * 2010-09-27 2012-03-29 Mark Henrik Sandstrom Application Load Adaptive Processing Resource Allocation
US8782645B2 (en) * 2011-05-11 2014-07-15 Advanced Micro Devices, Inc. Automatic load balancing for heterogeneous cores

Also Published As

Publication number Publication date
WO2012158753A1 (en) 2012-11-22
US8683468B2 (en) 2014-03-25
EP2710467B1 (en) 2017-11-22
CN103534686B (zh) 2017-07-11
US20120297163A1 (en) 2012-11-22
KR101559090B1 (ko) 2015-10-19
CN103534686A (zh) 2014-01-22
KR20140029480A (ko) 2014-03-10
EP2710467A1 (en) 2014-03-26
JP2014513853A (ja) 2014-06-05

Similar Documents

Publication Publication Date Title
JP5711853B2 (ja) 異種コアの自動カーネル移行
JP5859639B2 (ja) 異種コア用の自動負荷バランシング
US11954036B2 (en) Prefetch kernels on data-parallel processors
US11847508B2 (en) Convergence among concurrently executing threads
Yan et al. Alleviating irregularity in graph analytics acceleration: A hardware/software co-design approach
JP6159825B2 (ja) ハードウェアポインタを使用したsimdコア内での分岐ブランチに対するソリューション
Ubal et al. Multi2Sim: A simulation framework for CPU-GPU computing
Prabhu et al. Exposing speculative thread parallelism in SPEC2000
US20120331278A1 (en) Branch removal by data shuffling
Ezudheen et al. Parallelizing SystemC kernel for fast hardware simulation on SMP machines
JP2009259241A (ja) 汎用プロセッサによるリターゲティングされたグラフィックプロセッサ加速コードの実行
JP6236093B2 (ja) 並列パイプラインにおいてブランチを分岐するためのハードウェアおよびソフトウェアソリューション
Mikushin et al. KernelGen--The Design and Implementation of a Next Generation Compiler Platform for Accelerating Numerical Models on GPUs
US11934867B2 (en) Techniques for divergent thread group execution scheduling
US20090133022A1 (en) Multiprocessing apparatus, system and method
Du et al. Breaking the interaction wall: A DLPU-centric deep learning computing system
Ding et al. Multicore-aware code co-positioning to reduce WCET on dual-core processors with shared instruction caches
Ross et al. Scaling OpenSHMEM for Massively Parallel Processor Arrays
Leupers et al. Scalable Simulation for MPSoC Software and Architectures

Legal Events

Date Code Title Description
A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20140919

A621 Written request for application examination

Free format text: JAPANESE INTERMEDIATE CODE: A621

Effective date: 20140919

A871 Explanation of circumstances concerning accelerated examination

Free format text: JAPANESE INTERMEDIATE CODE: A871

Effective date: 20140919

A975 Report on accelerated examination

Free format text: JAPANESE INTERMEDIATE CODE: A971005

Effective date: 20141006

A131 Notification of reasons for refusal

Free format text: JAPANESE INTERMEDIATE CODE: A131

Effective date: 20141021

A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20150120

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: 20150210

A61 First payment of annual fees (during grant procedure)

Free format text: JAPANESE INTERMEDIATE CODE: A61

Effective date: 20150306

R150 Certificate of patent or registration of utility model

Ref document number: 5711853

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