TWI531974B - 管理巢狀執行串流的方法和系統 - Google Patents

Description

管理巢狀執行串流的方法和系統

本發明一般係關於電腦架構，尤其係關於管理巢狀執行串流的方法及系統。

在具有中央處理單元(CPU，central processing unit)以及圖形處理單元(GPU，graphics processing unit)的傳統計算系統中，該CPU決定哪一個特定計算任務由GPU執行並且用什麼順序執行。GPU計算任務通常包括橫跨並行資料集的高並行、高類似操作，例如影像或影像集。在傳統GPU執行模型內，該CPU利用選擇一對應的執行緒程式並且指示該GPU執行該執行緒程式的一組並行實例，初始一特定計算任務。在傳統GPU執行模型內，該CPU通常為可在該GPU上起始一執行緒程式執行的唯一實體。所有執行緒實例執行完成之後，該GPU需通知該CPU，並且等待該CPU發出另一個計算任務。通知該CPU並且等待下一個計算任務通常為一阻礙、序列化的操作，這讓該GPU內某些資源處於一暫時待命的狀態，因此降低整體系統效能。

在某些情況下，利用將順序計算任務佇存在一推送緩衝區內，這樣該GPU不用等待該CPU就可執行工作，如此可改善效能。當該CPU產生工作給該GPU的速度足夠快到該推送緩衝區內有工作等待給GPU執行時，不管該GPU是否可開始新任務，則包含固定資料流處理管線的計算任務可從此推送緩衝區模型中受益。不過，由該CPU啟動之資料相依計算任務仍舊保有GPU結果、CPU任務管理以及後續GPU任務執行之間的順序相依性。此問題的一種解決方案為提供一種機制，讓GPU執行緒程式佇放額外計算任務，不需要該CPU介入，並且等待這些計算任務完成。不過，這種方式有許多缺點。首先， CPU傳統上具有動態分配記憶體的手段，但是GPU並沒有。當該GPU啟動新計算任務時，將記憶體分配給該等計算任務，以儲存該任務執行期間存取的範圍以及參數資訊。在這種案例中，該GPU連結該CPU來分配記憶體給該新計算任務。然後，該GPU在佇存新任務之前，必須等待該CPU分配記憶體給該計算任務，效能因而可能降低。

第二，該CPU和GPU可發出新計算任務進入該推送緩衝區，可能發生僵持的情況。該CPU可佔據至該GPU的所有通訊通道，以便用於佇存新計算任務。然後，該GPU可佇存一存取該CPU的新計算任務，以便完成作業。在這種案例中，該CPU在釋出任何該等通訊通道之前，會等待一GPU任務完成，然而在同意透過一個該等受阻通訊通道將該任務存取至該CPU之前該GPU任務將無法被完成，因此造成僵持情況。

最後，佇存新計算任務並且從該推送緩衝區拉出任務來執行通常運用鎖定操作，確定任務已經依序執行，並且該推送緩衝區內的資訊已經正確保留與管理。雖然GPU執行類似的鎖定操作，但是鎖定操作本質上就緩慢。若該GPU利用鎖定操作同時佇列新任務，則對於系統效能會有負面衝擊。

如先前所例示，業界需要一種允許GPU更有效率佇存工作來執行之技術。

本發明的一個具體實施例揭示一種用於處理由執行緒第一群組所執行並且儲存在複數個任務中繼資料描述器佇列(TMDQ，task metadata descriptor queue)內複數個任務之電腦實施方法。該方法包含：接收包含在該等複數個任務內的一第一任務已經完成之通知；以及在一共同處理單元內決定所有任務是否都含在該等複數個任務之一子集內，並且關聯於已經執行的一第一TMDQ。若該子集內含的所有任務都尚未執行，則 該方法另包含啟動該等複數個任務內含的一第二任務。若該子集內含的所有任務都已執行，則該方法另包含更新一第一資料結構內與該第一TMDQ相關聯的一指標器；決定該等複數個任務內含的一第三任務已經佇存在該第一TMDQ內；以及啟動該第三任務。

本發明技術的一個優點為GPU可將計算任務佇存在任務佇列內，同時建立一任意數量的新任務佇列至任何任務巢狀階層，不用該CPU介入。由於在該CPU建立並佇存任務時該GPU不用等待，所以處理效率獲得增強。

在以下描述中，揭示許多特定細節以對本發明有更徹底之理解。但是，精通技術人士應該了解，在無一或多個該等特定細節之下還是可實施本發明。

系統概觀

第一圖為例示設置來實施本發明一或多個態樣的電腦系統100之方塊圖。電腦系統100包含一中央處理單元(CPU，central processing unit)102，以及通過一互連路徑通訊的一系統記憶體104，其中該路徑可包含一記憶體橋接器105。例如可為北橋晶片的記憶體橋接器105透過匯流排或其他通訊路徑106(例如超傳輸(HyperTransport)連結)，連接至一I/O(輸入/輸出)橋接器107。例如可為南橋晶片的I/O橋接器107接收來自一或多個使用者輸入裝置108(例如鍵盤、滑鼠)的使用者輸入，並透過通訊路徑106和記憶體橋接器105將該輸入轉送至CPU 102。一並行處理子系統112透過匯流排或第二通訊路徑113(例如周邊組件互連(PCI，Peripheral Component Interconnect)Express、加速圖形連接埠(Accelerated Graphics Port)或超傳輸連結)耦合至記憶體橋接器105；在一個具體實施例內，並行處理子系統112為傳遞畫素給顯示裝置110(例如 傳統陰極射線管或液晶監視器)的圖形子系統。系統磁碟114也連接至I/O橋接器107。一開關116提供I/O橋接器107與其他組件，像是網路配接器118以及許多外接卡120和121之間的連接。其他組件(未明確顯示)，包含萬用序列匯流排(USB，universal serial bus)或其他連接埠連接、CD(compact disc)光碟機、數位影音光碟(DVD，digital video disc)機、影片記錄裝置等等，也可連接至I/O橋接器107。第一圖內顯示的許多通訊路徑，包含特地指名的通訊路徑106和113，都可使用任何合適的協定來實施，例如PCI Express、AGP(加速圖形連接埠，Accelerated Graphics Port)、超傳輸或任何其他匯流排或點對點通訊協定，以及不同裝置之間的連接都可使用業界內已知的不同協定。

在一個具體實施例內，並行處理子系統112併入將圖形與視訊處理最佳化的電路，包含例如視訊輸出電路，並且構成一圖形處理單元(GPU，graphics processing unit)。在另一具體實施例內，並行處理子系統112併入將一般用途處理最佳化的電路，同時保留底層計算架構，本文內有更詳細描述。尚且在另一具體實施例內，並行處理子系統112可在單一子系統內合併一或多個其他系統元件，例如結合記憶體橋接器105、CPU 102和I/O橋接器107來形成一晶片上系統(SoC，system on chip)。

本文中顯示的該系統為例示，所以可進行改變與修改。包含橋接器的數量與配置、CPU 102的數量以及並行處理子系統112的數量這類連接拓撲可依照需求修改。例如在某些具體實施例內，系統記憶體104直接連接至CPU 102，而不是透過橋接器，並且其他裝置透過記憶體橋接器105以及CPU 102與系統記憶體104通訊。在其他替代拓撲中，並行處理子系統112連接至I/O橋接器107或直接連接至CPU 102，而不是連接至記憶體橋接器105。同樣在其他具體實施例中，I/O橋接器107和記憶體橋接器105可整合成為單一晶片，替代現有的 一或多個分散裝置。大型組件可包含二或多個CPU 102以及二或多個並行處理子系統112。本文中顯示的該等特定組件為選擇性；例如可支援任何數量的外接卡或周邊裝置。在某些具體實施例內，省略開關116，並且網路配接器118和外接卡120、121都直接連接至I/O橋接器107。

第二圖例示根據本發明一個具體實施例的並行處理子系統112。如所示，並行處理子系統112包含一或多個並行處理單元(PPU，parallel processing unit)202，每一個都耦合至本機並行處理(PP，parallel processing)記憶體204。一般來說，並行處理子系統包含數量為U的PPU，其中U1。(本文中多個類似物件都用參考號碼標示出該物件，並且用括號標示出該物件為相類似物件中的數目。)PPU 202以及並行處理記憶體204可使用一或多個積體電路裝置來實施，例如可程式處理器、特殊應用積體電路(ASIC，application specific integrated circuit)或記憶體裝置，或以任何其他技術可行方式。

請再次參閱第一圖及第二圖，在一些具體實施例內，並行處理子系統112內的某些或全部PPU 202為具有彩現管線的圖形處理器，其可設置成執行許多操作，這些操作有關透過記憶體橋接器105和第二通訊路徑113，從CPU 102及/或系統記憶體104供應的圖形資料產生畫素資料、與本機並行處理記憶體204(可用來當成圖形記憶體，包含例如傳統畫框緩衝區)互動來儲存與更新畫素資料、傳遞畫素資料給顯示裝置110等等。在某些具體實施例內，並行處理子系統112可包含操作當成圖形處理器的一或多個PPU 202，以及用於一般用途計算的一或多個其他PPU 202。該等PPU可一致或不同，並且每一PPU都可擁有一專屬的並行處理記憶體裝置或無專屬的並行處理記憶體裝置。並行處理子系統112內的一或多個PPU 202可輸出資料至顯示裝置110，或並行處理子系統112內的每一PPU 202都可輸出資料至一或多個顯示裝置110。

在操作上，CPU 102為電腦系統100的主要處理器，控制與協調其他系統組件的操作。尤其是，CPU 102發出指令控制PPU 202的操作。在某些具體實施例內，CPU 102將每一PPU 202的指令串流寫入至一資料結構(第一圖或第二圖內未明確顯示)，其可位於系統記憶體104內、並行處理記憶體204內或可存取CPU 102和PPU 202的其他儲存位置內。每一資料結構的指標已經寫入一推送緩衝區，開始在該資料結構內處理該指令串流。PPU 202讀取來自一或多個推送緩衝區的指令串流，然後非同步執行關於CPU 102的操作之指令。利用應用程式透過裝置驅動程式103來控制不同推送緩衝區的排程，指定每一推送緩衝區的執行優先順序。

此時請回頭參閱第二圖以及第一圖，每一PPU 202都包含一I/O(輸入/輸出)單元205，其透過連接至記憶體橋接器105(或在一個替代具體實施例內，直接至CPU 102)的通訊路徑113，與電腦系統100的剩餘組件通訊。PPU 202與電腦系統100剩餘組件的連接也可改變。在某些具體實施例內，並行處理子系統112實施成為可插入電腦系統100內擴充槽的外接卡。在其他具體實施例內，PPU 202可與例如記憶體橋接器105或I/O橋接器107這類匯流排橋接器整合在單一晶片上。在其他具體實施例內，PPU 202的某些或全部元件可與CPU 102整合在單一晶片上。

在一個具體實施例內，通訊路徑113為一PCI Express連結，其中分配專屬通道給每一PPU 202，如業界內所熟知。在此也可使用其他通訊路徑。一I/O單元205產生在通訊路徑113上傳輸的封包(或其他信號)，也從通訊路徑113上接收所有傳入封包(或其他信號)，將該等傳入封包導引至PPU 202的適當組件。例如：有關處理任務的指令可導引至主介面206，而有關記憶體操作的指令(例如讀取或寫入並行處理記憶體204)可導引至記憶體橫桿單元210。主介面206讀取每一推送緩衝 區，並將該推送緩衝區內儲存的該指令串流輸出至一前端212。

每一PPU 202均可包含一高並行處理架構。PPU 202(0)包含一處理叢集陣列230，其包含數量為C的一般處理叢集(GPC，general processing cluster)208，其中C1。每一GPC 208都可同時執行大量的(例如數百或數千)執行緒，其中每一執行緒都是一程式的實例。在許多應用當中，不同的GPC 208可分配用於處理不同種類的程式，或用於執行不同種類的計算。GPC 208的分配根據針對每一種程式或計算所關聯的工作負擔而改變。

GPC 208從任務/工作單元207內的工作分配單元當中接收要執行的處理任務，該工作分配單元接收指標來處理被編碼為任務中繼資料(TMD，task metadata)並儲存在記憶體內之任務。TMD的該指標包含在指令串流內，該指標係被儲存當成一推送緩衝區並由前端單元212從主介面206接收。可編碼為TMD的處理任務包含要處理的資料索引，以及定義如何處理該資料的狀態參數與指令(例如要執行哪個程式)。任務/工作單元207從前端212接收任務，並且確定在每一TMD指定的處理開始之前已經將GPC 208設置成有效狀態。一優先順序可指定給每一TMD，用來排定該處理任務的執行時間。處理任務也可從處理叢集陣列230接收。該TMD可選擇性地包含一參數，其控制該TMD加入處理任務清單(或處理任務指標清單)的頭部或尾部，藉此提供優先順序之上的另一控制等級。

記憶體介面214包含數量為D的分割單元215，這些單元每一個都直接耦合至一部分並行處理記憶體204，其中D 1。如所示，分割單元215的數量一般等於動態隨機存取記憶體(DRAM，dynamic random access memory)220的數量。在其他具體實施例內，分割單元215的數量可不等於記憶體裝置的數量。本領域之習知技藝者將可了解，DRAM 220可用其他合適的儲存裝置取代，並且可為一般傳統設計，因此省略其詳細說 明。像是畫框緩衝區或紋理地圖這類彩現目標可透過DRAM 220儲存，允許分割單元215並行寫入每一彩現目標的部分，以有效使用並行處理記憶體204的可用頻寬。

任一GPC 208都可處理寫入並行處理記憶體204內任一DRAM 220的資料。橫桿單元210設置成將每一GPC 208的輸出繞送至任意分割單元215的輸入或至其他GPC 208以進一步處理。GPC 208透過橫桿單元210與記憶體介面214通訊，來讀取或寫入許多外部記憶體裝置。在一個具體實施例內，橫桿單元210連接至記憶體介面214來與I/O單元205通訊，以及連接至本機並行處理記憶體204，藉此讓不同GPC 208內的處理核心與系統記憶體104或不在PPU 202本機上的其他記憶體通訊。在第二圖所示的具體實施例中，橫桿單元210直接連接I/O單元205。橫桿單元210可使用虛擬通道，以分隔GPC 208與分割單元215之間的流量串流。

GPC 208同樣可程式編輯來執行有關廣泛應用的處理任務，包含但不受限於線性與非線性資料傳輸、視訊及/或音訊資料篩選、模型化運算(例如套用實體規則來決定位置、速度以及其他物體屬性)、影像彩現運算(例如曲線細分著色、影點著色、幾何著色及/或畫素著色程式)等等。PPU 202可從系統記憶體104及/或本機並行處理記憶體204將資料傳輸進入內部(晶片上)記憶體、處理該資料並將結果資料寫回系統記憶體104及/或本機並行處理記憶體204，其中這種資料可由其他系統組件存取，包含CPU 102或另一並行處理子系統112。

一PPU 202可有任何數量的本機並行處理記憶體204，包含非本機記憶體，並且可在任何情況下使用本機記憶體和系統記憶體。例如：PPU 202可為統一記憶體架構(UMA，unified memory architecture)具體實施例內的圖形處理器。在這種具體實施例內，提供一些或無專屬圖形(並行處理)記憶體，並且PPU 202完全或幾乎完全使用系統記憶體。在UMA具體實施例 內，PPU 202可整合至橋接器晶片或處理器晶片，或提供當成分散式晶片，具有高速連結(例如PCI Express)透過橋接器晶片或其他通訊方式將PPU 202連接至系統記憶體。

如上述，任何數量的PPU 202都可包含在一並行處理子系統112內。例如：單一外接卡上可提供多個PPU 202，或多張外接卡可連接至通訊路徑113，或一或多個PPU 202可整合到一橋接器晶片上。多PPU系統內的PPU 202可彼此一致或不同。例如：不同的PPU 202可具有不同數量的處理核心、不同數量的本機並行處理記憶體等等。當存在多個PPU 202時，這些PPU可並行操作，以比單一PPU 202還要高產量的方式來處理資料。合併一或多個PPU 202的系統可在許多設置與外型因素之下實施，包含桌上型、膝上型或手持式個人電腦、伺服器、工作站、遊戲機、嵌入式系統等等。

多重並行任務排程

多重處理任務可在GPC 208上同時執行，並且一處理任務可在執行期間產生一或多個「子代」處理任務。任務/工作單元207接收該等任務，並且動態排定要由GPC 208執行的處理任務以及子代處理任務。

第三A圖為根據本發明的一個具體實施例，第二圖中任務/工作單元207的方塊圖。任務/工作單元207包含一任務管理單元300以及工作分配單元340。任務管理單元300根據執行優先順序等級來組織要排程的任務。針對每一優先順序等級，任務管理單元300儲存指標清單到對應至排程器表321內任務的TMD 322，其中該清單可用連結清單來實施。TMD 322可儲存在PP記憶體204或系統記憶體104內。任務管理單元300接受任務並且將該等任務儲存在排程器表321內的速率與任務管理單元300排定任務來執行的速率無關，因此任務管理單元300可在排定該等任務之前集中許多任務。然後根據優先順序資訊或使用其他技術，例如輪轉排程，來排定集中的任務。

工作分配單元340包含一任務表345，其中有插槽讓TMD 322針對執行的任務所佔用。任務管理單元300可排定任務在任務表345內有空插槽時執行。當無空插槽時，未佔用插槽的較高優先順序任務會驅逐佔用插槽的較低優先順序任務。當一任務遭到驅逐時，該任務會停止，並且若該任務尚未執行完畢，則該任務的指標會加入要排定的任務指標清單中，如此稍後會恢復該任務的執行。當在一任務執行期間產生子代處理任務時，將該子代任務的指標加入要排定的任務指標清單內。在一實施例中，由在處理叢集陣列230內執行的TMD 322可產生一子代任務。

與從前端212的任務/工作單元207所接收的任務不同，子代任務從處理叢集陣列230接收。子代任務並未插入推送緩衝區或傳輸至該前端。當子代任務已經產生或該子代任務的資料已經儲存在記憶體內時，並不會通知CPU 102。透過推送緩衝區提供的該等任務與子代任務間之另一差異在於，透過推送緩衝區提供的該等任務由該應用程式定義，而該等子代任務則在該等任務執行期間動態產生。

任務處理概觀

第三B圖為根據本發明的一個具體實施例，說明第二圖中PPU 202之一內之GPC 208的方塊圖。每一GPC 208都可設置成同時執行大量執行緒，其中「執行緒」一詞代表在特定輸入資料集上執行的特定程式之實例。在某些具體實施例內，單一指令、多重資料(SIMD，single-instruction,multiple-data)指令發行技術用於支援大量執行緒的並行執行，而不用提供多個獨立指令單元。在其他具體實施例內，單一指令、多重執行緒(SIMT，single-instruction,multiple-thread)技術用於支援大量一般同步執行緒的並行執行，使用共用指令單元，其設置成發出指令至每一GPC 208內的處理引擎集。不同於一SIMD執行區域，其中所有處理引擎一般都執行一致的指令，SIMT執 行允許不同執行緒更迅速遵循分散的執行路徑通過一已知執行緒程式。業界內精通技術人士將了解，SIMD處理區域代表一SIMT處理區域的功能子集。

透過將處理任務分配至串流多重處理器(SM，streaming multiprocessor)310的管線管理員305可有利地控制GPC 208的操作。管線管理員305也可設置成利用指定SM 310所輸出已處理資料的目的地，控制一工作分配橫桿330。

在一個具體實施例內，每一GPC 208都包含數量為M的SM 310，其中M1，每一SM 310都設置成處理一或多個執行緒群組。另外，每一SM 310可包含可管線化的一致功能執行單元集(例如執行單元與載入儲存單元，在第三C圖內顯示為執行單元302以及LSU 303)，允許在先前指令執行完成之前有新指令的發出，且其可以是功能執行單元的任何組合。在一個具體實施例內，該等功能單元支援許多種運算，包含整數與浮點演算(例如加法與乘法)、比較運算、布林運算(AND、OR、XOR)、位元位移和許多代數函數的計算(例如平面插值、三角函數、指數以及對數函數等等)；並且該相同功能單元硬體可用來執行不同運算。

這一系列指令傳輸至特定GPC 208構成一執行緒，如本文先前所定義，並且透過SM 310內並行處理引擎(未顯示)的同時執行之特定數量的執行緒之集合在此稱為「經線」或「執行緒群組」。如本文所使用，「執行緒群組」代表在不同輸入資料上同時執行相同程式的執行緒群組，其中該群組的一個執行緒係可被指派給SM 310內不同的處理引擎。一執行緒群組可包含數量比SM 310內處理引擎數量還要少的執行緒，在此案例中，某些處理引擎會在循環期間處理該執行緒群組時閒置。執行緒群組也可包含數量比SM 310內處理引擎數量還要多的執行緒，在此案例中，將在連續時脈循環上進行處理。因為每一SM 310都可同時支援最多G個執行緒群組，所以在任何已 知時間上GPC 208內都可執行最多G*M個執行緒群組。

此外，在SM 310內可同時啟用複數個相關執行緒群組(在不同執行階段內)。此執行緒群組的集合稱為「合作執行緒陣列」(CTA，cooperative thread array)或「執行緒陣列」。特定CTA的大小等於m*k，其中k為執行緒群組內同時執行的執行緒數量，通常為SM 310內並行處理引擎數量的整數倍數，並且m為SM 310內同時啟用的執行緒群組數量。CTA的大小一般由程式設計師以及該CTA可用的硬體資源數量(例如記憶體或暫存器)來決定。

每一SM 310都包含第一層(L1)快取(如第三C圖所示)，或使用SM 310之外對應L1快取內的空間，其用於執行負載與儲存操作。每一SM 310也要存取所有GPC 208之間共享的第二層(L2)快取，並且可用於在執行緒之間傳輸資料。最後，SM 310也要存取至晶片外「全域」記憶體，其可包含例如並行處理記憶體204及/或系統記憶體104。吾人了解，PPU 202之外的任何記憶體都可用來當成全域記憶體。此外，第一點五層(L1.5)快取335可包含在GPC 208內，設置成接收並固定由SM 310要求透過記憶體介面214從記憶體擷取的資料，這些資料包含指令、統一資料以及常數資料，並且將該要求的資料提供給SM 310。具有GPC 208內多個SM 310的具體實施例藉著共享共用指令和在L1.5快取335內快取的資料而得利。

每一GPC 208都可包含一記憶體管理單元(MMU，memory management unit)328，其設置成將虛擬位址映射至實體位址。在其他具體實施例內，MMU 328可位於記憶體介面214之內。MMU 328包含一組頁面表記錄(PTE，page table entry)，用於將一拼貼(tile)以及選擇性將一快取線索引的虛擬位址映射至實體位址。MMU 328可包含能夠位於多處理器SM 310或L1快取或GPC 208內的位址轉譯後備緩衝區(TLB，translation lookaside buffer)或快取。該實體位址經過處理來分配本機存取 的表面資料，以便於分割單元215之間有著有效要求交錯(request interleaving)。該快取線索引可用於決定一快取線的要求是否命中或未命中。

在圖形與計算應用當中，GPC 208可經過設置，如此每一SM 310都耦合至一紋理單元315，用於執行紋理映射操作，例如決定紋理樣本位置、讀取紋理資料以及篩選該紋理資料。紋理資料從內部紋理L1快取(未顯示)讀取，或在某些具體實施例內從SM 310內的該L1快取讀取，並且依照需求從所有GPC 208之間共享的L2快取、並行處理記憶體204或系統記憶體104擷取。每一SM 310都輸出處理過的任務至工作分配橫桿330，以便將該處理過的任務提供至另一GPC 208供進一步處理，或透過橫桿單元210將該處理過的任務儲存在L2快取、並行處理記憶體204或系統記憶體104內。preROP(預先光柵運算)325設置成從SM 310接收資料、將資料引導至分割單元215內的ROP單元，並且執行顏色混合、組織畫素顏色資料以及執行位址轉譯的最佳化。

吾人將了解，本文中顯示的該核心架構為例示，所以可進行改變與修改。任何數量的處理單元，例如SM 310或紋理單元315、preROP 325都可包含在GPC 208內。進一步如第二圖內所示，PPU 202可包含任意數量功能彼此類似的GPC 208，如此執行行為並不取決於哪一個GPC 208接收到特定處理任務。進一步，每一GPC 208都得利於與其他GPC 208無關的操作，運用分離並分散的處理單元、L1快取，來執行一或多個應用程式的任務。

精通此技術人士將了解，第一圖、第二圖、第三A圖和第三B圖內描述的架構並非用於限制本發明範疇，本說明書內的技術可在任何正確設置的處理單元上實施，在不悖離本發明範疇之下包含但不受限於一或多個CPU、一或多個多核心CPU、一或多個PPU 202、一或多個GPC 208、一或多個圖形 或特殊用途處理單元等等。

在本發明的具體實施例內，吾人想要使用一計算系統的PPU 202或其他處理器，運用執行緒陣列來執行一般用途計算。該執行緒陣列內的每一執行緒都可被指派一個獨一的執行緒識別碼(「執行緒ID」)，其可在該執行緒執行期間存取該執行緒。該執行緒ID可定義為一維度或多維度數值，控制該執行緒處理行為的許多態樣。例如：一執行緒ID可用於決定設定一執行緒的哪個輸入資料部分要處理及/或決定設定一執行緒的哪個輸出資料部分要產生或寫入。

每個執行緒指令的序列可包含至少一個指令，其定義該代表性執行緒與該執行緒陣列的一或多個其他執行緒之間的合作行為。例如：每一執行緒指令的順序可包含將在該順序內特定點上讓該代表執行緒運算執行中斷，直到一或多個其他執行緒到達該特定點上為止之指令、讓該代表執行緒將一或多個其他執行緒可存取的資料儲存至一共享記憶體內之指令、讓該代表執行緒根據其執行緒ID基本上讀取與更新儲存在一共享記憶體內一或多個該等其他執行緒已經存取過的資料之指令等等。該CTA程式也可包含一指令，計算該共享記憶體內所要讀取資料的位址，其中該位址為執行緒ID的函數。利用定義合適的功能並且提供同步技術，資料可利用CAT的一個執行緒寫入共享記憶體內的一已知位置，並且以可預測方式用相同CTA的不同執行緒從該位置當中讀取。因此，足以支援在執行緒之間共享的任何所要資料的模式，並且一CTA內的任何執行緒都可與相同CTA內任何其他執行緒共享資料。在CTA執行緒之間共享的資料內容(若有的話)由該CTA程式決定；如此吾人了解在使用CTA的特定應用當中，根據該CTA程式，CTA的執行緒彼此之間會或不會共享資料，並且在本文中可同時使用「CTA」與「執行緒陣列」等詞。

第三C圖為根據本發明的一個具體實施例，第三B圖中 SM 310的方塊圖。SM 310包含一指令L1快取370，其設置成透過L1.5快取335接收指令與常數。經線排程器與指令單元312接收來自指令L1快取370的指令與常數，並且根據該等指令與常數控制本機暫存檔304以及SM 310功能單元。SM 310功能單元包含N個執行(執行或處理)單元302以及P個載入儲存單元(LSU，load-store unit)303。

SM 310提供具備不同存取階層的晶片上(內部)資料儲存。特殊暫存器(未顯示)可由LSU 303讀取但是無法寫入，並且可用於儲存定義每一執行緒「位置」的參數。在一個具體實施例內，每一執行緒(或SM 310內每一執行單元302)可以對應一個特殊暫存器，其儲存一執行緒ID；每一執行緒ID暫存器都只能由個別一個執行單元302存取。特殊暫存器也可包含額外暫存器，可由執行TMD 322(或由所有LSU 303)所呈現相同處理任務的所有執行緒讀取，其儲存一CTA識別碼、該CTA維度、該CTA所屬網格的維度(或若TMD 322編碼一佇列任務而非一網格任務時的佇列位置)以及該CTA所被指派的TMD 322之識別碼。

若TMD 322為一網格TMD，則執行TMD 322會導致啟動並執行固定數量的CTA，來處理佇列525內所儲存的固定數量資料。CTA的數量依照網格寬度、高度與深度的乘積來指定。該固定數量的資料可儲存在TMD 322內，或TMD 322可儲存將由CTA處理的資料之指標。TMD 322也可儲存該CTA所執行程式的開始位址。

若TMD 322為佇列TMD，然後使用TMD 322的佇列功能，表示要處理的資料量並不需要固定。佇列記錄儲存指派給TMD 322由CTA處理的資料。該等佇列記錄也呈現執行緒執行期間由另一TMD 322產生的子代任務，藉此提供巢狀並行。一般來說，執行緒的執行或包含該執行緒的CTA會中止，直到子代任務執行完成為止。該佇列可儲存在TMD 322內， 或與TMD 322分開，在此案例中TMD 322儲存至該佇列的佇列指標。由該子代任務產生的資料可寫入該佇列，同時TMD 322代表已經執行的該子代任務。該佇列可實施為一圓形佇列，如此能夠儲存之資料總量並不受限於該佇列的大小。

屬於一網格的CTA具有暗示的網格寬度、高度和深度參數，指示該網格內個別CTA的位置。在初始化期間會寫入特殊暫存器，以回應透過前端212從裝置驅動程式103接收的命令，並且在一處理任務執行期間不會改變。前端212排程執行的每一處理任務。每一CTA都關聯於一特定TMD 322，以便同時執行一或多個任務。此外，單一GPC 208可同時執行多個任務。

一參數記憶體(未顯示)儲存可由相同CTA(或任何LSU 303)內的任何執行緒讀取但無法寫入的執行時間參數(常數)。在一個具體實施例內，裝置驅動程式103在導引SM 310開始執行使用這些參數的任務之前，提供參數給該參數記憶體。任何CTA內的任何執行緒(或SM 310內的任何執行單元302)都可透過記憶體介面214存取全域記憶體。全域記憶體的一部分可儲存在L1快取320內。

每一執行緒都使用本機暫存檔304當成伸展空間(scratch space)；每一暫存器都分配給一個執行緒專用，並且任何本機暫存檔304內的資料都只能由分配給該暫存器的該執行緒存取。本機暫存檔304可實施為實體上或邏輯上區分成P個通路的暫存檔，每一通路都具有某些數量的記錄(在此每一記錄都可儲存例如32位元字)。一個通路指派給該N個執行單元302以及P個載入儲存單元LSU 303的每一個，並且不同通路內的對應記錄可填入執行相同程式的不同執行緒之資料，來幫助SIMD執行。通路的不同部分可分配給該G個同時執行緒群組中不同的執行緒，如此本機暫存檔304內的一已知記錄只能由特定執行緒存取。在一個具體實施例內，本機暫存檔304內的 特定記錄保留給執行緒識別碼的儲存使用。此外，一統一L1快取375儲存N個執行單元302以及P個載入儲存單元LSU 303的每一通路之統一或常數值。

共享記憶體306可由單一CTA內的執行緒存取；換言之，共享記憶體306內任何位置都可由相同CTA內的任何執行緒(或SM 310內任何處理引擎)來存取。共享記憶體306可實施為具有互連的一共享暫存檔或共享晶片上快取記憶體，允許任何處理引擎從該共享記憶體內任何位置讀取或寫入。在其他具體實施例內，共享的狀態空間可映射至晶片外記憶體的每一CTA區域，並且快取儲存(cached)在L1快取320內。該參數記憶體可實施為該相同共享暫存檔或實施共享記憶體306的共享快取記憶體內的一指定區段，或實施為LSU 303具有唯讀存取的一個別共享暫存檔或晶片上快取記憶體。在一個具體實施例內，該參數記憶體的區域也用於儲存該CTA ID和任務ID，以及CTA和網格尺寸或佇列位置，用以實施該特殊暫存器的部分。SM 310內每一LSU 303都耦合至統一位址映射單元352，將提供用於載入與儲存統一記憶體空間內所指定指令的一位址轉換至每一分散記憶體空間內的一位址。因此，利用指定該統一記憶體空間內的一位址，可使用一指令來存取任何該本機、共享或全域記憶體空間。

每一SM 310內的L1快取320都可用於快取私用每一執行緒本機資料，並且也快取每一應用程式全域資料。在某些具體實施例內，該每一CTA共享資料可快取儲存在L1快取320內。LSU 303可透過一記憶體與快取互連380耦合至共享記憶體306以及L1快取320。

巢狀執行串流

第四圖例示根據本發明一個具體實施例，並行處理子系統112上的巢狀任務執行。如所示，CPU 102起始並行處理子系統112上任務420的執行。在任務420(0)完成之後，執行任務 420(1)。在任務420(1)完成之後，執行任務420(2)。在執行過程期間，任務420(1)喚起任務430(0)至430(2)，例如來計算任務420(1)所使用的中間結果。要維持適當指令執行順序，任務420(1)應該等待直到任務430完成才繼續。為了以這種方式等待，任務420(1)會受任務430上一個執行緒同步障礙之阻礙。每一任務420、430都可由一或更多執行緒、CTA或網格執行。雖然本說明書中將並行處理子系統112描述在一圖形處理單元(GPU)的範圍內，不過本說明書中描述的該等技術可實施於和CPU 102相關聯的任何共同處理單元之範圍內。

在此範例中，任務420(1)為任務430的親代，因此其就是任務420(1)的子代。雖然第四圖內只顯示一層親代-子代階層關係，實際上任意階層關係均可被實施。在一個具體實施例內，任務420和430每一都當成至少一個執行緒群組來執行，或第三B圖中SM 310內的至少一個CTA來執行。要在SM 310上執行具有親代與子代關係的執行緒程式，則應該實施三個系統元件，包含並行處理子系統112的硬體功能、並行處理子系統112的軟體執行時間功能以及用於程式編輯並行處理子系統112的語言支援架構。

支援一親代執行緒啟動並行處理子系統112內的一子代執行緒、CTA或網格所需的該硬體功能包含：從由SM 310產生的一要求啟動工作之新網格或CTA並且佇存來執行至任務/工作單元207、儲存SM 310的執行狀態、在SM 310內從該已儲存的執行狀態繼續執行，以及幫助一親代與子代任務之間的記憶體連貫性。需要支援處理子系統112內一親代執行緒啟動一子代執行緒、CTA或網格的該執行時間功能包含：啟動一新網格以回應來自SM 310內所執行之一執行緒的一要求、讓一親代執行緒在一子代執行緒群組上執行一執行緒同步障礙、確定該親代執行緒與該子代群組之間的記憶體一致性、排定已同步執行緒群組的工作與持續情況來保證正向計算進 度，以及確定親代執行緒與子代群組的正確執行語意。該語言支援架構包含指定從一親代執行緒啟動一子代執行緒程式，並且在該子代程式上執行一同步障礙之機制。

並行處理子系統112已經使用一執行緒導向程式編輯環境進行程式編輯，例如NVIDIA(tm)推出的CUDA(tm)程式編輯環境。在一個具體實施例內，該CUDA語言規格擴展成包含一子代啟動架構(「<<< >>>」)，以指定用於啟動一子代CUDA網格的細節。該子代啟動架構，本說明書內表示為「A<<<B>>>C」，包含一子代程式名稱(A)、網格組態參數(B)以及程式輸入參數(C)。該CUDA執行時間環境已經擴展，讓一親代執行緒在該子代CUDA網格上執行一同步障礙。雖然本發明以該CUDA程式編輯環境的範圍例示本發明具體實施例，不過精通技術人士了解，本說明書內傳授的技術可適用於任何並行程式編輯環境以及任何並行處理系統。也就是說，CUDA僅為例示，並不用於限制本發明的範疇或精神。

底下的表1例示運用示範CUDA程式內的子代啟動架構以及同步障礙。

在表1的範例中，執行緒程式「foo()」的實例使用一執行緒程式「A」，其中具有一指標器(*ptr)至由foo()分配的記憶體， 來啟動一子代網格。該子代網格內的執行緒可存取該已分配的記憶體。該親代執行緒foo()可在子代網格A完成之後繼續，由來自一阻擋同步障礙函數呼叫的返回來指示，該呼叫名為cudaThreadSynchronize()。

啟動至該GPU上的任務一般適合立即執行。因此缺少一種機制來確定依序執行一任務佇列內的任務，並行處理子系統112排定要執行的任何任務，不去考慮與先前已經啟動進入該相同任務佇列的任務上之相依性。執行次序可藉由如下所述一階層執行圖實施。

第五圖例示根據本發明一個具體實施例，包含相關聯任務中繼資料描述器佇列(TMDQ)以及任務的一階層執行圖。如所示，該階層執行圖包含巢狀深度0上的執行緒群組510、巢狀深度1上的TMDQ 512、任務520、530、540、一執行圖580，以及巢狀深度2上的一執行圖590。

巢狀深度0上的執行緒群組510包含由CPU 102建立並管理的執行緒。一執行緒群組包含執行緒的任何集合，包含一CTA，其中所有執行緒都存在於該相同巢狀深度上。一執行緒的巢狀深度為該執行緒階層之上的親代網格數量，例如：一CPU執行緒具有巢狀深度0，因為一CPU執行緒之上並無親代網格。若該CPU執行緒啟動一網格，則該網格稱為在巢狀深度1上。若巢狀深度1上該網格內的一執行緒啟動一新網格，則該新網格稱為在巢狀深度2上，以此類推。因為執行緒群組510內的該等執行緒都為CPU執行緒，所以每一個執行緒都在巢狀深度0上。

綜合第二圖的說明，TMDQ 512包含對於已知為任務的資料結構之指標器。每一TMDQ 512都指向屬於一或更多串流的任務。TMDQ(0)512(0)指向與一第一串流相關聯的任務520(0)，TMDQ(1)512(1)指向與一第二串流相關聯的任務530(0)和530(1)，TMDQ(2)512(2)則可指向與一第三串流相關聯的任 務534(0)、540(1)和540(2)。任何數量的TMDQ 512都可定義成每一TMDQ 512都對應到任意數量的任務。

任務520、530、540為資料結構，包含要由該GPU執行的一或更多命令。針對某一TMDQ如TMDQ 512，任務將以依序方式執行。舉例來說，當任務530(0)完成之後，才開始執行任務530(1)。類似地，任務540(0)完成之後，才開始執行任務540(1)，之後才開始執行任務540(2)。一旦任務開始啟動，則TMDQ 512前端的任務開始被執行。也就是說，一旦任務開始啟動，任務520(0)、530(0)和540(0)就立即執行。然而，不同TMDQ 512內的任務並無順序相依性，例如：任務530(1)可在任務540(1)之前、之後或同時執行。

巢狀深度1上的執行圖580為一執行緒群組，加上該相關聯的TMDQ和任務，這些已經由巢狀深度0上的任務之一者所啟動。任何任務都可啟動一或更多網格，而這些網格都位於比啟動該網格的該任務相關聯巢狀深度還要深一層之巢狀深度。如所示，存在於巢狀深度0上的任務540(1)在任務540(1)執行期間的某個時間上啟動執行圖580。執行圖580內每一任務與TMDQ的功能都與巢狀深度0上的任務與TMDQ實質上相同。當完成執行圖580內每一任務，並且已經完成任務540(1)內所有其他命令時，則開始執行任務540(2)。

巢狀深度2上的執行圖590為一執行緒群組，加上該相關聯的TMDQ和任務，這些已經由巢狀深度1上的任務之一者所啟動。執行圖590內每一任務與TMDQ的功能基本上都與較低巢狀階層上的任務與TMDQ實質上相同。當完成執行圖590內每一任務時，若一旦該啟動任務內所有其他命令都已經完成，則完成該啟動任務。如此，在任何網格內都保留執行順序，並且網格可為任意巢狀深度，而保留一串流內任務的執行順序。

一執行緒群組內的執行緒係可根據範圍來定義，其中該範 圍為可存取至相同串流和TMDQ資源的執行緒集合。該相同範圍內的執行緒可建立並共享TMDQ，只要該等執行緒都在相同巢狀深度上以及相同裝置上(GPU或CPU 102)。針對CPU執行緒，該範圍係定義為與該CUDA範圍相關聯的執行緒集合。針對GPU執行緒，該範圍可代表一合作執行緒陣列(CTA，Cooperative Thread Array)，或相同巢狀深度上存在的任何執行緒集合。

當新串流已經由一CPU執行緒建立時，CPU 102動態分配記憶體來支援該串流的管理。在該等串流任務完成之後接著該串流的摧毀(destroyed)時，CPU 102會釋放先前分配給該串流的記憶體。該GPU通常無法動態分配記憶體，因此該GPU預先分配範圍資料給可能同時執行的每一範圍。結果，與一GPU網格相關聯的一執行緒群組具有固定數量的TMDQ，這在該網格執行期間不會改變。使用該cudaStreamCreate()函數呼叫建立一GPU網格內的新串流。該函數呼叫回傳指向該網格內該預先分配TMDQ之一者的一整數索引。建立該串流並不需要動態分配記憶體，一旦已經完成一GPU串流內的所有任務，則使用一cudaStreamDestroy()函數呼叫摧毀該串流。因為並無記憶體動態分配給該GPU串流，cudaStreamDestroy()函數呼叫就不需要將記憶體釋放回佇池，因此簡單返回至該呼叫程式即可。

一旦已經建立一串流，則由該相關聯範圍內一或更多執行緒可將新任務啟動進入該串流。若一執行緒啟動一新任務進入目前無任務的一TMDQ內，則該新任務會在啟動之後立即執行。類似地，若一TMDQ內所有先前任務已經執行完畢，則啟動進入一TMDQ的一新任務會在啟動之後立即執行。另外，若一執行緒啟動一新任務進入具有尚未執行完成之一或更多等待中任務的一TMDQ內，則該新任務仍會被啟動進入該TMDQ，但是在等待中任務執行完畢之後才會開始執行該任 務。在兩案例中，該新任務透過不需要CPU 102干涉的非鎖定操作，依舊可以啟動進入該TMDQ。

第六圖例示根據本發明另一個具體實施例，包含相關聯TMDQ以及任務的一階層執行圖。如所示，該階層執行圖包含巢狀深度1上的執行緒群組610、巢狀深度2上的TMDQ 612、任務620、630、640、650、660、一執行圖680，以及巢狀深度3上的一執行圖690。除了以下詳述的以外，該階層執行圖的組件功能大體上如上面第五圖之描述。

如所示，執行緒群組610的每一TMDQ 612都具有一或更多等待中的任務。在一個範例中，關聯於串流670的任務620(0)已經啟動進入TMDQ 612(0)，但是關聯於串流675的任務660(0)則尚未啟動。關聯於一個串流的任務630已經啟動進入TMDQ(1)612(1)。類似地，關聯於一第二串流的任務640已經啟動進入TMDQ(2)612(2)、關聯於一第三串流的任務650已經啟動進入TMDQ(N)612(N)，並且所有中介的TMDQs 612也應該具有一或更多相關聯的任務。此時，執行緒群組610內的一執行緒嘗試建立一新串流675。不過，執行緒群組610具有巢狀深度1，並且關聯於該GPU。因為該GPU無法動態分配記憶體，因此無法建立一新TMDQ來容納新串流675。在這種案例中，關聯於新串流675的任務660會被啟動進入目前串流670所使用的TMDQ(0)。串流675啟動任務660(0)和660(1)進入TMDQ(0)612(0)。然後，串流670啟動任務620(1)進入TMDQ(0)612(0)。然後，串流675啟動任務660(2)進入TMDQ(0)612(0)。請注意，此方式可能會導致非必要的相依性。即使串流670和675彼此獨立，TMDQ的順序性質導致任務660(0)取決於完成任務620(0)、任務620(1)取決於完成任務660(1)、以此類推。雖然結果會降低效能，不過會正確保留串流670內任務620以及串流675內任務660的順序。

第七圖例示根據本發明一個具體實施例，包含與一執行緒 群組相關聯的參數與範圍資訊之一執行緒群組範圍720資料結構。如所示，執行緒群組範圍720包含一最後任務指標器740，用於該執行緒群組的每一TMDQ，以及包含一工作計數器750。

最後任務指標器740為該相關聯TMDQ內該最後任務的指標器。當一新任務已啟動進入一TMDQ時，最後任務指標器740透過一基本操作來更新，指出此時該新任務為該TMDQ內的該最後任務。底下的表2例示在示範CUDA程式內啟動一TMDQ的一新任務。

在表2的範例中，位址StreamEnd上的最後任務指標器740會由NewTask的指標器覆寫，並且最後任務指標器740內的先前值回到FormerStreamEnd。若FormerStreamEnd不為零(也就是FormerStreamEnd為一任務的指標器)，然後關聯於該任務的該StreamNext值更新成指向該等新啟動之任務。若FormerStreadEnd為零，然後該TMDQ內並無等待中的任務，則該新任務立即開始執行。

表2的範例在操作的關鍵區段內執行，如此避免其中一執行緒已經將一任務進入一串流，但是在啟動該新任務之前該執行緒已遭掃除的僵持情況。在這種案例中，若在該新任務完成之前不允許掃回該已掃除的執行緒，則可能發生僵持情況。不 過，因為該新任務尚未啟動，所以不會開始執行。

當一任務完成時，在並行處理子系統112上執行的一排程器讀取對應至關聯於該已完成任務的該TMDQ之該最後串流指標器。若該相關聯TMDQ的最後任務指標器740並未指向該已完成的任務，則該已完成的任務並不是該TMDQ內的最後任務。在此案例中，該排程器導致該TMDQ內的下一個任務被開始執行，如底下結合第八圖之說明。若該相關聯TMDQ的最後任務指標器740指向該已完成的任務，則該已完成的任務為該TMDQ內的最後任務。在此案例中，該排程器執行一基本比較與掃除，將最後任務指標器740設定為一空指標器，並且讀取目前最後任務指標器740內儲存的該值。該排程器執行「currentEnd=atomicCAS(&StreamEnd,finishedTask,NULL)」形式的函數呼叫，其中「StreamEnd」為該相關聯TMDQ的最後任務指標器740、「finishedTask」為該已完成任務的指標器並且「NULL」為該空指標器。該函數基本上回傳最後任務指標器740內儲存的該值，以該函數呼叫內的「currentEnd」所呈現。

若「currentEnd」的該值為該已完成任務的指標器，則該TMDQ內的所有任務已被完成，並且未啟動新任務。該排程器知道該串流內所有任務已經完成。若「currentEnd」的值並非該已完成任務的指標器，則已經啟動一新任務，並且已經更新執行緒群組範圍720來反應該新任務存在。在此案例中，該排程器讀取關聯於該已完成任務的該StreamNext指標器(以下說明)。若關聯於該已完成任務的該StreamNext指標器不為零，則該排程器開始執行位址StreamNext上的該任務。若「StreamNext」的值為該空指標器，則已經啟動一新任務，但是該任務狀態尚未更新來反應該新任務存在。在此案例中，該排程器監控StreamNext，直到該值從該空指標器改變成該新任務的指標器。然後，該排程器導致由StreamNext指引的該新 任務開始被執行。

第八圖例示根據本發明一個具體實施例，包含與一計算任務相關聯的參數之一任務狀態820資料結構。如所示，任務狀態820包含一任務識別碼(任務ID)840、一下一個串流指標器842、一執行緒群組範圍識別碼(執行緒群組範圍ID)844以及關聯於該任務的其他參數(未顯示)。

任務ID 840就是指向關聯於任務狀態820的該任務之唯一識別碼。當任務在一TMDQ上建立與啟動，會針對每一新任務建立任務狀態820。該任務ID可讓該排程器找出關聯於一已知任務狀態820的該任務。

下一個串流指標器842為該TMDQ內下一個任務的指標器。當一任務完成時，該排程器讀取該下一個串流指標器，決定到何處尋找該TMDQ內可開始執行的下一個任務。然後，該排程器導致由該下一個串流指標器842指出的該位址上之該任務能夠開始被執行。若該已完成任務為該TMDQ內的最後任務，則該下一個串流指標器842設定為空指標器。

執行緒群組範圍ID 820就是指向關聯於任務狀態820的該執行緒群組範圍720之唯一識別碼。當一任務完成時，該排程器讀取執行緒群組範圍ID 820，找出執行緒群組範圍720。然後排程器可執行相關聯的任務完成步驟，如此更新該工作計數器關閉一TMDQ與一範圍，如以上關聯於第七圖之描述。

精通技術人士將了解，本說明書中描述的該等技術僅為例示，可進行改變與修改。例如：該等描述技術足夠彈性運用在任何並行程式編輯環境以及任何並行處理系統內，而不管關聯於這種環境或系統的一GPU或其他共同處理器是否能動態分配記憶體。如此，不管該GPU預先分配記憶體是否關聯於一TMDQ，或是否需要動態分配記憶體至一TMDQ，都可運用該等描述的技術。

第九圖揭示根據本發明的一個具體實施例，用於處理一已 完成的計算任務之方法步驟流程圖。雖然已經結合第一圖至第八圖的系統來描述該等方法步驟，精通技術人士將了解，設置來執行該等方法步驟(以任何順序)的任何系統都在本發明範疇內。

方法900從步驟902開始，在此該排程器接收一計算任務已經完成的通知。在步驟904上，該排程器將關聯於該執行緒群組的一工作計數器遞減至該已完成任務所應有的計數。在步驟906上，該排程器決定該已完成任務是否為該相關聯TMDQ的最後任務，也就是該相關聯TMDQ內的所有任務都已經完成。例如：若關聯於該TMDQ的該串流末端指標器指向該已完成任務，則該排程器決定該已完成任務為該TMDQ的最後任務。若該已完成任務並非該TMDQ內最後任務，則方法900前往步驟908，在此該排程器開始執行該TMDQ內的下一個任務。然後方法900終止。

在步驟906上，若該已完成任務為該TMDQ內的最後任務，則方法900前往步驟910，在此該排程器使用一基本操作將關聯於該TMDQ的該串流末端指標器更新為一空指標器，反應出該TMDQ此時屬於被清空的狀態。在步驟912上，該排程器決定CPU 102或該執行緒群組是否剛啟動一新任務至該佇列。例如：當將該串流末端指標器基本上更新成一空指標器時，該排程器決定該指標器已經改變成不再指向該已完成任務。若CPU 102或執行緒群組未啟動一新任務，則方法900終止。

在步驟912上，若CPU 102或執行緒群組已經啟動一新任務，則方法900前往步驟914，在此該排程器等待該新任務佇存在該TMDQ內。例如：該排程器關聯於該TMDQ的該串流末端指標器更新成指向該新任務。在步驟916上，該排程器開始執行該新任務。然後方法900終止。

總結來說，所說明的技術提供一強化方式，讓GPU將新 計算任務佇存至TMDQ。尤其是，當建立一新TMDQ時已經預先分配記憶體給範圍資料，在此該記憶體包含資料空間，讓每一範圍可在該GPU上分開執行。在已經建立新TMDQ時，並且該CTA範圍並無可用項目給一新TMDQ時，該新TMDQ可與一現有TMDQ整合，在此該TMDQ內的計算任務包含來自每一原始TMDQ的任務。佇存至該TMDQ的新計算任務可立即執行，或可等待該TMDQ內的先前任務完成。在每一計算任務完成時執行一排程操作，以便保留任務的執行順序，而不用基本鎖定操作。

本發明技術的優點為GPU可將計算任務佇存在任務佇列內，同時建立一任意數量的新任務佇列至任何任務巢狀階層，不用該CPU的介入(如鎖定狀態)。由於在該CPU建立並佇存任務時該GPU不用等待，所以處理效率獲得增強。因為該GPU不需要來自該CPU的干涉，即使該CPU消耗至該GPU的所有通訊通道，也可避免僵持情況。計算任務的執行順序保留給該CPU和GPU所啟動的任務。

雖然上述都導引至本發明的具體實施例，在不悖離本發明基本領域之下可提供其他與進一步具體實施例。例如：本發明態樣可實現於硬體、軟體或軟體與硬體的組合之上，本發明的一個具體實施例可實施當成搭配電腦系統使用的程式產品。該程式產品的程式定義該等具體實施例(包含本文所述的方法)的功能，並且可包含在各種電腦可讀取儲存媒體上。例示的電腦可讀取儲存媒體包含但不受限於：(i)其上資訊永久儲存的不可抹寫儲存媒體(例如電腦內的唯讀記憶體裝置，例如CD-ROM光碟機可讀取的CD-ROM光碟、快閃記憶體、ROM晶片或任何一種固態非揮發性半導體記憶體)；以及(ii)其上儲存可變資訊的可抹寫儲存媒體(例如磁碟機或硬碟內的磁碟或任何一種固態隨機存取半導體記憶體)。這種電腦可讀取儲存媒體內當儲存具備本發明功能的電腦可讀取指令時，屬於本發 明的具體實施例。

因此，本發明的領域由下列申請專利範圍所決定。

100‧‧‧電腦系統

215‧‧‧分割單元

102‧‧‧中央處理單元

230‧‧‧處理叢集陣列

103‧‧‧裝置驅動程式

220‧‧‧動態隨機存取記憶體

104‧‧‧系統記憶體

105‧‧‧記憶體橋接器

300‧‧‧任務管理單元

106‧‧‧通訊路徑

302‧‧‧執行單元

107‧‧‧輸入/輸出橋接器

303‧‧‧載入儲存單元

108‧‧‧使用者輸入裝置

304‧‧‧本機暫存檔

110‧‧‧顯示裝置

305‧‧‧管線管理員

112‧‧‧並行處理子系統

306‧‧‧共享記憶體

113‧‧‧第二通訊路徑

310‧‧‧串流多重處理器

114‧‧‧系統磁碟

312‧‧‧經線排程器與指令單元

116‧‧‧開關

118‧‧‧網路配接器

315‧‧‧紋理單元

120-121‧‧‧外接卡

320‧‧‧第一層快取

202‧‧‧並行處理單元

321‧‧‧排程器表

204‧‧‧本機並行處理記憶體

322‧‧‧任務中繼資料

325‧‧‧預先光柵運算

205‧‧‧輸入/輸出單元

328‧‧‧記憶體管理單元

206‧‧‧主介面

330‧‧‧工作分配橫桿

207‧‧‧任務/工作單元

335‧‧‧第一點五層快取

208‧‧‧一般處理叢集

340‧‧‧工作分配單元

210‧‧‧記憶體橫桿單元

345‧‧‧任務表

212‧‧‧前端

352‧‧‧統一位址映射單元

214‧‧‧記憶體介面

370‧‧‧指令第一層快取

380‧‧‧記憶體與快取互連

630‧‧‧任務

420‧‧‧示範任務

640‧‧‧任務

430‧‧‧任務

650‧‧‧任務

510‧‧‧執行群組

660‧‧‧任務

512‧‧‧任務中繼資料描述器佇列

680‧‧‧執行圖

690‧‧‧執行圖

520‧‧‧任務

670‧‧‧串流

525‧‧‧佇列

675‧‧‧串流

530‧‧‧任務

720‧‧‧執行緒群組範圍

540‧‧‧任務

740‧‧‧最後任務指標器

580‧‧‧執行圖

750‧‧‧工作計數器

590‧‧‧執行圖

820‧‧‧任務狀態

610‧‧‧執行緒群組

840‧‧‧任務識別碼

612‧‧‧任務中繼資料描述器佇列

842‧‧‧下一個串流指標器

844‧‧‧執行緒群組範圍識別碼

620‧‧‧任務

所以，可以詳細瞭解本發明上述特徵之方式中，本發明的一更為特定的說明簡述如上，其可藉由參照到具體實施例來進行，其中一些例示於所附圖式中。但應注意所附圖式僅例示本發明的典型具體實施例，因此其並非要做為本發明之範圍的限制，本發明自可包含其它同等有效的具體實施例。

第一圖為例示設置來實施本發明一或多個態樣的電腦系統方塊圖；第二圖為根據本發明的一個具體實施例，用於第一圖中該電腦系統的一並行處理子系統之方塊圖；第三A圖為根據本發明的一個具體實施例，第二圖中前端的方塊圖；第三B圖為根據本發明的一個具體實施例，第二圖中並行處理單元之一者內一般處理叢集的方塊圖；第三C圖為根據本發明的一個具體實施例，第三B圖中該串流多重處理器一部分的方塊圖；第四圖例示根據本發明一個具體實施例，並行處理子系統上的巢狀任務執行；第五圖例示根據本發明一個具體實施例，包含相關聯任務中繼資料描述器佇列(TMDQ)以及任務的一階層執行圖；第六圖例示根據本發明另一個具體實施例，包含相關聯TMDQ以及任務的一階層執行圖；第七圖例示根據本發明一個具體實施例，包含與一執行緒群組相關聯的參數與範圍資訊之一執行緒群組範圍資料結構；第八圖例示根據本發明一個具體實施例，包含與一計算任務相關聯的參數之一任務狀態資料結構；以及 第九圖揭示根據本發明的一個具體實施例，用於處理一已完成的計算任務之方法步驟流程圖。

Claims (10)

1. 一種用於處理由一第一群組執行緒所執行支付數個任務並且儲存在複數個任務中繼資料描述器佇列(TMDQ，task metadata descriptor queue)內之電腦實施方法，該方法包括：接收包含在該等複數個任務內的一第一任務已經完成之通知；在一共同處理單元內決定所有任務是否都包含在該等複數個任務之一子集內，並且關聯於一第一TMDQ已經被執行；若該等複數個任務之該子集內含的所有任務都尚未執行，則啟動該等複數個任務內含的一第二任務；以及若該等複數個任務之該子集內含的所有任務都已執行，則：更新一第一資料結構內與該第一TMDQ相關聯的一指標器；決定該等複數個任務內含的一第三任務已經佇存在該第一TMDQ內；以及啟動該第三任務。
2. 一種用於處理由一第一群組執行緒所執行複數個任務並且儲存在複數個任務中繼資料描述器佇列(TMDQ，task metadata descriptor queue)內之子系統，包括：一任務管理單元，其設置成執行以下步驟：接收包含在該等複數個任務內的一第一任務已經完成之通知；在一共同處理單元內決定所有任務是否都包含在該等複數個任務之一子集內，並且關聯於一第一TMDQ已經執行；若該等複數個任務之該子集內含的所有任務都尚未執行，則啟動該等複數個任務內含的一第二任務；以及若該等複數個任務之該子集內含的所有任務都已執 行，則：更新一第一資料結構內與該第一TMDQ相關聯的一指標器；決定該等複數個任務內含的一第三任務已經佇存在該第一TMDQ內；以及啟動該第三任務。
3. 如申請專利範圍第2項之子系統，另包括在接收該第一任務已經完成的該通知之後，遞減與該等複數個任務內含的未完成任務數量相關之一計數器。
4. 如申請專利範圍第3項之子系統，另包括根據具有零值的該計數器，決定由該第一群組執行緒執行的該所有任務都已經完成。
5. 如申請專利範圍第2項之子系統，其中啟動該第三任務另包括等待更新關聯於該第三任務的一第二資料結構內一記憶體位置，以反應該第三任務的一所在位置。
6. 如申請專利範圍第2項之子系統，其中在該第一任務之前並且關聯於一第二群組執行緒所產生的一任務要求之下，已經建立該第一TMDQ。
7. 如申請專利範圍第2項之子系統，其中由在該第一任務之前並且關聯於一第二群組執行緒所產生的一任務，將該第一任務插入該第一TMDQ。
8. 如申請專利範圍第2項之子系統，其中該第一TMDQ包含對應至該第一群組執行緒內含的一第一執行緒的一第一組任務，並且一第二TMDQ包含對應至該第一群組執行緒內含的一第二執行緒的一第二組任務。
9. 如申請專利範圍第2項之子系統，其中該第一TMDQ包含對應至該第一群組執行緒內含的一第一執行緒的一第一組任務，以及對應至該第一群組執行緒內含的一第二執行緒的一第二組任務。
10. 如申請專利範圍第2項之子系統，其中決定該等複數個任 務的該子集內含之所有任務是否已經執行包含：決定至該第一TMDQ內所儲存一最後任務的一指標器是否指向該第一任務。