JP5859639B2 - 異種コア用の自動負荷バランシング - Google Patents

異種コア用の自動負荷バランシング Download PDF

Info

Publication number
JP5859639B2
JP5859639B2 JP2014510482A JP2014510482A JP5859639B2 JP 5859639 B2 JP5859639 B2 JP 5859639B2 JP 2014510482 A JP2014510482 A JP 2014510482A JP 2014510482 A JP2014510482 A JP 2014510482A JP 5859639 B2 JP5859639 B2 JP 5859639B2
Authority
JP
Japan
Prior art keywords
kernel
processor core
instructions
scheduling
runtime
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
JP2014510482A
Other languages
English (en)
Other versions
JP2014513373A (ja
JP2014513373A5 (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 JP2014513373A publication Critical patent/JP2014513373A/ja
Publication of JP2014513373A5 publication Critical patent/JP2014513373A5/ja
Application granted granted Critical
Publication of JP5859639B2 publication Critical patent/JP5859639B2/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/50Allocation of resources, e.g. of the central processing unit [CPU]
    • 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/5083Techniques for rebalancing the load in a distributed system

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)および/または他の特殊化された種類の処理コアを含んでもよい。アーキテクチャが異なる種類の複数のコアを含む場合、異種マルチコアアーキテクチャと呼ばれ得る。
現在、「スケジューラ」とも呼ばれ得る、オペレーティングシステム(OS)スケジューラまたはユーザレベルスケジューラは、例えばラウンドロビン方式等の種々の方式を用いて、異種マルチコアアーキテクチャを備えたコンピュータシステム上で作動する作業負荷をスケジューリングし得る。加えて、スケジューラは、コアの可用性に基づいて、これらの作業負荷をスケジューリングしてもよい。代替として、プログラマは、ランタイムシステムと組み合わせて、作業負荷をスケジューリングしてもよい。そのような場合、プログラマは、スケジューリングを行うために、ソフトウェアプラットフォームを利用し得る。例えば、OpenCL(登録商標)(Open Computing Language)フレームワークは、異種コンピュータ環境にわたってプログラミングをサポートし、異種計算用の低レベルアプリケーションプログラミングインターフェース(API)を含む。OpenCL(登録商標)フレームワーク(本明細書では、概して、「OpenCL」(登録商標)と呼ばれる)は、それぞれOpenCL(登録商標)デバイスと関連付けられる実行待ち行列を定義するために使用され得るC風の言語インターフェースを含む。OpenCL(登録商標)デバイスは、CPU、GPU、または異種マルチコアアーキテクチャ内の少なくとも1つのプロセッサコアを備える他のユニットであってもよい。OpenCL(登録商標)フレームワークでは、関数呼び出しは、OpenCL(登録商標)計算カーネル、または単に「計算カーネル」と呼ばれ得る。ソフトウェアプログラマは、実行待ち行列内の計算カーネルをスケジューリングしてもよい。計算カーネルは、1つ以上の計算の作業単位を生じさせるように、1つ以上のデータの記録と合致させられてもよい。各作業単位は、一意の識別子(ID)を有し得る。
上記のスケジューリングモデルは、スケジューリング方式とシステムリソースとの間に不一致がある場合に、移植性および性能が制限されるおそれがある。プログラマは、様々なシステム構成に及ぶアプリケーションの提供を試みつつ、移植性と効率性とを差し替えて提供するであろう。
複数の異種プロセッサコア間で作業単位の実行を効率的に自動スケジューリングするためのシステムおよび方法が考慮される。
一実施形態では、処理ノードは、第1のマイクロアーキテクチャを備える第1のプロセッサコアと、第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを備える第2のプロセッサコアとを含む。一実施形態では、第1のマイクロアーキテクチャは汎用マイクロアーキテクチャであり、第2のマイクロアーキテクチャは単一命令複数データ(SIMD)マイクロアーキテクチャである。処理ノードは、第1および第2のプロセッサコアのそれぞれに連結されたメモリを含む。メモリは、1つ以上の計算カーネルまたは関数呼び出しを備えたコンピュータプログラムを記憶する。コンパイラは、所与の関数呼び出しの命令を横断する際に、所与の関数呼び出しのランタイム前情報を計算するように構成されている。オペレーティングシステム(OS)内のスケジューラは、1つ以上のカーネルのそれぞれを、関連するデータのレコードと合致させることによって、1つ以上の作業単位を生じさせる。また、スケジューラは、計算されたランタイム前情報の少なくとも一部に基づいて、1つ以上の作業単位を第1または第2のプロセッサコアの何れかに割り当てる。加えて、スケジューラは、待機作業単位と同一のカーネルに対応する他の作業単位の動的なランタイム挙動に基づいて、待機作業単位のための最初の割り当てを、第1または第2のプロセッサコアの何れかから他方のプロセッサコアへ変更することが可能である。
以下の説明および図面を参照すると、これらおよび他の実施形態がさらに理解されるであろう。
異種マルチコアアーキテクチャを備えた例示的な処理ノードの一実施形態の一般化ブロック図である。 計算カーネルを定義するソースコードの一実施形態の一般化ブロック図である。 条件文を伴う計算カーネルを定義するソースコードの一実施形態の一般化ブロック図である。 ハードウェアリソースと計算カーネルとの間でスケジュールされた割り当ての一実施形態の一般化ブロック図である。 2種類のプロセッサコアのためのマイクロアーキテクチャの論理レイアウトの一実施形態の一般化ブロック図である。 汎用パイプライン実行フローの一実施形態の一般化ブロック図である。 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次元空間であってもよい。また、この次元空間は、インデックス空間と呼ばれてもよい。例えば、ソフトウェアアプリケーションは、画像ファイル等のデータの2次元(2D)アレイ上でデータ処理を行ってもよい。ソフトウェアアプリケーションは、2D画像のピクセルごとに、ソフトウェアプログラマによって開発されたアルゴリズムを実行してもよい。所与の計算カーネルは、インデックス空間(NDRange)にわたって呼び出されてもよい。
典型的には、コンパイル後、各計算カーネルの引数およびパラメータが設定される。加えて、関連メモリオブジェクトおよびバッファが作成される。計算カーネルの所与のインスタンスは、独自のソフトウェアスレッドとして実行されてもよい。インデックス空間内の所与の点における計算カーネルの所与のインスタンスは、「作業項目」と呼ばれてもよい。また、作業項目は、作業単位と呼ばれてもよい。作業単位は、2D画像の所与のピクセル(所与のインデックス)に対応するデータのレコード上で、計算カーネル内の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(登録商標)フレームワークは、並行して複数回、計算カーネルのインスタンスを呼び出し得る。JITコンパイル方法を用いると、これらのインスタンスは、後に呼び出されるように、ランタイムにコンパイルされる。計算カーネルへの召還(呼び出し)は、get_global_id(0)と名付けられた内部関数を呼び出すことによって取り出され得る、1つの関連する一意のID(作業単位ID)を有する。コード220における上記の例に関して、計算カーネル「Power2」は、INPUTアレイ内の各データ値に対して1回呼び出される。この場合、計算カーネル「Power2」は、10回呼び出される。したがって、10個の一意の作業単位IDが取り出される。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を備えてもよい。これらの計算カーネルの各々の内部に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は、アプリケーションの実行に利用可能なメモリが不十分であるときに、行動方針を判定し得る。前述のように、アプリケーションは、1つよりも多くの計算カーネルに分割されてもよく、システム410は、1つよりも多くのアプリケーションを実行していてもよい。したがって、並行して作動するいくつかの計算カーネルがあってもよい。スケジューラは、OS計算カーネル422を使用して、任意の時に、同時実行計算カーネルのどちらがプロセッサコア112および172に割り付けられるかを決定し得る。OS計算カーネル422は、タイムスライスと呼ばれる所与の時間量にわたって、1つ以上のコアを有し得るプロセッサのコア上で、プロセスが作動することを可能にし得る。オペレーティングシステム420内のスケジューラ424は、計算カーネルをコアに割り当てるための決定論理を備えてもよい。
一実施形態では、1つだけの計算カーネルが、任意の時に、ハードウェア計算ユニット412a〜412gおよび412h〜412rのうち何れか1つで実行することができる。これらのハードウェア計算ユニットは、関連データを用いて所与の作業単位の所与の命令の実行を取り扱うことが可能なハードウェアを備える。このハードウェアは、加算、乗算、ゼロ検出、ビット単位シフト、除算、ビデオグラフィックスおよびマルチメディア命令、またはプロセッサ設計の当業者に公知である他の演算を行うように構成された算術論理演算ユニットを含んでもよい。これらのハードウェア計算ユニットは、マルチスレッドプロセッサ内のハードウェアスレッド、SIMDマイクロアーキテクチャ内の並列ハードウェアカラム等を含んでもよい。
図4の鎖線は、割り当てを表し、必ずしも直接物理接続を表すとは限らない。したがって、例えば、ハードウェア計算ユニット412aは、作業単位442dを実行するように割り当てられてもよい。しかしながら、以降で(例えば、コンテキスト切り替え後に)、ハードウェア計算ユニット412aは、作業単位442hを実行するように割り当てられてもよい。一実施形態では、スケジューラ424は、ラウンドロビン方式を用いて、ハードウェア計算ユニット412a〜412rへ作業単位442a〜442qをスケジューリングしてもよい。代替として、スケジューラ424は、ラウンドロビン方式を用いて、コア112および172へ作業単位442a〜442qをスケジューリングしてもよい。所与のハードウェア計算ユニットへの所与の作業単位の割り当ては、関連プロセッサコアによって行われてもよい。別の実施形態では、スケジューラ424は、プロセッサコア112および172の可用性に基づいてスケジューリングを行ってもよい。さらに別の実施形態では、スケジューラ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は、他の命令によって消費されるリソースよりもはるかに大きくなり得る。実際に、計算集約的命令は、他の命令によるハードウェアリソースの使用を阻止し得る。
いくつかの計算集約的タスクは、汎用コア112内の共有リソースに圧力を加え得る。したがって、スループット損失が、計算集約的プロセスと、共有リソースを待つ他のプロセスとの両方に起こる。加えて、いくつかの命令は、共有リソース上で行われている計算をサポートするように、ダイ上の共有リソースおよび他のリソースを占有する。そのような長い待ち時間の命令は、長い待ち時間の間に、他のプロセスがいくつかのリソースを使用することを同時に阻止し得る。
ここで図7を参照すると、SIMDパイプライン実行フロー700の一実施形態を示す一般化ブロック図が示されている。命令702〜708が取り出され、関連データを伴うSIMDパイプラインに入力し得る。命令704は、分岐等の制御フロー転送命令であってもよい。命令706は、取られたパス内の第1の命令であってもよい。例えば、分岐命令704は、高レベル言語プログラム内のIF文と関連付けられてもよい。命令706は、高レベル言語プログラム内のTHEN文と関連付けられてもよい。命令708は、取られていないパス内の第1の命令であってもよい。命令708は、高レベル言語プログラム内のELSE文と関連付けられてもよい。
所与の行内の計算ユニットのそれぞれは、同一の計算ユニットであってもよい。これらの計算ユニットのそれぞれは、同一の命令に作用し得るが、異なるデータが異なる作業単位と関連付けられる。示されるように、いくつかの作業単位は、分岐命令704によって提供されるテストに合格し、他の作業単位はテストに不合格となる。SIMDコア172は、利用可能なパスのそれぞれを実行し、現在のパスを選択しなかった作業項目に対応する計算ユニット等の実行単位を選択的に無効にし得る。例えば、If−Then−Else構築文の実行中に、SIMDアーキテクチャの各列内に「Then」(パスA)および「Else」(パスB)パスを実行するように構成された実行単位がある。第1および第2の作業単位が実行を停止させ、第3の作業単位が継続中の実行を続けるのを待つ際に、並列実行の効率性が低減させられ得る。したがって、計算ユニットの全てが、分岐命令704の実行後に、所与の行内で動作中の計算ユニット710であるわけではない。多数の計算ユニットが所与のパイプ段階中で動作していない場合、SIMDコアの効率性およびスループットが低減させられる。
ここで図8を参照すると、静的情報を利用してプロセッサコアへ作業単位をスケジューリングするための方法800の一実施例が示されている。処理ノード110で具現化される構成要素、および上記の図4に示されたハードウェアリソース割り当ては、概して、方法800に従って動作し得る。論議の目的のために、この実施形態および以降で説明される方法の後続の実施形態におけるステップは、起こった順番に示される。しかしながら、いくつかのステップは、示されるのとは異なる順番で起こってもよく、いくつかのステップは、同時に行われてもよく、いくつかのステップは、他のステップと組み合わせられてもよく、いくつかのステップは、別の実施形態では欠けていてもよい。
ブロック802では、ソフトウェアプログラムまたはサブルーチンの場所が特定され、分析されてもよい。このフトウェアプログラムは、異種アーキテクチャ上でのコンパイルおよび実行のために書かかれてもよい。プログラムコードは、ソフトウェアアプリケーション、サブルーチン、動的なリンクされたライブラリまたはその他の任意の部分を参照し得る。パス名は、コマンドプロンプトにおいてユーザに入力されてもよく、ソースコードをコンパイルし始めるために、所与のディレクトリ場所またはその他から読み出されてもよい。プログラムコードは、C等の高水準言語や、OpenCL(登録商標)等のC風の言語等で、設計者によって書かれてもよい。一実施形態では、ソースコードは、静的にコンパイルされる。そのような実施形態では、ソースコードは、静的なフロントエンドコンパイル中に、中間表現(IR)に変換されてもよい。バックエンドコンパイルステップは、IRを機械コードに変換し得る。静的バックエンドコンパイルは、より多くの転換および最適化を行ってもよい。コンパイラは、プログラムコード内のカーネルを識別し得る。
ブロック804では、コンパイラは、カーネルの1つ以上の命令を読み取り、それらを分析し得る。条件文が識別された場合には(条件付きブロック806)、ブロック808において、条件文の数のカウントが増分されてもよい。条件文は、分岐等の制御フロー転送命令であってもよい。一実施形態では、別個のカウントが、例えば前方/後方分岐、直接/間接分岐、ジャンプ等の異なる種類の制御フロー転送命令のために維持されてもよい。コンパイラまたは他のツールは、分岐の方向、分岐の標的、またはメモリアクセス動作のアドレスを、静的に判定することが可能であり得る。しかしながら、一実施形態では、典型的には、実行中に関連データに行われる何らかの処理が、コンパイル中に行われてもよい。例えば、分岐の方向(取られる、または取られない)を判定する単純なテストが行われてもよい。コンパイルは「静的コンパイル」と呼ばれてもよいが、1つ以上の小さな動的演算が行われてもよい。また、このコンパイルは、「ランタイム前コンパイル」と呼ばれてもよい。この時に行われる動的ステップの別の例は、If−Then−Else構築文のTHEN、ELSEIFおよびELSEブロックのそれぞれにおいて実行する次の命令を識別することである。
メモリアクセス命令が識別された場合には(条件付きブロック810)、ブロック812において、対応するアクセスパターンが判定されてもよい。メモリアクセスは、順次、ストライド、直接、間接、グループ集合、散乱等であってもよい。再度、何らかの動的計算が、コンパイル中に、作業単位と関連付けられたデータを用いて行われてもよい。コンパイラは、異なるカテゴリのメモリアクセスのカウントを維持し得る。
一実施形態では、コード実行前に、静的バイナリ計測(static binary instrumentation)が行われてもよい。命令が計測に適格であるかどうかを判定するために、当該命令が点検されてもよい。計測は、測定およびエラーチェック分析が、分析ルーチンによる後続の実行において行われることを可能にする。加えて、プロファイリングデータが収集されてもよい。アプリケーションの性能は、例えばメモリプロファイル等の、結果として生じる作業単位の動的挙動の理解に基づいて増大させられてもよい。加えて、同一のカーネルに由来する完成作業単位の動的挙動に基づく、作業単位の動的スケジューリングが行われてもよい。静的コンパイル時間制御フローグラフおよびデータフローグラフが、ランタイム実行の前に、初期化された変数およびプログラム挙動を検出するために用いられてもよい。しかしながら、動的挙動は、さらなる情報を提供し得る。したがって、例えばロード/読取および記憶/書込動作等の制御フロー転送命令およびメモリアクセス命令が、少なくとも計測されてもよい。しかしながら、記憶する測定データの量および実行する分析量を低減させるために、フィルタリングが、所与の命令が別様に計測に適格であるときでさえも、計測された命令の数を削減させるために使用されてもよい。
命令が計測に適格ではない場合には(条件付きブロック814)、ブロック816において、計測段階中に、分析ルーチンは、インラインに配置され、または関数呼び出し内に存在してもよく、関数名は、計測された適格命令の前または後の何れかで、コード内にインラインに配置される。最後の命令に達した場合には(条件付きブロック818)、ブロック820で、スケジューラは、ランタイム前または静的情報に従って、異種アーキテクチャ内のコア112および172のうち対応する1つで実行するように、各作業単位をスケジューリングする。
異種マルチコアアーキテクチャで使用されるスケジューラ424は、ハードウェアリソースとコア内の組織化との間の合致、および作業単位の特性を優先させてもよい。例えば、低スレッドレベル並列性を伴うカーネルに対応する作業単位が、汎用プロセッサコア112上でスケジューリングされてもよい。
制御フロー転送命令の数が所与の閾値よりも大きい作業単位が、コア112上でスケジューリングされてもよい。代替として、関連データに基づく様々な方向を伴う比較的多数の制御フロー命令を備えるカーネルの作業単位が、コア112上でスケジュールに入れられてもよい。例えば、カーネルが多数の制御フロー転送命令を有するが、方向(取られる、取られない)が多数の作業単位の間で一致する場合には、作業単位は、SIMDコア172上でスケジューリングされてもよい。そうでなければ、制御フロー転送命令の命令が一貫していない、または様々である場合には、関連作業単位は、コア112上でスケジューリングされてもよい。
比較的多数のメモリアクセス命令が、順次またはストライド方式でメモリ位置のアクセスを行う場合には、対応する作業単位は、SIMDコア172上でスケジューリングされてもよい。比較的多数のメモリアクセス命令が、散乱または間接方式でメモリ位置のアクセスを行う場合には、対応する作業単位は、汎用コア112上でスケジューリングされてもよい。ランタイムにおいて、OpenCL(登録商標)コンパイラは、例えば汎用コア112およびSIMDコア172等のOpenCL(登録商標)デバイス種類ごとの複数のバージョンのカーネルを生成し得る。一実施例では、スケジューラ424は、SIMDコア172上で実行するように、所与のカーネルの最初の256個の作業単位をスケジューリングしてもよい。しかしながら、これらの作業単位の監視された動的挙動に基づいて、スケジューラ424は、汎用コア112へ、所与のカーネルの最後の16個の作業単位をスケジューリングしてもよい。
ここで図9を参照すると、動的情報を利用してプロセッサコアへ作業単位をスケジューリングするための方法900の一実施例が示されている。処理ノード110で具現化される構成要素および上記の図4に示されるハードウェアリソース割り当ては、概して、方法900に従って動作し得る。論議の目的で、この実施形態および以降で説明される方法の後続の実施形態におけるステップは、起こった順番に示される。しかしながら、いくつかのステップは、示されるのとは異なる順番で起こってもよく、いくつかのステップは、同時に行われてもよく、いくつかのステップは、他のステップと組み合わせられてもよく、いくつかのステップは、別の実施形態では欠けていてもよい。
ブロック902では、データの関連レコードが、所与のカーネルの各作業単位に割り当てられる。ブロック904では、スケジューラ424が、作業単位を異種コアへスケジューリングする。方法700は、スケジューリングを行うために使用されてもよい。ブロック906では、プロセッサコア112および172が、対応するスケジュールされた作業単位を実行する。ブロック908では、計測コードおよびツールが、実行作業単位の動的挙動を監視して収集する。収集されたデータは、1つ以上のテーブルに記憶されてもよい。1つ以上のテーブルの入力は、測定されている現在のシステムトポロジーを示すために、プロセッサコア識別子(ID)、カーネルIDおよび作業単位IDを利用し得る。
事象インデックスは、計測されたコードによって測定されている事象の種類を示し得る。実際の測定値が、比率値とともに記憶されてもよい。比率は、対応する頻度または割合の測定を含んでもよい。測定値および比率値が有効であるか否かを示すために、状態フィールドが使用されてもよい。1つ以上の構成可能な閾値が記憶されてもよい。一実施形態では、これらの閾値は、プログラム可能である。
スケジュールされた作業単位が、実行されるのを待っている場合には(条件付きブロック910)、ブロック912において、同一のカーネルに対応する任意の実行作業単位の監視された動的挙動が、分析されてもよい。ブロック914では、異種コアのうち1つが、所与の作業単位の効率的な実行のために好適であると判定される。例えば、作業単位あたりの命令の数が増加するにつれて、命令が汎用関数性に対応する可能性がより高い。したがって、測定された数が所与の閾値を通過するとき、汎用コア112は、待機作業単位を実行するのにより好適であり得る。加えて、取られた分岐の間の命令のカウントが使用されてもよい。
コード内の所与のループおよびループの数は、SIMDマイクロアーキテクチャを用いた効率的な実行を示し得る。所与の閾値を超える、実行された分岐の数および他の種類の制御フロー転送命令の数は、汎用コア112がより効率的な実行を提供することを示し得る。同様に、比較的多数のキャッシュミスは、汎用コア112が、作業単位を実行するのにSIMDコア172よりも効率的であり得ることを示し得る。実行された浮動小数点演算、実行されたグラフィックス処理動作、および書き込みバッファオーバーフローによるパイプライン失速の比較的多数は、SIMDコア172が、待機作業単位のためにより効率的な実行を提供することを示し得る。また、待機作業単位を実行する好ましいOpenCL(登録商標)デバイスの種類を判定するための実行時間が使用されてもよい。他のランタイム基準が可能であり、考慮される。加えて、その基準のそれぞれは、実行に好ましいOpenCL(商標)デバイスの種類を判定するために、全ての基準の集計式に使用される関連加重を有し得る。
ブロック916では、待機作業単位の効率的な実行のために上記で判定されるプロセッサコアと、以前にスケジュールされたプロセッサコアとの間で、比較が行われる。合致がある場合には(条件付きブロック918)、ブロック920において、スケジューラ424は、待機作業単位を、以前にスケジュールされたプロセッサコア上にスケジューリングする。合致がない場合には(条件付きブロック918)、ブロック922において、スケジューラ424は、対応するカーネルの動的挙動を利用して、待機作業単位を、上記の分析から見出されるプロセッサコア上にスケジューリングする。
上記の実施形態は、ソフトウェアを含み得ることを留意されたい。そのような実施形態では、方法および/または機構を実装するプログラム命令は、搬送され、またはコンピュータ可読媒体上に記憶されてもよい。プログラム命令を記憶するように構成された多数の種類の媒体が利用可能であり、ハードディスク、フロッピー(登録商標)ディスク、CD−ROM、DVD、フラッシュメモリ、プログラマブルROM(PROM)、ランダムアクセスメモリ(RAM)、および種々の他の形態の揮発性または不揮発性記憶装置を含む。一般的に言えば、コンピュータアクセス可能記憶媒体は、命令および/またはデータをコンピュータに提供するように、使用中にコンピュータによってアクセス可能な任意の記憶媒体を含んでもよい。例えば、コンピュータアクセス可能記憶媒体は、磁気または光学媒体、例えば、ディスク(固定またはリムーバブル)、テープ、CD−ROM若しくはDVD−ROM、CD−R、CD−RW、DVD−R、DVD−RWまたはBlu−Ray(登録商標)等の記憶媒体を含んでもよい。さらに、記憶媒体は、RAM(例えば、同期型ダイナミックRAM(SDRAM)、ダブルデータレート(DDR、DDR2、DDR3等)SDRAM、低出力DDR(LPDDR2等)SDRAM、Rambus DRAM(RDRAM)、スタティックRAM(SRAM)等)、ROM、フラッシュメモリ、ユニバーサルシリアルバス(USB)インターフェース等の周辺インターフェースを介してアクセス可能な不揮発性メモリ(例えば、フラッシュメモリ)等の、揮発性または不揮発性メモリ媒体を含んでもよい。記憶媒体は、微小電気機械システム(MEMS)、ならびにネットワークおよび/または無線リンク等の通信媒体を介してアクセス可能な記憶媒体を含んでもよい。
加えて、プログラム命令は、C等の高レベルプログラミング言語プログラム、またはVerilog、VHDL若しくはGDS IIストリームフォーマット(GDSII)等のデータベース形式等の設計言語(HDL)で、ハードウェア機能性の挙動レベル記述またはレジスタ転送レベル(RTL)記述を備えてもよい。場合によっては、記述は、合成ライブラリからのゲートのリストを備えるネットリストを生成するように記述を合成し得る合成ツールによって読み取られてもよい。ネットリストは、システムを備えるハードウェアの機能性も表す一式のゲートを備える。次いで、ネットリストは、マスクに適用される幾何学形状を表すデータセットを生成するように配置され、送られてもよい。次いで、マスクは、システムに対応する1つまたは複数の半導体回路を生産するために、種々の半導体製造ステップで使用されてもよい。代替として、コンピュータアクセス可能記憶媒体の命令は、所望に応じて、ネットリスト(合成ライブラリを伴う、または伴わない)またはデータセットであってもよい。加えて、命令は、Cadence(登録商標)、EVE(登録商標)およびMentor Graphics(登録商標)等のベンダからのハードウェアベース型のエミュレータによって、エミュレーションの目的で利用されてもよい。
上記の実施形態は、非常に詳細に説明されているが、上記の開示が完全に理解されると、多数の変形例および修正が、当業者に明白となるであろう。以下の請求項は、全てのそのような変形例および修正を包含するように解釈されることが意図される。

Claims (15)

  1. 第1の計算カーネルが第1の閾値以下の数の分岐命令を含むことを示すランタイム前静的情報であって、前記第1の計算カーネルのコンパイル中に判定されたランタイム前静的情報の少なくとも一部に基づいて、前記第1の計算カーネルを、複数のプロセッサコアのうち第1のマイクロアーキテクチャを有する第1のプロセッサコアにスケジューリングするステップと、
    前記第1の閾値よりも大きい数の分岐命令を含むことを示すランタイム前静的情報の少なくとも一部に基づいて、前記第1の計算カーネルを、複数のプロセッサコアのうち第2のマイクロアーキテクチャを有する第2のプロセッサコアにスケジューリングするステップと、
    前記第1のプロセッサコアまたは前記第2のプロセッサコアでの前記第1の計算カーネルのランタイム挙動であって、実行中の計算カーネルの挙動であるランタイム挙動に対応する、測定されたランタイム情報を受信するステップと、
    実行された分岐命令の数が第2の閾値よりも大きいことを示す前記測定されたランタイム情報の少なくとも一部に基づいて、前記第1の計算カーネルのスケジュールを、前記第1のプロセッサコアから、前記複数のプロセッサコアのうち前記第2のプロセッサコアに変更するステップと、
    を含む、異種マルチコアアーキテクチャにおいて作業単位をスケジューリングするための方法。
  2. 前記第1の計算カーネルのコンパイル中に、前記第1のプロセッサコアおよび第2のプロセッサコアの各々に対して、前記第1の計算カーネルの異なるバージョンのバイナリコードを生成するステップをさらに含む、請求項1に記載の方法。
  3. 前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャである、請求項1に記載の方法。
  4. 前記スケジューリングするステップは、
    第2の計算カーネルが第1の数の命令を含むことを、散乱メモリアクセスまたは間接的メモリアクセスを用いて判定するステップであって、前記散乱メモリアクセスは、不連続な記憶場所に対するアクセスであるステップと、
    前記第2の計算カーネルが第2の数の命令を含むことを、順次またはストライドメモリアクセスを用いて判定するステップと、
    命令の前記第1の数が命令の前記第2の数よりも大きいという判定に応じて、前記第2の計算カーネルを前記第2のプロセッサコアにスケジューリングするステップと、
    命令の前記第1の数が命令の前記第2の数よりも大きくないという判定に応じて、前記第2の計算カーネルを前記第1のプロセッサコアにスケジューリングするステップと、
    をさらに含む、請求項1に記載の方法。
  5. 2の計算カーネルが第3の閾値以下の数の第2の種類の命令を含むことを示すランタイム前静的情報であって、前記第2の計算カーネルのコンパイル中に判定されたランタイム前静的情報の少なくとも一部に基づいて、前記第2の計算カーネルを、前記複数のプロセッサコアのうち前記第2のプロセッサコアにスケジューリングするステップであって、前記第2の種類は、暗号、浮動小数点、ガーベジコレクションおよびビデオグラフィックスのうち少なくとも1つに対応する、ステップと、
    実行された前記第2の種類の命令の数が少なくとも第4の閾値であることを示す、測定されたランタイム情報の少なくとも一部に基づいて、前記第2の計算カーネルのスケジュールを、前記第2のプロセッサコアから前記第1のプロセッサコアに変更するステップであって、前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャである、ステップと、をさらに含む、請求項1に記載の方法。
  6. 前記第1のプロセッサコアにスケジューリングするステップは、
    前記第1の計算カーネルに含まれるメモリアクセスの少なくとも一部に基づいて、前記第1の計算カーネルを、前記第1のプロセッサコアにスケジューリングするステップを含む、請求項に記載の方法。
  7. 第1のマイクロアーキテクチャを備える第1のプロセッサコアと、
    前記第1のマイクロアーキテクチャとは異なる第2のマイクロアーキテクチャを備える第2のプロセッサコアと、
    スケジューラを備えるオペレーティングシステムと、を備え、
    前記スケジューラは、
    第1の計算カーネルが第1の閾値以下の数の分岐命令を含むことを示すランタイム前静的情報であって、第1の計算カーネルのコンパイル中に判定されたランタイム前静的情報の少なくとも一部に基づいて、前記第1の計算カーネルを前記第1のプロセッサコアにスケジューリングし、
    分岐命令の数が前記第1の閾値よりも大きいことを示すランタイム前静的情報の少なくとも一部に基づいて、前記第1の計算カーネルを、複数のプロセッサコアのうち第2のプロセッサコアにスケジューリングし、
    前記第1のプロセッサコアまたは前記第2のプロセッサコアでの前記第1の計算カーネルのランタイム挙動であって、実行中の計算カーネルの挙動であるランタイム挙動に対応する、測定されたランタイム情報を受信し、
    実行された分岐命令の数が第2の閾値よりも大きいことを示す前記測定されたランタイム情報の少なくとも一部に基づいて、前記第1の計算カーネルのスケジュールを、前記第1のプロセッサコアから前記第2のプロセッサコアに変更するように構成されている、
    異種マルチコアアーキテクチャを含むコンピュータシステム。
  8. コンパイラは、前記第1の計算カーネルのコンパイル中に、前記第1のプロセッサコアおよび第2のプロセッサコアの各々に対して、前記第1の計算カーネルの異なるバージョンのバイナリコードを生成するように構成されている、請求項7に記載のコンピュータシステム。
  9. 前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャである、請求項7に記載のコンピュータシステム。
  10. 前記スケジューラは、前記第1のプロセッサコアにスケジューリングすることを実行するために、
    第2の計算カーネルが第1の数の命令を含むことを、散乱メモリアクセスまたは間接的メモリアクセスを用いて判定し、
    前記第2の計算カーネルが第2の数の命令を含むことを、順次またはストライドメモリアクセスを用いて判定し、
    命令の前記第1の数が命令の前記第2の数よりも大きいという判定に応じて、前記第2の計算カーネルを前記第2のプロセッサコアにスケジューリングし、
    命令の前記第1の数が命令の前記第2の数よりも大きくないという判定に応じて、前記第2の計算カーネルを前記第1のプロセッサコアにスケジューリングするように構成されており
    前記散乱メモリアクセスは、不連続な記憶場所に対するアクセスである、
    請求項9に記載のコンピュータシステム。
  11. 前記スケジューラは、
    前記第2の計算カーネルが第3の閾値以下の数の第2の種類の命令を含むことを示すランタイム前静的情報であって、前記第2の計算カーネルのコンパイル中に判定されたランタイム前静的情報の少なくとも一部に基づいて、第2の計算カーネルを、前記複数のプロセッサコアのうち前記第2のプロセッサコアにスケジューリングすることであって、前記第2の種類は、暗号、浮動小数点、ガーベジコレクションおよびビデオグラフィックスのうち少なくとも1つに対応する、ことと、
    実行された前記第2の種類の命令の数が少なくとも第4の閾値であることを示す、測定されたランタイム情報の少なくとも一部に基づいて、前記第2の計算カーネルのスケジュールを、前記第2のプロセッサコアから前記第1のプロセッサコアに変更することであって、前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャである、ことと、
    を行うように構成されている、請求項7に記載のコンピュータシステム。
  12. 前記第1のプロセッサコアにスケジューリングすることは、
    前記第1の計算カーネルに含まれるメモリアクセスの少なくとも一部に基づいて、前記第1の計算カーネルを、前記第1のプロセッサコアにスケジューリングすることを含む、請求項11に記載のコンピュータシステム。
  13. 異種マルチコアアーキテクチャにおいて計算カーネルをスケジューリングするように構成されたプログラム命令を記憶するコンピュータ可読記憶媒体であって、
    前記プログラム命令は、
    第1の計算カーネルが第1の閾値以下の数の分岐命令を含むことを示すランタイム前静的情報であって、前記第1の計算カーネルのコンパイル中に判定されたランタイム前静的情報の少なくとも一部に基づいて、前記第1の計算カーネルを、複数のプロセッサコアのうち第1のマイクロアーキテクチャを有する第1のプロセッサコアにスケジューリングすることと、
    前記第1の閾値よりも大きい数の分岐命令を含むことを示すランタイム前静的情報の少なくとも一部に基づいて、前記第1の計算カーネルを、複数のプロセッサコアのうち第2のマイクロアーキテクチャを有する第2のプロセッサコアにスケジューリングすることと、
    前記第1のプロセッサコアまたは前記第2のプロセッサコアでの前記第1の計算カーネルのランタイム挙動であって、実行中の計算カーネルの挙動であるランタイム挙動に対応する、測定されたランタイム情報を受信することと、
    実行された分岐命令の数が第2の閾値よりも大きいことを示す前記測定されたランタイム情報の少なくとも一部に基づいて、前記第1の計算カーネルのスケジュールを、前記第1のプロセッサコアから、前記複数のプロセッサコアのうち前記第2のプロセッサコアに変更することと、
    を行うように実行可能である、
    コンピュータ可読記憶媒体。
  14. 前記第1のマイクロアーキテクチャは、単一命令複数データ(SIMD)マイクロアーキテクチャである、請求項13に記載のコンピュータ可読記憶媒体。
  15. 前記プログラム命令は、前記第1のプロセッサコアにスケジューリングすることを実行するために、
    第2の計算カーネルが第1の数の命令を含むことを、散乱メモリアクセスまたは間接的メモリアクセスを用いて判定し、
    前記第2の計算カーネルが第2の数の命令を含むことを、順次またはストライドメモリアクセスを用いて判定し、
    命令の前記第1の数が命令の前記第2の数よりも大きいという判定に応じて、前記第2の計算カーネルを前記第2のプロセッサコアにスケジューリングし、
    命令の前記第1の数が命令の前記第2の数よりも大きくないという判定に応じて、前記第2の計算カーネルを前記第1のプロセッサコアにスケジューリングするように実行可能であ
    前記散乱メモリアクセスは、不連続な記憶場所に対するアクセスである、
    請求項14に記載のコンピュータ可読記憶媒体。
JP2014510482A 2011-05-11 2012-05-11 異種コア用の自動負荷バランシング Active JP5859639B2 (ja)

Applications Claiming Priority (3)

Application Number Priority Date Filing Date Title
US13/105,250 US8782645B2 (en) 2011-05-11 2011-05-11 Automatic load balancing for heterogeneous cores
US13/105,250 2011-05-11
PCT/US2012/037433 WO2012155010A1 (en) 2011-05-11 2012-05-11 Automatic load balancing for heterogeneous cores

Publications (3)

Publication Number Publication Date
JP2014513373A JP2014513373A (ja) 2014-05-29
JP2014513373A5 JP2014513373A5 (ja) 2015-07-02
JP5859639B2 true JP5859639B2 (ja) 2016-02-10

Family

ID=46124772

Family Applications (1)

Application Number Title Priority Date Filing Date
JP2014510482A Active JP5859639B2 (ja) 2011-05-11 2012-05-11 異種コア用の自動負荷バランシング

Country Status (6)

Country Link
US (1) US8782645B2 (ja)
EP (1) EP2707797B1 (ja)
JP (1) JP5859639B2 (ja)
KR (1) KR101839544B1 (ja)
CN (1) CN103562870B (ja)
WO (1) WO2012155010A1 (ja)

Families Citing this family (73)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US8683468B2 (en) * 2011-05-16 2014-03-25 Advanced Micro Devices, Inc. Automatic kernel migration for heterogeneous cores
US10013731B2 (en) * 2011-06-30 2018-07-03 Intel Corporation Maximizing parallel processing in graphics processors
US9195501B2 (en) * 2011-07-12 2015-11-24 Qualcomm Incorporated Instruction culling in graphics processing unit
US8954017B2 (en) * 2011-08-17 2015-02-10 Broadcom Corporation Clock signal multiplication to reduce noise coupled onto a transmission communication signal of a communications device
US8707314B2 (en) * 2011-12-16 2014-04-22 Advanced Micro Devices, Inc. Scheduling compute kernel workgroups to heterogeneous processors based on historical processor execution times and utilizations
US9430807B2 (en) 2012-02-27 2016-08-30 Qualcomm Incorporated Execution model for heterogeneous computing
KR20130115574A (ko) * 2012-04-12 2013-10-22 삼성전자주식회사 단말기에서 태스크 스케줄링을 수행하는 방법 및 장치
US9626216B2 (en) * 2012-05-09 2017-04-18 Nvidia Corporation Graphics processing unit sharing between many applications
US9280395B2 (en) * 2012-05-30 2016-03-08 Intel Corporation Runtime dispatching among a heterogeneous group of processors
US8930956B2 (en) * 2012-08-08 2015-01-06 International Business Machines Corporation Utilizing a kernel administration hardware thread of a multi-threaded, multi-core compute node of a parallel computer
US20160162293A1 (en) * 2013-03-26 2016-06-09 Via Technologies, Inc. Asymmetric processor with cores that support different isa instruction subsets
US10423216B2 (en) 2013-03-26 2019-09-24 Via Technologies, Inc. Asymmetric multi-core processor with native switching mechanism
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
US9703696B1 (en) * 2013-09-11 2017-07-11 Altera Corporation Guided memory buffer allocation
KR102326945B1 (ko) * 2014-03-14 2021-11-16 삼성전자 주식회사 태스크 마이그레이션 방법 및 장치
US20170123775A1 (en) * 2014-03-26 2017-05-04 Empire Technology Development Llc Compilation of application into multiple instruction sets for a heterogeneous processor
US10133572B2 (en) * 2014-05-02 2018-11-20 Qualcomm Incorporated Techniques for serialized execution in a SIMD processing system
KR20150136345A (ko) * 2014-05-27 2015-12-07 삼성전자주식회사 태스크 그룹 전달 방법 및 이를 제공하는 전자 장치
US10291391B2 (en) * 2014-06-04 2019-05-14 Giesecke+Devrient Mobile Security Gmbh Method for enhanced security of computational device with multiple cores
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
KR102197874B1 (ko) 2014-09-01 2021-01-05 삼성전자주식회사 멀티-코어 프로세서를 포함하는 시스템 온 칩 및 그것의 쓰레드 스케줄링 방법
GB2524346B (en) * 2014-09-19 2016-12-21 Imagination Tech Ltd Separating Cores
JP6471456B2 (ja) * 2014-10-23 2019-02-20 日本電気株式会社 情報処理装置、情報処理方法および情報処理プログラム
WO2016068999A1 (en) 2014-10-31 2016-05-06 Hewlett Packard Enterprise Development Lp Integrated heterogeneous processing units
US9880953B2 (en) * 2015-01-05 2018-01-30 Tuxera Corporation Systems and methods for network I/O based interrupt steering
US9652817B2 (en) 2015-03-12 2017-05-16 Samsung Electronics Co., Ltd. Automated compute kernel fusion, resizing, and interleave
US9529950B1 (en) 2015-03-18 2016-12-27 Altera Corporation Systems and methods for performing profile-based circuit optimization using high-level system modeling
FR3035243B1 (fr) * 2015-04-20 2018-06-29 Commissariat A L'energie Atomique Et Aux Energies Alternatives Placement d'une tache de calcul sur un processeur fonctionnellement asymetrique
US9965343B2 (en) * 2015-05-13 2018-05-08 Advanced Micro Devices, Inc. System and method for determining concurrency factors for dispatch size of parallel processor kernels
US9983857B2 (en) 2015-06-16 2018-05-29 Architecture Technology Corporation Dynamic computational acceleration using a heterogeneous hardware infrastructure
US9501304B1 (en) 2015-06-16 2016-11-22 Architecture Technology Corporation Lightweight application virtualization architecture
US9626295B2 (en) * 2015-07-23 2017-04-18 Qualcomm Incorporated Systems and methods for scheduling tasks in a heterogeneous processor cluster architecture using cache demand monitoring
US11403099B2 (en) 2015-07-27 2022-08-02 Sony Interactive Entertainment LLC Backward compatibility by restriction of hardware resources
US9733978B2 (en) 2015-08-27 2017-08-15 Qualcomm Incorporated Data management for multiple processing units using data transfer costs
US9778961B2 (en) 2015-09-14 2017-10-03 Qualcomm Incorporated Efficient scheduling of multi-versioned tasks
US10360063B2 (en) 2015-09-23 2019-07-23 Qualcomm Incorporated Proactive resource management for parallel work-stealing processing systems
US9891926B2 (en) 2015-09-30 2018-02-13 International Business Machines Corporation Heterogeneous core microarchitecture
US9979656B2 (en) 2015-12-07 2018-05-22 Oracle International Corporation Methods, systems, and computer readable media for implementing load balancer traffic policies
US11550632B2 (en) * 2015-12-24 2023-01-10 Intel Corporation Facilitating efficient communication and data processing across clusters of computing machines in heterogeneous computing environment
KR101923210B1 (ko) * 2016-01-14 2018-11-28 서울대학교산학협력단 이종 멀티코어 프로세서를 활용한 암호화 처리 장치 및 암호화 처리 방법
KR102455675B1 (ko) * 2016-01-22 2022-10-17 주식회사 소니 인터랙티브 엔터테인먼트 하위 호환성을 위한 cpuid 스푸핑
JP6658136B2 (ja) * 2016-03-14 2020-03-04 コニカミノルタ株式会社 描画処理装置、画像処理装置、描画処理方法及び描画処理プログラム
US10025624B2 (en) * 2016-03-30 2018-07-17 International Business Machines Corporation Processing performance analyzer and process manager
US10303488B2 (en) * 2016-03-30 2019-05-28 Sony Interactive Entertainment Inc. Real-time adjustment of application-specific operating parameters for backwards compatibility
US11513805B2 (en) * 2016-08-19 2022-11-29 Wisconsin Alumni Research Foundation Computer architecture with synergistic heterogeneous processors
KR101828996B1 (ko) 2016-08-23 2018-02-13 울산과학기술원 이종 멀티프로세싱 환경에서 파이프라인 병렬화를 지원하는 동적 시스템 자원 할당 방법 및 장치
CN106777710A (zh) * 2016-12-22 2017-05-31 中国兵器装备集团自动化研究所 一种在fpga上实现的cuda内核的方法
CN107193631B (zh) * 2017-04-28 2019-09-13 华中科技大学 一种基于并行应用阶段检测的虚拟时间片调度方法及系统
DE102017109239A1 (de) * 2017-04-28 2018-10-31 Ilnumerics Gmbh Computerimplementiertes verfahren, computerlesbares medium und heterogenes rechnersystem
US11054883B2 (en) * 2017-06-19 2021-07-06 Advanced Micro Devices, Inc. Power efficiency optimization in throughput-based workloads
US10558418B2 (en) * 2017-07-27 2020-02-11 Advanced Micro Devices, Inc. Monitor support on accelerated processing device
KR102668340B1 (ko) * 2017-08-03 2024-05-22 넥스트 실리콘 리미티드 설정가능한 하드웨어 런타임 최적화
US10853044B2 (en) 2017-10-06 2020-12-01 Nvidia Corporation Device profiling in GPU accelerators by using host-device coordination
CN109697115B (zh) * 2017-10-20 2023-06-06 伊姆西Ip控股有限责任公司 用于调度应用的方法、装置以及计算机可读介质
CN109697121B (zh) * 2017-10-20 2023-05-05 伊姆西Ip控股有限责任公司 用于向应用分配处理资源的方法、设备和计算机可读介质
WO2019079994A1 (zh) * 2017-10-25 2019-05-02 华为技术有限公司 核心调度方法和终端
US10491524B2 (en) 2017-11-07 2019-11-26 Advanced Micro Devices, Inc. Load balancing scheme
CN109783141A (zh) * 2017-11-10 2019-05-21 华为技术有限公司 异构调度方法
US10719903B2 (en) * 2017-12-22 2020-07-21 International Business Machines Corporation On-the fly scheduling of execution of dynamic hardware behaviors
WO2019202371A1 (en) * 2018-04-16 2019-10-24 Badenhorst Emile A processor and a method of operating a processor
US20190370059A1 (en) * 2018-05-30 2019-12-05 Advanced Micro Devices, Inc. Multi-kernel wavefront scheduler
CN110716750A (zh) * 2018-07-11 2020-01-21 超威半导体公司 用于部分波前合并的方法和系统
DE102018212686A1 (de) * 2018-07-30 2020-01-30 Siemens Aktiengesellschaft Verfahren zum Betreiben einer Rechenanlage
US11188348B2 (en) * 2018-08-31 2021-11-30 International Business Machines Corporation Hybrid computing device selection analysis
US10798609B2 (en) 2018-10-16 2020-10-06 Oracle International Corporation Methods, systems, and computer readable media for lock-free communications processing at a network node
US10831535B2 (en) 2019-01-01 2020-11-10 International Business Machines Corporation Reducing minimum operating voltage through heterogeneous codes
CN110083469B (zh) * 2019-05-11 2021-06-04 广东财经大学 一种异构硬件组织运行统一内核方法及系统
US12008401B2 (en) * 2019-12-20 2024-06-11 Advanced Micro Devices, Inc. Automatic central processing unit (CPU) usage optimization
US11726823B2 (en) 2020-02-06 2023-08-15 Samsung Electronics Co., Ltd. Electronic device having heterogeneous processors and method of processing task using the heterogeneous processors
CN112199076B (zh) * 2020-10-10 2022-08-09 中国运载火箭技术研究院 一种飞行控制软件架构及其设计方法
KR20220094601A (ko) 2020-12-29 2022-07-06 삼성전자주식회사 스토리지 장치 및 그 구동 방법
US12008363B1 (en) 2021-07-14 2024-06-11 International Business Machines Corporation Delivering portions of source code based on a stacked-layer framework

Family Cites Families (22)

* 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
US6463582B1 (en) * 1998-10-21 2002-10-08 Fujitsu Limited Dynamic optimizing object code translator for architecture emulation and dynamic optimizing object code translation method
US6298477B1 (en) * 1998-10-30 2001-10-02 Sun Microsystems, Inc. Method and apparatus for selecting ways to compile at runtime
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
US7152170B2 (en) * 2003-02-20 2006-12-19 Samsung Electronics Co., Ltd. Simultaneous multi-threading processor circuits and computer program products configured to operate at different performance levels based on a number of operating threads and methods of operating
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
US8312455B2 (en) * 2007-12-19 2012-11-13 International Business Machines Corporation Optimizing execution of single-threaded programs on a multiprocessor managed by compilation
KR101366075B1 (ko) * 2007-12-20 2014-02-21 삼성전자주식회사 멀티코어 플랫폼에서의 태스크 이동 방법 및 장치
US8255669B2 (en) * 2008-01-30 2012-08-28 International Business Machines Corporation Method and apparatus for thread priority control in a multi-threaded processor based upon branch issue information including branch confidence information
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
JP2012511204A (ja) * 2008-12-08 2012-05-17 ケーピーアイティ クミンズ インフォシステムズ リミテッド リソースを最適化するためのタスク再編成方法
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
US8429635B2 (en) * 2009-10-28 2013-04-23 International Buisness Machines Corporation Controlling compiler optimizations
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
US20120079501A1 (en) * 2010-09-27 2012-03-29 Mark Henrik Sandstrom Application Load Adaptive Processing Resource Allocation

Also Published As

Publication number Publication date
EP2707797B1 (en) 2017-11-15
JP2014513373A (ja) 2014-05-29
KR20140027299A (ko) 2014-03-06
EP2707797A1 (en) 2014-03-19
KR101839544B1 (ko) 2018-03-16
CN103562870A (zh) 2014-02-05
US20120291040A1 (en) 2012-11-15
US8782645B2 (en) 2014-07-15
CN103562870B (zh) 2016-05-18
WO2012155010A1 (en) 2012-11-15

Similar Documents

Publication Publication Date Title
JP5859639B2 (ja) 異種コア用の自動負荷バランシング
JP5711853B2 (ja) 異種コアの自動カーネル移行
Yan et al. Alleviating irregularity in graph analytics acceleration: A hardware/software co-design approach
Stephenson et al. Flexible software profiling of gpu architectures
Wang et al. Dynamic thread block launch: A lightweight execution mechanism to support irregular applications on gpus
Ma et al. A memory access model for highly-threaded many-core architectures
US20120331278A1 (en) Branch removal by data shuffling
Dao et al. A performance model for GPUs with caches
US11934867B2 (en) Techniques for divergent thread group execution scheduling
Nélis et al. The variability of application execution times on a multi-core platform
Zhang et al. iMLBench: A machine learning benchmark suite for CPU-GPU integrated architectures
Voitsechov et al. Software-directed techniques for improved gpu register file utilization
Laso et al. CIMAR, NIMAR, and LMMA: Novel algorithms for thread and memory migrations in user space on NUMA systems using hardware counters
Tarakji et al. The development of a scheduling system GPUSched for graphics processing units
Salapura et al. Exploiting workload parallelism for performance and power optimization in Blue Gene
Panwar et al. Online performance projection for clusters with heterogeneous GPUs
Ukidave Architectural and Runtime Enhancements for Dynamically Controlled Multi-Level Concurrency on GPUs
Voudouris et al. Analysis and modeling of the timing bahavior of GPU architectures
Gong et al. PAQSIM: Fast Performance Model for Graphics Workload on Mobile GPUs
Ding et al. Multicore-aware code co-positioning to reduce WCET on dual-core processors with shared instruction caches
Liu et al. Two-level scratchpad memory architectures to achieve time predictability and high performance
Kiani et al. Rdmke: applying reuse distance analysis to multiple GPU kernel executions
Tendulkar et al. A runtime environment for real-time streaming applications on clustered multi-cores
Leggett et al. Parallelizing ATLAS reconstruction and simulation: issues and optimization solutions for scaling on multi-and many-CPU Platforms
Zope et al. A shared memory parallel block streaming model for irregular applications

Legal Events

Date Code Title Description
A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20140121

A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20150511

A621 Written request for application examination

Free format text: JAPANESE INTERMEDIATE CODE: A621

Effective date: 20150511

A871 Explanation of circumstances concerning accelerated examination

Free format text: JAPANESE INTERMEDIATE CODE: A871

Effective date: 20150511

A975 Report on accelerated examination

Free format text: JAPANESE INTERMEDIATE CODE: A971005

Effective date: 20150615

A131 Notification of reasons for refusal

Free format text: JAPANESE INTERMEDIATE CODE: A131

Effective date: 20150623

A521 Request for written amendment filed

Free format text: JAPANESE INTERMEDIATE CODE: A523

Effective date: 20150918

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

A61 First payment of annual fees (during grant procedure)

Free format text: JAPANESE INTERMEDIATE CODE: A61

Effective date: 20151216

R150 Certificate of patent or registration of utility model

Ref document number: 5859639

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