專利名稱：由通用處理器執(zhí)行重定目標(biāo)的圖形處理器加速代碼的制作方法
技術(shù)領(lǐng)域：
本發(fā)明的實(shí)施例大體上涉及編譯器程序，更具體地涉及一種應(yīng)用程序， 所述應(yīng)用程序被編寫用于由多核圖形處理器執(zhí)行，并且被重定目標(biāo)用于由 具有共享存儲(chǔ)器的通用處理器執(zhí)行。
背景技術(shù)：
現(xiàn)代圖像處理系統(tǒng)通常包括配置為以多線程方式執(zhí)行應(yīng)用程序的多核 圖像處理單元(GPU)。這種圖像處理系統(tǒng)還包括存儲(chǔ)器，該存儲(chǔ)器具有在 執(zhí)行線程之間共享的部分和每個(gè)線程專用的部分。
NVIDIA的CUDA (計(jì)算統(tǒng)一設(shè)備架構(gòu))技術(shù)提供了 一種C語言環(huán) 境，使得編程人員和開發(fā)人員能夠編寫解決復(fù)雜計(jì)算問題的軟件應(yīng)用程序， 這些復(fù)雜計(jì)算問題諸如視頻和音頻編碼、油氣勘探建模和醫(yī)學(xué)成像。應(yīng)用 程序配置為由多核GPU并行執(zhí)行，并且通常依靠多核GPU的特定特征。 由于同樣的特定特征在通用中央處理單元(CPU)中是不可用的，因此用 CUDA編寫的軟件應(yīng)用程序可能無法移植到通用CPU上運(yùn)行。
如前所述，在本領(lǐng)域中需要這樣一種技術(shù)，這種技術(shù)不需要編程人員 對(duì)應(yīng)用程序進(jìn)行修改就能讓用并行編程模型編寫的用于在多核GPU上執(zhí) 行的應(yīng)用程序能夠在通用CPU上運(yùn)行。

發(fā)明內(nèi)容
本發(fā)明的一個(gè)實(shí)施例給出了 一種用于配置通用處理器以執(zhí)行翻譯后的 應(yīng)用程序的方法。該方法包括接收翻譯后的應(yīng)用程序的步驟，所述翻譯后 的應(yīng)用程序是由用并行編程模型編寫的用于在多核圖像處理單元上執(zhí)行的應(yīng)用程序轉(zhuǎn)換而來的，和編譯翻譯后的應(yīng)用程序的步驟，以生成用于由通 用處理器執(zhí)行的編譯后的代碼?？捎脕韴?zhí)行編譯后的代碼的通用處理器內(nèi) 執(zhí)行內(nèi)核的數(shù)量被確定，并且將通用處理器配置成啟用該數(shù)量的執(zhí)行內(nèi)核。 編譯后的代碼被啟動(dòng)用于由包含該數(shù)量的執(zhí)行內(nèi)核的通用處理器來執(zhí)行。
在此披露的方法的一個(gè)優(yōu)勢(shì)在于用并行編程模型編寫的用于由多核 GPU執(zhí)行的應(yīng)用程序可以無需修改而移植到通用CPU上。為了能夠由通 用CPU執(zhí)行，將應(yīng)用程序中依賴于多核GPU的特定特征的部分由翻譯器 進(jìn)行轉(zhuǎn)換。將應(yīng)用程序劃分為同步獨(dú)立指令的區(qū)域。將這些指令分類為收 斂的或發(fā)散的，并且將這些區(qū)域之間共享的發(fā)散存儲(chǔ)器基準(zhǔn)進(jìn)行復(fù)制。在 由通用CPU執(zhí)行的過程中，插入線程循環(huán)以確保在各種線程之間能夠正確 地共享存儲(chǔ)器。


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