(57)【要約】  アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するデバイスに関する例が説明される。中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、アプリケーションのソースコードにおいて、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される。デバイスは、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコードを生成するために中間コードをコンパイルし、オブジェクトコードを記憶するように構成される。
2018517986.tif 000002
【特許請求の範囲】
【請求項１】
コンパイルする方法であって、
プロセッサにより、アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するステップであって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定される、ステップと、
前記プロセッサにより、前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルするステップと、
前記オブジェクトコードを記憶するステップと
を含む方法。
【請求項２】
前記プロセッサが前記オブジェクトコードを実行すべきであると判断するステップと、
前記プロセッサにより、前記オブジェクトコードを実行するステップと
をさらに含み、
コンパイルするステップは、前記プロセッサに前記下位レベルサブルーチンの複数のインスタンスを順次実行させる前記オブジェクトコードを生成するために、前記中間コードをコンパイルするステップを含む、請求項1に記載の方法。
【請求項３】
グラフィックス処理ユニット(GPU)またはデジタル信号プロセッサ(DSP)が前記オブジェクトコードを実行すべきであると判断するステップと、
前記GPUまたはDSPに前記オブジェクトコードを実行するよう命令するステップと
をさらに含み、
コンパイルするステップは、前記GPUまたはDSPに前記下位レベルサブルーチンの複数のインスタンスを並列に実行させる前記オブジェクトコードを生成するために、前記中間コードをコンパイルするステップを含む、請求項1に記載の方法。
【請求項４】
コンパイルするステップは、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別するステップと、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき回数を決定するステップと、
前記下位レベルサブルーチンが実行されるべき前記回数の前記決定に基づいて、前記オブジェクトコードを生成するステップと
を含む、請求項1に記載の方法。
【請求項５】
前記中間コードを受信するステップは、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信するステップを含む、請求項1に記載の方法。
【請求項６】
前記プロセッサにより、前記アプリケーションを利用するプログラムを実行するステップ
をさらに含み、
コンパイルするステップは、前記プログラムの実行中に前記中間コードをコンパイルするステップを含む、請求項1に記載の方法。
【請求項７】
コンパイルするためのデバイスであって、
メモリユニットと、
プロセッサを含む集積回路と
を含み、前記プロセッサは、
アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信することであって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定される、受信することと、
前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルすることと、
前記メモリユニットに前記オブジェクトコードを記憶することと
を行うように構成される、デバイス。
【請求項８】
前記プロセッサは、
前記プロセッサが前記オブジェクトコードを実行すべきであると判断することと、
前記オブジェクトコードを実行することと
を行うように構成され、
コンパイルするために、前記プロセッサは、前記プロセッサに前記下位レベルサブルーチンの複数のインスタンスを順次実行させる前記オブジェクトコードを生成するために、前記中間コードをコンパイルするように構成される、請求項7に記載のデバイス。
【請求項９】
グラフィックス処理ユニット(GPU)またはデジタル信号プロセッサ(DSP)のうちの少なくとも一方
をさらに含み、
前記プロセッサは、
前記GPUまたはDSPが前記オブジェクトコードを実行すべきであると判断することと、
前記GPUまたはDSPに前記オブジェクトコードを実行するよう命令することと
を行うように構成され、
コンパイルするために、前記プロセッサは、前記GPUまたはDSPに前記下位レベルサブルーチンの複数のインスタンスを並列に実行させる前記オブジェクトコードを生成するために、前記中間コードをコンパイルするように構成される、請求項7に記載のデバイス。
【請求項１０】
コンパイルするために、前記プロセッサは、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別することと、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき回数を決定することと、
前記下位レベルサブルーチンが実行されるべき前記回数の前記決定に基づいて、前記オブジェクトコードを生成することと
を行うように構成される、請求項7に記載のデバイス。
【請求項１１】
前記中間コードを受信するために、前記プロセッサは、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信するように構成される、請求項7に記載のデバイス。
【請求項１２】
前記プロセッサは、
前記アプリケーションを利用するプログラムを実行するように構成され、
コンパイルするために、前記プロセッサは、前記プログラムの実行中に前記中間コードをコンパイルするように構成される、請求項7に記載のデバイス。
【請求項１３】
ワイヤレス通信デバイスを含む、請求項7に記載のデバイス。
【請求項１４】
コンパイルするためのデバイスであって、
アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するための手段であって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定される、手段と、
前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルするための手段と、
前記オブジェクトコードを記憶するための手段と
を含むデバイス。
【請求項１５】
コンパイルするための前記手段を含むプロセッサが前記オブジェクトコードを実行すべきであると判断するための手段と、
前記オブジェクトコードを実行するための手段と
をさらに含み、
コンパイルするための前記手段は、前記プロセッサに前記下位レベルサブルーチンの複数のインスタンスを順次実行させる前記オブジェクトコードを生成するために、前記中間コードをコンパイルするための手段を含む、請求項14に記載のデバイス。
【請求項１６】
グラフィックス処理ユニット(GPU)またはデジタル信号プロセッサ(DSP)が前記オブジェクトコードを実行すべきであると判断するための手段と、
前記GPUまたはDSPに前記オブジェクトコードを実行するよう命令するための手段と
をさらに含み、
コンパイルするための前記手段は、前記GPUまたはDSPに前記下位レベルサブルーチンの複数のインスタンスを並列に実行させる前記オブジェクトコードを生成するために、前記中間コードをコンパイルするための手段を含む、請求項14に記載のデバイス。
【請求項１７】
コンパイルするための前記手段は、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別するための手段と、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき回数を決定するための手段と、
前記下位レベルサブルーチンが実行されるべき前記回数の前記決定に基づいて、前記オブジェクトコードを生成するための手段と
を含む、請求項14に記載のデバイス。
【請求項１８】
前記中間コードを受信するための前記手段は、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信するための手段を含む、請求項14に記載のデバイス。
【請求項１９】
前記アプリケーションを利用するプログラムを実行するための手段をさらに含み、
コンパイルするための前記手段は、前記プログラムの実行中に前記中間コードをコンパイルするための手段を含む、請求項14に記載のデバイス。
【請求項２０】
実行されると、1つまたは複数のプロセッサに、
アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信させる命令であって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定される、命令と、
前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルさせる命令と、
前記オブジェクトコードを記憶させる命令と
を含む非一時的コンピュータ可読記憶媒体。
【請求項２１】
前記1つまたは複数のプロセッサにコンパイルさせる前記命令は、前記1つまたは複数のプロセッサに、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別させる命令と、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき回数を決定させる命令と、
前記下位レベルサブルーチンが実行されるべき前記回数の前記決定に基づいて、前記オブジェクトコードを生成させる命令と
を含む、請求項20に記載のコンピュータ可読記憶媒体。
【発明の詳細な説明】
【技術分野】
【０００１】
本出願は、その内容全体が参照により本明細書に組み込まれる、2015年6月15日に出願され、「HIERARCHICAL SUB-ROUTINE INFORMATION IN INTERMEDIATE CODE」と題する米国特許出願第62/175,646号の優先権を主張する。
【０００２】
本開示は、コンパイラに関し、より詳細には、処理ユニット上で実行するポータブル中間コードを生成するコンパイラに関する。
【背景技術】
【０００３】
コンピューティングデバイス上のプロセッサは、高水準プログラミング言語で書かれたコード(すなわち、命令)をコンパイルすることから生成されたオブジェクトコードを実行する。オブジェクトコードを生成するために、コンパイラは最初に、高水準プログラミング言語で書かれたコードを中間コード(すなわち、中間言語によるコード)に変換する。同じもしくは何らかの他のシステム上の別のコンパイラ、または同じコンパイラは、中間言語によるコードを機械コードに変換し、次いで機械コードは、プロセッサまたはシステムにおける何らかの他のプロセッサによって実行される。
【発明の概要】
【課題を解決するための手段】
【０００４】
本開示は、高水準言語で書かれたアプリケーションにおける(たとえば、アプリケーションのソースコードにおける)サブルーチン間の階層的関係を識別する関数呼出しを中間コードに含めるための技法について説明する。たとえば、高水準言語で書かれたアプリケーションをコンパイルするコンパイラは、関数呼出しを中間コードに含めることができる。サブルーチン間の階層的関係は、第1のサブルーチンが第2のサブルーチンの完了に依存すること、および第2のサブルーチンが第1のサブルーチンよりも頻繁に実行されることを示す。この場合、第1のサブルーチンは、第2のサブルーチンよりも高いレベルにあると見なされる。
【０００５】
別のコンパイラ(または場合によっては同じコンパイラ)はさらに、中間コードをオブジェクトコードにコンパイルする。中間コードをオブジェクトコードにコンパイルするために、コンパイラ(たとえば、この他のコンパイラまたは同じコンパイラ)は、中間コードにおける関数呼出しを利用して、プロセッサの処理能力を活用するオブジェクトコードを生成し、プロセッサは、オブジェクトコードを実行することになる。このようにして、中間コードは、上位レベルサブルーチンと下位レベルサブルーチンとを分離する関数呼出しを含むので、中間コードは、異なるプロセッサ上で実行することに対して、よりポータブルであり、コンパイラは、オブジェクトコードが実行されるプロセッサの処理能力を最善の形で使用するように実行するためのサブルーチンの構成方法を決定することができる。
【０００６】
一例では、本開示は、コンパイルする方法について説明し、本方法は、プロセッサにより、アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するステップであって、中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、アプリケーションのソースコードにおいて、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される、ステップと、プロセッサにより、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコードを生成するために中間コードをコンパイルするステップと、オブジェクトコードを記憶するステップとを含む。
【０００７】
一例では、本開示は、コンパイルするためのデバイスについて説明し、本デバイスは、メモリユニットと、プロセッサを含む集積回路とを含み、プロセッサは、アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信することであって、中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、アプリケーションのソースコードにおいて、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される、受信することと、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコードを生成するために中間コードをコンパイルすることと、メモリユニットにオブジェクトコードを記憶することとを行うように構成される。
【０００８】
一例では、本開示は、コンパイルするためのデバイスについて説明し、本デバイスは、アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するための手段であって、中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、アプリケーションのソースコードにおいて、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される、手段と、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコードを生成するために中間コードをコンパイルするための手段と、オブジェクトコードを記憶するための手段とを含む。
【０００９】
一例では、本開示は、実行されると、1つまたは複数のプロセッサに、アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信させる命令であって、中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、アプリケーションのソースコードにおいて、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される、命令と、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコードを生成するために中間コードをコンパイルさせる命令と、オブジェクトコードを記憶させる命令とを含む非一時的コンピュータ可読記憶媒体について説明する。
【００１０】
1つまたは複数の例の詳細が、添付図面および以下の説明に記載される。他の特徴、目的、および利点は、説明および図面から、ならびに特許請求の範囲から明らかになろう。
【図面の簡単な説明】
【００１１】
【図１】本開示の態様を実装するための例示的なシステムを示すブロック図である。
【図２】本開示の態様を実装するように構成され得るデバイスを示すブロック図である。
【図３】図2の例示的なデバイスをさらに詳細に示すブロック図である。
【図４】本開示の態様を実装するための一例を示すフローチャートである。
【図５】本開示の態様を実装するための別の例を示すフローチャートである。
【発明を実施するための形態】
【００１２】
プロセッサ(たとえば、グラフィックス処理ユニット(GPU)または中央処理装置(CPU))は、プロセッサにいくつかの機能を実行させるプログラムの命令を実行する。アプリケーション開発者は、特定の規格のシンタックスを使用してプログラムを開発する。たとえば、アプリケーション開発者は、2つの非限定的な例としてOpenCL仕様または異種システムアーキテクチャ(HSA:Heterogeneous System Architecture)仕様のアプリケーションプログラミングインターフェース(API)によって規定された規格を使用し得る。説明しやすいように、本開示で説明する技法は、単に理解を助けるために、OpenCL仕様に関して説明する。ただし、本開示で説明する技法は、特定の規格に従って開発されたアプリケーションに限定されない。
【００１３】
OpenCL仕様は、アプリケーション開発者が、プロセッサにいくつかの機能を実行するよう命令する高水準プログラミング言語によるプログラムを書く(たとえば、アプリケーションのためのソースコードを書く)ことを可能にする。プログラムはその場合、GPU、CPU、フィールドプログラマブルゲートアレイ(FPGA)、デジタル信号プロセッサ(DSP)などのような様々な異なるプロセッサタイプ上で実行され得る。CPUは、ホストと呼ばれることがあり、これらの他の例示的なハードウェアユニットは、より一般的にデバイスと呼ばれることがある。本開示では、デバイスという用語は、CPU、GPU、FPGA、および/もしくはDSPを含むユニットを指すために、またはCPU、GPU、FPGA、および/もしくはDSPを指すために使用され得る。したがって、「デバイス」という用語は、GPU、FPGA、および/またはDSPなどの例に限定されると見なされるべきではなく、CPUならびにこれらの例示的な構成要素を格納するユニットも含む。
【００１４】
「プログラム」という用語は、プロセッサにいくつかの機能を実行するよう命令する高水準プログラミング言語で書かれた命令の任意のセットを指す。たとえば、「プログラム」という用語は、OpenCL仕様に従って開発されたプログラム、または任意の他の仕様を使用して開発されたプログラム、ならびにC++などの任意の高水準プログラミング言語で書かれたプログラムを含むように使用される。OpenCLをベースに構築された高水準プログラミングモデルである、高水準のプログラミング言語の別の例は、SYCLである。
【００１５】
プロセッサがプログラムの命令を実行するために、コンパイラは、プログラムの命令を中間言語(たとえば、アセンブリ言語)のコードに変換する。中間言語によるコードは、中間コードとも呼ばれる。この中間コードは、異なるプロセッサに対してかなりポータブルであり得る。たとえば、ユーザがデバイス(たとえば、モバイルデバイス、ラップトップ、デスクトップなど)に、実行するプログラムをダウンロードさせるとき、デバイスは、中間言語による中間コードをダウンロードし得る。
【００１６】
コンパイラはさらに、中間コードをダウンロードしたデバイス上で、プロセッサ(たとえば、CPU、GPU、FPGA、および/またはDSP)上で実行されるべきオブジェクトコードを生成するために、中間コードをコンパイルし得る。次いでプロセッサは、オブジェクトコードを実行し得る。
【００１７】
中間言語を生成するコンパイラは、フロントエンドコンパイラまたは高水準コンパイラ(high-level compiler)と呼ばれ、オブジェクトコードを生成するコンパイラは、バックエンドコンパイラまたは低水準コンパイラ(low-level compiler)と呼ばれる。フロントエンドコンパイラおよびバックエンドコンパイラは、異なるデバイス上で実行され得る。たとえば、フロントエンドコンパイラは、アプリケーション開発者のデバイス(または中間コードを格納するサーバ)上で、中間言語による中間コードを生成するために実行されることがあり、バックエンドコンパイラは、ユーザのデバイス上で、オブジェクトコードを生成するために実行されることがある。ある例では、バックエンドコンパイラは、オブジェクトコードを事前に記憶するのではなく、プログラムの実行中にオブジェクトコードを生成するオンザフライバックエンドコンパイラであり得る。
【００１８】
フロントエンドコンパイラおよびバックエンドコンパイラは、必ずしも異なるデバイス上で実行される必要はないことを理解されたい。たとえば、上記の例では、ユーザデバイスは、中間コードをダウンロードした。いくつかの他の例では、ユーザデバイスは、中間コードではなく、プログラムの命令をダウンロードし得る。そのような例では、ユーザデバイス上で実行されているコンパイラは、コンパイラのフロントエンドモジュールを介して中間コードを生成し、コンパイラのバックエンドモジュールを介してオブジェクトコードを生成し得る。本開示では、異なるコンパイラは、中間コードを生成し、オブジェクトコードを生成するもの(たとえば、異なるデバイス上で実行されているフロントエンドコンパイラおよびバックエンドコンパイラ)として説明されるが、本技法は、限定的なものと見なされるべきではない。
【００１９】
上記で説明したように、OpenCLは、プラットフォームにわたって(たとえば、異なるタイプのプロセッサ上で)実行されるプログラムを開発するための構造を規定する。場合によっては、開発者は、すべてのプロセッサタイプに適用可能であるわけではない不必要な動作を含んだOpenCLを使用したプログラムを作成している。たとえば、単一プログラム複数データ(SPMD:single-program-multiple-data)並列処理では、プロセッサは、同じプログラムの複数のインスタンスを実行し、各プログラムが異なるデータで動作し得る。OpenCLおよびHSA仕様は、SPMD並列処理を単一命令複数データ(SIMD:single-instruction-multiple-data)実行ユニットにマッピングするために有用である。ただし、そのようなマッピングでは、並列構造についての情報が失われ得る。
【００２０】
これに対処するために、OpenCL仕様は、「バリア」命令を可能にする。バリア命令は、プログラムのインスタンスに、プログラムのすべてのインスタンスが同期できる(すなわち、プログラムのより遅いインスタンスが追いつくことができる)ように実行を中止させる。同期は、プログラムのインスタンスが互いに通信することを可能にする。
【００２１】
たとえば、OpenCLワークグループの場合、インスタンス間で通信するために、開発者はバリアを使用し得る。これらのバリアは、独立したエンティティ間の通信動作として働く。ただし、場合によっては、これらのエンティティの多くは、同じスレッドにマッピングされてよく、そのように扱われることによって、より良く対処される。言い換えれば、異なる並列インスタンスを有するのではなく、直列実行があることが、より良い場合がある。場合によっては、OpenCLのツールチェーンは、エンティティが異なるスレッドにマッピングされるか、それとも同じスレッドにマッピングされるかを推測することが可能ではないことがある。
【００２２】
これに対処するために、SYCL高水準プログラミングモデルが、OpenCLをベースに構築され、並列処理の構造化形態を可能にした。SYCLでは、コードは、そのパフォーマンスがよりポータブルに(すなわち、複数の異なるタイプのデバイスにわたって使用可能に)なるように、より多くの情報を維持する。SYCLがそのような結果を達成する1つの方法は、階層構造を可能にすることによる。階層構造では、プログラムは、複数のサブルーチン(プログラムブロックまたはモジュールとも呼ばれる)を含む。サブルーチンのうちの1つは、別のサブルーチンに従属することがあり、サブルーチンの実行の頻度は異なり得る(たとえば、従属サブルーチンは、他方のサブルーチンが実行されるたびにN回実行され得る)。従属サブルーチンは、階層構造において下位レベルにあると見なされ、それが従属するサブルーチンは、階層構造において上位レベルにあると見なされ得る。
【００２３】
そのような階層構造の基礎的な例は、ネスト化forループ(nested for-loop)である。この例では、上位レベルforループが実行されるたびに、下位レベルforループがN回実行される。
【００２４】
いくつかの以前のケースでは、バリアは、異なるインスタンスのループが互いに追いつくことを可能にするために使用される。SYCLでは、バリアの両側でループ分割(loop fission)を達成しようとする代わりに、階層構造はそれを2つのループとして明示的に分解する。
【００２５】
ただし、そのような階層構造化プログラムが異なるプロセッサタイプ上で実行されるべき方法は、異なり得る。たとえば、超並列処理を実現するGPUでは、若干異なる入力およびより少ないコマンドにより並列に実行されるプログラムの多くのインスタンスを有することが有益であり得るが、超並列処理を実現しないCPUの場合、より多くのコマンドにより実行されるプログラムのより少ないインスタンスを有することが有益であり得る。
【００２６】
生じ得る1つの問題は、中間言語による中間コードを生成するためにプログラムをコンパイルするコンパイラ(たとえば、フロントエンドコンパイラ)が、プログラムが実行されることになるプロセッサタイプ(たとえば、CPUまたはGPU)を認識していないことがあることである。したがって、フロントエンドコンパイラは、中間言語による中間コードを生成する1つの方法を選択するが、この方法は、いくつかのプロセッサタイプには好適であるが、他のプロセッサタイプには好適ではないことがある。
【００２７】
たとえば、中間言語の一例は、SPIRVであり、SYCLを使用して書かれたプログラムは、SPIRVによるコードにコンパイルされる。場合によっては、階層構造を使用して開発されたプログラム(たとえば、SYCLを使用して書かれたプログラム)は、中間言語(たとえば、SPIRV)にコンパイルされるときに階層構造情報を失う。高水準コンパイラ(たとえば、フロントエンドコンパイラ)は、1つのマッピング(たとえば、プロセッサタイプに固有である中間コードを構成すること)を、そのポータブルマッピングをバックエンドコンパイラに任せる代わりに選択しなければならない。
【００２８】
プログラムが実行されることになるプロセッサのタイプを示す情報によりバックエンドコンパイラが構成され得るので、階層構造がさらにコンパイルされるべき方法をバックエンドコンパイラが決定できるようにすることが有益であり得る。いくつかの技法では、フロントエンドコンパイラは、中間コードにおける階層構造情報を失い、これは、バックエンドコンパイラが、オブジェクトコードが実行されるべきプロセッサタイプ向けにオブジェクトコードを特殊化するように中間コードをコンパイルすることができないことを意味する。
【００２９】
本開示で説明する技法によれば、フロントエンドコンパイラは、関数タイプの階層を追加し、中間言語(IL)による(たとえば、一例としてIL SPIRVによる)中間コードに対する命令をディスパッチするように構成され得る。バックエンドコンパイラは、階層構造のこの情報を利用して、中間コードをコンパイルし、プログラムが実行されるべきプロセッサの処理能力を活用するオブジェクトコードを生成することができる。
【００３０】
たとえば、階層構造を識別するために、フロントエンドコンパイラはサブルーチン(たとえば、プログラムブロック)内に、下位レベル従属サブルーチンを識別する関数を含め得る。下位レベルサブルーチンは、それが従属するサブルーチン(たとえば、上位レベルサブルーチン)よりも頻繁に実行され得る。一般に、フロントエンドコンパイラは、実行グラフにおいて何らかの実行ユニットを表すものとして関数を示し得る。実行グラフは、プログラムのサブルーチンが互いに関係する方法を識別する概念図であり得る。サブルーチンを識別するために使用される関数の例は、最新のSPMDワークアイテム、サブグループ全体、ワークグループ全体を含み、大局処理のためのディスパッチエントリポイントを含むこともある。
【００３１】
いくつかの例では、エントリポイントは、このスケール、すなわち、(最上位階層レベルから最下位階層レベルの一例として)Kernel、Workgroup、Subgroup、Work-itemにより示される。これらは、すべてのSPIR-VおよびHSAILカーネルにおいて現在、デフォルトであり得る。
【００３２】
SYCLによる並列処理のより細かい粒度(たとえば、下位階層レベル)にマッピングされるコード(たとえば、サブルーチン)のブロックの場合、呼び出し可能オブジェクト/ラムダ関数の本体として示される。これは、フロントエンドコンパイラが下位レベルサブルーチンを呼び出すための関数を上位レベルサブルーチンに含める1つの例示的な方法であり得る。本質的に、呼び出し可能オブジェクト/ラムダ関数として示すことで、並列処理のより細かい粒度は、それらを、インライン化され得る別個の関数にする。本開示では、インライン化は、プログラムにおいて情報が追加され得る関数を意味する。これの中間言語(IL)バージョンにより得られる中間コードは、関数境界が粒度遷移(granularity transition)を示すように、すべてのそのようなブロックを別個の関数にアウトライン化する(たとえば、分離する)。OpSGFunctionCall、OpWIFunctionCallなど、追加のディスパッチ関数が、より細かい粒度に遷移するように、そこで追加される。
【００３３】
いくつかの例では、フロントエンドコンパイラが、子サブグループのセット(たとえば、下位レベルサブルーチン)を開始するために並列ループ構成体も追加することが可能であり得る。これは、より動的な行為であり得る一方、いくつかのアーキテクチャは、かなり静的な階層を有する。静的な階層では、行為は、ワークグループ粒度からサブグループ粒度への移動が、ワークグループにおけるサブグループごとのサブグループ関数に入ることになる。同様に、サブグループからワークアイテムへの遷移は、サブグループにおけるワークアイテムごとのワークアイテム関数を呼び出すことになる。この例は、事実上、SIMDユニットでの均一状態から不均一状態への遷移である。そのような構造はまた、ハードウェアISA(命令セットアーキテクチャ)が異なる状態または異なる実行ユニットの概念を伴う場合に、当該ISAに埋め込まれ得る。スカラーユニットは、そのような状態変化を通じてベクトルユニットに遷移し得る。
【００３４】
たとえば、複数の下位レベルインスタンス(たとえば、下位レベルサブルーチン)がある例では、フロントエンドコンパイラは、下位レベルサブルーチンを一緒にSIMD命令にコンパイルし、単一のスレッドとして一緒に出し得る。SIMD構造において並列に実行されるスレッドの複数のインスタンスがあり得る。すべての場合において、フロントエンドコンパイラが階層を下るにつれて、同時スレッドの数が増加するとは限らず、代わりに、下位レベルは、単一のスレッドのデータ並列部分にマッピングされ得る。
【００３５】
したがって、いくつかの例では、サーバデバイス(たとえば、ユーザデバイスがアプリケーションをダウンロードするサーバデバイス)は、メモリユニットとプロセッサとを含む。プロセッサは、高水準プログラミング言語で書かれたプログラムのコードをコンパイルするためにコンパイラ(たとえば、フロントエンドコンパイラ)を実行するように構成され得る。この例では、コードをコンパイルすることは、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する、コンパイルすることから生成された中間言語による中間コードにおける情報を追加することを含み、下位レベルサブルーチンは、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行される。プロセッサは、メモリユニットに記憶する中間コードを出力し得る。メモリユニットは、プロセッサの内部またはプロセッサの外部にあり得る。
【００３６】
また、本開示では、デバイス(たとえば、ユーザデバイス)は、高水準プログラミング言語で書かれたプログラムのコードをコンパイルすることから生成された中間言語による中間コードを記憶するように構成されたメモリユニットを含む。中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行される。デバイスはまた、中間コードを受信し、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコードを生成するために中間言語による中間コードを(たとえば、バックエンドコンパイラにより)コンパイルし、オブジェクトコードを実行するように構成されたプロセッサを含む。
【００３７】
図1は、本開示の態様を実装するための例示的なシステム10を示すブロック図である。図示のように、システム10は、ホストデバイス12とユーザデバイス14とを含む。ホストデバイス12は、アプリケーションホスティングサーバ、デスクトップコンピュータ、モバイルコンピューティングデバイス、およびノートブック(たとえば、ラップトップ)コンピュータ、タブレットコンピュータを含む広範囲のデバイスを含み得る。ユーザデバイス14の例としては、デスクトップコンピュータ、ならびにモバイルコンピューティングデバイス、ノートブックコンピュータ、電話ハンドセット(いわゆる「スマート」フォン)、およびタブレットコンピュータなどのワイヤレス通信デバイスがある。ユーザデバイス14のさらなる例としては、セットトップボックス、テレビ、カメラ、ディスプレイデバイス、デジタルメディアプレーヤ、ビデオゲーミングコンソール、車内コンピュータなどがある。ユーザデバイス14のこれらの例のうちの1つまたは複数は、組み合わせられてよい(たとえば、ワイヤレス通信デバイス上で実行されているデジタルメディアプレーヤ)。
【００３８】
ユーザデバイス14は、チャネル16を介してホストデバイス12から中間コード22を受信し得る。チャネル16は、データをホストデバイス12からユーザデバイス14に移動させることが可能な1つまたは複数の媒体またはデバイスを含み得る。一例では、チャネル16は、ホストデバイス12がデータをリアルタイムでユーザデバイス14に直接送信することを可能にする、1つまたは複数の通信媒体を含み得る。この例では、ホストデバイス12は、ワイヤレス通信プロトコルなどの通信規格に従って、符号化されたビデオデータを変調することができ、変調されたビデオデータをユーザデバイス14に送信することができる。1つまたは複数の通信媒体は、無線周波数(RF)スペクトルまたは1つもしくは複数の物理伝送線路などの、ワイヤレスおよび/またはワイヤードの通信媒体を含み得る。1つまたは複数の通信媒体は、ローカルエリアネットワーク、ワイドエリアネットワーク、またはグローバルネットワーク(たとえば、インターネット)などの、パケットベースネットワークの一部を形成し得る。1つまたは複数の通信媒体は、ルータ、スイッチ、基地局、またはホストデバイス12からユーザデバイス14への通信を容易にする他の機器を含み得る。
【００３９】
別の例では、チャネル16は、ホストデバイス12によって生成された中間コード22を記憶する記憶媒体を含み得る。この例では、ユーザデバイス14は、ディスクアクセスまたはカードアクセスを介して記憶媒体にアクセスし得る。記憶媒体は、ブルーレイディスク、DVD、CD-ROM、フラッシュメモリ、またはデータを記憶するための他の適切なデジタル記憶媒体などの、ローカルにアクセスされる様々なデータ記憶媒体を含み得る。
【００４０】
さらなる例では、チャネル16は、ホストデバイス12によって生成された中間コード22を記憶するファイルサーバまたは別の中間記憶デバイスを含み得る。この例では、ユーザデバイス14は、ファイルサーバまたは他の中間記憶デバイスにおいて記憶された中間コード22に、ストリーミングまたはダウンロードを介してアクセスし得る。ファイルサーバは、データを記憶するとともにデータをユーザデバイス14に送信することが可能なタイプのサーバであり得る。例示的なファイルサーバは、(たとえば、ウェブサイト用の)ウェブサーバ、ファイル転送プロトコル(FTP)サーバ、ネットワーク接続ストレージ(NAS)デバイス、およびローカルディスクドライブを含む。
【００４１】
図示のように、ホストデバイス12は、メモリユニット18とプロセッサ24とを含む。プロセッサ24は、1つまたは複数のマイクロプロセッサ、デジタル信号プロセッサ(DSP)、特定用途向け集積回路(ASIC)、フィールドプログラマブルゲートアレイ(FPGA)、ディスクリート論理回路、ハードウェア、またはそれらの任意の組合せなどの、様々な適切な回路のいずれかとして実装され得る。メモリユニット18は、同期DRAM(SDRAM)を含むダイナミックランダムアクセスメモリ(DRAM)、磁気抵抗RAM(MRAM)、抵抗RAM(RRAM(登録商標))、または他のタイプのメモリデバイスなどの、様々なメモリデバイスのいずれかによって形成され得る。様々な例では、メモリユニット18は、プロセッサ24に対してオンチップであり得る。メモリユニット18がプロセッサ24のオフチップである例では、メモリユニット18の例としては、限定はしないが、ランダムアクセスメモリ(RAM)、電気消去可能プログラマブル読取り専用メモリ(EEPROM)、フラッシュメモリ、または所望のプログラムコードを命令および/もしくはデータ構造の形で搬送もしくは記憶するために使用され得るとともにコンピュータもしくはプロセッサによってアクセスされ得る他の媒体があり得る。
【００４２】
開発者は、実行されるべきアプリケーションを開発し、メモリユニット18上のソースコード20としてアプリケーションを記憶する。アプリケーションの例としては、OpenCLによるカーネルまたはOpenGLによるシェーダプログラムがあり、両方とも、グラフィックス処理ユニット(GPU)36などのGPU上で実行される。説明しやすいように、ソースコード20は、GPU36上で実行されるものとして説明される。ただし、本開示で説明する技法は、そのように限定されず、ソースコード20は、異なるプロセッサタイプ上で実行されるアプリケーションのためのものであってよい。
【００４３】
メモリユニット18は、ソースコード20および中間コード22を記憶する。ソースコード20を受信し、メモリユニット18に記憶する中間コード22を生成するフロントエンドコンパイラ26を、プロセッサ24は実行する。この例では、開発者は、SYCLなどの高水準言語によるソースコード20を開発し得る。開発者は、アプリケーションの複数のインスタンスが並列に実行され得るようなソースコード20を開発し得る。
【００４４】
ソースコード20を開発するための1つの方法は、階層的に構成されたワークグループ、サブグループ、およびワークアイテムをソースコード20に含めることである。たとえば、ワークグループは、1つもしくは複数のサブグループおよび/または1つもしくは複数のワークアイテムを含み、サブグループは、1つまたは複数のワークアイテムを含む。この例では、ワークグループは最上位レベルにあり、サブグループは中間レベルにあり、ワークアイテムは最下位レベルである。
【００４５】
説明しやすいように、ワークグループ、サブグループ、およびワークアイテムは、全般的にサブルーチンと呼ばれる。階層構造では、上位レベルサブルーチンは、下位レベルサブルーチンの実行の完了に依存し、下位レベルサブルーチンは、ソースコード20において、上位レベルサブルーチンよりも頻繁に実行されるものとして規定される。一例として、開発者は、ワークアイテムがN回実行されるべきであることをサブグループが規定するソースコード20を書くことができる。より詳細に説明するように、下位レベルサブルーチンは、ワークアイテムのすべてのインスタンスを並列に実行するのに十分な実行経路がある場合に、必ずしも常に上位レベルサブルーチンよりも頻繁に実行される必要はない。
【００４６】
本開示で説明する技法では、ホストデバイス12のプロセッサ24は、メモリユニット18からソースコード20を取り出し、プロセッサ24がメモリユニット18に再び記憶する中間コード22を生成するためにソースコード20をコンパイルするためにフロントエンドコンパイラ26を実行する。フロントエンドコンパイラ26の出力は、中間コード22であり、実行されるオブジェクトコードではないので、フロントエンドコンパイラ26は、「フロントエンド」コンパイラと呼ばれる。中間コード22の一例は、標準化ポータブル中間表現(SPIR:standardized portable intermediate representation)アセンブリ言語によるコードである。
【００４７】
中間コード22の利点は、中間コード22がポータブルであることである。たとえば、ユーザデバイス14などのデバイスは、中間コード22を取り出し、オブジェクトコード(たとえば、オブジェクトコード30)をローカルに生成するために中間コード22をコンパイルすることができる。このようにして、中間コード22は、様々なタイプのプロセッサに対して汎用的であり(したがってポータブルであり)得、オブジェクトコードは、オブジェクトコードが実行されるべきプロセッサタイプに基づいて、中間コード22からローカルに生成される。
【００４８】
一例として、ユーザデバイス14のプロセッサ32は、ユーザデバイス14にチャネル16を介して中間コード22を取り出させることができ、中間コード22をメモリユニット28に記憶することができる。プロセッサ32およびメモリユニット28の例は、プロセッサ24およびメモリユニット18に関して上記で説明したものと同様である。ただし、プロセッサ32およびメモリユニット28は、電力および空間が貴重なものではないデスクトップコンピュータまたはサーバまたはコンピューティングユニット上にあり得るプロセッサ24およびメモリユニット18と比較して、電力および空間が貴重なものであるハンドヘルドデバイス上にあり得るので、プロセッサ24およびメモリユニット18は、プロセッサ32およびメモリユニット28と比較して、より「ロバスト」(たとえば、より高速のプロセッサおよびより大きいメモリ)であり得る。本技法は、プロセッサ32およびメモリユニット28がプロセッサ24およびメモリユニット18ほどロバストではないことに限定されず、そのような説明は、単に例示の目的で与えられる。
【００４９】
プロセッサ32は中間コード22を取り出し、オブジェクトコード30を生成するためにバックエンドコンパイラ34を実行することができ、プロセッサ32はオブジェクトコード30をメモリユニット28に再び記憶する。いくつかの例では、プロセッサ32は、オブジェクトコード30によって表されるアプリケーションに実行を求めるプログラムの実行中にオブジェクトコード30を生成し得る。たとえば、プロセッサ32は、プログラム(たとえば、ビデオゲーム)を実行することがあり、ビデオゲームの実行中に、プロセッサ32は、アプリケーション(たとえば、頂点シェーダ)が実行されるべきであると判断することがある。この例では、中間コード22は、頂点シェーダのためのものであり得、プロセッサ32は、ビデオゲームの実行とともにオンザフライでオブジェクトコード30を生成するためにバックエンドコンパイラ34を実行し得る。本技法は、オンザフライでコンパイルすることに限定されると見なされるべきではなく、実行前にコンパイルすることも可能である。
【００５０】
図1に示すように、ユーザデバイス14はグラフィックス処理ユニット(GPU)36を含み、GPU36はオブジェクトコード30を実行する。ただし、いくつかの例で説明するように、プロセッサ32または何らかの他の処理ユニットがオブジェクトコード30を実行することがある。GPU36の例としては、限定はしないが、DSP、汎用マイクロプロセッサ、ASIC、FPGA、または他の等価な集積論理回路もしくはディスクリート論理回路がある。いくつかの例では、GPU36は、グラフィックス処理のために特別に設計された専用ハードウェアであり得る。たとえば、グラフィックス処理は、高速並列処理を必要とすることがあり、GPU36は、そのような高速並列処理のために特別に設計されることがある。GPU36が一般的な処理タスクなど、グラフィックス処理に加えてタスクを実行することが可能であり得る。したがって、GPU36は、一般的処理GPU(GPGPU:general processing GPU)と見なされ得る。本開示で説明する技法は、GPU36がグラフィックス関連タスクのみを実行する例またはGPU36がGPGPUである例に適用され得る。
【００５１】
上記で説明したように、オブジェクトコード30は、GPU36上で実行され得るが、代替的にプロセッサ32上で実行されてもよい。ただし、プロセッサ32およびGPU36の実行能力は異なることがあり、これは、オブジェクトコード30がプロセッサ32およびGPU36で別様に実行され、結果として実行を完了させるのに要する時間量が異なり得ることを意味する。
【００５２】
一例として、GPU36は、同じプログラムの複数のインスタンスを実行(たとえば、SPMD並列処理)できるようなSIMD処理を実現するように構成され得る。SPMD処理では、同じプログラムの複数のインスタンスが並列に(たとえば、同時に)実行されており、SIMD処理では、同じスレッドの複数のインスタンスが並列に(たとえば、同時に)実行されている。プログラムまたは命令のインスタンスの各々の場合の(たとえば、SPMDまたはSIMDの場合の)データは異なることがあるが、命令は同じである。
【００５３】
同じ命令(たとえば、スレッド)の複数のインスタンスが並列に実行され得るように、GPU36が複数のSIMDレーン(たとえば、32または64)を提供するという意味で、GPU36は特殊化され得る。一方、プロセッサ32は、並列処理を実現するように設計されていないことがあり、順次処理の方が適していることがある。たとえば、プロセッサ32は、並列に複数のインスタンスではなく、一度に命令の1つのインスタンスのみを実行することが可能であり得る。プロセッサ32およびGPU36の処理能力におけるそのような差のために、オブジェクトコード30は別様に実行され得る。
【００５４】
理解を助けるための基本的な例として、forループが32回ループスルーし(たとえば、for(int i = 0; i < 32; i++))、ループごとに、GPU36が特定のサブルーチンを実行すると仮定する。このforループを実施するための様々な方法があり得る。第1の方法は、forループを維持し、forループにサブルーチンを32回、順次実行させることである。第2の方法は、forループを「平坦化」し、サブルーチンの32個の並列実行をすることである。他の方法も存在する(たとえば、2まで数えるforループにより、16個のSIMDレーンでサブルーチンを実行する)。上記の例では、このforループを実行するための第1の方法は、プロセッサ32には好適であり、GPU36にはあまり適していない。このforループを実行するための第2の方法は、GPU36には好適であり、プロセッサ32にはあまり適していない。したがって、バックエンドコンパイラ34がこのforループのためのオブジェクトコード30の一部であるオブジェクトコードを生成する方法は、プロセッサ32またはGPU36がオブジェクトコード30を実行する際の効率性に影響する。
【００５５】
本開示で説明する技法によれば、プロセッサ24は、フロントエンドコンパイラ26を実行しており、中間コード22を生成する際の関数呼出しを含み得る。これらの関数呼出しは、ソースコード20の階層構造を識別し得る。そして、バックエンドコンパイラ34は、オブジェクトコード30を生成するために中間コード22をコンパイルする方法を決定するために、これらの関数呼出しを利用し得る。
【００５６】
たとえば、前述の例を続けると、ソースコード20におけるforループの命令を含むサブルーチンは、forループの一部であるサブルーチンよりも高いレベルにあると見なされ得る。これは、上位レベルサブルーチン(たとえば、forループの命令を含むもの)は、下位レベルサブルーチンのすべてのインスタンス(たとえば、ループされる命令)まで完了できないためであり、下位レベルサブルーチンはソースコード20において、上位レベルサブルーチンよりも多くの回数実行されるように実装される。
【００５７】
上記の例は、ソースコード20の階層構造の上位レベルサブルーチンおよび下位レベルサブルーチンを示すためにforループに基づいている。ただし、本技法はループに限定されず、他のサブルーチンにも及ぶ。たとえば、上位レベルサブルーチンは、下位レベルサブルーチンがN回実行されるべきであることを示す命令を含み得るが、必ずしもこれをforループとして定義するとは限らない。他のそのような例も存在し、本技法は上記の例に限定されると見なされるべきではない。
【００５８】
また、ソースコード20において、下位レベルサブルーチンは、上位レベルサブルーチンよりも多くの回数実行されるものと規定され得る。ただし、実行では、これが当てはまらないことがある。たとえば、上位レベルサブルーチンの命令が、下位レベルサブルーチンがN回実行されるべきであることを示し、GPU36上にN個のSIMDレーンが存在すると仮定する。この例では、バックエンドコンパイラ34は、下位レベルサブルーチンのためのコードを含む中間コード22をコピーし、その下位レベルサブルーチンコードを、上位レベルサブルーチンコードとともにN回コピーし得る。この例では、上位レベルサブルーチンおよび下位レベルサブルーチンは、両方ともN回実行されることになる。ただし、ソースコード20においては、下位レベルサブルーチンは、上位レベルサブルーチンよりも頻繁に実行されるように規定されている。
【００５９】
フロントエンドコンパイラ26が異なる階層的サブルーチンの間で線引きし得る様々な方法があり得る。上記で説明したように、OpenCLによれば、開発者は、ワークグループ、サブグループ、およびワークアイテム(たとえば、サブルーチンの異なる例)を伴うソースコード20を開発し得る。開発中、ワークグループ、サブグループ、およびワークアイテムの各々は、それぞれの識別値(たとえば、ワークグループ1、ワークグループ2、サブグループ1、サブグループ2、ワークアイテム1、ワークアイテム2など)を割り当てられ得る。フロントエンドコンパイラ26は、(たとえば、下位レベルサブルーチンが上位レベルサブルーチンに埋め込まれる場合)階層的境界を判断するためにサブルーチン識別値を利用し得る。特定のサブルーチンがどのレベルに属するかをフロントエンドコンパイラ26が判断し得る他の方法があり得、本技法は、上記の例に限定されると見なされるべきではない。
【００６０】
いくつかの例では、フロントエンドコンパイラ26は、異なるサブルーチンを分離することができ、サブルーチンの各レベルは、関数呼出しによって識別され得る。たとえば、最上位レベルサブルーチン(たとえば、ワークグループ)は、動作エントリポイント関数(OpEntryPointSKernel)または動作関数ワークグループ関数(OpFunctionWGもしくはOpWGFunctionCall)と呼ばれ得る。中間レベルサブルーチン(たとえば、サブグループ)は、動作関数サブグループ関数(OpFunctionSGまたはOpSGFunctionCall)と呼ばれ得る。最下位レベルサブルーチン(たとえば、ワークアイテム)は、動作関数ワークアイテム関数(OpFunctionWIまたはOpWIFunctionCall)と呼ばれ得る。最下位レベルサブルーチンは、単にOpFunctionとも呼ばれ得る。
【００６１】
ソースコード20全体を構文解析する際、プロセッサ24は、フロントエンドコンパイラ26の実行を介して、上位レベルサブルーチンにおいて、下位レベルサブルーチンのコードを、その下位レベルサブルーチンへの関数呼出しに置き換えることができる。一例として、サブグループ名は、SG1として識別され得、フロントエンドコンパイラ26は、サブグループSG1のためのコードを関数呼出し: OpSGFunctionCall(SG1)に置き換えることができる。別の例として、識別値WI1を有するワークアイテムがワークグループに埋め込まれ得る。この例では、ワークグループのためのソースコードの構文解析中、フロントエンドコンパイラ26は、ワークグループにおけるワークアイテムWI1のためのコードを関数呼出し: OpWGFunctionCall(WI1)に置き換えることができる。上記の例では、OpSGFunctionCall(SG1)は、SG1(たとえば、下位レベルサブルーチン)の実行を求めるワークグループサブルーチンにおける関数呼出しと見なされ得る。同様に、OpWGFunctionCall(WI1)は、WI1(たとえば、下位レベルサブルーチン)の実行を求めるワークグループサブルーチンまたはサブグループサブルーチンにおける関数呼出しと見なされ得る。
【００６２】
上記で説明したように、ユーザデバイス14は、チャネル16を介してホストデバイス12から中間コード22を取り出し得る。場合によっては、中間コード22は、サブルーチンが中間コード22に直接埋め込まれた、中間コードの連続的セットではない、分離されたサブルーチンのための中間コードを含み得る。たとえば、ユーザデバイス14は、中間コード22のための命令の1つの連続的セットではなく、中間コード22を取り出すことの一部として、ワークグループ、サブグループ、およびワークアイテムサブルーチンの各々のための中間コードを別個に取り出すことができる。サブルーチンは、別個のユニットに線引きされ得るが、上位レベルサブルーチンの各々は、1つまたは複数の下位レベルサブルーチンへの1つまたは複数の関数呼出しを含み得る。
【００６３】
メモリユニット28における中間コード22の取出しおよび記憶の後、プロセッサ32は、オブジェクトコード30を生成するために中間コード22をコンパイルするためにバックエンドコンパイラ34を実行する。バックエンドコンパイラ34は、中間コード22全体を構文解析し、下位レベルサブルーチンへの関数呼出しのあらゆるインスタンスで、下位レベルサブルーチンへの関数呼出しを平坦化するか、それとも主に順次方式で下位レベルサブルーチンを実行するかを決定するように構成され得る。上記で説明したように、関数呼出しを平坦化することは、複数の関数呼出しを並列に実行することを意味する。
【００６４】
一例として、プロセッサ32は、バックエンドコンパラ34がオブジェクトコード30を生成しているアプリケーションがプロセッサ32上で実行されるべきか、それともGPU36上で実行されるべきかを判断し得る。さらに、プロセッサ32は、プロセッサ32およびGPU36の処理能力を示す情報により事前構成され得る(またはランタイム中にそのような情報を受信し得る)。たとえば、プロセッサ32は、GPU36上のSIMDレーンの数(たとえば、GPU36が並列に実行できる命令の数)を示す情報により事前構成され得る(たとえば、プロセッサ32のオペレーティングシステム)。同様に、プロセッサ32は、プロセッサ32が並列に実行できる命令の数を示す情報により事前構成され得る。
【００６５】
このようにして、バックエンドコンパイラ34は、アプリケーションが実行されるべき処理ユニットの処理能力を活用するような方法で中間コード22をコンパイルし得る。たとえば、プロセッサ32が、アプリケーションがプロセッサ32上で実行されるべきであると判断した場合、バックエンドコンパイラ34は、中間コード22を構文解析することができ、上位レベルサブルーチンにおける同じ下位レベルサブルーチンの複数のインスタンスの実行(たとえば、同じ下位レベルサブルーチンへの複数の呼出し)を中間コード22が示す場合、プロセッサ32は並列処理能力を備えていないので、バックエンドコンパイラ34はこれらを、オブジェクトコード30において実行する順次命令として構成し得る。プロセッサ32が、アプリケーションがGPU36上で実行されるべきであると判断した場合、バックエンドコンパイラ34は、中間コード22を構文解析することができ、上位レベルサブルーチンにおける同じ下位レベルサブルーチンの複数のインスタンスの実行を中間コード22が示す場合、GPU36は並列処理能力を備えているので、バックエンドコンパイラ34はこれらを、オブジェクトコード30において実行する並列命令として構成し得る。
【００６６】
いくつかの例では、完全な並列実行に十分なSIMDレーンがGPU36にないことがある。これらの場合、バックエンドコンパイラ34は、できるだけ実行を並列化するが、必要な場合にループを使用することができる。一例として、32個の並列命令が処理される必要があるが、SIMDレーンが16個しかない場合、バックエンドコンパイラ34は、下位レベルサブルーチンの16個のインスタンスを並列に実行するが、下位レベルサブルーチンの32個のインスタンスすべてが実行されるように、2つの実行のループを含むことができる。
【００６７】
上記で説明したように、ワークアイテムは最下位レベルサブルーチンである。いくつかの例では、バックエンドコンパイラ34は、ワークアイテムへの関数呼出しを解析すると、ワークアイテムがスレッドごとに実施されるべきである(たとえば、ワークアイテムの単一の命令の複数のインスタンスが並列に実行される)と判断し得る。したがって、フロントエンドコンパイラ26は、スレッドレベルでワークアイテムサブルーチンをコンパイルすると見なされ得る。
【００６８】
バックエンドコンパイラ34は、サブグループへの関数呼出しを解析すると、サブグループがウェーブごとに実施される(たとえば、複数のスレッドが合わさって1つのウェーブを形成する)と判断し得る。たとえば、ウェーブごとに実施することは、ウェーブを構成するスレッドのための命令ならびにデータが同じであることを意味する。この例では、上位レベルサブルーチンは、下位レベルサブルーチンの多くの実行サイクルで一定であるので、サブグループサブルーチンは、サブグループサブルーチンが呼び出すワークアイテムサブルーチンのためのウェーブと見なされ得る。
【００６９】
いくつかの例では、フロントエンドコンパイラ26は、(たとえば、下位レベルサブルーチンが実行されるべき回数とコピーの数とが等しくなるように、下位レベルサブルーチンが実行されるべき回数に基づいて)下位レベルサブルーチンの複数のインスタンスをコピーし得る。フロントエンドコンパイラ26は、下位レベルサブルーチンのこれらの複数のインスタンスを一緒にSIMD命令にコンパイルし、下位レベルサブルーチンの複数のコピーを含むサブルーチンへの単一の関数呼出しを実現することができる。そのような例では、アプリケーションがGPU36上で実行されるべきである場合、バックエンドコンパイラ34は、並列実行のために下位レベルサブルーチンの複数のコピーを分離することができる。
【００７０】
上記の例は、例としてforループに関して説明している。ただし、本開示で説明する技法は、ベクトル演算など、他の場合にも適用可能であり得る。たとえば、Nビット変数は複数のMビットフロートと見なされ得る(たとえば、128ビット変数は4つの32ビットフロートである)。命令は、Nビット変数に対して実行されるべきプロセスを定めることができ、次いでこのプロセスは、Mビットフロートに適用され得る。
【００７１】
たとえば、128ビットSSE(ストリーミングSIMD拡張(streaming SIMD extension))を表すためにfloat4を使用すると、以下は、命令の例示的なセットである。
float a = 1.f;
float4 b = (float4)a; //これは、float4にわたって1.fを4回けた上げ(carry)している。
b.x = 2.f; //4つの値のうちの最初のもののみを更新する。
float c = sum(b); //bの各要素を追加する。cは、値5(2+1+1+1)を有することになる。
【００７２】
ただし、命令の上記セットは、命令が並列に実行され得るように、利用可能なSIMDレーンの数に基づいて分割され得る。たとえば、階層的コードでは、命令は、よりきめの細かい関数にbを含めることによって、bに演算を適用し得る。一例として、
float a = 1.f;
float b = a;
if(get_local_id(0) == 0){
b = 2.f;
}
c = cross-lane-sum(b)
【００７３】
上記のコードでは、4つのフロートの各々は、異なるレーンで演算される。たとえば、第1のレーン(たとえば、get_local_id(0) == 0の場合)では、bの値は2に等しく設定され、すべての他のレーンでは、bの値は(たとえば、float b = aのため)1に等しい。その場合、cは、5に等しい2+1+1+1となる、レーンにわたる合計に等しい。
【００７４】
このようにして、128ビットSSEのための命令の第1のセットは、きめの粗い関数と見なされ得、128ビットSSEのための命令の第2のセットは、きめの細かい関数と見なされ得る。場合によっては、2つのフォーマットを組み合わせることが可能であり得る。たとえば、きめの粗い関数がきめの細かい関数を呼び出す階層的コードがある。一例として、
float a = 1.f;
float4 b = (float4)a;
b = fine-update(b); //これは、中間言語(IL)におけるきめの粗い関数からきめの細かい関数への呼出し命令のうちの1つである。
float c = sum(b);
ここで、bは、ILにおけるきめの細かい関数であり、これは次のように表され得る。
void fine-update(float b){
if(get_local_id(0) == 0){
b = 2.f;
}
return b;
}
【００７５】
この例では、きめの細かいコードにおける32ビットフロートに直接マッピングされるきめの粗いコードにおけるfloat4変数(128ビットベクトル)がある。上記の例では、同じデータがベクトルコード(上記の例では一度に128ビット4要素)として、またはきめの細かいコードにおいて更新され得るように、クロスレーン(cross-lane)関数がきめの粗いコードに適用され得る。
【００７６】
したがって、いくつかの例では、ユーザデバイス14は、コンパイルするためのデバイスと見なされ得る。そのような例では、プロセッサ32は、アプリケーションのソースコード20をコンパイルすることから生成された中間コード22を受信するように構成され得る。説明したように、中間コード22は、フロントエンドコンパイラ26によるコンパイルから生成された情報を含み、中間コード22は、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する。下位レベルサブルーチンは、アプリケーションのソースコード20において、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される。プロセッサ32は、バックエンドコンパイラ34を介して、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコード30を生成するために中間コード22をコンパイルする。プロセッサ32は、メモリユニット28にオブジェクトコード30を記憶する。
【００７７】
一例では、プロセッサ32は、プロセッサ32がオブジェクトコード30を実行すべきであると判断することがあり、オブジェクトコード30を実行し得る。この例では、コンパイルするために、プロセッサ32は、バックエンドコンパイラ34を介して、プロセッサ32に下位レベルサブルーチンの複数のインスタンスを順次実行させるオブジェクトコード30を生成するために、中間コード22をコンパイルし得る。
【００７８】
別の例では、プロセッサ32は、GPU36がオブジェクトコード30を実行すべきであると判断することがあり、GPU36にオブジェクトコード30を実行するよう命令し得る。この例では、コンパイルするために、プロセッサ32は、バックエンドコンパイラ34を介して、GPU36に下位レベルサブルーチンの複数のインスタンスを並列に実行させるオブジェクトコード30を生成するために、中間コード22をコンパイルし得る。
【００７９】
バックエンドコンパイラ34が中間コード22をコンパイルするために、バックエンドコンパイラ34は、下位レベルサブルーチンの実行を求める上位レベルサブルーチンにおける関数呼出し(たとえば、OpFunctionSGまたはOpSGFunctionCall、OpFunctionWIまたはOpWIFunctionCallは、SG1またはWI1などのそれぞれの識別情報に基づいて下位レベルサブルーチンの実行を求める)を識別し得る。バックエンドコンパイラ34は、上位レベルサブルーチンから、(たとえば、下位レベルサブルーチンが実行されるべき回数を規定する上位レベルサブルーチンにおけるforループまたは命令がある場合)下位レベルサブルーチンが実行されるべき回数を決定し得る。
【００８０】
バックエンドコンパイラ34は、下位レベルサブルーチンが実行されるべき回数の決定に基づいて、オブジェクトコードを生成し得る。たとえば、プロセッサ32上で実行される場合、バックエンドコンパイラ34は、下位レベルサブルーチンが実行されるべき回数だけプロセッサ32に下位レベルサブルーチンを順次実行させるオブジェクトコード30を生成し得る。GPU36上で実行される場合、バックエンドコンパイラ34は、下位レベルサブルーチンが実行されるべき回数によって規定されたように、GPU36に複数のインスタンスを並列に実行させるオブジェクトコード30を生成し得る。
【００８１】
いくつかの例では、必ずしもすべてとは限らないが、バックエンドコンパイラ34は、オンザフライでオブジェクトコード30を生成し得る。たとえば、プロセッサ32は、オブジェクトコード30が生成されたアプリケーションを利用するプログラム(たとえば、ビデオゲームは、本開示で説明する技法に従ってオブジェクトコードが生成された頂点シェーダを利用する)を実行し得る。この例では、プロセッサ32は、バックエンドコンパイラ34を介して、プログラムの実行中に中間コード22をコンパイルする(たとえば、ビデオゲームの実行中に頂点シェーダのためのオブジェクトコードをコンパイルする)ことができる。
【００８２】
図2は、本開示で説明する1つまたは複数の例示的な技法を実施し得るデバイス14の一例を示すブロック図である。図2に示すように、また図1の例と同様に、デバイス14は、プロセッサ32(たとえば、中央処理装置(CPU))、GPU36、およびメモリユニット28を含み得る。デバイス14は、図1に示すものに加えて構成要素を含み得る。
【００８３】
その上、いくつかの例では、プロセッサ32およびGPU36は、単一の回路パッケージ内に格納された共通の集積回路として形成される(たとえば、共通のプロセッサとして形成される)ことがある。ただし、本開示の態様は、そのように限定されず、プロセッサ32およびGPU36の一方または両方は、別個の回路パッケージに格納された別個の集積回路であり得る。一般に、デバイス14は、本開示で説明する例示的な技法のうちの1つまたは複数を実行するように構成されたプロセッサ32を含む集積回路を含むものと見なされ得る。
【００８４】
メモリユニット28は、コンピュータ可読記憶媒体の一例であってもよい。たとえば、オブジェクトコード30および中間コード22を記憶することに加えて、メモリユニット28は、プロセッサ32およびGPU36に本開示において各々に起因する機能を実行させる命令を記憶し得る。メモリユニット28は、1つまたは複数のプロセッサ(たとえば、プロセッサ32またはGPU36)に様々な機能を実行させる命令を含むコンピュータ可読記憶媒体と見なされ得る。
【００８５】
メモリユニット28は、いくつかの例において、非一時的記憶媒体と見なされ得る。「非一時的」という用語は、記憶媒体が搬送波または伝搬信号において具現化されないことを示し得る。しかしながら、「非一時的」という用語は、メモリユニット28が移動不可能であることを意味すると解釈されるべきではない。一例として、メモリユニット28は、デバイス14から取り外され、別のデバイスに移動され得る。別の例として、メモリユニット28に実質的に類似する記憶デバイスが、デバイス14に挿入され得る。いくつかの例では、非一時的記憶媒体は、経時的に変化し得るデータを(たとえば、RAM内に)記憶することができる。
【００８６】
GPU36は、シェーダプロセッサ42および固定機能パイプライン44を含み得る。シェーダコアと呼ばれることがあるシェーダプロセッサ42は、プログラムのオブジェクトコードが実行されるGPU36のコアであり得る。たとえば、オブジェクトコード30は、シェーダプロセッサ42上で実行されてよく、いくつかの例では、シェーダプロセッサ42の複数の処理要素を介して並列方式で(たとえば、同時に)実行されてよい。一例として、バックエンドコンパイラ34が中間コード22をどのようにコンパイルすべきかを判断するためにプロセッサ32が利用し得るのは、シェーダプロセッサ42における処理要素の数である(たとえば、シェーダプロセッサ42の処理要素の数により、GPU36のSIMDレーンの数が決まる)。
【００８７】
固定機能パイプライン44は、固定機能を実行するハードウェアユニットを含み得る。固定機能パイプライン44およびシェーダプロセッサ42が合わさって、GPU36により処理するためのパイプラインを形成する。たとえば、処理の後、シェーダプロセッサ42は、固定機能パイプライン44の1つまたは複数のユニットがさらなる処理のために受信するデータを生成することができ、GPU36が最後の生成されたデータをメモリユニット28に出力するまで、シェーダプロセッサ42が(たとえば、頂点シェーダ、ジオメトリシェーダ、またはピクセルシェーダを介して)何らかの処理を間欠的に実行するのに伴って、各ユニットが次のユニットに出力する。
【００８８】
プロセッサ32は、メモリユニット28に記憶された、プログラム40などの1つまたは複数のプログラムを実行し得る。プログラム40の例としては、限定はしないが、ウェブブラウザ、ユーザインターフェース、電子メールアプリケーション、表計算アプリケーション、ワードプロセシングアプリケーション、グラフィックスオーサリングアプリケーション、ビデオゲーム、または表示用の閲覧可能オブジェクトを生成する他のアプリケーションがある。たとえば、プログラム40は、実行されると、ディスプレイ上に表示されるグラフィカルコンテンツを出力するビデオゲームであり得る。一例として、プログラム40は、GPU36またはプロセッサ32にオブジェクトコード30を実行させ得る(たとえば、プログラム40はビデオゲームであり、オブジェクトコード30は頂点シェーダのためのものであり、ビデオゲームはGPU36に頂点シェーダを実行させる)。いくつかのそのような例では、プロセッサ32は、プログラム40によって利用されるオブジェクトコード30を生成するために、中間コード22をオンザフライで(たとえば、プログラム40の実行中に)コンパイルし得る。
【００８９】
上記で説明したように、プロセッサ32は、中間コード22を受信し、バックエンドコンパイラ34を実行し得る。バックエンドコンパイラ34を介して、プロセッサ32は、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコード30を生成するために中間言語(たとえば、SPIR)による中間コード22をコンパイルし、最終的な実行のためにオブジェクトコード30を記憶し得る。プロセッサ32は、メモリユニット28にオブジェクトコード30を記憶し得る。メモリユニット28はプロセッサ32およびGPU36の外部にあるものとして示されているが、いくつかの例では、メモリユニット28は部分的に、プロセッサ32およびGPU36の中にあり得る。たとえば、プロセッサ32およびGPU36はそれぞれ、キャッシュメモリなどの内部メモリを(個別にまたは一緒に)含む。この内部メモリは、一般的な用語のメモリユニットの一部(たとえば、メモリユニット28の一部)であると見なされ得る。
【００９０】
いくつかの例では、バックエンドコンパイラ34は、アプリケーションが実行されるべきプロセッサタイプに基づいて、中間コード22をコンパイルし得る。たとえば、オブジェクトコード30がプロセッサ32上で実行されるべきである場合、バックエンドコンパイラ34は、プロセッサ32上での実行のために(たとえば、順次実行のために)オブジェクトコード30を最適化するために、中間コード22における情報を利用し得る。オブジェクトコード30がGPU36上で実行されるべきである場合、バックエンドコンパイラ34は、GPU36上での実行のために(たとえば、並列実行のために)オブジェクトコード30を最適化するために、中間コード22における情報を利用し得る。
【００９１】
上記で説明したように、中間コード22は、フロントエンドコンパイラ26を実行しているプロセッサ24を含むホストデバイス12によって生成され得る。しかしながら、例は必ずしもそのように限定されるとは限らない。いくつかの例では、プロセッサ32は、フロントエンドコンパイラ26およびバックエンドコンパイラ34を実行すること、またはフロントエンドコンパイラとバックエンドコンパイラの両方の機能を実行する単一のコンパイラを実行することがある。プロセッサ32は、そのような例では、中間コード22を生成し得る。
【００９２】
また図2に示すように、プロセッサ32は、グラフィックスドライバ38を実行するように構成され得る。グラフィックスドライバ38の機能は、ハードウェアもしくはファームウェアまたはハードウェア、ソフトウェアおよびファームウェアの組合せによって実行され得る。容易にするために、グラフィックスドライバ38は、プロセッサ32上で実行されているソフトウェアとして説明される。
【００９３】
グラフィックスドライバ38は、特定の規格に従って設計され得る。プロセッサ32上で実行されているグラフィックスドライバ38は、アプリケーションプログラミングインターフェース(API)(たとえば、OpenGLもしくはOpenCLまたは両方の組合せ)を実装するように構成され得る。そのような例では、シェーダプログラムまたはカーネル(たとえば、オブジェクトコード30が生成されるアプリケーション)は、グラフィックスドライバ38と同じAPIに従って構成され得る(たとえば、OpenCLによるカーネルおよびOpenGLによるシェーダプログラム)。図示されていないが、メモリユニット28は、プロセッサ32が実行のためにメモリユニット28から取り出すグラフィックスドライバ38のためのコードを記憶し得る。
【００９４】
グラフィックスドライバ38は、プロセッサ32とGPU36との間の通信を実行するように構成され得る。したがって、本開示では、プロセッサ32がGPU36に命令するか、またはさもなければGPU36と通信するものとして説明されるとき、プロセッサ32は、グラフィックスドライバ38を介してそのような機能を実行している。たとえば、GPU36がオブジェクトコード30を実行すべきであるとき、プロセッサ32は、グラフィックスドライバ38を介してGPU36に、オブジェクトコード30を実行する(たとえば、メモリユニット28からオブジェクトコード30を取り出し、オブジェクトコード30の命令を実行する)よう命令し得る。
【００９５】
図3は、図2の例示的なデバイスをさらに詳細に示すブロック図である。たとえば、上記のように、デバイス14の例としては、限定はしないが、モバイルワイヤレス電話、PDA、ビデオディスプレイを含むビデオゲーミングコンソール、モバイルビデオ会議ユニット、ラップトップコンピュータ、デスクトップコンピュータ、テレビセットトップボックスなどがある。
【００９６】
図3に示すように、デバイス14は、ディスプレイ46、プロセッサ32、GPU36、フレームバッファ54を含むメモリユニット28、トランシーバモジュール48、ユーザインターフェース50、およびディスプレイプロセッサ52を含み得る。プロセッサ32、GPU36、およびメモリユニット28は、図1に示すものと実質的に同様または同一であり得る。簡潔さのために、図1または図2に示されなかった構成要素で、図3に示す構成要素だけが、詳細に説明される。
【００９７】
デバイス14は、図3に示すように、明瞭さのために図3には示していない追加のモジュールまたはユニットを含むことができる。たとえば、デバイス14は、デバイス14がモバイルワイヤレス電話である例では電話通信をもたらし、デバイス14がメディアプレーヤである場合にはスピーカーをもたらすために、図3には示していない、スピーカーおよびマイクロフォンを含み得る。さらに、デバイス14に示す様々なモジュールおよびユニットは、デバイス14のすべての例において必要とは限らない。たとえば、ユーザインターフェース50およびディスプレイ46は、デバイス14が、デスクトップコンピュータ、または外部のユーザインターフェースもしくはディスプレイとインターフェースするために装備された他のデバイスである例ではデバイス14の外部にある場合がある。別の例として、ディスプレイ46がモバイルデバイスのタッチセンシティブディスプレイまたはプレゼンスセンシティブ(presence-sensitive)ディスプレイである例では、ユーザインターフェース50はディスプレイ46の一部であり得る。
【００９８】
ディスプレイ46は、液晶ディスプレイ(LCD)、陰極線管(CRT)ディスプレイ、プラズマディスプレイ、タッチセンシティブディスプレイ、プレゼンスセンシティブディスプレイ、または別のタイプのディスプレイデバイスを含み得る。ユーザインターフェース50の例は、限定はしないが、トラックボール、マウス、キーボード、および他のタイプの入力デバイスを含む。ユーザインターフェース50は、タッチスクリーンである場合もあり、ディスプレイ46の一部として組み込まれる場合がある。
【００９９】
トランシーバモジュール48は、(たとえば、図1のチャネル16を介した)デバイス14と別のデバイスまたはネットワークとの間のワイヤレス通信またはワイヤード通信を可能にするための回路を含み得る。トランシーバモジュール48は、1つまたは複数の変調器、復調器、増幅器、アンテナ、およびワイヤード通信またはワイヤレス通信のための他のそのような回路を含み得る。
【０１００】
ディスプレイプロセッサ52は、ディスプレイ46にグラフィックスデータを出力するように構成され得る。たとえば、ディスプレイプロセッサ52は、GPU36によって出力されたグラフィックデータを記憶するフレームバッファ54からデータを受信する。ディスプレイプロセッサ52は、取り出されたデータを出力し、グラフィックスデータを表示するようにディスプレイ46を制御する。
【０１０１】
図4は、本開示の態様を実装するための一例を示すフローチャートである。図4に示す例について、プロセッサ32に関して説明する。たとえば、ユーザデバイス14など、コンパイルするためのデバイスは、メモリユニット28および図4に示す例示的な技法を実行するように構成されたプロセッサ32を含む集積回路を含み得る。
【０１０２】
図4に示す例では、プロセッサ32は、アプリケーションのソースコード20をコンパイルすることから生成された中間コード22を受信し得る(60)。上記のように、また図4に示すように、中間コード22は、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、コンパイルすることから生成された情報を含み、下位レベルサブルーチンは、アプリケーションのソースコードにおいて、下位レベルサブルーチンを識別する上位レベルサブルーチンよりも頻繁に実行されるように規定される。
【０１０３】
いくつかの例では、プロセッサ32は、下位レベルサブルーチンのための中間コードとは別個の上位レベルサブルーチンのための中間コードを受信し得る。たとえば、中間コード22は、すべてのサブルーチンのための中間コードを含む順次コードではないことがある。そうではなく、中間コード22は、第1の関数の形式による上位レベルサブルーチンのための中間コードを含むことができ、第2の関数の形式による下位レベルサブルーチンのための中間コードを含むことができる。第1の関数は、実行する第2の関数を呼び出し得るが、プロセッサ32が第1の関数および第2の関数を受信するとき、プロセッサ32は第1の関数および第2の関数を別個に受信することがある。
【０１０４】
プロセッサ32は、バックエンドコンパイラ34を介して、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する情報に基づいて、オブジェクトコード30を生成するために中間コード22をコンパイルし得る(62)。一例として、バックエンドコンパイラ34は、下位レベルサブルーチンの実行を求める上位レベルサブルーチンにおける関数呼出しを識別し得る。バックエンドコンパイラ34は、上位レベルサブルーチンから、下位レベルサブルーチンが実行されるべき回数を決定し、下位レベルサブルーチンが実行されるべき回数の決定に基づいて、オブジェクトコード30を生成し得る。
【０１０５】
プロセッサ32は、オブジェクトコード30を記憶し得る(64)。たとえば、プロセッサ32は、最終的な実行のためにメモリユニット28にオブジェクトコード30を書き込み得る。たとえば、プロセッサ32またはGPU36のうちの一方がオブジェクトコード30を実行し得る(66)。
【０１０６】
一例として、プロセッサ32は、プロセッサ32がオブジェクトコード30を実行すべきであると判断することがあり、プロセッサ32は、オブジェクトコード30を実行し得る。この例では、コンパイルするために、プロセッサ32は、バックエンドコンパイラ34を介して、プロセッサ32に下位レベルサブルーチンの複数のインスタンスを順次実行させるオブジェクトコード30を生成するために、中間コード22をコンパイルし得る。別の例として、プロセッサ32は、GPU36がオブジェクトコード30を実行すべきであると判断することがあり、プロセッサ32は、グラフィックスドライバ38を介してGPU36に、シェーダプロセッサ42の1つまたは複数の処理要素上でオブジェクトコード30を実行するよう命令し得る。この例では、コンパイルするために、プロセッサ32は、バックエンドコンパイラ34を介して、GPU36に下位レベルサブルーチンの複数のインスタンスを並列に実行させるオブジェクトコード30を生成するために、中間コード22をコンパイルし得る。
【０１０７】
図5は、本開示の態様を実装するための別の例を示すフローチャートである。図示のように、プロセッサ24は、フロントエンドコンパイラ26を介して、上位レベルサブルーチンにおける下位レベルサブルーチンを識別する階層構造の情報を含む中間コード22を生成するために、ソースコード20をコンパイルし得る(70)。次いでプロセッサ24は、ユーザデバイス14による最終的な受信および実行のために、チャネル16を介して中間コード22を出力し得る(72)。たとえば、ユーザデバイス14は、下位レベルサブルーチンおよび上位レベルサブルーチンを識別する情報に基づいてオブジェクトコードを生成するために中間コードをコンパイルすること、オブジェクトコードを記憶すること、ならびにオブジェクトコードを実行することなど、図4に関して上記で説明した例示的な技法を実行し得る。
【０１０８】
上記で説明したように、中間コード22を生成する際、フロントエンドコンパイラ26は、ソースコード20における識別情報に基づいて、上位レベルサブルーチンと下位レベルサブルーチンとを区別し得る。フロントエンドコンパイラ26は、上位レベルサブルーチンにおける下位レベルサブルーチンを構文解析し、フロントエンドコンパイラ26は、上位レベルサブルーチンにおける下位レベルサブルーチンを、下位レベルサブルーチンへの関数呼出しに置き換える。これらの関数呼出しは、下位レベルサブルーチンの実行、およびオブジェクトコード30を生成するためにバックエンドコンパイラ34が依拠するものを求める。
【０１０９】
1つまたは複数の例では、説明した機能は、ハードウェア、ソフトウェア、ファームウェア、またはそれらの任意の組合せにおいて実装され得る。ソフトウェアにおいて実装される場合、機能は、非一時的コンピュータ可読媒体を備える製造品に1つまたは複数の命令またはコードとして記憶され得る。コンピュータ可読媒体は、コンピュータデータ記憶媒体、またはコンピュータプログラムを1つの場所から別の場所に転送することを容易にする任意の媒体を含む通信媒体を含むことができる。データ記憶デバイスは、本開示で説明する技法の実装のための命令、コードおよび/またはデータ構造を取り出すために、1つもしくは複数のコンピュータまたは1つもしくは複数のプロセッサによってアクセスされ得る、任意の利用可能な媒体であり得る。限定ではなく例として、そのようなコンピュータ可読媒体は、RAM、ROM、EEPROM、CD-ROMもしくは他の光ディスクストレージ、磁気ディスクストレージもしくは他の磁気ストレージデバイス、フラッシュメモリ、または所望のプログラムコードを命令もしくはデータ構造の形で搬送もしくは記憶するために使用され得るとともにコンピュータによってアクセスされ得る任意の他の媒体を含み得る。本明細書において使用される場合、ディスク(disk)およびディスク(disc)は、コンパクトディスク(disc)(CD)、レーザディスク(disc)、光ディスク(disc)、デジタル多用途ディスク(disc)(DVD)、フロッピーディスク(disk)、およびブルーレイディスク(disc)を含み、ディスク(disk)は、通常、データを磁気的に再生するが、ディスク(disc)は、レーザを用いて光学的にデータを再生する。上記の組合せはまた、コンピュータ可読媒体の範囲内に含まれるものとする。
【０１１０】
コードは、1つまたは複数のDSP、汎用マイクロプロセッサ、ASIC、FPGA、または他の等価の集積論理回路もしくはディスクリート論理回路などの、1つまたは複数のプロセッサによって実行され得る。加えて、いくつかの態様では、本明細書で説明する機能は、専用のハードウェアモジュールおよび/またはソフトウェアモジュール内で提供され得る。また、本技法は、1つまたは複数の回路または論理要素において完全に実装され得る。
【０１１１】
本開示の技法は、ワイヤレスハンドセット、集積回路(IC)、またはICのセット(たとえば、チップセット)を含む、様々なデバイスまたは装置において実装され得る。本開示では、開示される技法を実行するように構成されたデバイスの機能的側面を強調するために、様々な構成要素、モジュール、またはユニットが説明されているが、それらは、必ずしも異なるハードウェアユニットによる実現を必要とするとは限らない。むしろ、上記で説明したように、様々なユニットが、適切なソフトウェアおよび/またはファームウェアとともに、上記で説明した1つまたは複数のプロセッサを含めて、ハードウェアユニットにおいて組み合わせられ得るか、または相互動作可能なハードウェアユニットの集合によって与えられ得る。
【０１１２】
様々な例について説明した。これらおよび他の例は、以下の特許請求の範囲内に入る。
【符号の説明】
【０１１３】
10  システム
12  ホストデバイス
14  ユーザデバイス、デバイス
16  チャネル
18  メモリユニット
20  ソースコード
22  中間コード
24  プロセッサ
26  フロントエンドコンパイラ
28  メモリユニット
30  オブジェクトコード
32  プロセッサ
34  バックエンドコンパイラ
36  グラフィックス処理ユニット(GPU)
38  グラフィックスドライバ
40  プログラム
42  シェーダプロセッサ
44  固定機能パイプライン
46  ディスプレイ
48  トランシーバモジュール
50  ユーザインターフェース
52  ディスプレイプロセッサ
54  フレームバッファ
【図１】
2018517986.tif 000003
【図２】
2018517986.tif 000004
【図３】
2018517986.tif 000005
【図４】
2018517986.tif 000006
【図５】
2018517986.tif 000007
【手続補正書】
【提出日】平成29年12月18日(2017.12.18)
【手続補正１】
【補正対象書類名】特許請求の範囲
【補正対象項目名】全文
【補正方法】変更
【補正の内容】
【特許請求の範囲】
【請求項１】
コンパイルする方法であって、
プロセッサにより、アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するステップであって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定され、前記階層構造を識別する生成された前記情報は、前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを含み、前記中間コードを受信するステップは、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信するステップを含む、ステップと、
前記プロセッサにより、前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルするステップであって、コンパイルするステップは、前記オブジェクトコードが前記プロセッサ上で実行されるべきか、グラフィックス処理ユニット(GPU)上で実行されるべきか、それともデジタル信号プロセッサ(DSP)上で実行されるべきかの判断、および前記下位レベルサブルーチンが実行されるべき回数に基づいて、前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しを構成するステップを含む、ステップと、
前記オブジェクトコードを記憶するステップと
を含む方法。
【請求項２】
前記プロセッサが前記オブジェクトコードを実行すべきであると判断するステップと、
前記プロセッサにより、前記オブジェクトコードを実行するステップと
をさらに含み、
構成するステップは、前記プロセッサに前記下位レベルサブルーチンの複数のインスタンスを順次実行させる前記上位レベルサブルーチンにおける前記関数呼出しを構成するステップを含む、請求項1に記載の方法。
【請求項３】
前記GPUまたは前記DSPが前記オブジェクトコードを実行すべきであると判断するステップと、
前記GPUまたは前記DSPに前記オブジェクトコードを実行するよう命令するステップと
をさらに含み、
構成するステップは、前記GPUまたは前記DSPに前記下位レベルサブルーチンの複数のインスタンスを並列に実行させる前記上位レベルサブルーチンにおける前記関数呼出しを構成するステップを含む、請求項1に記載の方法。
【請求項４】
コンパイルするステップは、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別するステップと、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき前記回数を決定するステップと、
前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しの前記構成に基づいて、前記オブジェクトコードを生成するステップと
をさらに含む、請求項1に記載の方法。
【請求項５】
前記プロセッサにより、前記アプリケーションを利用するプログラムを実行するステップ
をさらに含み、
コンパイルするステップは、前記プログラムの実行中に前記中間コードをコンパイルするステップを含む、請求項1に記載の方法。
【請求項６】
コンパイルするためのデバイスであって、
メモリユニットと、
プロセッサを含む集積回路と
を含み、前記プロセッサは、
アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信することであって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定され、前記階層構造を識別する生成された前記情報は、前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを含み、前記中間コードを受信するために、前記プロセッサは、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信するように構成される、受信することと、
前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルすることであって、コンパイルするために、前記プロセッサは、前記オブジェクトコードが前記プロセッサ上で実行されるべきか、グラフィックス処理ユニット(GPU)上で実行されるべきか、それともデジタル信号プロセッサ(DSP)上で実行されるべきかの判断、および前記下位レベルサブルーチンが実行されるべき回数に基づいて、前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しを構成するように構成される、コンパイルすることと、
前記メモリユニットに前記オブジェクトコードを記憶することと
を行うように構成される、デバイス。
【請求項７】
前記プロセッサは、
前記プロセッサが前記オブジェクトコードを実行すべきであると判断することと、
前記オブジェクトコードを実行することと
を行うように構成され、
構成するために、前記プロセッサは、前記プロセッサに前記下位レベルサブルーチンの複数のインスタンスを順次実行させる前記上位レベルサブルーチンにおける前記関数呼出しを構成するように構成される、請求項6に記載のデバイス。
【請求項８】
前記GPUまたは前記DSP
をさらに含み、
前記プロセッサは、
前記GPUまたは前記DSPが前記オブジェクトコードを実行すべきであると判断することと、
前記GPUまたは前記DSPに前記オブジェクトコードを実行するよう命令することと
を行うように構成され、
構成するために、前記プロセッサは、前記GPUまたは前記DSPに前記下位レベルサブルーチンの複数のインスタンスを並列に実行させる前記上位レベルサブルーチンにおける前記関数呼出しを構成するように構成される、請求項6に記載のデバイス。
【請求項９】
コンパイルするために、前記プロセッサは、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別することと、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき前記回数を決定することと、
前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しの前記構成に基づいて、前記オブジェクトコードを生成することと
を行うように構成される、請求項6に記載のデバイス。
【請求項１０】
前記プロセッサは、
前記アプリケーションを利用するプログラムを実行するように構成され、
コンパイルするために、前記プロセッサは、前記プログラムの実行中に前記中間コードをコンパイルするように構成される、請求項6に記載のデバイス。
【請求項１１】
ワイヤレス通信デバイスを含む、請求項6に記載のデバイス。
【請求項１２】
コンパイルするためのデバイスであって、
アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信するための手段であって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定され、前記階層構造を識別する生成された前記情報は、前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを含み、前記中間コードを受信するための前記手段は、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信するための手段を含む、手段と、
前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルするための手段であって、コンパイルするための前記手段は、前記オブジェクトコードがプロセッサ上で実行されるべきか、グラフィックス処理ユニット(GPU)上で実行されるべきか、それともデジタル信号プロセッサ(DSP)上で実行されるべきかの判断、および前記下位レベルサブルーチンが実行されるべき回数に基づいて、前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しを構成するための手段を含む、手段と、
前記オブジェクトコードを記憶するための手段と
を含むデバイス。
【請求項１３】
コンパイルするための前記手段を含むプロセッサが前記オブジェクトコードを実行すべきであると判断するための手段と、
前記オブジェクトコードを実行するための手段と
をさらに含み、
構成するための前記手段は、前記プロセッサに前記下位レベルサブルーチンの複数のインスタンスを順次実行させる前記上位レベルサブルーチンにおける前記関数呼出しを構成するための手段を含む、請求項12に記載のデバイス。
【請求項１４】
前記GPUまたは前記DSPが前記オブジェクトコードを実行すべきであると判断するための手段と、
前記GPUまたは前記DSPに前記オブジェクトコードを実行するよう命令するための手段と
をさらに含み、
構成するための前記手段は、前記GPUまたは前記DSPに前記下位レベルサブルーチンの複数のインスタンスを並列に実行させる前記上位レベルサブルーチンにおける前記関数呼出しを構成するための手段を含む、請求項12に記載のデバイス。
【請求項１５】
コンパイルするための前記手段は、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別するための手段と、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき前記回数を決定するための手段と、
前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しの前記構成に基づいて、前記オブジェクトコードを生成するための手段と
を含む、請求項12に記載のデバイス。
【請求項１６】
前記アプリケーションを利用するプログラムを実行するための手段をさらに含み、
コンパイルするための前記手段は、前記プログラムの実行中に前記中間コードをコンパイルするための手段を含む、請求項12に記載のデバイス。
【請求項１７】
実行されると、1つまたは複数のプロセッサに、
アプリケーションのソースコードをコンパイルすることから生成された中間コードを受信させる命令であって、前記中間コードは、上位レベルサブルーチンにおける下位レベルサブルーチンの階層構造を識別する、前記コンパイルすることから生成された情報を含み、前記下位レベルサブルーチンは、前記アプリケーションの前記ソースコードにおいて、前記下位レベルサブルーチンを識別する前記上位レベルサブルーチンよりも頻繁に実行されるように規定され、前記階層構造を識別する生成された前記情報は、前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを含み、前記1つまたは複数のプロセッサに前記中間コードを受信させる前記命令は、前記1つまたは複数のプロセッサに、前記下位レベルサブルーチンのための中間コードとは別個の前記上位レベルサブルーチンのための中間コードを受信させる命令を含む、命令と、
前記上位レベルサブルーチンにおける前記下位レベルサブルーチンを識別する前記情報に基づいて、オブジェクトコードを生成するために前記中間コードをコンパイルさせる命令であって、前記1つまたは複数のプロセッサにコンパイルさせる前記命令は、前記1つまたは複数のプロセッサに、前記オブジェクトコードが前記1つまたは複数のプロセッサ上で実行されるべきか、グラフィックス処理ユニット(GPU)上で実行されるべきか、それともデジタル信号プロセッサ(DSP)上で実行されるべきかの判断、および前記下位レベルサブルーチンが実行されるべき回数に基づいて、前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しを構成させる命令を含む、命令と、
前記オブジェクトコードを記憶させる命令と
を含む非一時的コンピュータ可読記憶媒体。
【請求項１８】
前記1つまたは複数のプロセッサにコンパイルさせる前記命令は、前記1つまたは複数のプロセッサに、
前記下位レベルサブルーチンの実行を求める前記上位レベルサブルーチンにおける関数呼出しを識別させる命令と、
前記上位レベルサブルーチンから、前記下位レベルサブルーチンが実行されるべき前記回数を決定させる命令と、
前記下位レベルサブルーチンに対する前記上位レベルサブルーチンにおける前記関数呼出しの構成に基づいて、前記オブジェクトコードを生成させる命令と
を含む、請求項17に記載のコンピュータ可読記憶媒体。
【国際調査報告】
2018517986.tif 000008
2018517986.tif 000009
2018517986.tif 000010