JP6329274B2 - コンパイラ最適化のためのメモリ参照メタデータ - Google Patents

Description

本開示は、カーネルのソースコードをコンパイルすることに関し、より詳細には、メモリアクセスができるようにカーネルのソースコードをコンパイルするための技法に関する。

いわゆる異種コンピューティングアーキテクチャが推進されている。異種コンピューティングアーキテクチャでは、カーネルと呼ばれるプログラムが、フレームワークを使用してコンパイルされる場合があり、それによって、CPU(中央演算処理装置)、GPU(グラフィックス処理装置)、FPGA(フィールドプログラマブルゲートアレイ)などの様々な異なる種類のプロセッサがカーネルを実行する場合がある。異種コンピューティングをサポートする最近のフレームワークには、OpenCLフレームワークならびにDirectComputeフレームワークがある。

本開示では、コンパイル最適化に関するメタデータを生成するためにカーネルにおけるメモリ参照のメモリエイリアシングおよびメモリオーバーラップを検出するための技法について説明する。本開示の技法を実施するために、ジャストインタイムコンパイラ(JIT)などのコンパイラは、「カーネル」とも呼ばれるプログラムのソースコードをバイナリファイルとしてコンパイルする。コンパイラを実行するコンパイリングプロセッサは、実行時に(コンパイリングプロセッサがカーネルを実行するのに必要な引数を生成するときに)OpenCLなどの異種コンピューティングフレームワークを使用してカーネルをコンパイルする場合がある。本開示において説明する技法では、ドライバが、ターゲットプロセッサに、生成された引数を使用してカーネルを実行するよう命令する代わりに、引数を分析する。これらの引数は、まとめてバッファに渡され、カーネルを実行するターゲットプロセッサに渡される。ドライバ/ランタイムは、分析に基づいて、第1のメモリ参照と第2のメモリ参照との間の関係(たとえば、第1のメモリ参照のメモリ領域と第2のメモリ参照のメモリ領域が重複するかどうか、どの程度重複するか、など)を示すメタデータを生成する。

メモリ領域が同じではない場合、コンパイリングプロセッサは、コンパイラを使用してメタデータに基づき、ループアンローリングなどのより積極的なコンパイル技法を使用してカーネルを再コンパイルしてもよい。ドライバは、カーネルのメモリアクセスがどの程度重複するかを判定することができる場合もあり、メモリオーバーラップの量に基づいてより積極的な技法を使用してカーネルを再コンパイルしてもよい。このようにして、本開示の技法は、ジャストインタイムコンパイラを使用してコンパイルされたカーネルの実行性能を向上させる場合がある。

一例では、本開示では、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも一方によって、コンパイルされたカーネルのバイナリコードを実行するための引数を生成するステップと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、カーネル引数の第1のメモリ領域に対する第1のメモリ参照およびカーネル引数の第2のメモリ領域に対する第2のメモリ参照が同じメモリ領域を参照するものかどうかを判定するステップと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、判定に基づいて第1のメモリ参照および第2のメモリ参照に関連するメタデータを生成するステップとを含む方法について説明する。メタデータは、第1のメモリ領域と第2のメモリ領域との間の関係を示してもよい。この方法は、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、カーネルの第1および第2のメモリ参照が同じメモリ領域を参照するものではないと判定したことに応答して、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、コンパイラにメタデータに基づいてカーネルを再コンパイルさせるステップと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、ターゲットプロセッサに再コンパイルされたカーネルを実行するよう命令するステップとをさらに含む。

別の例では、本開示では、メモリとコンパイリングプロセッサとを含むデバイスであって、コンパイリングプロセッサが、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも一方によって、コンパイルされたカーネルのバイナリコードを実行するための引数を生成することと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、カーネル引数の第1のメモリ領域に対する第1のメモリ参照およびカーネル引数の第2のメモリ領域に対する第2のメモリ参照が同じメモリ領域を参照するものかどうかを判定することと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、判定に基づいて第1のメモリ参照および第2のメモリ参照に関連するメタデータを生成することとを行うように構成されるデバイスについて説明する。メタデータは、第1のメモリ領域と第2のメモリ領域との間の関係を示してもよく、コンパイリングプロセッサは、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、カーネルの第1および第2のメモリ参照が同じメモリ領域を参照するものではないと判定したことに応答して、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、コンパイラにメタデータに基づいてカーネルを再コンパイルさせることと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、ターゲットプロセッサに再コンパイルされたカーネルを実行するよう命令することとを行うようにさらに構成される。

別の例では、本開示では、命令を記憶する非一時的コンピュータ可読記憶媒体であって、命令が、実行されたときに、コンパイリングプロセッサに、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも一方によって、コンパイルされたカーネルのバイナリコードを実行するための引数を生成することと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、カーネル引数の第1のメモリ領域に対する第1のメモリ参照およびカーネル引数の第2のメモリ領域に対する第2のメモリ参照が同じメモリ領域を参照するものかどうかを判定することと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、判定に基づいて第1のメモリ参照および第2のメモリ参照に関連するメタデータを生成することとを行わせる非一時的コンピュータ可読記憶媒体について説明する。メタデータは、第1のメモリ領域と第2のメモリ領域との間の関係を示す。

コンパイリングプロセッサは、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、カーネルの第1および第2のメモリ参照が同じメモリ領域を参照するものではないと判定したことに応答して、コンパイリングプロセッサに、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、コンパイラにメタデータに基づいてカーネルを再コンパイルさせることと、コンパイリングプロセッサ上で実行されるコンパイラとランタイムとからなるグループのうちの少なくとも上記の一方によって、ターゲットプロセッサに再コンパイルされたカーネルを実行するよう命令することとを行わせる命令を実行するようにさらに構成されてもよい。

本開示の1つまたは複数の例の詳細を添付の図面および以下の説明に記載する。本開示の他の特徴、目的、および利点は、説明、図面、および特許請求の範囲から明らかになるであろう。

本開示の技法に従ってコンパイル最適化を補助するためにエイリアシング分析をサポートする例示的なコンピューティングデバイスを示すブロック図である。 本開示の技法に従ってカーネルを実行する場合があるプロセッサの1つまたは複数のシェーダコアの複数の処理要素を示す概念図である。 実行されたときに本開示の技法に従ってエイリアシングを生じさせることがあるコードを含むカーネルコードを示す概念図である。 本開示の技法に従って構成されたコンパイラが検出することができる場合があるエイリアシングの例を示す概念図である。 本開示の技法に従って構成されたコンパイラが検出することができる場合がある非重複メモリ参照の例を示す概念図である。 本開示の技法に従って構成されたドライバ/ランタイムが検出する場合がある重複メモリ参照を示す概念図である。 本開示の技法によるループアンローリングを示す概念図である。 本開示の技法によるコードリオーダリングを示す概念図である。 本開示の技法によるコードベクトル化を示す概念図である。 本開示の技法に従ってコンパイラ最適化を補助するためにコンパイラメタデータを生成するための例示的な方法を示す流れ図である。

上記において手短に説明したように、様々な異種コンピューティングフレームワークが現在開発中である。異種コンピューティングフレームワークのいくつかの例には、Khronosグループによって現在開発中のOpenCL(商標)フレームワークおよびMicrosoft(登録商標)によって現在開発中のDirectComputeフレームワークが含まれる。異種コンピューティングフレームワークは、CPU(中央演算処理装置)、GPU(グラフィックス処理装置)、FPGA(フィールドプログラマブルゲートアレイ)、DSP(デジタル信号プロセッサ)などの様々な異なる種類のプロセッサ上で単一のプログラムまたは「カーネル」を実行するのを可能にする。

カーネルを実行できるように準備するために、本開示ではコンパイリングプロセッサと呼ばれるプロセッサが、カーネルソースコードをコンパイルして、ターゲットプロセッサによって実行されるバイナリコードを生成する。ターゲットプロセッサは、コンパイリングプロセッサと同じであってもあるいは異なっていてもよい。コンパイリングプロセッサが使用するコンパイラの一例は、ジャストインタイムコンパイル(JIT)コンパイラと呼ばれる。JITコンパイラは、実行の前にソースコードをコンパイルすること(「事前」コンパイルとも呼ばれる)または事前に命令をまったくコンパイルしないこと(「インタープリテーション」と呼ばれる)ではなく、実行時(ランタイムとも呼ばれる)におけるソースコードのコンパイルを行う。

カーネルがコンパイルされた後、コンパイリングプロセッサは、ドライバおよびランタイムを介して、カーネルのコンパイルされたバイナリコードをターゲットプロセッサに転送する。カーネルはまた、ターゲットプロセッサ上でカーネルを実行するための引数のセットを実行時に受け入れ、コンパイリングプロセッサがまた、引数をターゲットプロセッサに転送する。カーネル引数は、バッファ、すなわち、引数に割り振られたメモリの領域を含む。たいていの場合、カーネルは、引数に対して動作する(すなわち、読取りまたは書込みを行う)コードセクションを含む。このように、引数は、カーネルの演算対象となる場合があるカーネル用のデータセットを含む。コンパイリングプロセッサのドライバ/ランタイムは、カーネルをターゲットプロセッサに転送した後、いくつかの例では実行時に引数をカーネルに供給する関数呼出しを実行する。カーネルが引数を受け取った後、ターゲットプロセッサはカーネルの実行を開始してもよい。

多くの例では、カーネルは、何らかのブール条件が満たされたとターゲットプロセッサが判定するかまたはある回数だけ繰り返されるまで実行されるループなどのコードセグメントを含む。コンパイラは、ループコードセクションを実行する性能を向上させるための様々な技法、たとえば、ループアンローリング、ならびにループコードセクションと非ループコードセクションとの両方の性能を向上させる場合があるコードリオーダリングおよび/またはベクトル化などの他の技法を使用することができる場合がある。

ループアンローリングは、コンパイラがループのいくつかの反復を拡張して演算、ループ終了試験などの、ループを制御する命令を減らすかもしくはなくし、および/またはループを実行する際のキャッシュ性能を向上させるための最適化プロセスである。コードリオーダリングは、コンパイラが一連の類似の命令をグループ化する(たとえば、まとめてロードまたは記憶する)ために使用する場合がある別の最適化である。コードリオーダリングは、場合によってはループコードセクションを実行する際にキャッシュ性能を向上させる場合がある。たとえば、コードリオーダリングは、いくつかのロード命令を(たとえば、ループ本体として)まとめて合体させることが、スカラー命令において使用されるオペランドのサイズの倍数であるキャッシュライン幅(以下により詳しく説明する)を有するシステム上の性能を向上させることがあるときに性能を向上させる場合がある。しかし、コンパイラがロードを安全に合体させることができるのは、コンパイルの前にロード/ストアバッファが互いにエイリアシングしていないとコンパイラが判定した場合に限られることがある。そうでない場合、並べ替えられたロード/ストア命令に起因してデータ破壊が生じる恐れがある。

ベクトル化は、コンパイラが、各々がオペランドの単一の対を一度に処理するいくつかのスカラー演算を含むソースコードを、オペランドの複数の対に対する1つの演算を一度に処理するベクトル命令に変換する場合がある別の最適化プロセスである。ベクトル化は、同じコードのスカラー実装形態に対して性能を向上させる場合がある並列化の形態である。ループアンローリング、コードリオーダリング、およびベクトル化について以下により詳しく説明する。

カーネルのコードセクションは、引数のメモリ領域を指す場合がある「ポインタ」とも呼ばれるメモリ参照を含んでもよい。たとえば、コードセクションは、カーネル引数の一部を参照する場合がある一連のメモリ参照(すなわち、カーネル引数に含まれるバッファのメモリ参照)を含んでもよい。カーネルは、引数バッファから値を読み出してもよく、引数バッファにデータを書き込んでもよい。

場合によっては、様々なメモリ参照、たとえば、異なる名前を有するポインタ変数が、メモリ内の同じデータ位置を参照する場合がある。様々な記号参照が同じメモリ領域を参照する状況を「エイリアシング」と呼ぶ。コンパイラは、静的分析またはその他の技法を使用してコンパイル時にエイリアシングを検出することを試みることがある。しかし、コンパイラは通常、ループコードセクション内のメモリ参照が参照するデータ(たとえば、カーネル引数)が実行時に供給されるときにループ内のメモリ参照のエイリアシングを検出することはできない。

コンパイラは、メモリ参照が同じメモリ領域を参照する(すなわち、メモリ参照の結果がエイリアシングになる)かどうかを明確に判定することができないときに、ループアンローリングおよびベクトル化などの最適化技法をループに対して実行できない場合がある。本開示の技法では、JITコンパイラが、カーネルループのメモリアクセスが同じメモリ領域を参照するかどうかを判定するのが可能になる場合がある。さらに、本開示の技法では、JITコンパイラが、メモリ参照間の関係に関するメタデータを生成し、生成されたメタデータに基づくベクトル化およびループアンローリングなどの最適化を使用してカーネルを再コンパイルするのが可能になる。

図1は、本開示の技法に従ってコンパイル最適化を補助するためにエイリアシング分析をサポートする例示的なコンピューティングデバイスを示すブロック図である。図1はコンピューティングデバイス2を含む。コンピューティングデバイス2は、パーソナルコンピュータ、デスクトップコンピュータ、ラップトップコンピュータ、コンピュータワークステーション、タブレットコンピューティングデバイス、ビデオゲームプラットフォームもしくはコンソール、ワイヤレス通信デバイス(たとえば、モバイル電話、セルラー電話、衛星電話、および/もしくはモバイル電話ハンドセット)、ポータブルビデオゲームデバイスまたは携帯情報端末(PDA)などのハンドヘルドデバイス、パーソナル音楽プレーヤ、ビデオプレーヤ、ディスプレイデバイス、テレビジョン、テレビジョンセットトップボックス、サーバ、中間ネットワークデバイス、メインフレームコンピュータ、またはグラフィカルデータを処理しおよび/もしくは表示する任意の他のタイプのデバイスを備えてもよい。

図1の例に示すように、コンピューティングデバイス2は、CPU16と、システムメモリ14と、グラフィックス処理ユニット(GPU)12と、ジャストインタイム(JIT)コンパイラ18と、ドライバ/ランタイム19とを含む。CPU16は様々なタイプのアプリケーションを実行してもよい。アプリケーションの例には、ウェブブラウザ、eメールアプリケーション、スプレッドシート、ビデオゲーム、表示用の可視オブジェクトを生成するアプリケーションなどが含まれる。1つまたは複数のアプリケーションを実行するための命令はシステムメモリ14内に記憶されてもよい。

CPU16はJITコンパイラ18を実行してもよい。したがって、CPU16は、一例として「コンパイリングプロセッサ」と呼ばれることがある。JITコンパイラ18は、CPU16によって実行されたときに、上述のように、OpenCLまたはDirectComputeなどの異種コンピューティングフレームワークを使用してカーネルのソースコードをコンパイルしてもよいコンパイラを備える。JITコンパイラ18は、ソースコードをターゲットプロセッサによって実行されるネイティブコードまたは中間コード(たとえば、バイトコード)としてコンパイルする。JITコンパイラ18は、「ランタイム」、すなわち、実行の前ではなく実行時にコンパイルを実行する。JITコンパイラ18は、OpenCLを使用したコンパイル時にclBuildProgram()関数を使用してコンパイルを実行してもよい。さらに、JITコンパイラ18は、カーネル20のデータアクセスパターンを分析して、ターゲットプロセッサGPU12上で実行される特定のファイバ(すなわち、スレッド)のデータアクセスが互いに依存していないかどうか、ならびに他の条件が成立するかどうかを判定するように構成されてもよい。

ドライバ/ランタイム19もJITコンパイラ18と相互作用して、カーネルソースコードをバイナリ命令またはバイトコード命令に変換する。ドライバ/ランタイム19は、ドライバを使用して、カーネルソースコード命令の、ターゲットプロセッサ(この例ではGPU12)用のネイティブコードまたはオブジェクトコードへのアーキテクチャ固有のコンパイルを実行してもよい。たとえば、ドライバ/ランタイム19は、ターゲットプロセッサに利用可能な特定のベクトル命令または実行リソースを認識してもよく、ターゲットプロセッサ上での実行性能を最適化するようにソースコードをネイティブコードにコンパイルしてもよい。いくつかの例では、たとえば、複数のターゲットプロセッサがある場合、たとえば、カーネルがCPU16およびGPU12上で実行される場合にいくつかの異なるドライバが存在してもよい。

カーネル20は、ターゲットプロセッサ、この例ではGPU12が実行することのできるネイティブコードまたはオブジェクトコード、たとえば、バイナリ命令で構成される。JITコンパイラ18はまた、GPU12の実行時実行を管理してもよい。CPU16は、カーネル20をGPU12が実行できるようにGPU12に送信してもよい。CPU16は、引数26を生成してもよく、引数26をさらに処理できるようにGPU12に転送してもよい。

CPU16は、引数26を割り振る前に、メモリの一領域である空きメモリバッファを引数26用に割り振る。バッファが割り振られた後、ドライバ/ランタイム19は引数26をバッファに格納する。引数26は、GPU12が処理することのできる複数のデータ値(たとえば、整数、浮動小数点値、オブジェクト、値の配列など)を含んでもよい。さらに、カーネル20の実行時に、GPU12は、引数26を格納するバッファにデータを出力として書き込んでもよい。出力されたデータは、出力引数を含んでもよく、GPU12はこれらの出力引数をCPU16に送り返してもよい。

CPU16がGPU12に転送する引数は、「入力引数」と呼ばれることもある。異種コンピューティングフレームワークがOpenCLフレームワークである例では、ドライバ/ランタイム19は、引数を生成して実行時にclSetKernelArg()関数に渡して(利用可能にして)もよい。clSetKernelArg()関数は、カーネル20を引数ならびにカーネル引数26のいずれかとして受け取り、実行を開始できるように引数をGPU12に転送する。

ドライバ/ランタイム19は、メモリを引数26用に割り振ることの一部として、カーネルに含まれるメモリ参照のいくつかまたはすべてに関連する引数26のアドレスおよびメモリ領域を判定する。メモリ参照は、特定のコードセクション、たとえば、「ループコードセクション」と呼ばれるループを含むコードセクションのメモリ参照であってもよい。ドライバ/ランタイム19は、判定されたメモリ領域に基づいて、カーネル20のループコードセクションまたは他のコードセクションのメモリ参照が引数26の同じメモリ領域を参照するものであるかどうかを解決する(すなわち、判定する)ことが可能であってもよい。

ドライバ/ランタイム19は、GPU12がカーネル20を実行するようにカーネル引数26を生成したことに応答して、カーネル20を実行してもよい。より詳細には、ドライバ/ランタイム19は、clEnqueueNDRangeKernel()関数を使用してターゲットプロセッサのGPU12にカーネル20をディスパッチしてもよい。実行時に、ドライバ/ランタイム19は引数26を分析し、カーネル20が引数26を受け取る。ドライバ/ランタイム19はまた、メモリ参照(たとえば、ポインタ)などを分析して、メモリ参照が引数26に割り振られたメモリ領域と同じメモリ領域を参照しているかどうかを判定する。ドライバ/ランタイム19は、メモリ参照と引数バッファを対として分析し、メモリ参照が同じメモリ領域を参照しているかどうかを判定してもよい。

ドライバ/ランタイム19は、メモリ参照が参照する引数26のメモリ領域間の関係に基づいてメモリ参照に関連するメタデータをさらに生成する。メタデータは、メモリ参照間の関係を示してもよい。たとえば、メタデータは、いくつかの非限定的な例として、重複メモリ参照のリスト、メモリ参照に関連するメモリ領域が重複しているかどうか、メモリ領域がどの程度重複しているか、重複箇所に含まれるバイトの数はいくつかを含んでもよい。

ドライバ/ランタイム19は、(もしあれば)生成されたメタデータをJITコンパイラ18に送り返す。ドライバ/ランタイム19は、メタデータに基づいて、2つのメモリ参照がまったく同じメモリ領域を共有するわけではないと判定したことに応答して、JITコンパイラ18に、ループアンローリング、コードリオーダリング、および/またはベクトル化などの様々な最適化を使用してカーネル20を再コンパイルさせてもよい。JITコンパイラ18は、生成されたメタデータに基づいてループアンローリング、コードリオーダリング、および/またはベクトル化のこれらの様々な最適化を適用してもよい。

コンパイリングプロセッサ、たとえばCPU16は、本開示の技法に従って、コンパイリングプロセッサ上で実行されるJITコンパイラ18とドライバ/ランタイム19とからなるグループのうちの少なくとも一方を使用して、コンパイルされたカーネル20のコード(たとえば、バイナリコードまたはオブジェクトコード)を実行するための引数26を生成するように構成されてもよい。JITコンパイラ18とドライバ/ランタイム19とからなるグループのうちの少なくとも上記の一方は、カーネル引数の第1のメモリ領域に対する第1のメモリ参照とカーネル引数の第2のメモリ領域に対する第2のメモリ参照が同じメモリ領域を参照しているかどうかを判定するようにさらに構成されてもよい。CPU16は、JITコンパイラ18とドライバ/ランタイム19とからなるグループのうちの少なくとも上記の一方によって、カーネルの第1および第2のメモリ参照が同じメモリ領域を参照するものではないと判定したことに応答して、CPU16上で実行されるJITコンパイラ18とドライバ/ランタイム19とからなるグループのうちの少なくとも上記の一方に、メタデータに基づいてカーネル20を再コンパイルさせることと、CPU16上で実行されるJITコンパイラ18とドライバ/ランタイム19とからなるグループのうちの少なくとも上記の一方によって、ターゲットプロセッサ、たとえば、GPU12に再コンパイルされたカーネル20を実行するよう命令することとを行うようにさらに構成されてもよい。

GPU12は、グラフィックスデータを処理するのに好適な大規模な並列処理を可能にする専用ハードウェアであってもよい。このようにして、CPU16は、GPU12によってよりうまく対処されるグラフィックス処理をオフロードする。CPU16は、特定のアプリケーション処理インターフェース(API)または異種コンピューティングフレームワークに従ってGPU12と通信してもよい。そのようなAPIの例には、Microsoft(登録商標)によるDirectX(登録商標) APIおよびKhronosグループによるOpenGL(登録商標)が含まれ、異種コンピューティングフレームワークの例には、MicrosoftによるDirectCompute、KhronosグループによるOpenCL(商標)が含まれる。しかし、本開示の態様は、上述のAPIおよびフレームワークに限定されず、他のタイプのAPIに拡張されてもよい。

CPU16およびGPU12の例には、限定はしないが、デジタル信号プロセッサ(DSP)、汎用マイクロプロセッサ、特定用途向け集積回路(ASIC)、フィールドプログラマブルゲートアレイ(FPGA)、または他の等価の集積論理回路もしくはディスクリート論理回路が含まれる。いくつかの例では、GPU12は、グラフィックス処理に適した大規模な並列処理機能をGPU12に付与する集積論理回路および/またはディスクリート論理回路を含む専用ハードウェアであってもよい。場合によっては、GPU12は、汎用処理を含んでもよく、汎用GPU(GPGPU)と呼ばれることもある。本開示において説明する技法は、GPU12がGPGPUである例に適用可能である。

システムメモリ14は、1つまたは複数のコンピュータ可読記憶媒体を備えてもよい。システムメモリ14の例には、限定はしないが、ランダムアクセスメモリ(RAM)、読取り専用メモリ(ROM)、電気的消去可能プログラマブル読取り専用メモリ(EEPROM)、フラッシュメモリ、あるいは所望のプログラムコードを命令および/またはデータ構造の形で伝送または格納するのに使用することができ、かつコンピュータまたはプロセッサによってアクセスすることができる任意の他の媒体が含まれる。

いくつかの態様では、システムメモリ14は、CPU16および/またはGPU12に本開示ではCPU16およびGPU12に帰属する機能を実行させる命令を含んでもよい。したがって、システムメモリ14は、1つまたは複数のプロセッサ、たとえば、CPU16およびGPU12に様々な機能を実行させる命令を含むコンピュータ可読記憶媒体であってもよい。

システムメモリ14は、いくつかの例では、非一時的記憶媒体とみなされることがある。「非一時的」という用語は、記憶媒体が搬送波では具体化されないこと、またはいくつかの例では、非一時的コンピュータ可読記憶媒体が経時的に変化する可能性があるデータを(たとえば、RAMに)格納する場合があることを示すことがある。

CPU16は、JITコンパイラ18およびドライバ/ランタイム19を使用して、ソースコードをGPGPUアプリケーション用のネイティブコード(たとえば、コマンドおよびデータ)またはバイトコードとしてコンパイルしてもよい。例示的なGPGPUデータおよびコマンドには、光線追跡アプリケーション、物理シミュレーション用のコマンドおよびシーンデータ、または任意の他のタイプのGPGPUカーネル用のデータが含まれる。GPGPUアプリケーション、たとえば、カーネル20は、DirectXなどのグラフィックスAPIまたはOpenGLを使用するか、あるいはOpen Compute Language(OpenCL)、またはOpenCompute、またはDirectComputeなどのより汎用的なコンピュートAPIを使用してコンパイルされてもよい。CPU16は、カーネル20用のデータを処理のためにコマンドバッファに送ってもよい。様々な例では、コマンドバッファは、システムメモリ14の一部であってもまたはGPU12の一部であってもよい。いくつかの例では、CPU16は、カーネル20のコマンドおよびデータをGPU12がPCI-Expressバスなどの専用バスまたは別の汎用シリアルバスもしくはパラレルバスを介して処理できるように伝送してもよい。

GPU12は、コマンドバッファに記憶されたカーネル20の演算を実行するために処理パイプラインを実装してもよい。処理パイプラインは、GPU12上で実行されるソフトウェアまたはファームウェアによって定義されるような関数を実行することと、非常に特殊な関数を実行するように配線接続された固定関数ユニットによって関数を実行することとを含む。カーネル20を実行できるように固定関数ユニットをバイパスすることが可能であってもよく、またはカーネル20を実行するために固定関数ユニットを使用してもよい。

カーネル20は、GPU12の1つまたは複数の処理要素(「シェーダコア」または“PE”とも呼ばれる)上で実行されてもよい。ユーザがシェーダを任意の他のプロセッサと同様に所望のタスクを任意の考えられる方法で実行するようにプログラムすることができるので、シェーダコア22は、ユーザによる関数の融通性に富んだ使用を可能にする。しかし、固定関数ユニットは、固定関数ユニットがタスクを実行する方法に関して配線接続される。したがって、固定関数ユニットでは、関数の融通性に富んだ使用が可能にならない場合がある。本開示の技法は、GPUシェーダコア22上での、カーネル20などのカーネルの実行を参照する。

CPU16が、グラフィカルシーンをレンダリングすることまたはコマンドバッファに対してカーネルを実行することに関連するデータおよび/またはコマンドを送った後、GPU12は、GPU12のパイプラインを通してコマンドの実行を開始する。GPU12のスケジューラ24はスレッドを作成し、スレッドがカーネルに関連する作業の基本ユニットを実行する。スケジューラ24は、スレッドをシェーダコア22の特定の処理要素に割り当てる。

図2は、本開示の技法に従ってカーネルを実行する場合があるプロセッサの1つまたは複数のシェーダコアの複数の処理要素を示す概念図である。図2は、GPU12またはCPU16の一部を示す。GPU12は複数の処理要素42A〜42N(PE42)を含み、これらの処理要素がカーネル、たとえばカーネル20の一部を実行してもよい。いくつかの例では、PE42上で実行される場合があるカーネル20の部分は「ワープ」または「作業ユニット」と呼ばれることがある。PE42は、シェーダコア22(図1)のうちの1つまたは複数の一部であってもよい。ワープまたは作業ユニットは、「ファイバ」とも呼ばれるスレッドのグループを含んでもよく、GPUスケジューラ24は各スレッドを複数の処理要素、たとえばPE42に実行できるように割り当ててもよい。図2の各PEは、特定の時間に(たとえば、並列実行が可能なように同時に)複数のデータ値に対して、ベクトル命令などの単一の命令を実行することができる単一命令多重データ(SIMD)ユニットを備えてもよい。PE42は、単一の浮動小数点値に対する単一の演算などの、単一のデータ値に対する単一の命令の実行をサポートしてもよい。

図2は、GPU12のスケジューラがPE42に実行できるように割り当てる命令44も含む。いくつかの例では、命令44は、コマンドバッファに格納されてもよい。命令44には、各PEが実行するように構成されたカーネルの命令のセットを含んでもよい。プログラムカウンタ(PC)50は、PE42のうちの1つまたは複数が実行する現在の命令を示す。PE42上での命令の実行が終了した後、PC50の値がカーネル20の次の命令のアドレスに増分されてもよい。図2はレジスタ46も含む。レジスタ46A〜46N(レジスタ46)は、複数のデータ値または単一の値を保持することのできる汎用レジスタであってもよい。レジスタ46は、「バンク化」されてもよく、すなわち、特定のPE用のデータをロードし格納してもよい。一例として、レジスタ46Aは、PE42A用のデータの格納に限定されてもよく、他のPE用のデータのロードまたは格納を行わなくてもよい。レジスタ46の各々は、PE42のうちの1つにデータを供給しおよび/またはPE42のうちの1つからデータを供給されてもよく、次いで、PE42がそのデータを処理してもよい。

PE42、命令44、レジスタ46、キャッシュ48、およびPC50は、GPU12のシェーダコア22のコアまたは一部を含んでもよい。様々な例では、ワープ40は、ジオメトリシェーダ、ピクセルシェーダ、および/または頂点シェーダなどのシェーダの一部を含んでもよく、これらの部分は、GPU12のグラフィックスパイプラインの一部であってもよく、またはカーネル20などのカーネルの一部を含んでもよい。いくつかの例では、GPU12は、ワープによって生成された結果をさらなる処理ができるようにパイプラインの別の段階に供給してもよい。

図2はまた、キャッシュ48も含む。キャッシュ48は、頻繁にアクセスされる命令およびデータを実行時に高速の取出しおよび格納ができるように格納する小型メモリである。キャッシュ48は、単一のキャッシュとして示されているが、複数のキャッシュレベルおよび/または別個のキャッシュを表してもよい。上述のように、カーネル20の実行時には、GPU12は、PC50の値によって示されるアドレスに位置する命令44のうちの1つを取り出す。GPU12は次いで、PE42にPC50のアドレスに格納された命令を実行させる。PC50は、いくつかの例ではレジスタであってもよい。

GPU12は、不必要に低速である、PC50のアドレスにおける命令のシステムメモリからのフェッチを行うのではなく、キャッシュ48をチェックして、キャッシュ48が現在、次に実行すべき命令を含んでいるかどうかを判定する。命令を格納するキャッシュ48の部分を命令キャッシュ(「Iキャッシュ」)と呼ぶ。次に実行すべき命令がキャッシュ48に格納されている場合は、「キャッシュヒット」と呼ばれ、GPU12はキャッシュされている命令をロードし実行する。次に実行すべき命令がキャッシュ48に格納されていない場合は「キャッシュミス」と呼ばれ、GPU12は、あるより低速のメモリ、たとえば、システムメモリ14から次に実行される命令をロードする。

GPU12は、メモリアドレスに格納されたデータ値(たとえば、オペランド)を必要とする命令(たとえば、加算、乗算、ロード、ストア)の実行時に、まず、オペランドがレジスタ、たとえばレジスタ46のうちの1つに格納されているかどうかを判定する。要求されたデータ値がレジスタ46に格納されていない場合、GPU12は、データキャッシュ(「dキャッシュ」)と呼ばれるデータ値を保持するキャッシュ48の部分に対してデータ値へのアクセスを試みる。データ値がキャッシュ48内に格納されている場合、GPU12は、要求されたデータ値をキャッシュ48からロードする。それ以外の場合、GPU12は、要求されたデータ値をより低速のメモリ、たとえば、システムメモリ14からロードしなければならない。同様に、命令がPE42にデータ値をメモリに格納し直させるかまたはデータ値を修正させる場合、キャッシュ48は、その値をキャッシュ48自体に格納してもよく、それによって、そのデータ値は、再び書き込まれるかまたは読み出される場合、どのレジスタ46にも格納されていない場合に、キャッシュ48から高速に取り出されるかまたはキャッシュ48に上書きされる。

GPU12は、データをキャッシュ「ライン」と呼ばれる固定サイズブロックでキャッシュ48との間で転送する。キャッシュ48は、数百または数千個の異なるラインを格納する容量を有してもよい。各ラインは、特定のメモリアドレスに関連付けられ、複数のデータバイトを格納してもよい。たとえば、キャッシュ48の各ラインは、一例として64バイトのデータを格納してもよい。各ラインに格納されるバイトの数をキャッシュ「幅」と呼ぶ。キャッシュ48が64バイトのデータを記憶することができるラインを有する例では、キャッシュ48のキャッシュ幅は64バイトである。キャッシュ幅は、以下により詳しく説明するように、コードリオーダリング最適化技法の性能に影響を与えることがある。

キャッシュ48からデータを取り出すロード演算の間、GPU12は、取り出されたキャッシュデータをレジスタ46のうちの1つまたは複数あるいは図示されていない他のレジスタにロードしてもよい。命令の実行時に、PE42は、レジスタ46から1つまたは複数のデータ値を読み出してもよい。PE42は、データ値に対して1つまたは複数の演算を実行し、新しい値をレジスタ46に格納し直してもよい。PE42は、分岐、ジャンプ、gotoなどのフロー制御命令を実行してもよい。しかし、1つのPC50しかないので、PE42は、ある特定の所与の時間にはPC50によって示される命令44のうちの1つしか実行できない。

GPU12などのプロセッサは、大量のベクトルレジスタおよびベクトル命令を有してもよい。したがって、JITコンパイラ18などのコンパイラは、ベクトル化などの最適化を使用してアプリケーションをコンパイルすることができ、ベクトル命令をサポートするかまたはGPU12などのSIMDアーキテクチャを有するプロセッサのスループットまたは実行性能を向上させる場合がある。

より詳細には、GPU12は、図2に示すシェーダコアと同様の数百個または数千個のシェーダコアを含んでもよい。各シェーダコアは、ベクトル命令を実行することが可能であってもよい。複数のオペランドを有するベクトル命令を実行すると、ベクトル命令ではなくスカラー命令を含む最適化されていないコードに対して性能が大幅に改善される場合がある。さらに、実行性能の向上は、ベクトル命令を実行することができる多数のSIMDコアを有するアーキテクチャ上ではより程度が高くなることがある。その理由としては、汎用性のより高いプロセッサでは、ベクトル命令を実行することができるレジスタおよび/またはコアの数が限られている場合があることが挙げられる。

図3Aは、実行されたときに本開示の技法に従ってエイリアシングを生じさせることがあるコードを含むカーネルコードを示す概念図である。図3Aの例はカーネルコード80を含む。カーネルコード80は、ライン82、84、86、および88を含む。

カーネルコード80のライン82はcompute_output関数である。ライン82のcompute_output関数は、カーネルが実行を開始するときにターゲットプロセッサ(たとえば、GPU12)が呼び出す関数である。この関数は、compute_output関数が、ドライバ/ランタイム19がカーネル20の実行を開始するために使用するプログラムエントリポイントであるという点で、Cプログラミング言語における"int main()"関数とほぼ同等である。ターゲットプロセッサがCPU16である場合、Cランタイムライブラリは、ドライバ/ランタイム19の実行時ランタイム構成要素を含む場合がある。GPU12がターゲットプロセッサである場合、ドライバ/ランタイム19のドライバ構成要素はランタイムを含んでもよい。compute_output関数は4つの入力引数、(1)inputImage、(2)global_cdf、(3)outputImage、および(4)local_cdfを含む。inputImageは入力引数のバッファへのポインタである。outputImageは、カーネルが実行を終了したときに、出力引数を含むバッファへのポインタである。引数global_cdfおよびlocal_cdfは、値の配列へのポインタである。ライン84は、実行されたときに、GPU12に変数を割り振らせ初期設定させる複数のステートメントを表す場合がある。一例として、ライン84を実行すると、PE42はinputImage[i]の値を初期設定しロードするなどである。

ライン86は、ループ初期化ステートメントである。ループ初期化ステートメントは、ループが固定反復回数にわたって反復することを示す。ループは、変数"start_offset"に等しい開始インデックスiにおいて反復を開始し、各反復が実行を終了したときにiを1だけ増分させる。各ループ反復が完了したときに、GPU12は、ブール条件"i<final_offset"が依然として真であるかどうかをチェックする。GPU12は、iの値が値"final_offset"以上になったときにループの実行を停止する。

各ループ反復内で、GPU12はoutputImageの値を、local_cdf[inputImage[i]]の値に等しいoutputImage[i]として示されるインデックスiに設定する。local_cdfは、この例では、inputImage[i]の値によってインデックス付けされる配列である。inputImage[i]は、変数iによってインデックス付けされ、GPU12は、変数iをループ反復ごとに増分させる。

上述のように、outputImageとinputImageはどちらもメモリ参照である。outputImageおよびinputImageへのポインタがメモリ内の同じ領域を参照する場合がある(すなわち、outputImageとinputImageとの間にエイリアシングまたは部分的なエイリアシングが生じる)。outputImageとinputImageがメモリ内の互いに異なる領域または重複する領域を参照する場合もある(すなわち、outputImageとinputImageとの間にエイリアシングが生じない)。JITコンパイラ18は、inputImageとoutputImageとの間にエイリアシングが生じない(すなわち、まったく同じメモリ領域を参照するわけではない)かどうかを判定することができない場合、ベクトル化、コードリオーダリング、および/またはループアンローリングなどの特定のコンパイラ最適化を使用することができない場合がある。

図3Bは、本開示の技法に従って構成されたドライバ/ランタイムが検出することができる場合があるエイリアシングの例を示す概念図である。しかし、コンパイラはエイリアシングに関して最適化することができない。図3Bの例は、GPU12がメモリに格納する場合があるバッファ100を示す。一例として、図3Bのポインタ、outputImageおよびinputImageはバッファ100の一部を参照する場合がある。図3Bの例では、バッファ100はメモリアドレス0x800(16進)から開始する。

この例では、inputImageとoutputImageはどちらも、バッファ100内に格納された単一のエントリ(たとえば、単一のオブジェクト、変数など)を参照する。すなわち、この例では、inputImageとoutputImageとの間にまったく同じメモリ領域に関するエイリアシングが生じ、これはクロスハッチングによって示されている。ドライバ/ランタイム19は、inputImageとoutputImageが同じメモリ領域を参照することを検出することができてもよい。inputImageとoutputImageが同じメモリ領域を参照するので、JITコンパイラ18は、ループアンローリングおよび/またはベクトル化などの最適化を実行することができない。

ドライバ/ランタイム19は、図3Bに示すように2つのメモリ参照が同じメモリ領域を参照していることを検出したことに応答して、メタデータを生成しない場合がある。さらに、JITコンパイラ18はカーネル20を再コンパイルしない場合がある。すなわち、JITコンパイラは、図3Cおよび図3Dに示す場合に関してカーネル20を再コンパイルすることがある。したがって、JITコンパイラ18は、図3Cおよび図3Dにも示すようにいずれのコード最適化も実行しない場合がある。

図3Cは、本開示の技法に従って構成されたドライバ/ランタイムが検出することができる場合がある非重複メモリ参照の例を示す概念図である。図3Cは、図3Bに示すバッファと同じバッファであるバッファ120を示す。バッファ120は同様に、図3Bのバッファ100と同じメモリアドレスの0x800から開始する。

図3Cでは、inputImageおよびoutputImageは、バッファ120の2つの異なるメモリ領域を参照するメモリ参照である。inputImageが参照するメモリ領域は、水平ハッチングによって示されている。outputImageが参照するメモリ領域は、垂直ハッチングによって示されている。JITコンパイラ18は、カーネルコード80、より詳細には、ライン86および88を実行する前に、iの値に関わらず、inputImage[i]とoutputImage[i]がループの同じ反復時に同じメモリ領域を参照することはないと判定する場合がある。

実行時に、ドライバ/ランタイム19は、inputImage[i]およびoutputImage[i]の初期値に基づき、かつinputImage[i]およびoutputImage[i]のメモリアドレスがループ86における反復過程にわたって収束しないことに基づいて、inputImage[i]とoutputImage[i]が同じメモリ領域を参照することはないと判定することができてもよい。言い換えれば、inputImageおよびoutputImageの参照されるインデックスは、常に同じインデックス値iによって参照され、GPU12はインデックス値iを単調増加させる。

ドライバは、メモリ参照inputImageとoutputImageが同じメモリ領域を参照しないと判定したことに応答して、inputImageとoutputImageとの間の関係を示すメタデータを生成してもよい。このメタデータは、inputImageおよびoutputImageに関連するメモリ領域が重複せず、一例として2つのエントリによって分離されることを示す場合がある。メタデータは、inputImageおよびoutputImageに関連する領域のサイズならびにinputImageとoutputImageとの間のバイト数を示すこともある。JITコンパイラ18は、メタデータを生成した後、ドライバ/ランタイム19からメタデータを受け取り、以下により詳しく説明するように様々な最適化を適用することによってメタデータに基づいてカーネル20を再コンパイルしてもよい。

図3Dは、本開示の技法に従って構成されたドライバ/ランタイムが検出する場合がある重複メモリ参照を示す概念図である。図3Dは、引数バッファ、たとえば、引数26(図1)であってもよいバッファ130を含む。バッファ130は、この例ではアドレス0x800から開始する。バッファ130は、バッファ130の周囲の矩形内の別個の矩形として示されている複数のデータ値を含む。

前述の例と同様に、inputImageおよびoutputImageは、バッファ130の領域を参照するメモリ参照である。この例では、inputImageおよびoutputImageが参照する領域が重複しているが、全体的に重複しているわけではない。inputImageのみに関連するメモリ領域は、水平ハッチングによって示された矩形によって示されている。outputImageのみに関連するメモリ領域は、垂直ハッチングによって示された矩形によって示されている。inputImageとoutputImageとの両方によって参照される重複メモリ領域は、クロスハッチングによって示された矩形によって示されている。

実行時に、ドライバは、inputImageメモリ参照およびoutputImageメモリ参照が同じメモリ領域を参照しているかどうかを判定する。この例では、inputImageとoutputImageが重複しているが、同じメモリ領域を参照してはいない。ドライバ/ランタイム19は、inputImageとoutputImageが重複しているが同一ではないことを検出し、JITコンパイラ18用のメタデータを生成する。このメタデータは、各領域の開始アドレスおよび終了アドレスなどの、inputImageおよびoutputImageに関連する領域に関する情報を示す場合がある。このメタデータは、重複領域のサイズならびに重複領域の開始アドレスおよび/または終了アドレスなどの、重複領域に関する情報をさらに含んでもよい。JITコンパイラ18は、ドライバ/ランタイム19によって生成されたメタデータを受け取り、本開示による最適化技法を適用することによってカーネル20を再コンパイルしてもよい。

図4Aは、本開示の技法によるループアンローリングを示す概念図である。図4Aは、概して図3Aに示すカーネルコード80に対応するコードセクション140を含む。図4Aの例では、ドライバ/ランタイム19および/またはJITコンパイラ18は、図3Cおよび図3Dに示すように、メモリ参照inputImageおよびoutputImageは同じメモリ領域を参照しないと判定していてもよい。inputImageとoutputImageが同じメモリ領域を参照しないので、JITコンパイラ18はカーネルコード80に対してループアンローリングを実行している。ライン142〜150は、1つの反復を4つの反復としてアンローリングした結果を示す。

図3Aのライン86および88は、単一の反復を実行し、各反復の後で変数iを1だけ増分させることを示すが、ライン142のアンローリングされたループは各反復後にiを4だけ増分させる。ライン144は、local_cdf[inputImage[i]]の値をoutputImage[i]に割り当てる。ライン146は、local_cdf[inputImage[i+1]]の値をoutputImage[i+1]に割り当てる。ライン148は、local_cdf[inputImage[i+2]]の値をoutputImage[i+2]に割り当て、ライン150は、local_cdf[inputImage[i+3]]の値をoutputImage[i+3]に割り当てる。ライン144〜150の結果は、local_cdf[inputImage[i+x]]の出力をoutputImage[i+x]の対応する値に割り当てた結果であり、この場合、x[0...3]である。したがって、ライン142〜150に示すアンローリングされたループコードセクションは、図3Aのライン86〜88の4つの反復と同じ効果を有する。

コードセクション140のループアンローリングは、図3Aのループコードセクション80と比較していくつかの利点を有する場合がある。第1の利点は、割当ての各々を特定の順序に並べることによって、JITコンパイラ18および/またはドライバ/ランタイム19は、特定の順序に並べられていないコードセクションと比較してターゲットプロセッサ、たとえばGPU12上でより高いキャッシュ性能を実現することができる場合があることである。

たとえば、GPU12は、ライン144を実行した後、inputImageおよびoutputImageに関連するメモリ領域のデータのいくつかまたはすべてをキャッシュ、たとえばキャッシュ48に格納していてもよい。命令を実行するのに必要なデータがレジスタ、たとえばレジスタ46に格納されていない場合、このデータについてキャッシュ、たとえばキャッシュ48にアクセスすることが必要になる場合がある。より詳細には、GPU12は、inputImageおよびoutputImageのエントリ、たとえば、inputImage[i+1]、[i+2]など、ならびにoutputImage[i+1]、[i+2]などをキャッシュ48に格納してもよい。inputImageおよびoutputImageのエントリがGPU12のキャッシュに格納されている場合、GPU12は、ライン144〜150のinputImageおよびoutputImageの参照されたインデックスのデータについて、より低速のメモリにアクセスするのではなく、キャッシュに高速にアクセスすることができる場合がある。

さらに、コードセクション140がアンローリングされるとき、inputImage[i,i+1,i+2...]およびoutputImage[i,i+1,など]の値が単一のキャッシュラインに格納される場合がある。対照的に、アンローリングされないとき、inputImageおよびoutputImage[i]の値は異なるキャッシュラインに格納される場合がある。ループアンローリングの結果として行われることがある、inputImageのすべての値の、1回のキャッシュ読取りにおける単一のキャッシュラインからの取出しは、アンローリングされたコードを実行する際に行われることがある複数のキャッシュ読取りの実行よりも高速である場合がある。

データについて、より低速のシステムメモリ、たとえばシステムメモリ14ではなくGPU12のキャッシュにアクセスすることによって、ライン86〜88と比較してライン142〜150のループの方が実行性能が向上する場合がある。いくつかの例では、GPU12は、ライン144〜150間に依存性がなく、inputImageまたはoutputImageの値がカーネル20において前に算出された値に依存すると仮定して、たとえば、スーパースカラー実行をサポートするプロセッサ、またはSIMDプロセッサ上でライン144〜150を並列に実行することができる場合もある。

図4Aのコードセクション140に示すようなループアンローリングは、キャッシュ性能を向上させることに加えて、GPU12がループに関連するブール条件を評価する回数、ならびにGPU12が各ループ反復を終了した後に実行するジャンプの回数も低減させる。図3Aのコードセクション80と比較して、ライン142〜150のコードは、ブール条件、すなわちライン142の"i<final_offset"が真であるかどうかを評価する前に反復当たり4つのラインを実行する。対照的に、コードセクション80は、ライン82のブール条件が真であるかどうかを評価する前に1ラインのみを実行する。したがって、GPU12がライン142のブール条件を評価する回数は、コードセクション80と比較して低減される。

GPU12は、ライン142〜150のループの反復を完了した後、ブール条件"i<final_offset"が依然として真であると判定した場合、ライン150からジャンプしてライン144に戻る。コードセクション140において、GPU12は、4つのラインを実行した後ジャンプする。GPU12は、コードセクション80を実行する際、各反復後にジャンプする。したがって、コードセクション80と比較して、コードセクション140のアンローリングされたコードは、ブール条件の評価とGPU12が実行するジャンプの回数との両方を低減させ、コードセクション140の実行性能を向上させる場合がある。

図4Bは、本開示の技法によるコードリオーダリングを示す概念図である。図4Bはコードセクション160を含み、コードセクション160は、ライン162、164、166、168、170、および172をさらに含む。上述のように、ドライバ/ランタイム19および/またはJITコンパイラ18は、同じメモリ領域へのメモリ参照のエイリアシングが生じているかどうかを判定してもよい。図4Aに関して上記において説明したように、JITコンパイラ18は、特定のコードセクションにおいてメモリエイリアシングが生じていないと判定したドライバ/ランタイム19からメタデータを受け取ったことに応答して、図4Aに示すループアンローリングなどの特定の最適化を実行してもよい。

特定のコードセクションにおけるメモリ参照が同じメモリ領域を参照していないと判定したことに応答してJITコンパイラ18および/またはドライバ/ランタイム19が実行する場合がある別の最適化は、図4Bに示すコードリオーダリングである。コード160は概して、図4Bのアンローリングされたコードの並べ替えられたアセンブリ言語表現に相当してもよい。JITコンパイラ18および/またはドライバ/ランタイム19は、非ループコードセクションならびにループコードセクションにコードリオーダリングを適用してもよい。図4Bにおいて、JITコンパイラ18は、ロードおよびストアのすべてがグループ化されるように図4Aのロードおよびストアを並べ替えている。

ライン162および164はロード命令であり、JITコンパイラ18および/またはドライバ/ランタイム19がこれらのロード命令をグループ化している。図4Aにおいて、ライン144などのラインは、複数のロード命令およびストア命令を含む。たとえば、JITコンパイラ18は、ライン144を実行するために、3つの別個の命令を生成する場合がある。第1の命令はロード命令であってもよく、この命令は、inputImage[i]が参照するメモリ位置の値をr0として示されたレジスタにロードする。第2の命令はロード命令であってもよく、この命令は、local_cdf[inputImage[i]]の値をロードし、ロードされた値を同じレジスタr0に格納し、それによってr0の前の値を上書きする。ライン144に含まれる最後の命令はストア命令であってもよく、この命令は、r0の値をoutputImage[i]が参照するメモリに格納する。

ライン162〜172は、ライン144〜150を含む命令に対して並べ替えられたロード命令およびストア命令を示す。ライン162において、アセンブリコードは、GPU12に、(ロード命令"ldg"を使用して)inputImage[i]が参照するメモリ領域の値をレジスタr0にロードさせる。同様に、ライン164は、GPU12に、メモリ参照inputImage[i+1]が参照する値をレジスタr1にロードさせる。ライン162および164とライン166との間に生じる場合があり、簡潔のために図示されていない以後の命令には、GPU12にinputImageが参照するメモリ領域のデータをレジスタにロードさせる追加のロード命令を含めてもよい。

ライン166、168、および簡潔のために図示されていない他のラインでは、JITコンパイラ18はバッファlocal_cdfからのロードをグループ化している。ライン166は、local_cdf[r0]のコンテンツ、すなわち、インデックスr0における配列local_cdfからメモリのコンテンツをロードし、local_cdf[r0]のコンテンツをレジスタr0に格納し、それによってr0のコンテンツを上書きするロード命令を含む。同様に、ライン168の命令は、GPU12に、メモリ参照local_cdfが、現在レジスタr1に格納されている値によって示されるインデックスにおける参照するコンテンツをレジスタr1に格納させる。したがって、実行時に、168の命令はGPU12にr1の前の値を上書きさせる。ライン168とライン170との間に生じ、簡潔のために図示されていない他の命令には、同様に、実行されたときに、GPU12にlocal_cdf[rx]からデータをロードさせる命令を含めてもよく、ここでxはある整数である。

JITコンパイラ18は、コードセクション140の命令のリオーダリングの一部として、ストア命令もグループ化する。この一例として、リオーダリングの後、JITコンパイラ18はライン170および172をグループ化している。ライン170は、r0のコンテンツをメモリの位置outputImage[i]に格納するストア命令を含む。同様に、ライン172は、実行されたときに、GPU12に、レジスタr1の値を、メモリの、outputImage[i+1]が参照する位置に格納させる。簡潔のために図示されていない他の命令は、実行されたときに、同様に、GPU12に、レジスタ、たとえばレジスタrx(xは整数である)の値をメモリの位置outputImage[i+x]に格納させる場合がある。

ロードおよびストアを並べ替えると、図3Aのコード80と比較してコード160の実行性能が向上する場合がある。より詳細には、ロードおよびストアを並べ替えると、キャッシュライン幅に応じて、特定の場合に性能が向上することがある。たとえば、コードリオーダリングは、いくつかのロード命令をまとめて合体することが、スカラー命令において使用されるオペランドのサイズの倍数であるキャッシュライン幅を有するシステム上の性能を向上させることがあるときに実行性能を向上させる場合がある。

図4Cは、本開示の技法によるコードベクトル化を示す概念図である。図4Cはコードセクション180を含み、コードセクション180は、ライン182、184、および186をさらに含む。上述のように、JITコンパイラ18および/またはコンパイラ/ドライバ19は、コードセクションにおけるメモリ参照が同じメモリ領域を参照していないと判定したことに応答して、図4Aに示すループアンローリングなどの特定の最適化を実行してもよい。JITコンパイラ18および/またはコンパイラ/ドライバ19は、ループコードセクションのメモリ参照に関する情報を含むドライバ/ランタイム19からのメタデータに基づいてそのコードセクションをベクトル化するように構成される。

ベクトル化は、コンパイラ(たとえば、JITコンパイラ18)および/またはドライバ/ランタイム19が、各々が単一のオペランドを有する複数のスカラー命令を、複数のオペランドを有する単一のベクトル命令として組み合わせるプロセスである。ベクトル化は、特定のコードセクションを完了するためにプロセッサが実行する必要のある命令の数を低減させ、ならびに固有のハードウェア機能を利用してデータをシステムメモリ14とGPU12との間を移動させることによって実行性能を向上させる並列化の一形態である。図4Cのコードセクション180の例では、JITコンパイラ18は、図4Bに示すようにロードおよびストアを並べ替えてもよい。JITコンパイラ18は、ロードおよびストアを並べ替えた後、図4Cに示すように、同様の命令の各グループをベクトル化してもよい。

ライン182において、JITコンパイラ18は、複数のロード(ldg)命令を単一のベクトル化された命令として組み合わせている。ベクトル化された命令は、実行されたときに、インデックス[i]〜[i+3]におけるinputImageをレジスタr0〜r3にロードする。同様に、ライン184において、JITコンパイラ18は、ライン166、168などの複数のロード命令を、local_cdf[r0-r3]の値をレジスタr0〜r3にロードする単一のベクトル化されたロード命令として組み合わせる。さらに、ライン186において、JITコンパイラ18は、ライン170〜172のストア("stg"命令)を、レジスタr0〜r3の値をoutputImage[i]〜outputImage[i+3]に格納する単一のベクトル化されたストア命令として組み合わせている。

JITコンパイラ18および/またはドライバ/ランタイム19は、図4Bおよび図4Cに示すように命令を並べ替えるかまたはベクトル化するために、任意の依存性に配慮しなければならない。依存性は、ステートメントまたは命令間の実行順序制約を生成する関係である。一例として、ステートメントS2よりも前に別のステートメントS1を実行しなければならない場合にステートメントS2にはステートメントS1に依存性がある。JITコンパイラ18および/またはドライバ/ランタイム19は、依存性によってベクトル化および/またはコードリオーダリングが妨げられるかどうかを判定するために、ドライバ/ランタイム19から取得されるメタデータに基づいて、本開示の技法によるコードのリオーダリングまたはベクトル化の前に依存性分析を実行してもよい。

図5は、本開示の技法に従ってコンパイラ最適化を補助するためにコンパイラメタデータを生成するための例示的な方法を示す流れ図である。概して、図6の方法が、コンパイリングプロセッサ、たとえばCPU16を実行するJITコンパイラ18および/またはドライバ/ランタイム19ならびにターゲットプロセッサ、たとえばGPU12からなるグループのうちの少なくとも一方によって実施されてもよいことを理解されたい。いくつかの例では、ターゲットプロセッサとコンパイリングプロセッサは同じであってもよい。さらに、複数のコンパイリングプロセッサおよび/またはターゲットプロセッサがあってもよい。

図5の方法では、コンパイリングプロセッサ、たとえばCPU16は、ドライバ/ランタイム19および/またはJITコンパイラ18を使用して、コンパイルされたカーネル20のバイナリコードまたはバイトコードを実行するための引数(たとえば、カーネル引数26)を生成する(202)。ドライバ/ランタイム19および/またはJITコンパイラ18はさらに、カーネル引数26の第1のメモリ領域に対する第1のメモリ参照およびカーネル引数26の第2のメモリ領域に対する第2のメモリ参照がカーネル引数26の同じメモリ領域を参照しているかどうかを判定する(204)か、または図3B、図3C、および図3Dに示すような考えられる関係の他の例を判定する。

CPU16は、ドライバ/ランタイム19および/またはJITコンパイラ18を使用して、第1のメモリ参照および第2のメモリ参照に関連するメタデータを生成する(204)。メタデータは、第1のメモリ領域と第2のメモリ領域との間の重複領域などの、第1のメモリ領域と第2のメモリ領域との間の関係を示す。メタデータは、第1のメモリ領域と第2のメモリ領域との間の重複のバイト数をさらに含む場合がある。いくつかの例では、メタデータは、メモリ重複領域の開始アドレスとメモリ重複領域の終了アドレスとを含んでもよい。図5に関して説明した例が、例示のみを目的としてメモリ参照の単一の対を参照していることを理解されたい。ドライバ/ランタイム19および/またはJITコンパイラ18は、カーネル引数26のメモリ参照のすべての対に関するメタデータを導出してもよい。

CPU16上で実行されるJITコンパイラ18は、ドライバ/ランタイム19を使用して、第1および第2のメモリ参照がカーネル引数26の同じメモリ領域を参照していないと判定したことに応答して、CPU16に、JITコンパイラ18を使用してメタデータに基づいてカーネル20を再コンパイルさせてもよい(208)。最後に、ターゲットプロセッサ、たとえばGPU12は、再コンパイルされたカーネルを実行してもよい(210)。いくつかの例では、ドライバ/ランタイム19および/またはJITコンパイラ18は、メタデータに基づいて、第1および第2のメモリ参照が同じメモリ領域を参照していないと判定してもよく、この情報を使用して最適化によってカーネル20を再コンパイルしてもよい。

いくつかのさらなる例では、CPU16は、カーネル20の第1のメモリ参照および第2のメモリ参照が同じメモリ領域を参照しているかどうかを判定するために、ドライバ/ランタイム19を使用して、第1および第2のメモリ参照を含むカーネル20のループコードセクションを判定してもよい。JITコンパイラ18は、カーネルを再コンパイルするために、ドライバ/ランタイム19および/またはJITコンパイラ18によって生成されたメタデータに基づいてループコードセクションをアンローリングしてもよい。JITコンパイラ18はまた、カーネルを再コンパイルするために、ループコードセクションのロード演算およびストア演算のうちの少なくとも一方を並べ替えるかまたはループコードセクションの複数のスカラー命令を生成されたメタデータに基づいて少なくとも1つのベクトル命令としてベクトル化してもよい。様々な例では、JITコンパイラ18は、Microsoft DirectComputeおよび/またはKhronos GroupによるOpenCLなどの異種フレームワークを使用してカーネル20を再コンパイルしてもよい。

本開示において説明した技法は、少なくとも部分的に、ハードウェア、ソフトウェア、ファームウェア、またはそれらの任意の組合せで実装することができる。たとえば、前述の技法の様々な態様は、1つまたは複数のマイクロプロセッサ、デジタル信号プロセッサ(DSP)、特定用途向け集積回路(ASIC)、フィールドプログラマブルゲートアレイ(FPGA)、または任意の他の同等の集積論理回路もしくはディスクリート論理回路、ならびにそのような構成要素の任意の組合せを含む、1つまたは複数のプロセッサ内に実装されてもよい。「プロセッサ」または「処理回路」という用語は概して、前述の論理回路のいずれかを指すか、あるいは前述の論理回路のいずれかと他の論理回路または処理を実行する個別ハードウェアなどの任意の他の同等の回路との組合せを指す場合がある。

そのようなハードウェア、ソフトウェア、およびファームウェアは、本開示において説明する様々な動作および機能をサポートするために同一のデバイス内に実装されてもあるいは別個のデバイス内に実装されてもよい。さらに、前述のユニット、モジュール、または構成要素はいずれも、まとめて実装されてもよく、あるいは離散しているが相互運用可能である論理デバイスとして別個に実装されてもよい。様々な特徴をモジュールまたはユニットとして記述することは、様々な機能態様を強調することを目的としたものであり、そのようなモジュールまたはユニットを別個のハードウェア構成要素またはソフトウェア構成要素によって実現しなければならないことを必ずしも意味しない。むしろ、1つまたは複数のモジュールまたはユニットに関連する機能が、別個のハードウェア構成要素、ファームウェア構成要素、および/またはソフトウェア構成要素によって実行されても、あるいは共通のハードウェア構成要素もしくはソフトウェア構成要素または別個のハードウェア構成要素もしくはソフトウェア構成要素内に集積されてもよい。

本開示において説明する技法はまた、命令を記憶するコンピュータ可読記憶媒体などのコンピュータ可読媒体に格納されても、あるいはコンピュータ可読媒体において具体化または符号化されてもよい。コンピュータ可読媒体に埋め込まれるかまたはコンピュータ可読媒体において符号化された命令は、たとえば、命令が1つまたは複数のプロセッサによって実行されるときには、これらのプロセッサに、本明細書において説明する技法を実行させてもよい。コンピュータ可読記憶媒体には、ランダムアクセスメモリ(RAM)、読取り専用メモリ(ROM)、プログラマブル読取り専用メモリ(PROM)、消去可能プログラマブル読取り専用メモリ(EPROM)、電子的消去可能プログラマブル読取り専用メモリ(EEPROM)、フラッシュメモリ、ハードディスク、CD-ROM、フロッピーディスク、カセット、磁気媒体、光メディア、または他の有形のコンピュータ可読記憶媒体を含めてもよい。

コンピュータ可読媒体には、上記に記載したような、有形の記憶媒体に対応する、コンピュータ可読記憶媒体を含めてもよい。コンピュータ可読媒体にはまた、たとえば、通信プロトコルに従って、ある場所から別の場所へのコンピュータプログラムの転送を容易にする任意の媒体を含む通信媒体を含めてもよい。このようにして、「コンピュータ可読媒体」という句は、一般に、(1)非一時的な有形のコンピュータ可読記憶媒体、および(2)一時的な信号もしくは搬送波などの無形のコンピュータ可読通信媒体に対応してもよい。

様々な態様および例について説明した。しかし、以下の特許請求の範囲から逸脱せずに本開示の構造または技法に修正を施すことが可能である。

14 システムメモリ
16 CPU
19 ドライバ/ランタイム
20 カーネル
22 シェーダコア
26 引数
48 キャッシュ
80 コードセクション
82 ライン
86 ループ
100 バッファ
130 バッファ
140 コードセクション
142 ライン
144 ライン
160 コード
166 ライン
168 ライン
180 コードセクション

Claims (30)

1. カーネルを実行できるようにコンパイルする方法であって、
   コンパイリングプロセッサ上で実行されるコンパイラによって、カーネルをコンパイルするステップと、
   前記コンパイリングプロセッサ上で実行される前記コンパイラとドライバとからなるグループのうちの少なくとも一方によって、前記コンパイルされたカーネルのコードを実行するためのカーネル引数を生成するステップと、
   前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記カーネル引数の第1のメモリ領域に対する第1のメモリ参照と前記カーネル引数の第2のメモリ領域に対する第2のメモリ参照とが同じメモリ領域を参照しているかどうかを判定するステップと、
   前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記判定に基づいて前記第1のメモリ参照および前記第2のメモリ参照に関連するメタデータを生成するステップであって、前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の関係を示す、ステップと、
   前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を対象としていないと判定したことに応答して、
   前記コンパイリングプロセッサ上で実行される前記コンパイラによって、前記メタデータに基づいて前記カーネルを再コンパイルするステップと、
   前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、ターゲットプロセッサに、前記再コンパイルされたカーネルを実行するよう命令するステップと
   を含む、方法。
2. 前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定するステップは、前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのループコードセクションを判定するステップをさらに含み、
   前記カーネルを再コンパイルするステップは、前記メタデータに基づいて前記ループコードセクションをアンローリングするステップと、前記アンローリングされたループコードセクションをコンパイルするステップとを含む、
   請求項1に記載の方法。
3. 前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定するステップは、前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのコードセクションを判定するステップをさらに含み、
   前記カーネルを再コンパイルするステップは、前記コードセクションの前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、前記メタデータに基づいて前記コードセクションのロード演算およびストア演算のうちの少なくとも一方を並べ替えるステップをさらに含む、
   請求項1に記載の方法。
4. 前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定するステップは、前記コンパイリングプロセッサ上で実行される前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのコードセクションを判定するステップをさらに含み、
   前記カーネルを再コンパイルするステップは、前記コードセクションの前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、前記メタデータに基づいて前記コードセクションの複数のスカラー命令を少なくとも1つのベクトル命令へとベクトル化するステップをさらに含む、
   請求項1に記載の方法。
5. 前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の重複領域をさらに示す、請求項1に記載の方法。
6. 前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の重複のバイト数を含む、請求項5に記載の方法。
7. 前記メタデータは、前記メモリ重複領域の開始アドレスおよび前記メモリ重複領域の終了アドレスのうちの少なくとも一方を含む、請求項5に記載の方法。
8. 前記コンパイリングプロセッサは中央演算処理ユニット(CPU)を備え、前記ターゲットプロセッサはグラフィックス処理ユニット(GPU)を備える、請求項1に記載の方法。
9. 前記コンパイラは、異種コンピューティングフレームワークを使用して前記カーネルを再コンパイルする、請求項1に記載の方法。
10. 前記カーネル引数は、前記引数用に割り振られたメモリのバッファ領域を含む、請求項1に記載の方法。
11. メモリと、
    コンパイリングプロセッサであって、
    前記コンパイリングプロセッサのコンパイラによって、カーネルをコンパイルすることと、
    前記コンパイリングプロセッサの前記コンパイラとドライバとからなるグループのうちの少なくとも一方によって、前記コンパイルされたカーネルのコードを実行するためのカーネル引数を生成することと、
    前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記カーネル引数の第1のメモリ領域に対する第1のメモリ参照と前記カーネル引数の第2のメモリ領域に対する第2のメモリ参照とが同じメモリ領域を参照しているかどうかを判定することと、
    前記コンパイリングプロセッサのコンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記判定に基づいて前記第1のメモリ参照および前記第2のメモリ参照に関連するメタデータを生成することであって、前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の関係を示す、生成することと、
    前記コンパイリングプロセッサのコンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、
    前記コンパイリングプロセッサの前記コンパイラによって、前記メタデータに基づいて前記カーネルを再コンパイルすることと、
    前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、ターゲットプロセッサに、前記再コンパイルされたカーネルを実行するよう命令することと
    を行うように構成される、コンパイリングプロセッサと
    を備える、デバイス。
12. 前記コンパイリングプロセッサは、前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定するために、前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのループコードセクションを判定するようにさらに構成され、
    前記コンパイリングプロセッサは、前記カーネルを再コンパイルするために、前記メタデータに基づいて前記ループコードセクションをアンローリングし、かつ前記アンローリングされたループコードセクションをコンパイルするようにさらに構成される、
    請求項11に記載のデバイス。
13. 前記コンパイリングプロセッサは、前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定するために、前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのコードセクションを判定するようにさらに構成され、
    前記コンパイリングプロセッサは、前記カーネルを再コンパイルするために、前記コードセクションの前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、前記メタデータに基づいて前記コードセクションのロード演算およびストア演算のうちの少なくとも一方を並べ替えるようにさらに構成される、
    請求項11に記載のデバイス。
14. 前記コンパイリングプロセッサは、前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定するために、前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのコードセクションを判定するようにさらに構成され、
    前記コンパイリングプロセッサは、前記カーネルを再コンパイルするために、前記コードセクションの前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、前記メタデータに基づいて前記コードセクションの複数のスカラー命令を少なくとも1つのベクトル命令へとベクトル化するようにさらに構成される、
    請求項11に記載のデバイス。
15. 前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の重複領域を示す、請求項11に記載のデバイス。
16. 前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の重複のバイト数を含む、請求項15に記載のデバイス。
17. 前記メタデータは、前記メモリ重複領域の開始アドレスおよび前記メモリ重複領域の終了アドレスのうちの少なくとも一方を含む、請求項15に記載のデバイス。
18. 前記コンパイリングプロセッサは中央演算処理ユニット(CPU)を備え、前記ターゲットプロセッサはグラフィックス処理ユニット(GPU)を備える、請求項11に記載のデバイス。
19. 前記コンパイラは、異種コンピューティングフレームワークを使用して前記カーネルを再コンパイルする、請求項11に記載のデバイス。
20. 前記カーネル引数は、前記引数用に割り振られたメモリのバッファ領域を含む、請求項11に記載のデバイス。
21. 記憶された命令を含む非一時的コンピュータ可読記憶媒体であって、前記命令は、実行されたときに、コンパイリングプロセッサに、
    前記コンパイリングプロセッサのコンパイラによって、カーネルをコンパイルすることと、
    前記コンパイリングプロセッサの前記コンパイラとドライバとからなるグループのうちの少なくとも一方によって、前記コンパイルされたカーネルのコードを実行するためのカーネル引数を生成することと、
    前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記カーネル引数の第1のメモリ領域に対する第1のメモリ参照と前記カーネル引数の第2のメモリ領域に対する第2のメモリ参照とが同じメモリ領域を参照しているかどうかを判定することと、
    前記コンパイリングプロセッサのコンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記判定に基づいて前記第1のメモリ参照および前記第2のメモリ参照に関連するメタデータを生成することであって、前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の関係を示す、生成することと、
    前記コンパイリングプロセッサのコンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記カーネル引数の前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、
    前記コンパイリングプロセッサの前記コンパイラによって、前記メタデータに基づいて前記カーネルを再コンパイルすることと、
    前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、ターゲットプロセッサに、前記再コンパイルされたカーネルを実行するよう命令することと
    を行わせる、非一時的コンピュータ可読記憶媒体。
22. 前記コンパイリングプロセッサに、前記カーネル引数の前記第1のメモリ参照と前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定させる前記命令は、実行されたときに、前記コンパイリングプロセッサに、前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのループコードセクションを判定させる命令をさらに含み、
    前記コンパイリングプロセッサに、前記カーネルを再コンパイルさせる前記命令は、実行されたときに、前記コンパイリングプロセッサに、前記メタデータに基づいて前記ループコードをアンローリングさせ、かつ前記アンローリングされたループコードセクションをコンパイルさせる命令をさらに含む、
    請求項21に記載の非一時的コンピュータ可読記憶媒体。
23. 前記コンパイリングプロセッサに、前記カーネル引数の前記第1のメモリ参照と前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定させる命令は、実行されたときに、前記コンパイリングプロセッサに、前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのコードセクションを判定させる命令をさらに含み、
    前記コンパイリングプロセッサに、前記カーネルを再コンパイルさせる前記命令は、実行されたときに、前記コンパイリングプロセッサに、前記コードセクションの前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、前記メタデータに基づいて前記コードセクションのロード演算およびストア演算のうちの少なくとも一方を並べ替えさせる命令をさらに含む、
    請求項21に記載の非一時的コンピュータ可読記憶媒体。
24. 前記コンパイリングプロセッサに、前記カーネル引数の前記第1のメモリ参照と前記第2のメモリ参照が前記同じメモリ領域を参照しているかどうかを判定させる前記命令は、実行されたときに、前記コンパイリングプロセッサに、前記コンパイリングプロセッサの前記コンパイラと前記ドライバとからなる前記グループのうちの前記少なくとも一方によって、前記第1のメモリ参照および前記第2のメモリ参照を含む前記カーネルのコードセクションを判定させる命令をさらに含み、
    前記コンパイリングプロセッサに、前記カーネルを再コンパイルさせる前記命令は、実行されたときに、前記コンパイリングプロセッサに、前記コードセクションの前記第1のメモリ参照および前記第2のメモリ参照が前記同じメモリ領域を参照していないと判定したことに応答して、前記メタデータに基づいて前記コードセクションの複数のスカラー命令を少なくとも1つのベクトル命令へとベクトル化させる命令をさらに含む、
    請求項21に記載の非一時的コンピュータ可読記憶媒体。
25. 前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の重複領域を示す、請求項21に記載の非一時的コンピュータ可読記憶媒体。
26. 前記メタデータは、前記第1のメモリ領域と前記第2のメモリ領域との間の重複のバイト数を含む、請求項25に記載の非一時的コンピュータ可読記憶媒体。
27. 前記メタデータは、前記メモリ重複領域の開始アドレスおよび前記メモリ重複領域の終了アドレスのうちの少なくとも一方を含む、請求項25に記載の非一時的コンピュータ可読記憶媒体。
28. 前記コンパイリングプロセッサは中央演算処理ユニット(CPU)を備え、前記ターゲットプロセッサはグラフィックス処理ユニット(GPU)を備える、請求項21に記載の非一時的コンピュータ可読記憶媒体。
29. 前記コンパイラは、異種コンピューティングフレームワークを使用して前記カーネルを再コンパイルする、請求項21に記載の非一時的コンピュータ可読記憶媒体。
30. 前記カーネル引数は、前記引数用に割り振られたメモリのバッファ領域を含む、請求項21に記載の非一時的コンピュータ可読記憶媒体。

Images