本發(fā)明涉及編譯內(nèi)核的源代碼，且更具體地說涉及用于編譯用于存儲器存取的內(nèi)核的源代碼的技術(shù)。

背景技術(shù)：

已存在朝向所謂的異構(gòu)計(jì)算架構(gòu)的趨勢。在異構(gòu)計(jì)算架構(gòu)中，稱為內(nèi)核的程序可使用框架來進(jìn)行編譯，使得例如CPU(中央處理單元)、GPU(圖形處理單元)、FPGA(現(xiàn)場可編程門陣列)等多種不同類型的處理器可執(zhí)行所述內(nèi)核。最近的支持異構(gòu)計(jì)算的框架包含OpenCL框架以及DirectCompute框架。

技術(shù)實(shí)現(xiàn)要素：

本發(fā)明描述用于檢測內(nèi)核中的存儲器參考的存儲器混疊和存儲器重疊以便產(chǎn)生用于編譯優(yōu)化的元數(shù)據(jù)的技術(shù)。為了執(zhí)行本發(fā)明的技術(shù)，例如及時(shí)編譯器(JIT)等編譯器將程序的源代碼(也被稱作“內(nèi)核”)編譯為二進(jìn)制文件。執(zhí)行編譯器的編譯處理器可在運(yùn)行時(shí)間(當(dāng)編譯處理器產(chǎn)生執(zhí)行內(nèi)核所需要的自變量時(shí))使用例如OpenCL等異構(gòu)計(jì)算框架編譯所述內(nèi)核。在本發(fā)明中描述的技術(shù)中，驅(qū)動程序分析在緩沖器中一起傳遞的將傳遞到將執(zhí)行內(nèi)核的目標(biāo)處理器的自變量，而不是指示目標(biāo)處理器使用所產(chǎn)生自變量執(zhí)行內(nèi)核。基于所述分析，驅(qū)動程序/運(yùn)行時(shí)間產(chǎn)生指示第一存儲器參考與第二存儲器參考之間的關(guān)系(例如，第一存儲器參考和第二存儲器參考的存儲器區(qū)是否重疊、重疊達(dá)何種范圍等)的元數(shù)據(jù)。

如果存儲器區(qū)是不相同的，那么編譯處理器可使用編譯器基于元數(shù)據(jù)且使用例如循環(huán)展開等更積極的編譯技術(shù)而重新編譯內(nèi)核。驅(qū)動程序也可以能夠確定內(nèi)核的存儲器存取重疊達(dá)何種范圍，且可基于存儲器重疊量使用更積極的技術(shù)重新編譯內(nèi)核。以此方式，本發(fā)明的技術(shù)可增加使用及時(shí)編譯器編譯的內(nèi)核的執(zhí)行性能。

在一個(gè)實(shí)例中，本發(fā)明描述一種方法，其包括：通過由在編譯處理器上執(zhí)行的編譯器和運(yùn)行時(shí)間組成的群組中的至少一者產(chǎn)生用于執(zhí)行經(jīng)編譯內(nèi)核的二進(jìn)制代碼的自變量；通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者確定對所述內(nèi)核自變量的第一存儲器區(qū)的第一存儲器參考和對所述內(nèi)核自變量的第二存儲器區(qū)的第二存儲器參考是否參考同一存儲器區(qū)；通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者基于所述確定而產(chǎn)生與所述第一存儲器參考和所述第二存儲器參考相關(guān)聯(lián)的元數(shù)據(jù)。所述元數(shù)據(jù)可指示所述第一存儲器區(qū)與所述第二存儲器區(qū)之間的關(guān)系。所述方法進(jìn)一步包含響應(yīng)于通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者確定所述內(nèi)核的第一和第二存儲器參考并不參考所述同一存儲器區(qū)而進(jìn)行以下操作：通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者致使編譯器基于所述元數(shù)據(jù)而重新編譯所述內(nèi)核；以及通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者指示目標(biāo)處理器執(zhí)行所述經(jīng)重新編譯的內(nèi)核。

在另一實(shí)例中，本發(fā)明描述一種裝置，其包含存儲器和編譯處理器，所述編譯處理器經(jīng)配置以：通過由在所述編譯處理器上執(zhí)行的編譯器和運(yùn)行時(shí)間組成的群組中的至少一者產(chǎn)生用于執(zhí)行經(jīng)編譯內(nèi)核的二進(jìn)制代碼的自變量；通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者確定對所述內(nèi)核自變量的第一存儲器區(qū)的第一存儲器參考和對所述內(nèi)核自變量的第二存儲器區(qū)的第二存儲器參考是否參考同一存儲器區(qū)；通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者基于所述確定而產(chǎn)生與所述第一存儲器參考和所述第二存儲器參考相關(guān)聯(lián)的元數(shù)據(jù)。所述元數(shù)據(jù)可指示所述第一存儲器區(qū)與所述第二存儲器區(qū)之間的關(guān)系，且響應(yīng)于通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者確定所述內(nèi)核的第一和第二存儲器參考并不參考所述同一存儲器區(qū)，所述編譯處理器進(jìn)一步經(jīng)配置以：通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者致使編譯器基于所述元數(shù)據(jù)而重新編譯所述內(nèi)核；以及通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者指示目標(biāo)處理器執(zhí)行所述經(jīng)重新編譯的內(nèi)核。

在另一實(shí)例中，本發(fā)明描述一種存儲指令的非暫時(shí)性計(jì)算機(jī)可讀存儲媒體，所述指令當(dāng)執(zhí)行時(shí)致使編譯處理器進(jìn)行以下操作：通過由在所述編譯處理器上執(zhí)行的編譯器和運(yùn)行時(shí)間組成的群組中的至少一者產(chǎn)生用于執(zhí)行經(jīng)編譯內(nèi)核的二進(jìn)制代碼的自變量；通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者確定對所述內(nèi)核自變量的第一存儲器區(qū)的第一存儲器參考和對所述內(nèi)核自變量的第二存儲器區(qū)的第二存儲器參考是否參考同一存儲器區(qū)；通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者基于所述確定而產(chǎn)生與所述第一存儲器參考和所述第二存儲器參考相關(guān)聯(lián)的元數(shù)據(jù)。所述元數(shù)據(jù)指示第一存儲器區(qū)與第二存儲器區(qū)之間的關(guān)系，且

響應(yīng)于通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者確定所述內(nèi)核的第一和第二存儲器參考并不參考所述同一存儲器區(qū)，所述編譯處理器可進(jìn)一步經(jīng)配置以執(zhí)行致使所述編譯處理器進(jìn)行以下操作的指令：通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者致使編譯器基于所述元數(shù)據(jù)而重新編譯所述內(nèi)核；以及通過由在所述編譯處理器上執(zhí)行的所述編譯器和所述運(yùn)行時(shí)間組成的所述群組中的所述至少一者指示目標(biāo)處理器執(zhí)行所述經(jīng)重新編譯的內(nèi)核。

在附圖和下文描述中陳述本發(fā)明的一或多個(gè)實(shí)例的細(xì)節(jié)。本發(fā)明的其它特征、目標(biāo)和優(yōu)點(diǎn)將從所述描述和圖式以及權(quán)利要求書而顯而易見。

附圖說明

圖1是說明根據(jù)本發(fā)明的技術(shù)的支持混疊分析以輔助編譯優(yōu)化的實(shí)例計(jì)算裝置的框圖。

圖2是說明根據(jù)本發(fā)明的技術(shù)的可執(zhí)行內(nèi)核的處理器的一或多個(gè)著色器核心的多個(gè)處理元件的概念圖。

圖3A是說明根據(jù)本發(fā)明的技術(shù)的包含當(dāng)執(zhí)行時(shí)可造成混疊的代碼的內(nèi)核代碼的概念圖。

圖3B是說明根據(jù)本發(fā)明的技術(shù)配置的編譯器可能夠檢測的混疊的實(shí)例的概念圖。

圖3C是說明根據(jù)本發(fā)明的技術(shù)配置的編譯器可能夠檢測的不重疊存儲器參考的實(shí)例的概念圖。

圖3D是說明根據(jù)本發(fā)明的技術(shù)配置的驅(qū)動程序/運(yùn)行時(shí)間可檢測的重疊存儲器參考的概念圖。

圖4A是說明根據(jù)本發(fā)明的技術(shù)的循環(huán)展開的概念圖。

圖4B是說明根據(jù)本發(fā)明的技術(shù)的代碼重排序的概念圖。

圖4C是說明根據(jù)本發(fā)明的技術(shù)的代碼向量化的概念圖。

圖5是根據(jù)本發(fā)明的技術(shù)用于產(chǎn)生編譯器元數(shù)據(jù)以輔助編譯器優(yōu)化的實(shí)例方法的流程圖。

具體實(shí)施方式

如上文簡要地描述，當(dāng)前在開發(fā)各種異構(gòu)計(jì)算框架。異構(gòu)計(jì)算框架的一些實(shí)例包含當(dāng)前由Khronos集團(tuán)開發(fā)的OpenCL^TM框架以及當(dāng)前由開發(fā)的DirectCompute框架。異構(gòu)計(jì)算框架允許單個(gè)程序或“內(nèi)核”在多種不同處理器上執(zhí)行，例如CPU(中央處理單元)、GPU(圖形處理單元)、FPGA(現(xiàn)場可編程門陣列)、DSP(數(shù)字信號處理器)等。

為了準(zhǔn)備內(nèi)核用于執(zhí)行，在本發(fā)明中稱為編譯處理器的處理器編譯內(nèi)核源代碼以產(chǎn)生將由目標(biāo)處理器執(zhí)行的二進(jìn)制代碼。目標(biāo)處理器可為同一處理器或不同于目標(biāo)處理器。編譯處理器使用的編譯器的一個(gè)實(shí)例稱為及時(shí)編譯(JIT)編譯器。JIT編譯器在執(zhí)行時(shí)間(也被稱作運(yùn)行時(shí)間)編譯源代碼，而不是在執(zhí)行之前編譯(有時(shí)稱為“提前”編譯)或根本不需要先前編譯的指令(稱為“解譯”)。

一旦內(nèi)核已經(jīng)編譯，編譯處理器便經(jīng)由驅(qū)動程序和運(yùn)行時(shí)間將內(nèi)核的經(jīng)編譯二進(jìn)制代碼轉(zhuǎn)移到目標(biāo)處理器。內(nèi)核還在運(yùn)行時(shí)間接受自變量集合以用于在目標(biāo)處理器上執(zhí)行所述內(nèi)核，編譯處理器也將所述自變量集合轉(zhuǎn)移到目標(biāo)處理器。內(nèi)核自變量包括緩沖器，即為自變量分配的存儲器的區(qū)域。在大多數(shù)情況下，內(nèi)核包含對自變量操作(即，對其讀取或從其寫入)的代碼區(qū)段。以此方式，自變量包括內(nèi)核可對其操作的用于內(nèi)核的數(shù)據(jù)集。在將內(nèi)核轉(zhuǎn)移到目標(biāo)處理器之后，編譯處理器的驅(qū)動程序/運(yùn)行時(shí)間執(zhí)行函數(shù)調(diào)用，其在一些實(shí)例中在運(yùn)行時(shí)間將自變量提供到內(nèi)核。一旦內(nèi)核已接收自變量，目標(biāo)處理器便可開始內(nèi)核的執(zhí)行。

在許多情況下，內(nèi)核包含代碼段，例如循環(huán)，其執(zhí)行直到目標(biāo)處理器確定一些布爾型條件已滿足或達(dá)到某一次數(shù)的迭代。編譯器可能夠采用各種技術(shù)以改善執(zhí)行循環(huán)代碼區(qū)段的性能，例如循環(huán)展開以及其它技術(shù)，例如代碼重排序，和/或可改善循環(huán)和非循環(huán)代碼區(qū)段的執(zhí)行的向量化。

循環(huán)展開是一種優(yōu)化過程，編譯器通過所述過程擴(kuò)展循環(huán)的若干迭代以減少或消除控制循環(huán)的指令，例如算術(shù)操作、循環(huán)結(jié)束測試，和/或當(dāng)執(zhí)行循環(huán)時(shí)改善高速緩沖存儲器執(zhí)行性能。代碼重排序是另一優(yōu)化，編譯器可使用其來分組一系列類似指令(例如，一起加載或存儲)。代碼重排序可當(dāng)在某些情況下執(zhí)行循環(huán)代碼區(qū)段時(shí)改善高速緩沖存儲器性能。舉例來說，當(dāng)一起聚結(jié)許多加載指令(例如，在循環(huán)主體內(nèi))可改善具有高速緩沖存儲器線寬的系統(tǒng)(下文更詳細(xì)地論述)上的性能時(shí)，代碼重排序可改善性能，所述高速緩沖存儲器線寬是在標(biāo)量指令中使用的操作數(shù)的大小的倍數(shù)。然而，如果編譯器在編譯之前確定加載/存儲緩沖器并不彼此混疊，那么編譯器聚結(jié)負(fù)載可僅為安全的。否則，由于經(jīng)重排序的加載/存儲指令，可發(fā)生數(shù)據(jù)損壞。

向量化是另一優(yōu)化過程，編譯器通過所述過程可將包含若干標(biāo)量操作的源代碼轉(zhuǎn)換為向量指令，所述若干標(biāo)量操作中的每一者每次處理單個(gè)對的操作數(shù)，所述向量指令每次處理多對操作數(shù)上的一個(gè)操作。向量化是一種形式的并行度，其可改善相對于同一代碼的標(biāo)量實(shí)施方案的性能。下文更詳細(xì)地描述循環(huán)展開、代碼重排序以及向量化。

內(nèi)核的代碼區(qū)段可含有存儲器參考，也被稱作“指標(biāo)”，其可參考自變量的存儲器區(qū)域。舉例來說，代碼區(qū)段可包含可參考內(nèi)核自變量的部分的一系列存儲器參考(即，對包含在內(nèi)核自變量中的緩沖器的存儲器參考)。內(nèi)核可從自變量緩沖器讀取值，且還可將數(shù)據(jù)寫入到自變量緩沖器。

在一些情況下，不同存儲器參考(例如具有不同名稱的指針變量)可參考存儲器中的同一數(shù)據(jù)位置。其中不同符號參考參考同一存儲器區(qū)的情形稱為“混疊”。編譯器可嘗試在編譯時(shí)間使用靜態(tài)分析或其它技術(shù)檢測混疊。然而，當(dāng)循環(huán)代碼區(qū)段中的存儲器參考所參考的數(shù)據(jù)(例如，內(nèi)核自變量)是在運(yùn)行時(shí)間供應(yīng)時(shí)，編譯器通常不能夠檢測循環(huán)中的存儲器參考的混疊。

當(dāng)編譯器不能夠明確地確定存儲器參考是否參考同一存儲器區(qū)(即存儲器參考導(dǎo)致混疊)時(shí)，編譯器可能不能夠?qū)ρh(huán)執(zhí)行優(yōu)化技術(shù)，例如循環(huán)展開和向量化。本發(fā)明的技術(shù)可使得JIT編譯器能夠確定內(nèi)核循環(huán)的存儲器存取是否參考同一存儲器區(qū)。另外，本發(fā)明的技術(shù)使得JIT編譯器能夠產(chǎn)生關(guān)于存儲器參考之間的關(guān)系的元數(shù)據(jù)，且基于所產(chǎn)生元數(shù)據(jù)使用例如向量化和循環(huán)展開等優(yōu)化來重新編譯內(nèi)核。

圖1是說明根據(jù)本發(fā)明的技術(shù)的支持混疊分析以輔助編譯優(yōu)化的實(shí)例計(jì)算裝置的框圖。圖1包含計(jì)算裝置2。計(jì)算裝置2可包括個(gè)人計(jì)算機(jī)、桌上型計(jì)算機(jī)、膝上型計(jì)算機(jī)、計(jì)算機(jī)工作站、平板計(jì)算裝置、視頻游戲平臺或控制臺、無線通信裝置(例如，移動電話、蜂窩式電話、衛(wèi)星電話及/或移動電話手持機(jī))、手持式裝置(例如，便攜式視頻游戲裝置或個(gè)人數(shù)字助理(PDA))、個(gè)人音樂播放器、視頻播放器、顯示裝置、電視、電視機(jī)頂盒、服務(wù)器、中間網(wǎng)絡(luò)裝置、主機(jī)計(jì)算機(jī)，或處理及/或顯示圖形數(shù)據(jù)的任何其它類型的裝置。

如圖1的實(shí)例中所說明，計(jì)算裝置2包含CPU 16、系統(tǒng)存儲器14、圖形處理單元(GPU)12、及時(shí)(JIT)編譯器18以及驅(qū)動程序/運(yùn)行時(shí)間19。CPU 16可執(zhí)行各種類型的應(yīng)用程序。應(yīng)用程序的實(shí)例包含網(wǎng)絡(luò)瀏覽器、電子郵件應(yīng)用程序、電子數(shù)據(jù)表、視頻游戲、產(chǎn)生用于顯示的可觀看對象的應(yīng)用程序及類似應(yīng)用程序。用于執(zhí)行一或多個(gè)應(yīng)用程序的指令可存儲在系統(tǒng)存儲器14內(nèi)。

CPU 16還可執(zhí)行JIT編譯器18。因此，CPU 16可出于實(shí)例的目的稱為“編譯處理器”。JIT編譯器18包括當(dāng)由CPU 16執(zhí)行時(shí)可使用如上文所描述的例如OpenCL或DirectCompute等異構(gòu)計(jì)算框架編譯內(nèi)核的源代碼的編譯器。JIT編譯器18將源代碼編譯為原生代碼或中間代碼(例如，字節(jié)代碼)用于由目標(biāo)處理器執(zhí)行。JIT編譯器18在“運(yùn)行時(shí)間”執(zhí)行編譯，即在執(zhí)行時(shí)而不是在執(zhí)行之前執(zhí)行編譯。JIT編譯器18可當(dāng)使用OpenCL編譯時(shí)使用clBuildProgram()函數(shù)執(zhí)行編譯。另外，JIT編譯器18可經(jīng)配置以分析內(nèi)核20的數(shù)據(jù)存取模式以確定在目標(biāo)處理器GPU 12上執(zhí)行的某些纖維(即線程)的數(shù)據(jù)存取是否是獨(dú)立的以及其它條件是否保持。

驅(qū)動程序/運(yùn)行時(shí)間19還與JIT編譯器18交互以將內(nèi)核源代碼翻譯為二進(jìn)制指令或字節(jié)代碼指令。驅(qū)動程序/運(yùn)行時(shí)間19可使用驅(qū)動程序來執(zhí)行內(nèi)核源代碼指令到用于目標(biāo)處理器(在此實(shí)例中為GPU 12)的原生或目標(biāo)代碼的架構(gòu)特定編譯。舉例來說，驅(qū)動程序/運(yùn)行時(shí)間19可知道可用于目標(biāo)處理器的特定向量指令或執(zhí)行資源，且可以優(yōu)化目標(biāo)處理器上的執(zhí)行性能的方式將源代碼編譯為原生代碼。在一些實(shí)例中，例如如果存在多個(gè)目標(biāo)處理器，例如如果內(nèi)核將在CPU 16和GPU 12上執(zhí)行，那么可存在不同驅(qū)動程序。

內(nèi)核20由目標(biāo)處理器(在此實(shí)例中為GPU 12)能夠執(zhí)行的原生或目標(biāo)代碼構(gòu)成，例如二進(jìn)制指令。JIT編譯器18還可管理GPU 12的運(yùn)行時(shí)間執(zhí)行。CPU 16可將內(nèi)核20傳輸?shù)紾PU 12用于執(zhí)行。CPU 16還可產(chǎn)生自變量26，CPU 16可將其轉(zhuǎn)移到GPU 12用于進(jìn)一步處理。

在分配自變量26之前，CPU 16為自變量26分配自由存儲器緩沖器，其為存儲器的區(qū)。一旦緩沖器已經(jīng)被分配，驅(qū)動程序/運(yùn)行時(shí)間19便在緩沖器中存儲自變量26。自變量26可包括GPU 12能夠處理的多個(gè)數(shù)據(jù)值(例如，整數(shù)、浮點(diǎn)值、對象、值的陣列等)。另外，在內(nèi)核20的執(zhí)行期間，GPU 12可將數(shù)據(jù)寫入到存儲自變量26的緩沖器作為輸出。輸出的數(shù)據(jù)可包括GPU 12可轉(zhuǎn)移返回至CPU 16的輸出自變量。

CPU 16轉(zhuǎn)移到GPU 12的自變量可被稱為“輸入自變量”。在其中異構(gòu)計(jì)算框架是OpenCL框架的實(shí)例中，驅(qū)動程序/運(yùn)行時(shí)間19可產(chǎn)生自變量且在運(yùn)行時(shí)間傳遞(使得可用)到clSetKernelArg()函數(shù)。clSetKernelArg()函數(shù)接收內(nèi)核20作為自變量以及內(nèi)核自變量26中的任一者，且將自變量轉(zhuǎn)移到GPU 12以使得執(zhí)行可開始。

作為為自變量26分配存儲器的部分，驅(qū)動程序/運(yùn)行時(shí)間19確定與包含在內(nèi)核中的存儲器參考中的一些或全部相關(guān)聯(lián)的自變量26的地址和存儲器區(qū)。存儲器參考可為特定代碼區(qū)段的存儲器參考，例如包含循環(huán)的代碼區(qū)段，稱為“循環(huán)代碼區(qū)段”?；谒_定的存儲器區(qū)，驅(qū)動程序/運(yùn)行時(shí)間19可能夠解析(即，確定)內(nèi)核20的循環(huán)代碼區(qū)段或其它代碼區(qū)段的存儲器參考是否參考自變量26的同一存儲器區(qū)。

響應(yīng)于產(chǎn)生內(nèi)核自變量26用于GPU 12執(zhí)行內(nèi)核20，驅(qū)動程序/運(yùn)行時(shí)間19可執(zhí)行內(nèi)核20。更確切地說，驅(qū)動程序/運(yùn)行時(shí)間19可使用clEnqueueNDRangeKernel()函數(shù)將內(nèi)核20分派到目標(biāo)處理器。在運(yùn)行時(shí)間，驅(qū)動程序/運(yùn)行時(shí)間19分析自變量26，內(nèi)核20接收所述自變量。驅(qū)動程序/運(yùn)行時(shí)間19還分析存儲器參考(例如，指針)等以確定存儲器參考是否參考為自變量26分配的存儲器區(qū)的同一存儲器區(qū)。驅(qū)動程序/運(yùn)行時(shí)間19可以成對方式分析存儲器參考和自變量緩沖器以確定存儲器參考是否參考同一存儲器區(qū)。

驅(qū)動程序/運(yùn)行時(shí)間19進(jìn)一步基于存儲器參考所參考的自變量26的存儲器區(qū)之間的關(guān)系產(chǎn)生與存儲器參考相關(guān)聯(lián)的元數(shù)據(jù)。所述元數(shù)據(jù)可指示存儲器參考之間的關(guān)系。舉例來說，作為一些非限制性實(shí)例，所述元數(shù)據(jù)可包含重疊存儲器參考的列表、與存儲器區(qū)相關(guān)聯(lián)的存儲器區(qū)是否重疊、存儲器區(qū)重疊的范圍，以及所述重疊包括多少字節(jié)。

驅(qū)動程序/運(yùn)行時(shí)間19將所產(chǎn)生元數(shù)據(jù)(如果存在)提供回到JIT編譯器18。響應(yīng)于基于元數(shù)據(jù)確定兩個(gè)存儲器參考并不共享確切同一存儲器區(qū)，驅(qū)動程序/運(yùn)行時(shí)間19可致使JIT編譯器18可使用各種優(yōu)化來重新編譯內(nèi)核20，例如循環(huán)展開、代碼重排序和/或向量化。JIT編譯器18可基于所產(chǎn)生元數(shù)據(jù)應(yīng)用循環(huán)展開代碼重排序和/或向量化的這些各種優(yōu)化。

根據(jù)本發(fā)明的技術(shù)，編譯處理器(例如CPU 16)可經(jīng)配置以使用在編譯處理器上執(zhí)行的由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19組成的群組中的至少一者產(chǎn)生用于執(zhí)行經(jīng)編譯內(nèi)核20的代碼(例如，二進(jìn)制代碼或目標(biāo)代碼)的自變量26。由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19組成的群組中的所述至少一者可進(jìn)一步經(jīng)配置以確定對內(nèi)核自變量的第一存儲器區(qū)的第一存儲器參考和對內(nèi)核自變量的第二存儲器區(qū)的第二存儲器參考是否參考同一存儲器區(qū)。響應(yīng)于通過由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19組成的群組中的所述至少一者確定所述內(nèi)核的第一和第二存儲器參考并不參考同一存儲器區(qū)，CPU 16可進(jìn)一步經(jīng)配置以：借助在CPU 16上執(zhí)行的由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19組成的群組中的所述至少一者致使由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19組成的群組中的所述至少一者基于元數(shù)據(jù)重新編譯內(nèi)核20，且通過在CPU 16上執(zhí)行的由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19組成的群組中的所述至少一者指示目標(biāo)處理器(例如GPU 12)執(zhí)行經(jīng)重新編譯的內(nèi)核20。

GPU 12可為允許大規(guī)模并行處理的專用硬件，所GPU 12可為允許大規(guī)模并行處理的專用硬件，所述硬件對于處理圖形數(shù)據(jù)非常適合。以此方式，CPU 16卸載由GPU 12更好地處置的圖形處理。CPU 16可根據(jù)特定應(yīng)用處理接口(API)或異構(gòu)計(jì)算框架與GPU 12通信。此些API的實(shí)例包含的API以及Khronos集團(tuán)的異構(gòu)計(jì)算框架的實(shí)例包含微軟的DirectCompute、Khronos集團(tuán)的OpenCL^TM。然而，本發(fā)明的各方面不限于上述API和框架，且可延伸到其它類型的API。

CPU 16及GPU 12的實(shí)例包含(但不限于)數(shù)字信號處理器(DSP)、通用微處理器、專用集成電路(ASIC)、現(xiàn)場可編程邏輯陣列(FPGA)或其它等效集成或離散邏輯電路。在一些實(shí)例中，GPU 12可為包含集成及/或離散邏輯電路的專用硬件，所述電路為GPU 12提供適合于圖形處理的大規(guī)模并行處理能力。在一些情況下，GPU 12還可包含通用處理，且可被稱作通用GPU(GPGPU)。本發(fā)明中描述的技術(shù)適用于GPU 12為GPGPU的實(shí)例。

系統(tǒng)存儲器14可包括一或多個(gè)計(jì)算機(jī)可讀存儲媒體。系統(tǒng)存儲器14的實(shí)例包含(但不限于)隨機(jī)存取存儲器(RAM)、只讀存儲器(ROM)、電可擦除可編程只讀存儲器(EEPROM)、快閃存儲器或可用以載運(yùn)或存儲呈指令及/或數(shù)據(jù)結(jié)構(gòu)形式的所要程序代碼及可由計(jì)算機(jī)或處理器存取的任何其它媒體。

在一些方面中，系統(tǒng)存儲器14可包含致使CPU 16及/或GPU 12執(zhí)行在本發(fā)明中歸于CPU 16及GPU 12的功能的指令。因此，系統(tǒng)存儲器14可為包括使一或多個(gè)處理器(例如，CPU 16和GPU 12)執(zhí)行各種功能的指令的計(jì)算機(jī)可讀存儲媒體。

在一些實(shí)例中，系統(tǒng)存儲器14可被視為非暫時(shí)性儲存媒體。術(shù)語“非暫時(shí)性”可指示存儲媒體不是以載波或?qū)嵗齺眢w現(xiàn)，非暫時(shí)性存儲媒體可存儲可隨時(shí)間改變和數(shù)據(jù)(例如，存儲在RAM中)。

使用JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19，CPU 16可將源代碼編譯為用于GPGPU應(yīng)用程序的原生代碼(例如，命令和數(shù)據(jù))或字節(jié)代碼。實(shí)例GPGPU數(shù)據(jù)和命令包含用于射線跟蹤應(yīng)用、物理學(xué)模擬的命令和場景數(shù)據(jù)，或用于任何其它類型的GPGPU內(nèi)核的數(shù)據(jù)。GPGPU應(yīng)用程序(例如，內(nèi)核20)還可使用圖形API(例如DirectX或OpenGL)或使用更通用的計(jì)算API(例如開放計(jì)算語言(OpenCL)或OpenCompute或DirectCompute)來編譯。CPU 16可將用于內(nèi)核20的數(shù)據(jù)發(fā)射到命令緩沖器以供處理。在各種實(shí)例中，命令緩沖器可為系統(tǒng)存儲器14的部分或GPU 12的部分。在一些實(shí)例中，CPU 16可經(jīng)由專用總線(例如PCI-Express總線或另一通用串行或并行總線)將用于GPU 12的內(nèi)核20的命令和數(shù)據(jù)發(fā)射到程序。

為執(zhí)行命令緩沖器中存儲的內(nèi)核20的操作，GPU 12可實(shí)施處理管線。處理管線包含執(zhí)行如由執(zhí)行于GPU 12上的軟件或固件定義的功能，及由經(jīng)硬接線以執(zhí)行極特定功能的固定功能單元執(zhí)行功能?？捎锌赡芾@過用于內(nèi)核20的執(zhí)行的固定功能單元，或者內(nèi)核20的執(zhí)行可使用固定功能單元。

內(nèi)核20可在GPU 12的一或多個(gè)處理元件(也被稱作“著色器核心”或“PE”)上執(zhí)行。著色器核心22為用戶提供功能靈活性，因?yàn)橛脩艨蓪χ鬟M(jìn)行編程以用任何可設(shè)想方式執(zhí)行所需任務(wù)，如同任何其它處理器。然而，固定功能單元經(jīng)針對固定功能單元執(zhí)行任務(wù)的方式而硬接線。因此，固定功能單元可不提供大量功能靈活性。本發(fā)明的技術(shù)是針對例如內(nèi)核20等內(nèi)核在GPU著色器核心22上的執(zhí)行。

一旦CPU 16將與再現(xiàn)圖形場景或執(zhí)行內(nèi)核相關(guān)聯(lián)的數(shù)據(jù)和/或命令傳輸?shù)矫罹彌_器，GPU 12便通過GPU 12的管線開始命令的執(zhí)行。GPU 12的調(diào)度器24產(chǎn)生線程，所述線程執(zhí)行與內(nèi)核相關(guān)聯(lián)的工作的基本單元。調(diào)度器24將線程指派到著色器核心22的特定處理元件。

圖2是說明根據(jù)本發(fā)明的技術(shù)的可執(zhí)行內(nèi)核的處理器的一或多個(gè)著色器核心的多個(gè)處理元件的概念圖。圖2說明GPU 12或CPU 16的部分。GPU 12包含多個(gè)處理元件42A-42N(PE 42)，其可執(zhí)行內(nèi)核(例如內(nèi)核20)的部分。在一些實(shí)例中，可在PE 42上執(zhí)行的內(nèi)核20的部分可被稱為“線程束”或“工作單元”。PE 42可為著色器核心22(圖1)中的一或多者的一部分。線程束或工作單元可包括線程的群組，也被稱作“纖維”，GPU調(diào)度器24可將其指派到多個(gè)處理元件(例如，PE 42)用于執(zhí)行。圖2的每一PE可包括單指令多數(shù)據(jù)(SIMD)單元，其能夠在特定時(shí)間(例如，在同時(shí)以用于并行執(zhí)行)對多個(gè)數(shù)據(jù)值執(zhí)行單個(gè)指令，例如向量指令。PE 42還可支持對單一數(shù)據(jù)值執(zhí)行單一指令，例如對單一浮點(diǎn)值的單一操作。

圖2還包含GPU 12的調(diào)度器指派PE 42以供執(zhí)行的指令44。在一些實(shí)例中，指令44可存儲于命令緩沖器中。指令44可包含每一PE經(jīng)配置以經(jīng)配置以執(zhí)行的內(nèi)核的指令集。程序計(jì)數(shù)器(PC)50指示PE 42中的一或多者將執(zhí)行的當(dāng)前指令。在指令結(jié)束在PE 42上執(zhí)行之后，PC 50的值可遞增到內(nèi)核20的下一個(gè)指令的地址。圖2還包含寄存器46。寄存器46A-46N(寄存器46)可為能夠保持多個(gè)數(shù)據(jù)值或單個(gè)值的通用寄存器。寄存器46可被“儲備”，即，可加載和存儲用于特定PE的數(shù)據(jù)。作為一實(shí)例，寄存器46A可限于存儲用于PE 42A的數(shù)據(jù)，且可不加載或存儲用于其它PE的數(shù)據(jù)。寄存器46中的每一者可將數(shù)據(jù)供應(yīng)到PE 42中的一者和/或供應(yīng)來自其的數(shù)據(jù)，PE 42可隨后處理所述數(shù)據(jù)。

PE 42、指令44、寄存器46、高速緩沖存儲器48和PC 50可包括GPU 12的著色器核心22的核心或部分。在各種實(shí)例中，線程束40可包括著色器的部分，例如幾何著色器、像素著色器和/或頂點(diǎn)著色器，其可為GPU 12的圖形管線的部分或包括例如內(nèi)核20等內(nèi)核的部分。在一些實(shí)例中，GPU 12可將由線程束產(chǎn)生的結(jié)果饋送到管線的另一級中用于額外處理。

圖2還包含高速緩沖存儲器48。高速緩沖存儲器48是存儲頻繁存取的指令和數(shù)據(jù)以用于在執(zhí)行期間的快速檢索和存儲的小存儲器。雖然說明為單個(gè)高速緩沖存儲器，但高速緩沖存儲器48可表示多個(gè)高速緩沖存儲器層級和/或單獨(dú)的高速緩沖存儲器。如上文所描述，在內(nèi)核20的執(zhí)行期間，GPU 12檢索位于由PC 50的值指示的地址的指令44中的一者。GPU 12隨后致使PE 42執(zhí)行存儲在PC 50的所述地址處的指令，所述PC 50在一些實(shí)例中可為寄存器。

GPU 12檢查高速緩沖存儲器48以確定高速緩沖存儲器48當(dāng)前是否包含將執(zhí)行的下一指令，而不是從系統(tǒng)存儲器獲取在PC 50的地址處的指令，這將是不必要地緩慢的。高速緩沖存儲器48的存儲指令的部分稱為指令高速緩沖存儲器(“I高速緩沖存儲器”)。如果將執(zhí)行的下一指令存儲在高速緩沖存儲器48中，稱為“高速緩沖存儲器中”，那么GPU 12加載且執(zhí)行經(jīng)高速緩沖存儲的指令。如果將執(zhí)行的下一指令不存儲在高速緩沖存儲器48中，稱為“高速緩沖存儲器未中”，那么GPU 12從某個(gè)較慢的存儲器、例如從系統(tǒng)存儲器14加載下一指令用于執(zhí)行。

在需要存儲在存儲器地址處的數(shù)據(jù)值(例如，操作數(shù))的指令執(zhí)行(例如，相加、相乘、加載、存儲等)期間，GPU 12首先確定所述操作數(shù)是否存儲在寄存器(例如，寄存器46中的一者)內(nèi)。如果所請求數(shù)據(jù)值不存儲在寄存器46中，那么GPU 12嘗試從高速緩沖存儲器48的保持?jǐn)?shù)據(jù)值的部分存取所述數(shù)據(jù)值，所述部分稱為數(shù)據(jù)高速緩沖存儲器(“d高速緩沖存儲器”)。如果數(shù)據(jù)值存儲在高速緩沖存儲器48內(nèi)，那么GPU 12從高速緩沖存儲器48加載所請求數(shù)據(jù)值。否則，GPU 12必須從較慢的存儲器(例如系統(tǒng)存儲器14)加載所請求數(shù)據(jù)值。類似地，如果指令致使PE 42存儲或修改回到存儲器中的數(shù)據(jù)值，那么高速緩沖存儲器48可存儲所述值到高速緩沖存儲器48以使得如果其被再次寫入或讀取，那么在所述數(shù)據(jù)值不存儲在寄存器46中的一者的情況下從高速緩沖存儲器48快速檢索或?qū)Ω咚倬彌_存儲器48快速覆寫所述數(shù)據(jù)值。

GPU 12以稱為高速緩沖存儲器“線”的固定大小塊對高速緩沖存儲器48和從高速緩沖存儲器48轉(zhuǎn)移數(shù)據(jù)。高速緩沖存儲器48可具有存儲數(shù)百或數(shù)千不同線的容量。每一線與特定存儲器地址相關(guān)聯(lián)，且可存儲多個(gè)字節(jié)的數(shù)據(jù)。舉例來說，作為一個(gè)實(shí)例，高速緩沖存儲器48的每一線可存儲64個(gè)字節(jié)的數(shù)據(jù)。存儲在每一線中的字節(jié)的數(shù)目稱為高速緩沖存儲器“寬度”。在其中高速緩沖存儲器48具有可存儲64個(gè)字節(jié)的數(shù)據(jù)的線的實(shí)例中，高速緩沖存儲器48的高速緩沖存儲器寬度是64字節(jié)。高速緩沖存儲器寬度可影響代碼重排序優(yōu)化技術(shù)的性能，如下文將更詳細(xì)地論述。

在從高速緩沖存儲器48檢索數(shù)據(jù)的加載操作期間，GPU 12可將檢索的高速緩沖存儲器數(shù)據(jù)加載到寄存器46中的一或多者中或未描繪的其它寄存器中。在指令的執(zhí)行期間，PE 42可從寄存器46讀取一或多個(gè)數(shù)據(jù)值。PE 42可對所述數(shù)據(jù)值執(zhí)行一或多個(gè)操作，且將新值存儲回到寄存器46。PE 42可執(zhí)行流控制指令，例如分支、跳轉(zhuǎn)、去往等。然而因?yàn)榇嬖趩我籔C 50，所以PE 42可在給定時(shí)間在一個(gè)特定處僅執(zhí)行由PC 50指示的指令44中的一者。

例如GPU 12的處理器可具有廣泛量的向量寄存器和向量指令。因此，可使用例如向量化等優(yōu)化來編譯應(yīng)用程序的例如JIT編譯器18等編譯器可增加支持向量指令或具有SIMD架構(gòu)的處理器(例如GPU 12)的處理量或執(zhí)行性能。

更確切地說，GPU 12可包含類似于圖2中說明的那些著色器核心的數(shù)百或數(shù)千個(gè)著色器核心。每一著色器核心可能夠執(zhí)行向量指令。執(zhí)行具有多個(gè)操作數(shù)的向量指令可極大地改善相對于含有標(biāo)量指令而不是向量指令的未經(jīng)優(yōu)化代碼的性能。此外，執(zhí)行性能增加可在具有能夠執(zhí)行向量指令的較大數(shù)目的SIMD核心的架構(gòu)上較大，因?yàn)檩^多通用處理器可具有有限數(shù)目的能夠執(zhí)行向量指令的寄存器和/或核心。

圖3A是說明根據(jù)本發(fā)明的技術(shù)的包含當(dāng)執(zhí)行時(shí)可造成混疊的代碼的內(nèi)核代碼的概念圖。圖3A的實(shí)例包含內(nèi)核代碼80。內(nèi)核代碼80包含行82、84、86和88。

內(nèi)核代碼80的行82是compute_output函數(shù)。行82的compute_output函數(shù)是當(dāng)內(nèi)核開始執(zhí)行時(shí)目標(biāo)處理器(例如，GPU 12)調(diào)用的函數(shù)。其大致等效于C編程語言中的“int main()”函數(shù)，因?yàn)閏ompute_output函數(shù)是驅(qū)動程序/運(yùn)行時(shí)間19用來開始內(nèi)核20的執(zhí)行的程序進(jìn)入點(diǎn)。如果目標(biāo)處理器是CPU 16，那么C運(yùn)行時(shí)間庫可包括驅(qū)動程序/運(yùn)行時(shí)間19的運(yùn)行時(shí)間組件。如果GPU 12是目標(biāo)處理器，那么驅(qū)動程序/運(yùn)行時(shí)間19的驅(qū)動程序組件可包括運(yùn)行時(shí)間。compute_ouput函數(shù)包含四個(gè)輸入自變量：(1)inputImage，(2)global_cdf，(3)outputImage，和(4)local_cdf。inputImage是到輸入自變量的緩沖器的指針。outputImage是當(dāng)內(nèi)核結(jié)束執(zhí)行時(shí)到緩沖器的指針，將包含輸出自變量。自變量global_cdf和local_cdf是到值的陣列的指針。行84可表示當(dāng)執(zhí)行時(shí)致使GPU 12分配且初始化變量的多個(gè)語句。作為一實(shí)例，執(zhí)行行84可致使PE 42初始化且加載inputImage[i]等的值。

行86是循環(huán)初始化語句。循環(huán)初始化語句指示循環(huán)重復(fù)固定數(shù)目的迭代。循環(huán)以等于變量“start_offset”的起始索引i開始迭代，且當(dāng)每一迭代結(jié)束執(zhí)行時(shí)將i遞增一。在每一循環(huán)迭代的完成時(shí)，GPU 12檢查以查看布爾型條件“i<final_offset”是否仍為真。當(dāng)i的值等于或大于值“final_offset”時(shí)，GPU 12停止執(zhí)行循環(huán)。

在每一循環(huán)迭代內(nèi)，GPU 12將outputImage的值設(shè)定于索引i，表示為等于local_cdf[inputImage[i]]的值的outputImage[i]。Local_cdf是陣列，其在此實(shí)例中以inputImage[i]的值作索引。inputImage[i]又以GPU 12隨著每一循環(huán)迭代而遞增的變量i作索引。

如上文所論述，outputImage和inputImage都是存儲器參考。有可能到outputImage和inputImage的指針可指代存儲器中的同一區(qū)(即outputImage和inputImage混疊或部分地混疊)。也可能outputImage和inputImage可指代存儲器中的不同區(qū)或重疊區(qū)(即outputImage和inputImage并不混疊)。如果JIT編譯器18不能夠確定inputImage和outputImage是否并不混疊(即并不參考確切同一存儲器區(qū))，那么編譯器可不能夠使用某些編譯器優(yōu)化，例如向量化、代碼重排序和/或循環(huán)展開。

圖3B是說明根據(jù)本發(fā)明的技術(shù)配置的驅(qū)動程序/運(yùn)行時(shí)間可能夠檢測的混疊的實(shí)例的概念圖。然而，編譯器可能不能夠針對混疊進(jìn)行優(yōu)化。圖3B的實(shí)例說明GPU 12可在存儲器中存儲的緩沖器100。出于實(shí)例的目的，來自圖3B的指針outputImage和inputImage可參考緩沖器100的部分。在圖3B的實(shí)例中，緩沖器100開始于存儲器地址0x800(十六進(jìn)制)。

在此實(shí)例中，inputImage和outputImage都參考存儲在緩沖器100內(nèi)的單個(gè)條目(例如，單個(gè)對象、變量等)。即，在此實(shí)例中，inputImage和outputImage混疊到確切同一存儲器區(qū)，其以交叉散列指示。驅(qū)動程序/運(yùn)行時(shí)間可能夠檢測到inputImage和outputImage參考同一存儲器區(qū)。因?yàn)閕nputImage和outputImage參考同一存儲器區(qū)，所以JIT編譯器18不能夠執(zhí)行優(yōu)化，例如循環(huán)展開和/或向量化。

響應(yīng)于如圖3B中所說明檢測到兩個(gè)存儲器參考參考同一存儲器區(qū)，驅(qū)動程序/運(yùn)行時(shí)間19可不產(chǎn)生任何元數(shù)據(jù)。另外，JIT編譯器18可不重新編譯內(nèi)核20，因?yàn)镴IT編譯器可針對圖3C到3D中說明的情況來進(jìn)行。因此，JIT編譯器18可不執(zhí)行也如圖3C到3D中說明的代碼優(yōu)化中的任一者。

圖3C是說明根據(jù)本發(fā)明的技術(shù)配置的驅(qū)動程序/運(yùn)行時(shí)間可能夠檢測的不重疊存儲器參考的實(shí)例的概念圖。圖3C說明作為與圖3B中所說明的同一緩沖器的緩沖器120。緩沖器120類似地與圖3B的緩沖器100開始于同一存儲器地址0x800。

在圖3C中，inputImage和outputImage是參考緩沖器120的兩個(gè)不同存儲器區(qū)的存儲器參考。inputImage參考的存儲器區(qū)由水平散列指示。outputImage參考的存儲器區(qū)由垂直散列指示。在內(nèi)核代碼80且更具體地說行86和88的執(zhí)行之前，JIT編譯器18不管i的值如何都確定在循環(huán)的同一迭代期間inputImage[i]和outputImage[i]將不參考同一存儲器區(qū)。

在運(yùn)行時(shí)間期間，驅(qū)動程序/運(yùn)行時(shí)間19可能夠基于inputImage[i]和outputImage[i]的初始值且基于inputImage[i]和outputImage[i]的存儲器地址在迭代通過循環(huán)86的過程中并不收斂的事實(shí)而確定inputImage[i]和outputImage[i]并不參考同一存儲器區(qū)。換句話說，inputImage和outputImage的所參考索引始終由GPU 12單調(diào)增加的同一索引值i所參考。

響應(yīng)于確定存儲器參考inputImage和outputImage并不參考同一存儲器區(qū)，驅(qū)動程序可產(chǎn)生指示inputImage與outputImage之間的關(guān)系的元數(shù)據(jù)。作為一實(shí)例，元數(shù)據(jù)可指示與inputImage和outputImage相關(guān)聯(lián)的存儲器區(qū)不重疊，且通過兩個(gè)條目分離。元數(shù)據(jù)還可指示與inputImage和ouptutImage相關(guān)聯(lián)的區(qū)的大小，以及inputImage與outputImage之間的字節(jié)數(shù)目。在產(chǎn)生元數(shù)據(jù)之后，JIT編譯器18可從驅(qū)動程序/運(yùn)行時(shí)間19接收元數(shù)據(jù)，且通過應(yīng)用各種優(yōu)化基于元數(shù)據(jù)重新編譯內(nèi)核20，如下文更詳細(xì)描述。

圖3D是說明根據(jù)本發(fā)明的技術(shù)配置的驅(qū)動程序/運(yùn)行時(shí)間可檢測的重疊存儲器參考的概念圖。圖3D包含緩沖器130，其可為自變量緩沖器，例如自變量26(圖1)。在此實(shí)例中緩沖器130開始于地址0x800。緩沖器130包含多個(gè)數(shù)據(jù)值，其說明為緩沖器130的包含矩形內(nèi)的單獨(dú)矩形。

如先前實(shí)例中，inputImage和outputImage是存儲器參考，其參考緩沖器130的區(qū)。在此實(shí)例中，inputImage和outputImage參考的區(qū)重疊，但不是完全重疊。僅與inputImage相關(guān)聯(lián)的存儲器區(qū)以水平散列矩形指示。僅與outputImage相關(guān)聯(lián)的存儲器區(qū)以垂直散列矩形指示。由inputImage和outputImage兩者參考的重疊存儲器區(qū)以交叉影線矩形指示。

在運(yùn)行時(shí)間，驅(qū)動程序確定inputImage和outputImage存儲器參考是否參考同一存儲器區(qū)。在此實(shí)例中，inputImage和outputImage重疊，但并不參考同一存儲器區(qū)。驅(qū)動程序/運(yùn)行時(shí)間19檢測到inputImage和outputImage重疊但不相同，且產(chǎn)生用于JIT編譯器18的元數(shù)據(jù)。所述元數(shù)據(jù)可指示關(guān)于與inputImage和outputImage相關(guān)聯(lián)的區(qū)的信息，例如每一區(qū)的起始和結(jié)束地址。所述元數(shù)據(jù)可進(jìn)一步包含關(guān)于重疊區(qū)的信息，例如重疊區(qū)的大小以及重疊區(qū)的起始和/或結(jié)束地址。JIT編譯器18接收由驅(qū)動程序/運(yùn)行時(shí)間19產(chǎn)生的元數(shù)據(jù)，且可通過應(yīng)用根據(jù)本發(fā)明的優(yōu)化技術(shù)而重新編譯內(nèi)核20。

圖4A是說明根據(jù)本發(fā)明的技術(shù)的循環(huán)展開的概念圖。圖4A包含代碼區(qū)段140，其一般對應(yīng)于圖3A中說明的內(nèi)核代碼80。在圖4A的實(shí)例中，驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18可能已確定存儲器參考inputImage和outputImage并不參考同一存儲器區(qū)，如圖3C和3D中所說明。因?yàn)閕nputImage和outputImage并不參考同一存儲器區(qū)，所以JIT編譯器18已對內(nèi)核代碼80執(zhí)行循環(huán)展開。行142到150說明將一個(gè)迭代展開為四個(gè)迭代的結(jié)果。

圖3A的行86和88說明執(zhí)行單個(gè)迭代且在每一迭代之后將變量i遞增一，而行142的展開循環(huán)在每一迭代之后將i遞增四。行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的結(jié)果是將local_cdf[inputImage[i+x]]的輸出指派給outputImage[i+x]的對應(yīng)值，其中x[0…3]。因此，當(dāng)執(zhí)行時(shí)，行142到150中說明的展開循環(huán)代碼區(qū)段具有與圖3A的行86到88的四個(gè)迭代相同的效果。

代碼區(qū)段140的循環(huán)展開可具有相對于圖3A的循環(huán)代碼區(qū)段80的若干益處。第一優(yōu)點(diǎn)是通過一個(gè)接一個(gè)地對指派中的每一者排序，JIT編譯器18和/或驅(qū)動程序/運(yùn)行時(shí)間19可能夠?qū)崿F(xiàn)相對于未經(jīng)排序代碼區(qū)段的在目標(biāo)處理器(例如GPU 12)上的更好的高速緩沖存儲器性能。

舉例來說，在執(zhí)行行144之后，GPU 12可能已在高速緩沖存儲器(例如高速緩沖存儲器48)中存儲與inputImage和outputImage相關(guān)聯(lián)的存儲器區(qū)的數(shù)據(jù)中的一些或全部。如果執(zhí)行指令所需要的數(shù)據(jù)不存儲在寄存器(例如寄存器46)中，那么可能需要從高速緩沖存儲器(例如高速緩沖存儲器48)存取數(shù)據(jù)。更確切地說，GPU 12可在高速緩沖存儲器48中存儲inputImage和outputImage的條目，例如inputImage[i+1]、[i+2]等，以及outputImage[i+1]、[i+2]等。如果inputImage和outputImage的條目存儲在GPU 12的高速緩沖存儲器中，那么GPU 12可能夠從高速緩沖存儲器快速存取行144到150的inputImage和outputImage的參考索引的數(shù)據(jù)，而不是從較慢的存儲器存取所述參考索引。

另外，當(dāng)代碼區(qū)段140展開時(shí)，inputImage[i，i+1，i+2…]和outputImage[i，i+1，等]的值可存儲在單個(gè)高速緩沖存儲器線中。相比之下當(dāng)不展開時(shí)，inputImage和outputImage[i]的值可存儲在不同高速緩沖存儲器線中?？捎裳h(huán)展開引起的單個(gè)高速緩沖存儲器讀取中從單個(gè)高速緩沖存儲器線檢索inputImage的全部值相對于當(dāng)執(zhí)行展開代碼時(shí)可導(dǎo)致的執(zhí)行多次高速緩沖存儲器讀取來說可更快。

從GPU 12的高速緩沖存儲器存取數(shù)據(jù)而不是從較慢的系統(tǒng)存儲器(例如系統(tǒng)存儲器14)存取數(shù)據(jù)可相對于行86到88增加執(zhí)行行142到150的循環(huán)的性能。在一些實(shí)例中，GPU 12也可以能夠例如在支持超標(biāo)量執(zhí)行的處理器或假定行144到150之間不存在相依性的SIMD處理器上并行地執(zhí)行行144到150，其中inputImage或outputImage的值取決于內(nèi)核20中先前計(jì)算的值。

除改善高速緩沖存儲器性能之外，如圖4A的代碼區(qū)段140中所說明的循環(huán)展開也減少GPU 12評估與循環(huán)相關(guān)聯(lián)的布爾型條件的次數(shù)，以及GPU 12在完成每一循環(huán)迭代之后執(zhí)行的跳轉(zhuǎn)的次數(shù)。與圖3A的代碼區(qū)段80相比，在評估行142的布爾型條件“i<final_offset”是否為真之前，行142到150的代碼每迭代執(zhí)行四行。代碼區(qū)段80相比之下在評估行82的布爾型條件是否為真之前僅執(zhí)行一行。因此，GPU 12評估行142的布爾型條件的次數(shù)相對于代碼區(qū)段80減少。

在GPU 12完成行142到150的循環(huán)的迭代之后，且如果GPU 12確定布爾型條件“i<final_offset”仍為真，那么GPU 12從行150跳轉(zhuǎn)回到行144。在代碼區(qū)段140中，GPU 12在執(zhí)行四行之后執(zhí)行跳轉(zhuǎn)。當(dāng)執(zhí)行代碼區(qū)段80時(shí)，GPU 12在每一迭代之后跳轉(zhuǎn)。因此，相對于代碼區(qū)段80，代碼區(qū)段140的展開代碼減少了布爾型條件的評估以及GPU 12執(zhí)行的跳轉(zhuǎn)的數(shù)目，其可改善執(zhí)行代碼區(qū)段140的執(zhí)行性能。

圖4B是說明根據(jù)本發(fā)明的技術(shù)的代碼重排序的概念圖。圖4B包含代碼區(qū)段160，其進(jìn)一步包含行162、164、166、168、170和172。如上文所論述，驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18可確定對存儲器的參考是否混疊到存儲器的同一區(qū)。如上文相對于圖4A所論述，JIT編譯器18可響應(yīng)于從驅(qū)動程序/運(yùn)行時(shí)間19接收到確定特定代碼區(qū)段中不存在存儲器混疊的元數(shù)據(jù)而執(zhí)行某些優(yōu)化，例如圖4A中說明的循環(huán)展開。

JIT編譯器18和/或驅(qū)動程序/編譯器19可響應(yīng)于確定特定代碼區(qū)段中的存儲器參考并不參考同一存儲器區(qū)而執(zhí)行的另一優(yōu)化是圖4B說明的代碼重排序。代碼160可一般對應(yīng)于圖4B的展開代碼的經(jīng)重排序匯編語言表示。JIT編譯器18和/或驅(qū)動程序/運(yùn)行時(shí)間19可將代碼重排序應(yīng)用于非循環(huán)以及循環(huán)代碼區(qū)段。在圖4B中，JIT編譯器18已經(jīng)重排序圖4A的加載和存儲以使得全部加載和存儲分組在一起。

行162和164是JIT編譯器18和/或驅(qū)動程序/編譯器19已經(jīng)分組在一起的加載指令。在圖4A中，例如行144的行包含多個(gè)加載和存儲指令。舉例來說，為了執(zhí)行行144，JIT編譯器18可產(chǎn)生三個(gè)單獨(dú)指令。第一指令可為加載指令，其將來自inputImage[i]參考的存儲器位置的值加載到表示為r0的寄存器中。第二指令可為加載指令，其加載local_cdf[inputImage[i]]的值且將加載的值存儲到同一寄存器r0中，進(jìn)而覆寫r0的先前值。包含在行144中的最終指令可為存儲指令，其將來自r0的值存儲到outputImage[i]參考的存儲器中。

行162到172說明相對于包括行144到150的指令經(jīng)重排序的加載和存儲指令。在行162中，組合代碼指示GPU 12將來自inputImage[i]參考的存儲器區(qū)的值加載(使用加載指令“l(fā)dg”)到寄存器r0中。類似地，行164致使GPU 12將存儲器參考inputImage[i+1]參考的值加載到寄存器r1中。可在行162和164之后但在行166之前發(fā)生且出于簡潔起見未說明的后續(xù)指令可包含致使GPU 12將來自inputImage參考的存儲器區(qū)的數(shù)據(jù)加載到寄存器中的額外加載指令。

在行166、168和為簡潔起見未說明的其它行中，JIT編譯器18已經(jīng)將從緩沖器的加載local_cdf分組在一起。行166包含加載local_cdf[r0]的內(nèi)容(即在索引r0處來自陣列l(wèi)ocal_cdf的存儲器的內(nèi)容)且將local_cdf[r0]的內(nèi)容存儲到寄存器r0中進(jìn)而覆寫r0的內(nèi)容的加載指令。類似地，行168的指令致使GPU 12在由當(dāng)前存儲在寄存器r1中的值指示的索引處將存儲器參考local_cdf參考的內(nèi)容存儲到寄存器r1中。因此，在執(zhí)行時(shí)168的指令致使GPU 12覆寫r1的先前值。在行168之后且在行170之前發(fā)生且出于簡潔起見未說明的其它指令可類似地包含當(dāng)執(zhí)行時(shí)致使GPU 12從local_cdf[rx]加載數(shù)據(jù)的指令，其中x是某個(gè)整數(shù)。

作為重排序代碼區(qū)段140的指令的部分，JIT編譯器18還將存儲指令分組在一起。作為此一實(shí)例，在重排序之后，JIT編譯器18已將行170和172分組在一起行170包含在位置outputImage[i]將r0的內(nèi)容存儲到存儲器中的存儲指令。類似地，行172當(dāng)執(zhí)行時(shí)致使GPU 12在outputImage[i+1]參考的位置將寄存器r1的值存儲到存儲器中。出于簡潔起見未說明的其它指令當(dāng)執(zhí)行時(shí)可類似地致使GPU 12在位置outputImage[i+x]將寄存器(例如寄存器rx，其中x是整數(shù))的值存儲到存儲器。

重排序加載和存儲可相對于圖3A的代碼80改善執(zhí)行代碼160的性能。更確切地說，重排序加載和存儲可在某些情況下取決于高速緩沖存儲器線寬度而改善性能。舉例來說，當(dāng)將許多加載指令聚結(jié)在一起可改善具有作為標(biāo)量指令中使用的操作數(shù)大小的倍數(shù)的高速緩沖存儲器線寬的系統(tǒng)上的性能時(shí)，代碼重排序可改善執(zhí)行性能。

圖4C是說明根據(jù)本發(fā)明的技術(shù)的代碼向量化的概念圖。圖4C包含代碼區(qū)段180，其進(jìn)一步包含行182、184和186。如上文所論述，JIT編譯器18和/或編譯器/驅(qū)動程序19可響應(yīng)于確定代碼區(qū)段中的存儲器參考并不參考同一存儲器區(qū)而執(zhí)行某些優(yōu)化，例如圖4A中說明的循環(huán)展開。JIT編譯器18和/或編譯器/驅(qū)動程序19經(jīng)配置以基于來自驅(qū)動程序/運(yùn)行時(shí)間19的包含與代碼區(qū)段的存儲器參考相關(guān)的信息的元數(shù)據(jù)而向量化循環(huán)代碼區(qū)段。

向量化是其中編譯器(例如，JIT編譯器18)和/或驅(qū)動程序/運(yùn)行時(shí)間19將各自具有單個(gè)操作數(shù)的多個(gè)標(biāo)量指令組合為具有多個(gè)操作數(shù)的單個(gè)向量指令的過程。向量化是一種形式的并行度，其通過減少處理器需要執(zhí)行以完成特定代碼區(qū)段的指令數(shù)目以及通過利用固有硬件能力在系統(tǒng)存儲器14與GPU 12之間移動數(shù)據(jù)而改善執(zhí)行性能。在圖4C的代碼區(qū)段180的實(shí)例中，JIT編譯器18可如圖4B中所說明重排序加載和存儲。一旦JIT編譯器18已經(jīng)重排序加載和存儲，JIT編譯器18便接著可向量化類似指令的群組，如圖4C中所說明。

在行182中，JIT編譯器18已經(jīng)將多個(gè)加載(ldg)指令組合為單個(gè)向量化指令。當(dāng)執(zhí)行時(shí)，向量化指令將索引[i]到[i+3]處的inputImage加載到寄存器r0到r3。類似地，在行184中，JIT編譯器18將行166、168等的多個(gè)加載指令組合為單個(gè)向量化加載指令，其將local_cdf[r0-r3]的值加載到寄存器r0到r3中。并且，在行186中，JIT編譯器18已將行170到172的存儲(“stg”指令)組合為單個(gè)向量化存儲指令，其將寄存器r0到r3的值存儲到outputImage[i]到outputImage[i+3]中。

為了如圖4B和4C中所說明重排序或向量化指令，JIT編譯器18和/或驅(qū)動程序/運(yùn)行時(shí)間19必須注意任何相依性。相依性是產(chǎn)生語句或指令之間的執(zhí)行次序約束的關(guān)系。作為一實(shí)例，如果S1必須在S2之前執(zhí)行，那么存在語句S2對另一語句S1的相依性。為了確定相依性是否禁止向量化和/或代碼重排序，JIT編譯器18和/或驅(qū)動程序/運(yùn)行時(shí)間19可基于從驅(qū)動程序/運(yùn)行時(shí)間19獲得的元數(shù)據(jù)在根據(jù)本發(fā)明的技術(shù)重排序或向量化代碼之前執(zhí)行相依性分析。

圖5是說明根據(jù)本發(fā)明的技術(shù)用于產(chǎn)生編譯器元數(shù)據(jù)以輔助編譯器優(yōu)化的實(shí)例方法的流程圖。應(yīng)大體上理解，圖6的方法可通過由JIT編譯器18和驅(qū)動程序/運(yùn)行時(shí)間19執(zhí)行編譯處理器(例如，CPU 16)和目標(biāo)處理器(例如，GPU 12)組成的群組中的至少一者執(zhí)行。在一些實(shí)例中，目標(biāo)處理器和編譯處理器可為相同的。另外，可存在多于一個(gè)編譯處理器和/或目標(biāo)處理器。

在圖5的方法中，編譯處理器(例如CPU 16)使用驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18以產(chǎn)生自變量(例如，內(nèi)核自變量26)用于執(zhí)行經(jīng)編譯內(nèi)核20的二進(jìn)制代碼或字節(jié)代碼(200)。驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18進(jìn)一步確定對內(nèi)核自變量26的第一存儲器區(qū)的第一存儲器參考和對內(nèi)核自變量26的第二存儲器區(qū)的第二存儲器參考是否參考內(nèi)核自變量26的同一存儲器區(qū)(202)，或如圖3B、3C和3D中所說明的可能關(guān)系的其它實(shí)例。

CPU 16使用驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18以產(chǎn)生與第一存儲器參考和第二存儲器參考相關(guān)聯(lián)的元數(shù)據(jù)(204)。所述元數(shù)據(jù)指示第一存儲器區(qū)與第二存儲器區(qū)之間的關(guān)系，例如第一存儲器區(qū)與第二存儲器區(qū)之間的重疊區(qū)。所述元數(shù)據(jù)可進(jìn)一步包含第一和第二存儲器區(qū)之間的重疊的字節(jié)數(shù)目。在一些實(shí)例中，所述元數(shù)據(jù)可包含存儲器重疊的起始地址以及存儲器重疊區(qū)的結(jié)束地址。應(yīng)理解，相對于圖5描述的實(shí)例僅出于實(shí)例的目的而參考單個(gè)對的存儲器參考。驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18可針對內(nèi)核自變量26的所有對的存儲器參考導(dǎo)出元數(shù)據(jù)。

響應(yīng)于使用驅(qū)動程序/運(yùn)行時(shí)間19確定第一和第二存儲器參考并不參考內(nèi)核自變量26的同一存儲器區(qū)，在CPU 16上執(zhí)行的JIT編譯器18可致使CPU 16基于所述元數(shù)據(jù)使用JIT編譯器18重新編譯內(nèi)核20(206)。最后，目標(biāo)處理器(例如GPU 12)可執(zhí)行經(jīng)重新編譯的內(nèi)核(210)。在一些實(shí)例中，驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18可基于元數(shù)據(jù)確定第一和第二存儲器參考并不參考同一存儲器區(qū)，且可使用此信息借助優(yōu)化來重新編譯內(nèi)核20。

在一些額外實(shí)例中，為了確定內(nèi)核20的第一存儲器參考和第二存儲器參考是否參考同一存儲器區(qū)，CPU 16可使用驅(qū)動程序/運(yùn)行時(shí)間19以確定內(nèi)核20的包含第一和第二存儲器參考的循環(huán)代碼區(qū)段。并且，為了重新編譯內(nèi)核，JIT編譯器18可基于由驅(qū)動程序/運(yùn)行時(shí)間19和/或JIT編譯器18產(chǎn)生的元數(shù)據(jù)而展開所述循環(huán)代碼區(qū)段。為了重新編譯內(nèi)核，JIT編譯器18還可基于所產(chǎn)生元數(shù)據(jù)而重排序所述循環(huán)代碼區(qū)段的加載操作和存儲操作和存儲操作中的至少一者或?qū)⑺鲅h(huán)代碼區(qū)段的多個(gè)標(biāo)量指令向量化為至少一個(gè)向量指令。在各種實(shí)例中，JIT編譯器18可使用例如微軟DirectCompute和/或Khronos集團(tuán)的OpenCL等異構(gòu)框架來重新編譯內(nèi)核20。

本發(fā)明中所描述的技術(shù)可至少部分實(shí)施于硬件、軟件、固件或其任何組合中。舉例來說，所描述技術(shù)的各種方面可實(shí)施于一或多個(gè)處理器中，包含一或多個(gè)微處理器、數(shù)字信號處理器(DSP)、專用集成電路(ASIC)、現(xiàn)場可編程門陣列(FPGA)，或任何其它等效集成或離散邏輯電路，以及此等組件的任何組合。術(shù)語“處理器”或“處理電路”可一般指代前述邏輯電路中的任一者(單獨(dú)或結(jié)合其它邏輯電路)或例如執(zhí)行處理的離散硬件的任何其它等效電路。

此硬件、軟件和固件可實(shí)施于同一裝置內(nèi)或單獨(dú)裝置內(nèi)以支持本發(fā)明中所描述的各種操作和功能。另外，所描述的單元、模塊或組件中的任一者可一起或單獨(dú)作為離散但可互操作邏輯裝置而實(shí)施。將不同特征描繪為模塊或單元意圖強(qiáng)調(diào)不同功能方面且未必暗示此類模塊或單元必須由單獨(dú)硬件或軟件組件實(shí)現(xiàn)。而是，與一或多個(gè)模塊或單元相關(guān)聯(lián)的功能性可由單獨(dú)硬件、固件及/或軟件組件執(zhí)行，或集成到共用或單獨(dú)硬件或軟件組件內(nèi)。

本發(fā)明中所描述的技術(shù)也可存儲、體現(xiàn)或編碼于計(jì)算機(jī)可讀媒體(例如，存儲指令的計(jì)算機(jī)可讀存儲媒體)中。嵌入或編碼于計(jì)算機(jī)可讀媒體中的指令可致使一或多個(gè)處理器執(zhí)行本文中所描述的技術(shù)(例如，當(dāng)由一或多個(gè)處理器執(zhí)行指令時(shí))。計(jì)算機(jī)可讀存儲媒體可包含隨機(jī)存取存儲器(RAM)、只讀存儲器(ROM)、可編程只讀存儲(PROM)、可擦除可編程只讀存儲器(EPROM)、電可擦除可編程只讀存儲器(EEPROM)、快閃存儲器、硬盤、CD-ROM、軟盤、卡盒、磁性媒體、光學(xué)媒體或其它有形計(jì)算機(jī)可讀存儲媒體。

計(jì)算機(jī)可讀媒體可包含計(jì)算機(jī)可讀存儲媒體，其對應(yīng)于例如上文所列的有形存儲媒體的有形存儲媒體。計(jì)算機(jī)可讀媒體也可包括通信媒體，其包含促進(jìn)計(jì)算機(jī)程序從一個(gè)地點(diǎn)到另一地點(diǎn)的傳送(例如，根據(jù)通信協(xié)議)的任何媒體。以此方式，短語“計(jì)算機(jī)可讀媒體”大體上可對應(yīng)于(1)非暫時(shí)性有形計(jì)算機(jī)可讀存儲媒體，和(2)例如暫時(shí)性信號或載波等非有形計(jì)算機(jī)可讀通信媒體。

已描述各種方面和實(shí)例。然而，可在不脫離所附權(quán)利要求書的范圍的情況下對本發(fā)明的結(jié)構(gòu)或技術(shù)作出修改。