專利名稱：：為了用于由通用處理器執(zhí)行而對應(yīng)用程序重定目標(biāo)的制作方法為了用于由通用處理器執(zhí)行而對應(yīng)用程序重定目標(biāo)相關(guān)專利申請的交叉引用本申請要求于2008年4月9日提交的美國臨時申請序列號為61/043，708(律師案巻號No.NVDA/SC-08-0007-US0)的優(yōu)先權(quán)，在此通過4爰引的方式納入本說明書。才支術(shù)領(lǐng)域本發(fā)明的實施例大體上涉及編譯器程序，更具體地涉及一種翻譯器，所述翻譯器用于將為了由多核圖像處理器執(zhí)行而編寫的應(yīng)用程序重定目標(biāo)，以便由具有共享存儲器的通用處理器執(zhí)行。
背景技術(shù)：
：現(xiàn)代圖像處理系統(tǒng)通常包括配置為以多線程方式執(zhí)行應(yīng)用程序的多核圖像處理單元(GPU)。這種圖像處理系統(tǒng)還包括存儲器，該存儲器具有在執(zhí)行線程之間共享的部分和每個線程專用的部分。NVIDIA的CUDA(計算統(tǒng)一設(shè)備架構(gòu))技術(shù)提供了一種C語言環(huán)境，使得編程人員和開發(fā)人員能夠編寫解決復(fù)雜計算問題的軟件應(yīng)用程序，這些復(fù)雜計算問題諸如視頻和音頻編碼、油氣勘探建模和醫(yī)學(xué)成像。應(yīng)用程序配置為由多核GPU并行執(zhí)行，并且通常依靠多核GPU的特定特4正。由于同樣的特定特征在通用中央處理單元(CPU)中是不可用的，因此用CUDA編寫的軟件應(yīng)用程序可能無法移植到通用CPU上運行。如前所述，在本領(lǐng)域中需要這樣一種技術(shù)，這種技術(shù)不需要編程人員對應(yīng)用程序進行修改就能讓用并行編程模型編寫的用于在多核GPU上執(zhí)行的應(yīng)用程序能夠在通用CPU上運行。
發(fā)明內(nèi)容本發(fā)明的一個實施例給出了一種將用并行編程模型編寫的程序翻:^,為用于由通用CPU執(zhí)行的方法。該方法包括接收用并行編程模型編寫的用于由多核圖像處理單元執(zhí)行的應(yīng)用程序的步驟，和將該應(yīng)用程序劃分成同步獨立指令的區(qū)域以創(chuàng)建劃分的應(yīng)用程序的步驟。將一個循環(huán)插入到劃分的應(yīng)用程序中的至少一個區(qū)域附近，以創(chuàng)建用于由通用處理器執(zhí)行的翻譯后的應(yīng)用程序。該循環(huán)在合作線程陣列維度上進行迭代，該線程陣列維度與由位于多核圖像處理單元內(nèi)的并行處理器同時執(zhí)行的多個線程相對應(yīng)。在此披露的方法的一個優(yōu)勢在于用并行編程模型編寫的用于由多核GPU執(zhí)行的應(yīng)用程序可以無需修改而移植到通用CPU上。為了能夠由通用CPU執(zhí)行，將應(yīng)用程序中依賴于多核GPU的特定特征的部分由翻譯器進行轉(zhuǎn)換。將應(yīng)用程序劃分為同步獨立指令的區(qū)域。將這些指令分類為收斂的或發(fā)散的，并且將這些區(qū)域之間共享的發(fā)散存儲器基準(zhǔn)進行復(fù)制。在由通用CPU執(zhí)行的過程中，插入線程循環(huán)以確保在各種線程之間能夠正確地共享存儲器。為了詳細地理解本發(fā)明的上述特征，對于以上筒要說明的發(fā)明，將參照實施例進行更為具體的描述，其中對部分實施例結(jié)合附圖進行了說明。然而，需要注意的是，附圖中示出的只是本發(fā)明代表性的實施例，因此不能認為附圖限制了本發(fā)明的范圍，本發(fā)明可以適用于其他同樣有效的實施例。圖l是說明計算機系統(tǒng)的框圖2是根據(jù)本發(fā)明的一個實施例對計算機系統(tǒng)進行說明的框圖3A是根據(jù)本發(fā)明的一個實施例，將用于由多核圖像處理單元執(zhí)行而編寫的代碼翻譯為由通用處理器執(zhí)行的代碼的方法步驟的流程圖3B是根據(jù)本發(fā)明的一個實施例，說明翻譯為劃分后的代碼的輸入代碼的概念圖3C是根據(jù)本發(fā)明的一個實施例，說明翻譯為優(yōu)化后的代碼的輸入代碼的概念圖；以及圖4是根據(jù)本發(fā)明的一個實施例，說明由通用處理器^丸行翻譯后的代碼的方法步驟的流程圖。具體實施例方式解。然而，對于本領(lǐng)域:術(shù)人員來說顯而易見的i,'即使缺少其中二個或多個這些具體細節(jié)也可以實施本發(fā)明。在其他例子中，為了避免引起與本發(fā)明的混淆，對一些公知的特征沒有進行描述。5圖l示出了配置為用于執(zhí)行用CUDA編寫的代碼的計算機系統(tǒng)100的框圖。計算機系統(tǒng)100包括CPU102和系統(tǒng)存儲器104，兩者通過包含存儲器橋105的總線路徑互相通信。存儲器橋105例如可以是北橋芯片，通過總線或其他通信路徑106(例如，超傳輸鏈接)與I/O(輸7v/輸出)橋107相連接。I/O橋107例如可以是南橋芯片，從一個或多個用戶輸入裝置108(例如，鍵盤、鼠標(biāo))接收用戶輸入，并將該輸入通過路徑106和存儲器橋105轉(zhuǎn)發(fā)給CPU102。多線程處理子系統(tǒng)112通過總線或其他通信路徑113(例如，PCIExpress、加速圖像端口(AGP)或者超傳輸鏈接)與存儲器橋105相連接。在一個實施例中，多線程處理子系統(tǒng)112是將像素輸送到顯示裝置110(例如，傳統(tǒng)的CRT或基于LCD的顯示器)的圖像子系統(tǒng)。系統(tǒng)盤114還與I/O橋107相連接。開關(guān)116為I/O橋107和諸如網(wǎng)絡(luò)適配器118以及各種外插卡120和121的其他部件之間提供了連接。其他部件(圖中沒有示出)包括USB或其他端口連接、CD驅(qū)動器、DVD馬區(qū)動器、電影刻錄裝置及類似的部件，也可以與1/0橋107相連接。在圖1PCI(外設(shè)部件互連)、PCIExpress(PCI-E)、AGP(加速圖像端口)、超傳輸或其他任何一種總線或點對點通信協(xié)議，并且不同裝置之間的連接可以4吏用本領(lǐng)域已知的不同協(xié)議。CPU102作為計算機系統(tǒng)IOO的控制處理器來運行，管理和協(xié)調(diào)其他系統(tǒng)部件的工作。特別是，CPU102發(fā)出控制多線程處理子系統(tǒng)112中的并行處理器134工作的命令。在一些實施例中，CPU102將用于并行處理器134的命令流寫入到命令緩沖器(圖中未示出)，該命令緩沖器可以位于系統(tǒng)存儲器104中、子系統(tǒng)存儲器138中、或者是CPU102和并行處理器134都可以訪問的其他存儲位置中。并行處理器134從命令緩沖器中讀出命令流，并且相對于CPU102的工作異步地執(zhí)行這些命令。系統(tǒng)存儲器104包括操作系統(tǒng)的執(zhí)行映像、裝置驅(qū)動器103和配置為用于由多線程處理子系統(tǒng)112執(zhí)行的CUDA代碼101。CUDA代碼101包含了意在多線程處理子系統(tǒng)112上執(zhí)行的編程指令。在本文的描述中，代碼指的是任何計算機代碼、指令、和/或可以用處理器執(zhí)行的函數(shù)。例如，在各種實施例中，所述代碼可以包含C代碼、C十+代碼等等。在一個實施例中，所述代碼可以包括一種計算機語言的擴展語言(例如，C、C++的擴展等等)。操作系統(tǒng)提供了用于管理和協(xié)調(diào)計算機系統(tǒng)IOO工作的詳細指令。裝置驅(qū)動器103提供了用于管理和協(xié)調(diào)多線程處理子系統(tǒng)112,特別是并行處理器134工作的詳細指令。另外，裝置驅(qū)動器103可以提供編譯功能，用來生成特別針對并行處理器134進行優(yōu)化的機器代碼。裝置驅(qū)動器103可以結(jié)合由NVIDIA公司提供的CUD八tm框架來提供。在一個實施例中，多線程處理子系統(tǒng)112包含了一個或多個并行處理器134,該并行處理器134可以例如用一個或多個集成電路裝置來實現(xiàn)，所述集成電路裝置例如是可編程處理器、專用集成電路(ASICs)。并行處理器134可以包括針對圖像和視頻處理進行優(yōu)化的電路，例如包括^L頻輸出電路和圖像處理單元(GPU)。在另一個實施例中，多線程處理子系統(tǒng)112可以和一個或多個其他的系統(tǒng)元件集成，例如存儲器橋105、CPU102和I/O橋107,以形成片上系統(tǒng)(SoC)。一個或多個并行處理器134可以將數(shù)據(jù)輸出到顯示裝置110上，或者每一個并行處理器34可以將數(shù)據(jù)輸出到一個或多個顯示裝置110上。并行處理器134有利地實現(xiàn)包括一個或多個處理內(nèi)核的高度并行處理器，每一個處理內(nèi)核能夠同時執(zhí)行大量的線程，其中每一個線程是一例程序，比如代碼101。并行處理器134能夠被編程以執(zhí)行與各種廣泛的應(yīng)用相關(guān)的處理任務(wù)，所述應(yīng)用包括但不限于線性和非線性數(shù)據(jù)變換、視頻和/或音頻數(shù)據(jù)的過濾、建模運算(例如，應(yīng)用物理法則來確定物體的位置、速度和其他屬性)、圖像渲染操作(例如，鑲嵌著色器、頂點著色器、幾何著色器和/或像素著色器編程)等等。并行處理器134可以將數(shù)據(jù)從系統(tǒng)存儲器104和/或本地子系統(tǒng)存儲器138傳輸?shù)奖镜?片上)存儲器，對數(shù)據(jù)進行處理，并將結(jié)果數(shù)據(jù)寫回到系統(tǒng)存儲器104和/或子系統(tǒng)存儲器138中，在那里這些數(shù)據(jù)能夠由包括CPU102或另一多線程處理子系統(tǒng)112的其他系統(tǒng)部件來訪問。并行處理器134可以設(shè)置有任意數(shù)量的子系統(tǒng)存儲器138，也可以不包括子系統(tǒng)存儲器138，并且可以使用子系統(tǒng)存儲器138和系統(tǒng)存儲器104的任意組合。例如，在統(tǒng)一存儲器架構(gòu)(UMA)的實施例中，并行處理器134可以是圖像處理器。在這些實施例中，將會設(shè)置極少的甚至不設(shè)置任何專用子系統(tǒng)存儲器138，并行處理器134將只使用或者幾乎只使用系統(tǒng)存儲器104。在UMA實施例中，并行處理器134可以被集成到橋芯片或處理器芯片中，或者是設(shè)置為具有高速鏈接(例如PCI-E)的分離的芯片，該芯片通過橋芯片或其他通信裝置將并行處理器134連接到系統(tǒng)存儲器104。7如上所述，多線程處理子系統(tǒng)112可以包括任意數(shù)量的并行處理器134。例如，多個并行處理器134可以設(shè)置在單個的外插卡上，或者多個外插卡可以與通信路徑113相連，或者一個或多個并行處理器134可以被集成到橋芯片中。當(dāng)存在有多個并行處理器134時，那些并行處理器134可以以高于單個并行處理器134可能達到的數(shù)據(jù)吞吐量來并行工作處理數(shù)據(jù)。包含有一個或多個并行處理器134的系統(tǒng)可以實現(xiàn)為各種配置和形式，包括桌上型電腦、筆記本電腦、或是手持個人計算機、服務(wù)器、工作站、游戲控制臺、嵌入式系統(tǒng)等等。在并行處理器134的一些實施例中，使用單指令多數(shù)據(jù)(SIMD)指令發(fā)送技術(shù)來支持大量線程的并行執(zhí)行，而無需提供多個獨立指令單元。在其他實施例中，使用單指令多線程(SIMT)技術(shù)來支持大量大體上同步化的線程的并行執(zhí)行。不同于其中所有處理引擎通常執(zhí)行相同指令的SIMD執(zhí)行方案，SIMT"l丸行允許不同的線程更易于通過給定的線程程序來跟隨發(fā)散的執(zhí)行路徑。本領(lǐng)域技術(shù)人員能夠理解的是，SIMD處理方案代表的是SIM丁處理方案的功能性子集。并行處理器134中的功能性單元支持多種運算，包括整數(shù)和浮點計算(例如，加法和乘法)、比較運算、布爾運算(AND、OR、XOR)、移位和各種代數(shù)函數(shù)的計算(例如，平面插值、三角、指數(shù)和對數(shù)函數(shù)等等)。傳輸?shù)轿挥诓⑿刑幚砥?34的處理內(nèi)核(未示出)中的特定處理單元(未示出)的指令序列構(gòu)成了線程，如在此前所定義的，在跨越處于一個處理內(nèi)核中的多個處理單元同時執(zhí)行的一定數(shù)量的線程集合在這里被稱作"線程組"。如在此所使用的，"線程組"指的是一組對不同的輸入數(shù)據(jù)執(zhí)行同一程序的線程，該組里的每一個線程被分配給處理內(nèi)核中不同的處理單元。線程組可以包含比處理單元的數(shù)量更少的線程，在此情況下，在該線程組正在進行處理的周期內(nèi)，一些處理單元將處于空閑狀態(tài)。線程組也可以包含比處理單元的數(shù)量更多的線程，在此情況下，處理將發(fā)生在多個時鐘周期上。由于每一個處理內(nèi)核能夠同時支持多達G個線程組，因此遵循在任意給定的時間上，多達GxM個線程組可以在處理內(nèi)核中執(zhí)4亍，這里M是并行處理器134內(nèi)的處理內(nèi)核的數(shù)量。另外，處理內(nèi)核中的多個相關(guān)的線程稱為"合作線程陣列，，("CTA，，)。CTA的尺寸通常由編程人員和CTA可用的硬件資源的數(shù)量來決定，所述硬件資源諸如存儲器或寄存器。CUDA編程模型反映的是GPU加速器的系統(tǒng)架構(gòu)。每一個線程都有專用的本地地址空間，并且每一個CTA共享的地址空間被用來在CTA中的線程之間傳遞數(shù)據(jù)。處理內(nèi)核還訪問片外"全局"存儲器，該存儲器可以例如包括子系統(tǒng)存儲器138和/或系統(tǒng)存儲器104。CUDA應(yīng)用程序的主機部分用常規(guī)的方法和工具來編譯，而內(nèi)核函數(shù)指定CTA處理。在最高層，CUDA存儲器模型將主機和設(shè)備存儲器空間分開，這樣主機代碼和內(nèi)核代碼就僅能直接訪問它們各自的存儲器空間。API(應(yīng)用編程接口)函數(shù)允許在主機和設(shè)備存儲器空間之間進行數(shù)據(jù)復(fù)制。在CUDA編程模型的共享存儲器CPU執(zhí)行過程中，控制CPU線程能夠在沒有潛在數(shù)據(jù)竟?fàn)幍那闆r下與并行CTA并行執(zhí)行。主機存儲器空間由C編程語言定義，并且設(shè)備存儲器空間被指定為全局的、恒定的、本地的、共享的和紋理的。所有線程可以訪問這些全局的、恒定的和紋理的存儲器空間。正如前面已經(jīng)解釋過的，對本地空間的訪問限于單個線程，并且對共享空間的訪問限于在CTA內(nèi)的線程。這種存儲器模型鼓勵在對于低等待時間的訪問時使用小存儲器空間，并且鼓勵對通常有更長等待時間的大存儲器空間進行明智地使用。CUDA壽呈序，比如代碼101,通常會纟皮組織成以一維、二維或三維形式(例如x、y和z)的一組同步或異步執(zhí)行的CTA。三元索引唯一地識別線程塊中的線程。線程塊自身由隱含定義的二元變量來區(qū)分。這些索引的范圍在運行時間時被定義，并且運行時間環(huán)境檢查這些索S1是否符合任何硬件限制。每一個CTA可以和其他CTA—起由并行處理器134并行執(zhí)行。很多CTA可以和每一個執(zhí)行一個或多個CTA的并行處理器134并行運行。運行時間環(huán)境負責(zé)根據(jù)要求來同步或異步地管理CUDA代碼101的執(zhí)行。CTA內(nèi)的線程通過使用共享存儲器和被稱為synchthreads()的壁壘同步圖元互相通信并同步。CUDA確保在線程塊內(nèi)的線程將同時存活，并為線程塊內(nèi)的線程提供了架構(gòu)以執(zhí)行快速壁壘同步和本地數(shù)據(jù)共享。在(由一維或多維定義的)CTA內(nèi)不同的線程塊對于其創(chuàng)建、執(zhí)行或退隱沒有順序的要求。另外，不允許并行CTA訪問系統(tǒng)調(diào)用，包括1/0。CUDA編程模型僅僅執(zhí)行并行CTA之間的全局同步，并為CTA中的塊之間的有限通信提供內(nèi)稟原子操作(intrinsicatomicoperations)。每個線程的主體被稱作內(nèi)核，用CUDA來指定，其在標(biāo)準(zhǔn)C中可以用存儲器模型注釋和壁壘同步圖元來表示。CUDA程序的語義是，每一個內(nèi)核由CTA內(nèi)的全部線程以按照由壁壘同步圖元所暗示的存儲器排序的順序來執(zhí)行。特別是，在壁壘同步圖元之前發(fā)生的CTA內(nèi)的全部共享存儲器基準(zhǔn)必須在壁壘同步圖元之后發(fā)生的任何共享存儲器基準(zhǔn)之前完成。內(nèi)核代碼中的每一個壁壘同步圖元實例在概念上代表了單獨的邏輯壁壘，并應(yīng)當(dāng)被處理為靜態(tài)的。當(dāng)CUDA線程可以采取不同的構(gòu)建分支時，在如果-否則(if-else)構(gòu)建的兩條路徑上都調(diào)用壁壘同步圖元是非法的。雖然線程塊內(nèi)的所有線程都會到達其中一個同步圖元，但它們代表的是單獨的壁壘，每個都要求要么全部線程到達壁壘，要么沒有線程到達壁壘。因此，這樣的內(nèi)核不會正確地執(zhí)行。更普遍地，如果同步圖元是包含在對于線程塊中不同的線程表現(xiàn)不一樣的任何控制流程構(gòu)建內(nèi)的話，則CUDA代碼不能確保被正確地執(zhí)行。圖2是根據(jù)本發(fā)明的一個實施例對計算機系統(tǒng)200進行說明的框圖。計算機系統(tǒng)200包括CPU202和系統(tǒng)存儲器204,兩者通過包含存儲器橋.205的總線路徑互相通信。存儲器橋205例如可以是北橋芯片，與I/O(輸入/輸出)橋107通過總線或其他通信路徑106(例如，超傳輸鏈接)相連接。CPU202產(chǎn)生輸出以在顯示裝置210(例如，常規(guī)的CRT或基于LCD的顯示器)上顯示。多線程處理子系統(tǒng)112不包括在計算機系統(tǒng)200中，并且CUDA代碼101沒有被改寫用于由通用處理器來執(zhí)行，所述通用處理器比如CPU202。CUDA代碼101被改寫用于由多線程處理子系統(tǒng)112執(zhí)行，并且使用翻i斧器220進行翻譯以產(chǎn)生翻譯后的代碼201,該翻譯后的代碼201不包括壁壘同步圖元。為了讓CPU202運行由代碼IOI表示的程序，代碼101首先必須要被翻譯成代碼201。然后該翻譯后的代碼可以由編譯器225編譯成用于由CPU202來執(zhí)行。編譯器225可以執(zhí)行專門針對于CPU202的優(yōu)化。翻譯代碼指的是將由第一計算機語言編寫的代碼轉(zhuǎn)換成由第二計算機語言編寫的代碼。編譯代碼指的是將由一種計算機語言(例如，源代碼)編寫的代碼轉(zhuǎn)換成由另一種計算機語言(例如，目標(biāo)代碼)編寫的代碼。翻譯器220結(jié)合圖3A進行描述，編譯器225結(jié)合圖4進行描述。編譯器225可以包含在裝置驅(qū)動器203中，該裝置驅(qū)動器203配置為在代碼101、代碼201和CPU202之間接口。將運行時間環(huán)境227配置成實現(xiàn)編譯后的代碼的功能，例如，輸入和輸出、存儲器管理等等。運行時間環(huán)境227還啟動編譯后的代碼用于由CPU202執(zhí)行。翻譯器220執(zhí)行優(yōu)化轉(zhuǎn)換，以便將跨越CUDA線程組的細粒度線程的操作串連成單個的CPU線程，而運行時間環(huán)境227將線程組調(diào)度為由CPU202并行處理的工作單元。10的這種可移植性的首要障礙是并行性的粒度。常規(guī)的CPU不支持單個CUDACTA所要求的上百個硬件線程的環(huán)境。因此，在通用CPU上實現(xiàn)CUDA編程模型的系統(tǒng)的首要目標(biāo)是要將任務(wù)級的并行性分配給可用的CPU內(nèi)核。同時，該系統(tǒng)必須將一個任務(wù)中的微線程整合成單個的CPU線程，以防止過度調(diào)度的開銷和頻繁的內(nèi)核間同步化。圖3A是根據(jù)本發(fā)明的一個實施例，將用于由多核圖像處理單元(例如多線程處理子系統(tǒng)112)執(zhí)行而編寫的代碼101翻譯為由通用處理器(例如CPU202)執(zhí)行的代碼201的方法步驟的流程圖。為了保存用于代碼101中的壁壘同步圖元語義，將翻譯器220配置為執(zhí)行如圖3A所示的一個或多個步驟。通過將圍繞在壁壘同步圖元周圍的代碼101進行劃分，翻譯器220將并行的線程"展開"，降低了共享狀態(tài)的使用，改進了用于存儲器訪;處i器執(zhí)行。在^^改變其目標(biāo)定為由多線程處理子系統(tǒng)、、U2執(zhí)行的CUDA代碼101的情況下，用CUP202來執(zhí)行代碼201有可能獲得很好的執(zhí)行效果。編譯器225可以利用由CPU202提供的向量指令容量，并且在為了執(zhí)行而編譯代碼201時進行優(yōu)化。在步驟300中，翻譯器220接收到用于由多核GPU執(zhí)行而編寫的代碼101,所述多核GPU例如多線程處理子系統(tǒng)112或者包括一個或多個并行處理器134的處理器，所述代碼101例如是CUDA代碼101。在步驟300中接收到的代碼可以表示為由邊連接的基本塊節(jié)點組成的控制流程圖。每個基本塊指定了由目標(biāo)環(huán)境(例如CPU202)執(zhí)行的才栗作。在步驟305中，翻譯器220將圍繞在壁壘同步圖元周圍的代碼101進行劃分，以生成劃分后的代碼。劃分后的代碼在圖3B和圖3C中示出，并且所述劃分過程結(jié)合那些圖來進行描述。同步劃分是一個代碼區(qū)域，在該區(qū)域內(nèi)的操作順序完全由該劃分內(nèi)的基本塊的控制流和數(shù)據(jù)流屬性來決定。劃分具有這樣的屬性，即線程循環(huán)能夠在一個劃分的周圍被插入以運行并行線程。通過將每個同步線程圖元替換成邊，將基本塊節(jié)點分成不同的劃分區(qū)域，所述控制流程圖可以用來生成同步劃分控制流程圖。在步驟310中，劃分后的代碼被分類，以便每一個語句被識別為收斂的或發(fā)散的。劃分后的代碼可以包括表達式和語句。表達式是一種計算，該計算可以包含常量、隱式線程身份(threadID)和由編程人員創(chuàng)建命名的變量，但是沒有副作用或賦值。簡單的語句定義為導(dǎo)致單一賦值的計算表ii達式。通用語句也能代表壁壘、控制流有條件的或循環(huán)構(gòu)建、或者是一系列語句塊。CTA的維度x、y和z通過代碼進行傳播以決定每個操作是否依賴于一個或多個CTA維度。引用了維度x、y和/或z中的線程身份(線程標(biāo)識符)的搡作被認為是發(fā)散的，這是因為引用了CTA維度的線程可以與執(zhí)行期間同一個CTA內(nèi)的其他線程所不同。例如，依賴于線程身份x(threadID.x)的操作對于x維度是發(fā)散的。另一個不依賴于線程身份x(threadID.x)的操作在x維度中是收斂的。發(fā)散的語句對于每一個它們所引用的CTA維度都要求有線程循環(huán)。在步驟315,使用分類信息來對劃分后的代碼進行性能優(yōu)化，從而產(chǎn)生優(yōu)化后的代碼。例如，一個劃分內(nèi)的指令可以被重新排序以融合才喿作，這樣具有同一類別的操作被放在一組，并能落入在步驟325中插入的同一個線程循環(huán)內(nèi)。操作按照這樣的方式排序，即在其變量向量中具有較少線程身份維度的操作放在依賴于較多線程身份維度的操作之前。這種重新排序是有效的，因為一個語句必須具有一個變量向量，該變量向量是它所依賴的語句的變量向量的超集。因此其變量向量中只有一個維度的語句不能依賴于任何在其變量向量中有不同的維度或有多于一個維度的語句。在步驟320中，根據(jù)需要將優(yōu)化后的代碼中的線程-本地存儲器基準(zhǔn)提升至陣列基準(zhǔn)，以保證每一個對象實例有唯一的位置來存放一個值。特別的，從一個分區(qū)輸送到另一個分區(qū)的數(shù)據(jù)需要進行復(fù)制，這樣使其對每一個分區(qū)都是可用的。符合下面條件之一的變量將被提升至陣列基準(zhǔn)具有交叉分區(qū)依賴性的本地變量(在一個分區(qū)中被賦值并且在另一個分區(qū)中被引用)。在步驟320中，翻譯器220將線程-本地存儲器基準(zhǔn)提升至陣列基準(zhǔn)。表1所示的程序包含同步化壁壘圖元和發(fā)散基準(zhǔn)。表1—global—voidfUnction(){intleftlndex,rightlndex;SharedMem[threadldX.x]=...;〃將值存儲到共享存儲器中l(wèi)eftlndex=...threadld.x..;rightlndex=...threadld.x;—synchthreads();=…(SharedMem[leftlndex]+SharedMem[rightlndex])/2.0;將表1所示的程序劃分成在同步線程圖元前的第一分區(qū)和在同步線程圖元后的第二分區(qū)。第二分區(qū)包括基準(zhǔn)(leftlndex和rightlndex),該基準(zhǔn)在第一分區(qū)中進行計算并依賴于CTA維度。如果不提升發(fā)散基準(zhǔn)，則第二分區(qū)將無法正確地使用由第一分區(qū)的最后一次迭代所計算的值。第二分區(qū)應(yīng)該使用為第一分區(qū)的threaded,x的每一次相應(yīng)的迭代所計算的值。為確保計算是正確的，將發(fā)散基準(zhǔn)按表2所示的進行提升。表2voidfunction(){for(inttid一x=0;tid—x<dimblock.X;tid—x++){SharedMem[tid.x]=...;〃將值存入存儲器中l(wèi)eftlndexArray[tid—x]=...threadld.x.rightlndexArray[tid一x]=…threadld.x;for(inttid一x=0;tid一x<dimblock.X;tid—x++){=…(SharedMem[leftlndexArray[tid—x]]+SharedMem[rightlndexArray[tid一x]])/20;在步驟325中，為在其變量向量中包含線程身份維度的那些語句生成線程循環(huán)。使用自適應(yīng)循環(huán)嵌套來同時評估相當(dāng)于循環(huán)互換、循環(huán)分裂和循環(huán)不變式消除的轉(zhuǎn)換，以獲得最佳的冗余消除效果。嵌套循環(huán)在線程身份元組的每一個維度的值上動態(tài)地生成，以最佳地適應(yīng)應(yīng)用，而不是假定特定的循環(huán)嵌套并在該嵌套的基礎(chǔ)之上對應(yīng)用進行評估。當(dāng)語句在步驟315中被排序之后，可以僅在一些語句周圍對線程身份維度生成循環(huán)，這些語句在其變量向量中包含該維度。為了消除循環(huán)開銷，翻譯器220可以融合相鄰的語句組，在這些相鄰的語句組中一個語句組所具有的變量向量是另一個語句組的變量向量的子集。圖3B是根據(jù)本發(fā)明的一個實施例，說明被翻譯成為劃分后的代碼350的輸入代碼101的概念圖。輸入代碼330^皮配置為用于由多線程處理子系統(tǒng)112執(zhí)行，并且包括由同步壁壘指令336進行分隔的代碼序列331和332。CTA中的所有線程將在其中任一線程開始執(zhí)行代碼序列332之前完成代碼序列331的執(zhí)行。翻譯器220將輸入代碼330進行劃分以產(chǎn)生劃分后的代碼350,其中分區(qū)351包括由代碼序列331代表的指令，分區(qū)352包括由代碼序列332代表的指令。當(dāng)劃分后的代碼350由本身不支持同步壁壘指令的通用處理器執(zhí)行時，線程循環(huán)353^C插入到分區(qū)352周圍，以確保同步語義得到維持。此例中，代碼分區(qū)351包含收斂基準(zhǔn)，分區(qū)352可以包含發(fā)散基準(zhǔn)。因此，線程循環(huán)353在分區(qū)352周圍被插入。在圖3A的步驟325中，翻譯器220將線程循環(huán)(比如線程循環(huán)353)插入到經(jīng)過優(yōu)化的代碼中，以生成被翻譯成由CPU202執(zhí)4亍的代碼201。每一個分區(qū)可以具有為每一個CTA維度插入的線程循環(huán)。同步分區(qū)和線程循環(huán)插入的例子在表3和4中示出。表3示出的程序被翻譯為表4示出的程序。<table>tableseeoriginaldocumentpage14</column></row><table>表3的程序使用了明確的同步來確保存儲器在CTA中各個線程之間的正確共享。翻譯器220將程序劃分為兩個分區(qū)，每一個分區(qū)依賴于xCTA維度。因此，線程循環(huán)在這兩個分區(qū)的每一個周圍被插入以確保翻譯后的程序按正確的順序來執(zhí)行操作。<table>tableseeoriginaldocumentpage14</column></row><table>=…(SharedMem[tid—x]+SharedMem[tid—x—1])/2.0;用于將程序翻譯為由通用處理器來執(zhí)行的更加簡單的才支術(shù)是給每個CTA維度插入明確的線程循環(huán)，這樣就不必為在同一分區(qū)內(nèi)的基準(zhǔn)而決定維度的依賴性。例如，表5示出的程序被翻譯為如表6所示的程序。注意，表5中插入的一個或多個線程循環(huán)可能不是必須的，這是因為不用決定維度的依賴性就可以生成程序。表5—global—voidfonction(){Sharedl=...=Sharedl表6voidfunction()(for(inttid一x=0;tid_x<dimblock.X;tidx+十){for(inttid_y=0;tid_y<dimblock.Y;tid_y++){for(inttid—z=0;tid_z<dimblock.Z;tid—z++){Sharedl=...=Sharedl圖3C是才艮據(jù)本發(fā)明的一個實施例，說明翻"^為優(yōu)化后的代碼360的輸入代碼333的概念圖。將輸入代碼333配置為用于由多線程處理子系統(tǒng)112執(zhí)行，并且包括由同步壁壘指令335進行分隔的代碼序列334和338。CTA中的所有線程會在其中任一線程開始執(zhí)行代碼序列338之前完成代碼其中分區(qū)361包括由代碼序列334代表的指令，分區(qū)362、364和365包括由代碼序列338代表的指令。15分區(qū)362包括在第一CTA維度上發(fā)散的第一部分指令。分區(qū)364包括收斂的第二部分指令。分區(qū)365包括在第二CTA維度上發(fā)散的第三部分指令。當(dāng)劃分后的代碼360由本身不支持同步壁壘指令的通用處理器執(zhí)行時，線程循環(huán)363被插入到分區(qū)362周圍以確保同步語義得到維持。線程循環(huán)363在第一CTA維度上進行迭代。線程循環(huán)366在分區(qū)365周圍被插入以便在第二CTA維度上進行迭代。表7示出了CUDA內(nèi)核的例子，表8示出了用于由通用處理器執(zhí)行的CUDA內(nèi)核的翻譯。該示例性內(nèi)核與一系列小矩陣相乘。每一個線程塊計算系列中一個小矩陣的相乘，而每一個線程為其塊計算結(jié)果矩陣的一個元素。表7示例的CUDA內(nèi)核(1)一global—small_mm—list(float*A—list,float*B—list,，constintsize){(2)floatsum;(3)intmatrix—start,col，row,out—index,i;(4)martrix—start=blockldx.x*size*size;(5)col=matrix—start+threadIDx.x;(6)row—matrix—start+threadldx.y*size);(7)sum=0.0;(8)for(i=0;i<size;i++)(9)sum+=A—list[row+i]*B_list[col+(i*size)];//在重寫輸入數(shù)據(jù)之前進行同步(10)—syncthread();(11)out—index=matrix—start+(threadldx.y*size)+threadldx.x;(12)A—list[out一index]=sum;要注意的是表7中位于第(9)行的語句具有變量向量(x,y),因為col(列)依賴于x維度且row(行)依賴于y維度。z維度從來不^f吏用，所以沒有循環(huán)被插入在z上進行迭代。通常的成本分析技術(shù)可以用來決定如在表7中示出的示例性內(nèi)核中的語句5和6這樣的情況。由于每一個僅依賴于一個線程身份維度，因此選擇x和y索引循環(huán)中的任一嵌套順序?qū)⑵仁拐Z句的冗余執(zhí)行，或者是分區(qū)的主循環(huán)嵌套之外的冗余循環(huán)。表8翻譯后的CUDA內(nèi)核(1)一global—small—mm—list(float*A—list,float*B—list,,constintsize)(2)floatsum[];(3)intmatrix一start[],col[]，row[],out—index,i;(4)matrix一start[threadID]=blocklDx.x*size*size;for(threadID.x=0;threadID.x<blockDim.x;threadID.x++){(5)col[threadID]=matrix—start+threadIDx.x;for(threadID.y=0;threadID.y<blockDim.y;threadID.y++){(6)row[threadID]=matrix_start[threadID]+(threadIDx.y*size);(7)sum[threadID]=0.0;(8)for(i[threadID]=0;i<size;i++)(9)sum[threadID]+=A—list[讓[threadID〗+i]*BJist[col[threadID]+(i*size)];(10)for(threadID.x=0;threadID.x<blockDim.x;threadID.x++){for(threadID.y=0;threadID.y<blockDim.y;threadID.y++){(11)out—index=matrix一start[threadID]+(threadID.y*size)+threadID.x;(12)A—list[out一index]=sum[threadID];圖4是根據(jù)本發(fā)明的一個實施例，由通用處理器(比如CPU202)執(zhí)行翻譯后的代碼201的方法步驟的流程圖。在步驟400中，編譯器225將翻譯后的代碼201進行編譯，選擇性地執(zhí)行針對CPU的優(yōu)化，以生成編譯后的代碼。在步驟405中，在CPU202中可用的執(zhí)行內(nèi)核400的數(shù)量由設(shè)備驅(qū)動器203決定。為了提高性能，翻譯后的代碼201可以自動地進行縮17放以用于在可用4丸行內(nèi)核上執(zhí)行。在步驟410中，運行時間環(huán)境227或裝置驅(qū)動器203將CPU202配置為啟用執(zhí)行翻譯后的代碼201的該數(shù)量的執(zhí)4亍內(nèi)核。運行時間環(huán)境227可以生成多個操作系統(tǒng)(OS)運行時間線程，該運行時間線程能夠由環(huán)境變量控制。默認狀態(tài)下，系統(tǒng)中內(nèi)核的數(shù)量可以用作OS運行時間線程的數(shù)量。在步驟410中，可以對將要啟動的CUDA線程的數(shù)量進行評估，并且在統(tǒng)計上劃分為運行時間線程的數(shù)量。每一個運行時間線程順序地執(zhí)行一部分編譯后的代碼，并在壁壘上等待。當(dāng)所有的運行時間線程都到達了壁晝，則CTA完成。在步驟415中，運行時間環(huán)境227或裝置驅(qū)動器203啟動編譯后的代碼用于由CPU202執(zhí)行。翻譯器220、編譯器225以及運行時間環(huán)境227被用來將CUDA應(yīng)用程序轉(zhuǎn)換為由通用CPU執(zhí)行的代碼。CUDA編程模型支持批量的同步任務(wù)并行，其中每個任務(wù)由細粒度SPMD線程組成。CUDA編程模型的使用已經(jīng)被限制于那些愿意為由GPU執(zhí)行而編寫特殊代碼的程序員。這些特殊代碼可以轉(zhuǎn)換為由通用CPU來執(zhí)行，并不需要程序員來重新編寫CUDA應(yīng)用程序。由CUDA支持的三個關(guān)鍵的抽象概念是SPMD線程塊、壁壘同步和共享的存儲器。翻譯器220將遍及CUDA線程塊的細粒度線程串行化為單一的CPU線程，并且執(zhí)行優(yōu)化變形以轉(zhuǎn)換CUDA應(yīng)用程序。雖然前面所述的是本發(fā)明的實施例，在不背離本發(fā)明基本范圍的前提下可以設(shè)計出本發(fā)明更多的實施例。例如，本發(fā)明的某些方面可以由^^更件或軟件來實現(xiàn)，或者是由硬件與軟件結(jié)合在一起來實現(xiàn)。本發(fā)明的一個實施例可以實現(xiàn)為計算機系統(tǒng)所使用的程序產(chǎn)品。程序產(chǎn)品的程序?qū)嵤├墓δ?包括在此描述的方法)進行定義，并且能夠被包含在各種各樣的計算機可讀存儲介質(zhì)內(nèi)。說明性的計算機可讀存儲介質(zhì)包括但不限于(i)信息在其上永久保存的非可寫存儲介質(zhì)(例如，計算機內(nèi)的只讀存儲裝置，如可被CD-ROM驅(qū)動器讀出的CD-ROM盤、閃存、ROM芯片或者任意類型的固態(tài)非易失性半導(dǎo)體存儲器)；以及(ii)其上存儲有可改變的信息的可寫存儲介質(zhì)(例如，磁盤驅(qū)動器內(nèi)的軟盤或硬盤驅(qū)動器或任意類型的固態(tài)隨機存取半導(dǎo)體存儲器)。當(dāng)攜帶有引導(dǎo)本發(fā)明的功能的計算機可讀指令時，這樣的計算機可讀存儲介質(zhì)就是本發(fā)明的實施例。因此，本發(fā)明的范圍由以下的權(quán)利要求來界定。18權(quán)利要求1.一種計算系統(tǒng)，所述計算系統(tǒng)被配置用于將應(yīng)用程序翻譯為由通用處理器來執(zhí)行，所述計算系統(tǒng)包括處理器，所述處理器被配置用于執(zhí)行翻譯器；和系統(tǒng)存儲器，所述系統(tǒng)存儲器與所述處理器連接并被配置用于存儲所述翻譯器、第一應(yīng)用程序和第二應(yīng)用程序，所述第一應(yīng)用程序使用并行編程模型來編寫以用于在多核圖像處理單元上執(zhí)行，所述第二應(yīng)用程序被配置用于由通用處理器執(zhí)行，以及所述翻譯器被配置用于接收所述第一應(yīng)用程序；將所述第一應(yīng)用程序劃分為同步獨立指令的區(qū)域以生成劃分后的應(yīng)用程序；以及在至少其中一個劃分后的應(yīng)用程序周圍插入循環(huán)，以生成所述第二應(yīng)用程序，其中所述循環(huán)在合作線程陣列維度上進行迭代，所述合作線程陣列維度對應(yīng)于由多核圖像處理單元內(nèi)的并行處理器同時執(zhí)行的多個線程。2.根據(jù)權(quán)利要求1的計算系統(tǒng)，其特征在于，所述翻譯器進一步被配置用于識別所述應(yīng)用程序內(nèi)的同步壁壘指令。3.根據(jù)權(quán)利要求2的計算系統(tǒng)，其特征在于，劃分后的應(yīng)用程序的第一區(qū)域包括位于所述同步壁壘指令之前的指令，劃分后的應(yīng)用程序的第二區(qū)域包括位于所述同步壁壘指令之后的指令。4.根據(jù)權(quán)利要求3的計算系統(tǒng)，其特征在于，所述插入循環(huán)的步驟包括在所述劃分后的應(yīng)用程序的第一區(qū)域周圍插入第一循環(huán)，以確保在所述合作線程陣列中的任一線程開始執(zhí)行所述劃分后應(yīng)用程序的第二區(qū)區(qū)域的執(zhí)行。5.根據(jù)權(quán)利要求2的計算系統(tǒng)，其特征在于，所述第一應(yīng)用程序表示為包括由邊連接的基本塊節(jié)點的控制流程圖。6.根據(jù)權(quán)利要求5的計算系統(tǒng)，其特征在于，所述翻譯器進一步被配置為將所述同步壁壘指令替換成邊，以便將所述基本塊節(jié)點分隔成與第一區(qū)域?qū)?yīng)的第一基本塊節(jié)點和與第二區(qū)域?qū)?yīng)的第二基本塊節(jié)點。7.根據(jù)權(quán)利要求1的計算系統(tǒng)，其特征在于，所述翻譯器進一步被配置用于將所述劃分后的應(yīng)用程序進行分類，以便識別每個語句相對于所述合作線程陣列維度是收斂的還是發(fā)散的。8.根據(jù)權(quán)利要求1的計算系統(tǒng)，其特征在于，所述翻譯器進一步被配置用于在所述劃分后的應(yīng)用程序的至少一個區(qū)域周圍插入額外的循環(huán)，以生成所述第二應(yīng)用程序，其中所述額外的循環(huán)在不同的合作線程陣列維度上進行迭代。9.根據(jù)權(quán)利要求1的計算系統(tǒng)，其特征在于，所述第一應(yīng)用程序是CUDA(計算統(tǒng)一設(shè)備架構(gòu))應(yīng)用程序。10.根據(jù)權(quán)利要求1的計算系統(tǒng)，其特征在于，進一步包括運行時間庫，所述運行時間庫將所述通用處理器配置成能使所述通用處理器內(nèi)的多個執(zhí)行內(nèi)核執(zhí)行所述第二應(yīng)用程序。全文摘要本發(fā)明公開了一種為了用于由通用處理器執(zhí)行而對應(yīng)用程序重定目標(biāo)。本發(fā)明的一個實施例提出了一種技術(shù)，用于將使用并行編程模型編寫的用于在多核圖像處理單元(GPU)上執(zhí)行的應(yīng)用程序翻譯成用于由通用中央處理單元(CPU)執(zhí)行。所述應(yīng)用程序中依賴于多核GUP的特定特征的部分由翻譯器轉(zhuǎn)換成由通用CPU執(zhí)行。所述應(yīng)用程序被劃分為同步獨立指令區(qū)域。所述指令被分類為收斂的或發(fā)散的，并且在區(qū)域之間共享的發(fā)散存儲器基準(zhǔn)被復(fù)制。插入線程循環(huán)，以確保在由通用CPU執(zhí)行期間各種線程之間存儲器的正確共享。文檔編號G06F9/45GK101556544SQ200910117898公開日2009年10月14日申請日期2009年4月9日優(yōu)先權(quán)日2008年4月9日發(fā)明者巴斯蒂安·約翰·馬特烏斯·阿特斯,維諾德·格羅夫,賈揚特·B·科爾希,邁克爾·墨菲,道格拉斯·塞勒,鮑里斯·柏林申請人:輝達公司