本發(fā)明總地涉及計(jì)算機(jī)架構(gòu)，并且，更具體地，涉及用于管理嵌套（nested）執(zhí)行流的方法和系統(tǒng)。

背景技術(shù)：
在具有中央處理單元(CPU)和圖形處理單元(GPU)二者的常規(guī)計(jì)算系統(tǒng)中，CPU確定由GPU實(shí)施哪些具體計(jì)算任務(wù)以及以什么次序?qū)嵤?。GPU計(jì)算任務(wù)典型地包括跨并行數(shù)據(jù)集的高度并行、高度類似的操作，該并行數(shù)據(jù)集諸如圖像或圖像集。在常規(guī)GPU執(zhí)行模型中，CPU通過選擇相應(yīng)的線程程序并且指導(dǎo)GPU執(zhí)行線程程序的并行實(shí)例集來發(fā)起特定計(jì)算任務(wù)。在常規(guī)GPU執(zhí)行模型中，CPU經(jīng)常是可在GPU上發(fā)起線程程序的執(zhí)行的僅有的實(shí)體。在所有線程實(shí)例完成執(zhí)行后，GPU必須通知CPU并等待將由CPU所發(fā)出的另一個(gè)計(jì)算任務(wù)。通知CPU并等待下一個(gè)計(jì)算任務(wù)是使GPU內(nèi)的某些資源暫時(shí)閑置的典型的阻塞型、序列化操作，從而降低整體系統(tǒng)性能。在某些場景中可通過在入棧緩沖區(qū)中對順序的計(jì)算任務(wù)進(jìn)行排隊(duì)，GPU可從該入棧緩沖區(qū)中拉取工作用于執(zhí)行而不用等待CPU，從而改進(jìn)性能。當(dāng)CPU能夠足夠快地生成用于GPU的工作，使每當(dāng)GPU能夠開始新任務(wù)時(shí)工作均在入棧緩沖區(qū)內(nèi)掛起（pending）時(shí)，包括固定數(shù)據(jù)流處理管線的計(jì)算任務(wù)從該入棧緩沖區(qū)模型中受益。然而，依賴于數(shù)據(jù)的計(jì)算任務(wù)仍在GPU結(jié)果、CPU任務(wù)管理、以及后續(xù)的必須由CPU來啟動(dòng)的GPU任務(wù)執(zhí)行之間存在順序依賴性。解決該問題的一個(gè)辦法是提供用于GPU線程程序的機(jī)制來對附加的計(jì)算任務(wù)進(jìn)行排隊(duì)而不要求來自CPU的干預(yù)，并等待那些計(jì)算任務(wù)的完成。然而，這種方法有幾個(gè)缺點(diǎn)。首先，常規(guī)地，CPU具有動(dòng)態(tài)地分配存儲(chǔ)器的裝置，但GPU沒有。當(dāng)新計(jì)算任務(wù)由GPU所啟動(dòng)時(shí)，計(jì)算任務(wù)被分配到存儲(chǔ)器以存儲(chǔ)在任務(wù)執(zhí)行期間所訪問的上下文和參數(shù)信息。在這種情況下，GPU使CPU分配存儲(chǔ)器用于新計(jì)算任務(wù)。然后，在對新任務(wù)進(jìn)行排隊(duì)之前，GPU等待CPU給計(jì)算任務(wù)分配存儲(chǔ)器，從而降低性能。其次，在CPU和GPU均能夠啟動(dòng)新計(jì)算任務(wù)進(jìn)入入棧緩沖區(qū)之處可能發(fā)生死鎖的情況。CPU可能出于對新計(jì)算任務(wù)進(jìn)行排隊(duì)的目的而占據(jù)對GPU的所有通信信道。然后GPU可能對為完成而訪問CPU的新計(jì)算任務(wù)進(jìn)行排隊(duì)。在這種情況下，CPU在釋放任意通信信道前等待GPU任務(wù)以完成，同時(shí)GPU任務(wù)不能完成直到任務(wù)被許可經(jīng)由所阻塞的通信信道之一訪問CPU為止，這導(dǎo)致死鎖。最后，對新計(jì)算任務(wù)進(jìn)行排隊(duì)和從入棧緩沖區(qū)拉取任務(wù)用于執(zhí)行典型地利用鎖定操作以確保任務(wù)順序地執(zhí)行以及確保正確地保存和管理入棧緩沖區(qū)中的信息。盡管GPU實(shí)施類似的鎖定操作，但鎖定操作本身是緩慢的。如果GPU在對新任務(wù)進(jìn)行排隊(duì)的同時(shí)采用鎖定操作，那么系統(tǒng)性能將受到負(fù)面的影響。如前所述，本領(lǐng)域中所需要的是允許GPU更有效率地對工作進(jìn)行排隊(duì)用于執(zhí)行的技術(shù)。

技術(shù)實(shí)現(xiàn)要素：
本發(fā)明的一個(gè)實(shí)施例闡述用于處理正由第一組線程所執(zhí)行并存儲(chǔ)在多個(gè)任務(wù)元數(shù)據(jù)描述符隊(duì)列(TMDQ)內(nèi)的多個(gè)任務(wù)的計(jì)算機(jī)實(shí)現(xiàn)的方法。方法包括接收包括在多個(gè)任務(wù)中的第一任務(wù)已完成的通知，以及在協(xié)處理單元內(nèi)確定是否包括在多個(gè)任務(wù)的子集中并與第一TMDQ相關(guān)聯(lián)的所有任務(wù)已經(jīng)執(zhí)行。如果并非包括在多個(gè)任務(wù)的子集中的所有任務(wù)已經(jīng)執(zhí)行，那么方法進(jìn)一步包括啟動(dòng)包括在多個(gè)任務(wù)中的第二任務(wù)。如果包括在多個(gè)任務(wù)的子集中的所有任務(wù)已經(jīng)執(zhí)行，那么方法進(jìn)一步包括更新與第一TMDQ相關(guān)聯(lián)的第一數(shù)據(jù)結(jié)構(gòu)中的指針、確定將在第一TMDQ中對包括在多個(gè)任務(wù)中的第三任務(wù)進(jìn)行排隊(duì)、以及啟動(dòng)第三任務(wù)。所公開的技術(shù)的一個(gè)優(yōu)點(diǎn)是GPU使能以在任務(wù)隊(duì)列內(nèi)對計(jì)算任務(wù)進(jìn)行排隊(duì)，還將任意數(shù)目的新任務(wù)隊(duì)列創(chuàng)建到任何任意嵌套級(jí)別，而不用CPU干預(yù)。在當(dāng)CPU創(chuàng)建任務(wù)并對任務(wù)進(jìn)行排隊(duì)的同時(shí)GPU不進(jìn)行等待的情況下，提高了處理效率。附圖說明因此，可以詳細(xì)地理解本發(fā)明的上述特征，并且可以參考實(shí)施例得到對如上面所簡要概括的本發(fā)明更具體的描述，其中一些實(shí)施例在附圖中示出。然而，應(yīng)當(dāng)注意的是，附圖僅示出了本發(fā)明的典型實(shí)施例，因此不應(yīng)被認(rèn)為是對其范圍的限制，本發(fā)明可以具有其他等效的實(shí)施例。圖1是示出配置為實(shí)現(xiàn)本發(fā)明的一個(gè)或多個(gè)方面的計(jì)算機(jī)系統(tǒng)的框圖；圖2是根據(jù)本發(fā)明的一個(gè)實(shí)施例的、用于圖1的計(jì)算機(jī)系統(tǒng)的并行處理子系統(tǒng)的框圖；圖3A是根據(jù)本發(fā)明的一個(gè)實(shí)施例的、圖2的前端的框圖；圖3B是根據(jù)本發(fā)明的一個(gè)實(shí)施例的、圖2的并行處理單元之一內(nèi)的通用處理集群的框圖；圖3C是根據(jù)本發(fā)明的一個(gè)實(shí)施例的、圖3B的流多處理器的一部分的框圖；圖4示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、并行處理子系統(tǒng)上的嵌套任務(wù)執(zhí)行；圖5示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、包括相關(guān)聯(lián)的任務(wù)元數(shù)據(jù)描述符隊(duì)列(TMDQ)和任務(wù)的層次執(zhí)行圖；圖6示出根據(jù)本發(fā)明的另一個(gè)實(shí)施例的、包括相關(guān)聯(lián)的TMDQ和任務(wù)的層次執(zhí)行圖；圖7示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、包括與線程組相關(guān)聯(lián)的參數(shù)和上下文信息的線程組上下文數(shù)據(jù)結(jié)構(gòu)；圖8示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、包括與計(jì)算任務(wù)相關(guān)聯(lián)的參數(shù)的任務(wù)狀況數(shù)據(jù)結(jié)構(gòu)；以及圖9是根據(jù)本發(fā)明的一個(gè)實(shí)施例的、用于處理已完成計(jì)算任務(wù)的方法步驟的流程圖。具體實(shí)施方式在下面的描述中，將闡述大量的具體細(xì)節(jié)以提供對本發(fā)明更透徹的理解。然而，本領(lǐng)域的技術(shù)人員應(yīng)該清楚，本發(fā)明可以在沒有一個(gè)或多個(gè)這些具體細(xì)節(jié)的情況下得以實(shí)施。系統(tǒng)概述圖1為示出了配置為實(shí)現(xiàn)本發(fā)明的一個(gè)或多個(gè)方面的計(jì)算機(jī)系統(tǒng)100的框圖。計(jì)算機(jī)系統(tǒng)100包括經(jīng)由可以包括存儲(chǔ)器橋105的互連路徑通信的中央處理單元（CPU）102和系統(tǒng)存儲(chǔ)器104。存儲(chǔ)器橋105可以是例如北橋芯片，經(jīng)由總線或其他通信路徑106（例如超傳輸（HyperTransport）鏈路）連接到I/O（輸入/輸出）橋107。I/O橋107，其可以是例如南橋芯片，從一個(gè)或多個(gè)用戶輸入設(shè)備108（例如鍵盤、鼠標(biāo)）接收用戶輸入并且經(jīng)由通信路徑106和存儲(chǔ)器橋105將該輸入轉(zhuǎn)發(fā)到CPU102。并行處理子系統(tǒng)112經(jīng)由總線或第二通信路徑113（例如外圍部件互連（PCI）Express、加速圖形端口或超傳輸鏈路）耦連到存儲(chǔ)器橋105；在一個(gè)實(shí)施例中，并行處理子系統(tǒng)112是將像素傳遞到顯示設(shè)備110（例如傳統(tǒng)的基于陰極射線管或液晶顯示器的監(jiān)視器）的圖形子系統(tǒng)。系統(tǒng)盤114也連接到I/O橋107。交換器116提供I/O橋107與諸如網(wǎng)絡(luò)適配器118以及各種插卡120和121的其他部件之間的連接。其他部件（未明確示出），包括通用串行總線（USB）或其他端口連接、壓縮磁盤（CD）驅(qū)動(dòng)器、數(shù)字視頻光盤（DVD）驅(qū)動(dòng)器、膠片錄制設(shè)備及類似部件，也可以連接到I/O橋107。圖1所示的各種通信路徑包括具體命名的通信路徑106和113可以使用任何適合的協(xié)議實(shí)現(xiàn)，諸如PCI-Express、AGP（加速圖形端口）、超傳輸或者任何其他總線或點(diǎn)到點(diǎn)通信協(xié)議，并且如本領(lǐng)域已知的，不同設(shè)備間的連接可使用不同協(xié)議。在一個(gè)實(shí)施例中，并行處理子系統(tǒng)112包含經(jīng)優(yōu)化用于圖形和視頻處理的電路，包括例如視頻輸出電路，并且構(gòu)成圖形處理單元（GPU）。在另一個(gè)實(shí)施例中，并行處理子系統(tǒng)112包含經(jīng)優(yōu)化用于通用處理的電路，同時(shí)保留底層（underlying）的計(jì)算架構(gòu)，本文將更詳細(xì)地進(jìn)行描述。在又一個(gè)實(shí)施例中，可以將并行處理子系統(tǒng)112與一個(gè)或多個(gè)其他系統(tǒng)元件集成在單個(gè)子系統(tǒng)中，諸如結(jié)合存儲(chǔ)器橋105、CPU102以及I/O橋107，以形成片上系統(tǒng)（SoC）。應(yīng)該理解，本文所示系統(tǒng)是示例性的，并且變化和修改都是可能的。連接拓?fù)洌虻臄?shù)目和布置、CPU102的數(shù)目以及并行處理子系統(tǒng)112的數(shù)目，可根據(jù)需要修改。例如，在一些實(shí)施例中，系統(tǒng)存儲(chǔ)器104直接連接到CPU102而不是通過橋，并且其他設(shè)備經(jīng)由存儲(chǔ)器橋105和CPU102與系統(tǒng)存儲(chǔ)器104通信。在其他替代性拓?fù)渲校⑿刑幚碜酉到y(tǒng)112連接到I/O橋107或直接連接到CPU102，而不是連接到存儲(chǔ)器橋105。而在其他實(shí)施例中，I/O橋107和存儲(chǔ)器橋105可能被集成到單個(gè)芯片上而不是作為一個(gè)或多個(gè)分立設(shè)備存在。大型實(shí)施例可以包括兩個(gè)或兩個(gè)以上的CPU102以及兩個(gè)或兩個(gè)以上的并行處理子系統(tǒng)112。本文所示的特定部件是可選的；例如，任何數(shù)目的插卡或外圍設(shè)備都可能得到支持。在一些實(shí)施例中，交換器116被去掉，網(wǎng)絡(luò)適配器118和插卡120、121直接連接到I/O橋107。圖2示出了根據(jù)本發(fā)明一個(gè)實(shí)施例的并行處理子系統(tǒng)112。如所示的，并行處理子系統(tǒng)112包括一個(gè)或多個(gè)并行處理單元（PPU）202，每個(gè)并行處理單元202都耦連到本地并行處理（PP）存儲(chǔ)器204。通常，并行處理子系統(tǒng)包括U個(gè)PPU，其中U≥1。（本文中，類似對象的多個(gè)實(shí)例需要時(shí)以標(biāo)識(shí)對象的參考數(shù)字和標(biāo)識(shí)實(shí)例的括號(hào)中的數(shù)字來表示。）PPU202和并行處理存儲(chǔ)器204可使用一個(gè)或多個(gè)集成電路設(shè)備來實(shí)現(xiàn)，諸如可編程處理器、專用集成電路（ASIC）或存儲(chǔ)器設(shè)備，或者以任何其他技術(shù)可行的方式來實(shí)現(xiàn)。再參考圖1以及圖2，在一些實(shí)施例中，并行處理子系統(tǒng)112中的一些或所有PPU202是具有渲染管線的圖形處理器，其可以配置為實(shí)施與下述相關(guān)的各種操作：經(jīng)由存儲(chǔ)器橋105和第二通信路徑113從CPU102和/或系統(tǒng)存儲(chǔ)器104所供應(yīng)的圖形數(shù)據(jù)生成像素?cái)?shù)據(jù)，與本地并行處理存儲(chǔ)器204（可被用作圖形存儲(chǔ)器，包括例如常規(guī)幀緩沖區(qū)（buffer））交互以存儲(chǔ)和更新像素?cái)?shù)據(jù)，傳遞像素?cái)?shù)據(jù)到顯示設(shè)備110等等。在一些實(shí)施例中，并行處理子系統(tǒng)112可包括一個(gè)或多個(gè)作為圖形處理器而操作的PPU202以及一個(gè)或多個(gè)用于通用計(jì)算的其他PPU202。這些PPU可以是同樣的或不同的，并且每個(gè)PPU可具有專用并行處理存儲(chǔ)器設(shè)備或不具有專用并行處理存儲(chǔ)器設(shè)備。并行處理子系統(tǒng)112中的一個(gè)或多個(gè)PPU202可輸出數(shù)據(jù)到顯示設(shè)備110，或者并行處理子系統(tǒng)112中的每個(gè)PPU202可輸出數(shù)據(jù)到一個(gè)或多個(gè)顯示設(shè)備110。在操作中，CPU102是計(jì)算機(jī)系統(tǒng)100的主處理器，控制和協(xié)調(diào)其他系統(tǒng)部件的操作。具體地，CPU102發(fā)出控制PPU202的操作的命令。在一些實(shí)施例中，CPU102寫入用于每個(gè)PPU202的命令流到數(shù)據(jù)結(jié)構(gòu)中（在圖1或圖2中未明確示出），該數(shù)據(jù)結(jié)構(gòu)可位于系統(tǒng)存儲(chǔ)器104、并行處理存儲(chǔ)器204、或CPU102和PPU202都可訪問的其他存儲(chǔ)位置中。將指向每個(gè)數(shù)據(jù)結(jié)構(gòu)的指針寫到入棧緩沖區(qū)（pushbuffer）以發(fā)起對數(shù)據(jù)結(jié)構(gòu)中的命令流的處理。PPU202從一個(gè)或多個(gè)入棧緩沖區(qū)讀取命令流，然后相對于CPU102的操作異步地執(zhí)行命令?？梢越?jīng)由設(shè)備驅(qū)動(dòng)程序103由應(yīng)用程序?yàn)槊總€(gè)入棧緩沖區(qū)指定執(zhí)行優(yōu)先級(jí)以控制對不同入棧緩沖區(qū)的調(diào)度。現(xiàn)在返回參考圖2和圖1，每個(gè)PPU202包括經(jīng)由連接到存儲(chǔ)器橋105（或者，在一個(gè)替代性實(shí)施例中，直接連接到CPU102）的通信路徑113與計(jì)算機(jī)系統(tǒng)100的其余部分通信的I/O（輸入/輸出）單元205。PPU202到計(jì)算機(jī)系統(tǒng)100的其余部分的連接也可以變化。在一些實(shí)施例中，并行處理子系統(tǒng)112可實(shí)現(xiàn)為可插入到計(jì)算機(jī)系統(tǒng)100的擴(kuò)展槽中的插卡。在其他實(shí)施例中，PPU202可以和諸如存儲(chǔ)器橋105或I/O橋107的總線橋集成在單個(gè)芯片上。而在其他實(shí)施例中，PPU202的一些或所有元件可以和CPU102集成在單個(gè)芯片上。在一個(gè)實(shí)施例中，通信路徑113是PCIExpress鏈路，如本領(lǐng)域所知的，其中專用通道被分配到每個(gè)PPU202。也可以使用其他通信路徑。I/O單元205生成用于在通信路徑113上傳送的包（或其他信號(hào)），并且還從通信路徑113接收所有傳入的包（或其他信號(hào)），將傳入的包引導(dǎo)到PPU202的適當(dāng)部件。例如，可將與處理任務(wù)相關(guān)的命令引導(dǎo)到主機(jī)接口206，而將與存儲(chǔ)器操作相關(guān)的命令（例如，對并行處理存儲(chǔ)器204的讀取或?qū)懭耄┮龑?dǎo)到存儲(chǔ)器交叉開關(guān)單元210。主機(jī)接口206讀取每個(gè)入棧緩沖區(qū)，并且將存儲(chǔ)在入棧緩沖區(qū)中的命令流輸出到前端212。有利地，每個(gè)PPU202都實(shí)現(xiàn)高度并行處理架構(gòu)。如詳細(xì)示出的，PPU202（0）包括處理集群陣列230，該陣列230包括C個(gè)通用處理集群（GPC）208，其中C≥1。每個(gè)GPC208能夠并發(fā)執(zhí)行大量的（例如，幾百或幾千）線程，其中每個(gè)線程是程序的實(shí)例（instance）。在各種應(yīng)用中，可分配不同的GPC208用于處理不同類型的程序或用于執(zhí)行不同類型的計(jì)算。GPC208的分配可以取決于因每種類型的程序或計(jì)算所產(chǎn)生的工作量而變化。GPC208從任務(wù)/工作單元207內(nèi)的工作分布單元接收所要執(zhí)行的處理任務(wù)。工作分布單元接收指向編碼為任務(wù)元數(shù)據(jù)（TMD）并存儲(chǔ)在存儲(chǔ)器中的處理任務(wù)的指針。指向TMD的指針包括在存儲(chǔ)為入棧緩沖區(qū)并由前端單元212從主機(jī)接口206接收的命令流中?？梢跃幋a為TMD的處理任務(wù)包括所要處理的數(shù)據(jù)的索引，以及定義數(shù)據(jù)將被如何處理（例如，什么程序?qū)⒈粓?zhí)行）的狀態(tài)參數(shù)和命令。任務(wù)/工作單元207從前端212接收任務(wù)并確保在每一個(gè)TMD所指定的處理發(fā)起前，將GPC208配置為有效狀態(tài)?？梢詾槊總€(gè)TMD指定用來調(diào)度處理任務(wù)的執(zhí)行的優(yōu)先級(jí)。還可從處理集群陣列230接收處理任務(wù)?？蛇x地，TMD可包括控制將TMD添加到處理任務(wù)列表（或指向處理任務(wù)的指針的列表）的頭部還是尾部的參數(shù)，從而提供除優(yōu)先級(jí)以外的另一級(jí)別的控制。存儲(chǔ)器接口214包括D個(gè)分區(qū)單元215，每個(gè)分區(qū)單元215直接耦連到并行處理存儲(chǔ)器204的一部分，其中D≥1。如所示的，分區(qū)單元215的數(shù)目一般等于動(dòng)態(tài)隨機(jī)存取存儲(chǔ)器（DRAM）220的數(shù)目。在其他實(shí)施例中，分區(qū)單元215的數(shù)目也可以不等于存儲(chǔ)器設(shè)備的數(shù)目。本領(lǐng)域的技術(shù)人員應(yīng)該理解DRAM220可以用其他合適的存儲(chǔ)設(shè)備來替代并且可以是一般常規(guī)的設(shè)計(jì)。因此省略了詳細(xì)描述。諸如幀緩沖區(qū)或紋理映射圖的渲染目標(biāo)可以跨DRAM220加以存儲(chǔ)，這允許分區(qū)單元215并行寫入每個(gè)渲染目標(biāo)的各部分以有效地使用并行處理存儲(chǔ)器204的可用帶寬。任何一個(gè)GPC208都可以處理要被寫到并行處理存儲(chǔ)器204內(nèi)的任何DRAM220的數(shù)據(jù)。交叉開關(guān)單元210配置為路由每個(gè)GPC208的輸出到任何分區(qū)單元215的輸入或到另一個(gè)GPC208用于進(jìn)一步處理。GPC208通過交叉開關(guān)單元210與存儲(chǔ)器接口214通信，以對各種外部存儲(chǔ)器設(shè)備進(jìn)行讀取或?qū)懭搿Ｔ谝粋€(gè)實(shí)施例中，交叉開關(guān)單元210具有到存儲(chǔ)器接口214的連接以和I/O單元205通信，以及到本地并行處理存儲(chǔ)器204的連接，從而使得在不同GPC208內(nèi)的處理內(nèi)核能夠與系統(tǒng)存儲(chǔ)器104或?qū)τ赑PU202而言非本地的其他存儲(chǔ)器通信。在圖2所示的實(shí)施例中，交叉開關(guān)單元210直接與I/O單元205連接。交叉開關(guān)單元210可使用虛擬信道來分開GPC208與分區(qū)單元215之間的業(yè)務(wù)流。另外，GPC208可被編程以執(zhí)行與種類繁多的應(yīng)用相關(guān)的處理任務(wù)，包括但不限于，線性和非線性數(shù)據(jù)變換、視頻和/或音頻數(shù)據(jù)過濾、建模操作（例如，應(yīng)用物理定律以確定對象的位置、速率和其他屬性）、圖像渲染操作（例如，曲面細(xì)分（tessellation）著色器、頂點(diǎn)著色器、幾何著色器、和/或像素著色器程序）等等。PPU202可將數(shù)據(jù)從系統(tǒng)存儲(chǔ)器104和/或本地并行處理存儲(chǔ)器204轉(zhuǎn)移到內(nèi)部（片上）存儲(chǔ)器中，處理該數(shù)據(jù)，并且將結(jié)果數(shù)據(jù)寫回到系統(tǒng)存儲(chǔ)器104和/或本地并行處理存儲(chǔ)器204，其中這樣的數(shù)據(jù)可以由其他系統(tǒng)部件訪問，所述其他系統(tǒng)部件包括CPU102或另一個(gè)并行處理子系統(tǒng)112。PPU202可配備有任何容量（amount）的本地并行處理存儲(chǔ)器204，包括沒有本地存儲(chǔ)器，并且可以以任何組合方式使用本地存儲(chǔ)器和系統(tǒng)存儲(chǔ)器。例如，在統(tǒng)一存儲(chǔ)器架構(gòu)（UMA）實(shí)施例中，PPU202可以是圖形處理器。在這樣的實(shí)施例中，將不提供或幾乎不提供專用的圖形（并行處理）存儲(chǔ)器，并且PPU202會(huì)以排他或幾乎排他的方式使用系統(tǒng)存儲(chǔ)器。在UMA實(shí)施例中，PPU202可集成到橋式芯片中或處理器芯片中，或作為具有高速鏈路（例如，PCIExpress）的分立芯片提供，所述高速鏈路經(jīng)由橋式芯片或其他通信手段將PPU202連接到系統(tǒng)存儲(chǔ)器。如上所示，在并行處理子系統(tǒng)112中可以包括任何數(shù)目的PPU202。例如，可在單個(gè)插卡上提供多個(gè)PPU202、或可將多個(gè)插卡連接到通信路徑113、或可將一個(gè)或多個(gè)PPU202集成到橋式芯片中。在多PPU系統(tǒng)中的PPU202可以彼此同樣或不同。例如，不同的PPU202可能具有不同數(shù)目的處理內(nèi)核、不同容量的本地并行處理存儲(chǔ)器等等。在存在多個(gè)PPU202的情況下，可并行操作那些PPU從而以高于單個(gè)PPU202所可能達(dá)到的吞吐量來處理數(shù)據(jù)。包含一個(gè)或多個(gè)PPU202的系統(tǒng)可以以各種配置和形式因素來實(shí)現(xiàn)，包括臺(tái)式電腦、筆記本電腦或手持式個(gè)人計(jì)算機(jī)、服務(wù)器、工作站、游戲控制臺(tái)、嵌入式系統(tǒng)等等。多個(gè)并發(fā)任務(wù)調(diào)度可以在GPC208上并發(fā)執(zhí)行多個(gè)處理任務(wù)并且處理任務(wù)在執(zhí)行期間可以生成一個(gè)或多個(gè)“子”處理任務(wù)。任務(wù)/工作單元207接收任務(wù)并動(dòng)態(tài)調(diào)度處理任務(wù)和子處理任務(wù)用于由GPC208執(zhí)行。圖3A為根據(jù)本發(fā)明一個(gè)實(shí)施例的圖2的任務(wù)/工作單元207的框圖。任務(wù)/工作單元207包括任務(wù)管理單元300和工作分布單元340。任務(wù)管理單元300基于執(zhí)行優(yōu)先級(jí)級(jí)別來組織所要調(diào)度的任務(wù)。對于每個(gè)優(yōu)先級(jí)級(jí)別，任務(wù)管理單元300將指向與任務(wù)相對應(yīng)的TMD322的指針的列表存儲(chǔ)在調(diào)度器表321中，其中所述列表可以實(shí)現(xiàn)為鏈表?？梢詫MD322存儲(chǔ)在PP存儲(chǔ)器204或系統(tǒng)存儲(chǔ)器104中。任務(wù)管理單元300接受任務(wù)并將任務(wù)存儲(chǔ)在調(diào)度器表321中的速度與任務(wù)管理單元300調(diào)度任務(wù)用于執(zhí)行的速度是解耦的。因此，任務(wù)管理單元300可以在調(diào)度任務(wù)之前收集數(shù)個(gè)任務(wù)。之后可以基于優(yōu)先級(jí)信息或使用其他技術(shù)諸如輪叫調(diào)度來調(diào)度所收集的任務(wù)。工作分布單元340包括具有槽的任務(wù)表345，每個(gè)槽可以被用于正在執(zhí)行的任務(wù)的TMD322所占用。當(dāng)任務(wù)表345中有空閑槽時(shí)，任務(wù)管理單元300可以調(diào)度任務(wù)用于執(zhí)行。當(dāng)沒有空閑槽時(shí)，未占用槽的較高優(yōu)先級(jí)任務(wù)可以驅(qū)逐占用槽的較低優(yōu)先級(jí)任務(wù)。當(dāng)任務(wù)被驅(qū)逐時(shí)，該任務(wù)被停止，并且如果該任務(wù)的執(zhí)行沒有完成，則將指向該任務(wù)的指針添加到所要調(diào)度的任務(wù)指針的列表以使得任務(wù)的執(zhí)行稍后恢復(fù)。當(dāng)生成子處理任務(wù)時(shí)，在任務(wù)的執(zhí)行期間，將指向該子任務(wù)的指針添加到所要調(diào)度的任務(wù)指針的列表?？梢杂稍谔幚砑宏嚵?30中執(zhí)行的TMD322生成子任務(wù)。不同于由任務(wù)/工作單元207從前端212接收的任務(wù)，子任務(wù)從處理集群陣列230接收。子任務(wù)不被插入入棧緩沖區(qū)或傳送到前端。當(dāng)生成子任務(wù)或?qū)⒂糜谧尤蝿?wù)的數(shù)據(jù)存儲(chǔ)在存儲(chǔ)器中時(shí)不通知CPU102。通過入棧緩沖區(qū)提供的任務(wù)與子任務(wù)之間的另一個(gè)區(qū)別是通過入棧緩沖區(qū)提供的任務(wù)由應(yīng)用程序來定義而子任務(wù)是在任務(wù)執(zhí)行期間動(dòng)態(tài)生成的。任務(wù)處理概述圖3B為根據(jù)本發(fā)明一個(gè)實(shí)施例的在圖2的PPU202之一內(nèi)的GPC208的框圖。每個(gè)GPC208可配置為并行執(zhí)行大量線程，其中術(shù)語“線程”是指在特定輸入數(shù)據(jù)集上執(zhí)行的特定程序的實(shí)例。在一些實(shí)施例中，單指令、多數(shù)據(jù)（SIMD）指令發(fā)出技術(shù)用于在不提供多個(gè)獨(dú)立指令單元的情況下支持大量線程的并行執(zhí)行。在其他實(shí)施例中，單指令、多線程（SIMT）技術(shù)用于使用配置為向GPC208中的每一個(gè)內(nèi)的處理引擎集發(fā)出指令的公共指令單元來支持大量一般來說同步的線程的并行執(zhí)行。不同于所有處理引擎通常都執(zhí)行同樣指令的SIMD執(zhí)行機(jī)制，SIMT執(zhí)行通過給定線程程序允許不同線程更容易跟隨分散執(zhí)行路徑。本領(lǐng)域普通技術(shù)人員應(yīng)該理解SIMD處理機(jī)制代表SIMT處理機(jī)制的功能子集。經(jīng)由將處理任務(wù)分布到流多處理器（SM）310的管線管理器305來有利地控制GPC208的操作。管線管理器305還可配置為通過為由SM310所輸出的處理數(shù)據(jù)指定目的地來控制工作分布交叉開關(guān)330。在一個(gè)實(shí)施例中，每個(gè)GPC208包括M個(gè)SM310，其中M≥1，每個(gè)SM310配置為處理一個(gè)或多個(gè)線程組。另外，如本領(lǐng)域已知的，每個(gè)SM310有利地包括可以管線化的同樣功能執(zhí)行單元集（例如執(zhí)行單元和加載-存儲(chǔ)單元—在圖3C中示出為Exec單元302和LSU303），其允許在前一個(gè)指令完成之前發(fā)出新指令。可提供功能執(zhí)行單元的任何組合。在一個(gè)實(shí)施例中，功能單元支持各種各樣的操作，包括整數(shù)和浮點(diǎn)運(yùn)算（例如加法和乘法）、比較操作、布爾操作（AND、OR、XOR）、移位和各種代數(shù)函數(shù)的計(jì)算（例如平面插值、三角函數(shù)、指數(shù)函數(shù)和對數(shù)函數(shù)等等）；以及相同功能單元硬件可均衡地用來實(shí)施不同的操作。如本文之前所定義的，傳送到特定GPC208的一系列指令構(gòu)成線程，并且跨SM310內(nèi)的并行處理引擎（未示出）的某一數(shù)目的并發(fā)執(zhí)行線程的集合在本文中稱為“線程束（warp）”或“線程組”。如本文所使用的，“線程組”是指對不同輸入數(shù)據(jù)并發(fā)執(zhí)行相同程序的一組線程，所述組的一個(gè)線程被指派到SM310內(nèi)的不同處理引擎。線程組可以包括比SM310內(nèi)的處理引擎數(shù)目少的線程，在這種情況下一些處理引擎在該線程組正在被處理的周期期間處于閑置狀態(tài)。線程組還可以包括比SM310內(nèi)的處理引擎數(shù)目多的線程，在這種情況下處理在連續(xù)的時(shí)鐘周期內(nèi)發(fā)生。因?yàn)槊總€(gè)SM310可以并發(fā)支持多達(dá)G個(gè)線程組，結(jié)果是在任何給定時(shí)間在GPC208中可以執(zhí)行多達(dá)G*M個(gè)線程組。此外，多個(gè)相關(guān)線程組可以在SM310內(nèi)同時(shí)活動(dòng)（在執(zhí)行的不同階段）。該線程組集合在本文中稱為“協(xié)作線程陣列”（“CTA”）或“線程陣列”。特定CTA的大小等于m*k，其中k是線程組中并發(fā)執(zhí)行線程的數(shù)目并且通常是SM310內(nèi)的并行處理引擎數(shù)目的整數(shù)倍，以及m是SM310內(nèi)同時(shí)活動(dòng)的線程組的數(shù)目。CTA的大小一般由編程者以及可用于CTA的硬件資源諸如存儲(chǔ)器或寄存器的容量來確定。每個(gè)SM310包括一級(jí)（L1）高速緩存（圖3C所示）或使用用于實(shí)施加載和存儲(chǔ)操作的SM310外部的相應(yīng)L1高速緩存中的空間。每個(gè)SM310都還有權(quán)訪問在所有GPC208之間共享并且可用于在線程之間轉(zhuǎn)移數(shù)據(jù)的二級(jí)（L2）高速緩存。最后，SM310還有權(quán)訪問片外“全局”存儲(chǔ)器，所述“全局”存儲(chǔ)器可以包括例如并行處理存儲(chǔ)器204和/或系統(tǒng)存儲(chǔ)器104。應(yīng)該理解，PPU202外部的任何存儲(chǔ)器可用作全局存儲(chǔ)器。此外，一點(diǎn)五級(jí)（L1.5）高速緩存335可以包括在GPC208內(nèi)，其配置為接收并保持由SM310所請求的經(jīng)由存儲(chǔ)器接口214從存儲(chǔ)器獲取的數(shù)據(jù)，包括指令、一致（uniform）數(shù)據(jù)和常數(shù)數(shù)據(jù)，并將所請求的數(shù)據(jù)提供給SM310。在GPC208中具有多個(gè)SM310的實(shí)施例有利地共享了高速緩存在L1.5高速緩存335中的公共指令和數(shù)據(jù)。每個(gè)GPC208可以包括配置為將虛擬地址映射到物理地址中的存儲(chǔ)器管理單元（MMU）328。在其他實(shí)施例中，MMU328可以駐留在存儲(chǔ)器接口214內(nèi)。MMU328包括用于將虛擬地址映射到像素塊（tile）的物理地址的頁表?xiàng)l目（PTE）集和可選地包括高速緩存行索引。MMU328可以包括地址轉(zhuǎn)換后備緩沖區(qū)（TLB）或可以駐留在多處理器SM310或L1高速緩存或GPC208內(nèi)的高速緩存。物理地址經(jīng)處理以分布表面數(shù)據(jù)訪問位置來允許高效請求在分區(qū)單元215之間交錯(cuò)。高速緩存行索引可用于確定用于高速緩存行的請求是命中還是未命中。在圖形和計(jì)算應(yīng)用中，GPC208可配置為使得每個(gè)SM310耦連到用于實(shí)施紋理映射操作例如確定紋理樣本位置、讀出紋理數(shù)據(jù)以及過濾該紋理數(shù)據(jù)的紋理單元315。從內(nèi)部紋理L1高速緩存（未示出）或者在一些實(shí)施例中從SM310內(nèi)的L1高速緩存讀出紋理數(shù)據(jù)并根據(jù)需要從在所有GPC208之間共享的L2高速緩存、并行處理存儲(chǔ)器204或系統(tǒng)存儲(chǔ)器104中獲取紋理數(shù)據(jù)。為了將所處理的任務(wù)提供給另一個(gè)GPC208用于進(jìn)一步處理或?yàn)榱私?jīng)由交叉開關(guān)單元210將所處理的任務(wù)存儲(chǔ)在L2高速緩存、并行處理存儲(chǔ)器204或系統(tǒng)存儲(chǔ)器104中，每個(gè)SM310將所處理的任務(wù)輸出到工作分布交叉開關(guān)330。preROP（預(yù)光柵操作）325配置為從SM310接收數(shù)據(jù)、將數(shù)據(jù)引導(dǎo)到分區(qū)單元215內(nèi)的ROP單元以及針對顏色混合實(shí)施優(yōu)化、組織像素顏色數(shù)據(jù)和實(shí)施地址轉(zhuǎn)譯。應(yīng)該理解本文所述的內(nèi)核架構(gòu)是示例性的并且變化和修改都是可能的。任何數(shù)目的處理單元例如SM310或紋理單元315、preROP325可以包括在GPC208內(nèi)。進(jìn)一步地，如圖2所示，PPU202可以包括任何數(shù)目的GPC208，所述GPC208有利地在功能上彼此相似以使得執(zhí)行行為不取決于哪個(gè)GPC208接收特定處理任務(wù)。進(jìn)一步地，每個(gè)GPC208有利地使用分開且各異的處理單元、L1高速緩存來獨(dú)立于其他GPC208操作以為一個(gè)或多個(gè)應(yīng)用程序執(zhí)行任務(wù)。本領(lǐng)域普通技術(shù)人員應(yīng)該理解圖1、2、3A和3B所描述的架構(gòu)決不限制本發(fā)明的范圍并且在不脫離本發(fā)明范圍的情況下本文所教導(dǎo)的技術(shù)可以在任何經(jīng)適當(dāng)配置的處理單元上實(shí)現(xiàn)，所述處理單元包括但不限于一個(gè)或多個(gè)CPU、一個(gè)或多個(gè)多核CPU、一個(gè)或多個(gè)PPU202、一個(gè)或多個(gè)GPC208、一個(gè)或多個(gè)圖形或?qū)Ｓ锰幚韱卧鹊取Ｔ诒景l(fā)明的實(shí)施例中，使用計(jì)算系統(tǒng)的PPU202或其他處理器來使用線程陣列執(zhí)行通用計(jì)算是可取的。為線程陣列中的每個(gè)線程指派在線程的執(zhí)行期間對于線程可訪問的唯一的線程標(biāo)識(shí)符（“線程ID”）?？杀欢x為一維或多維數(shù)值的線程ID控制線程處理行為的各方面。例如，線程ID可用于確定線程將要處理輸入數(shù)據(jù)集的哪部分和/或確定線程將要產(chǎn)生或?qū)戄敵鰯?shù)據(jù)集的哪部分。每線程指令序列可包括定義線程陣列的代表性線程和一個(gè)或多個(gè)其他線程之間的協(xié)作行為的至少一個(gè)指令。例如，每線程指令序列可能包括在序列中的特定點(diǎn)處暫停用于代表性線程的操作執(zhí)行直到諸如其他線程的一個(gè)或多個(gè)到達(dá)該特定點(diǎn)的時(shí)間為止的指令、用于代表性線程將數(shù)據(jù)存儲(chǔ)在其他線程的一個(gè)或多個(gè)有權(quán)訪問的共享存儲(chǔ)器中的指令、用于代表性線程原子地讀出和更新存儲(chǔ)在其他線程的一個(gè)或多個(gè)基于它們的線程ID有權(quán)訪問的共享存儲(chǔ)器中的數(shù)據(jù)的指令等等。CTA程序還可以包括計(jì)算數(shù)據(jù)將從其讀出的共享存儲(chǔ)器中的地址的指令，該地址是線程ID的函數(shù)。通過定義合適的函數(shù)并提供同步技術(shù)，可以以可預(yù)測的方式由CTA的一個(gè)線程將數(shù)據(jù)寫入共享存儲(chǔ)器中的給定位置并由同一個(gè)CTA的不同線程從該位置讀出數(shù)據(jù)。因此，數(shù)據(jù)在線程之間共享的任何期望模式可以得到支持，以及CTA中的任何線程可以與同一個(gè)CTA中的任何其他線程共享數(shù)據(jù)。如果存在數(shù)據(jù)在CTA的線程之間的共享，則其范圍由CTA程序確定；因此，應(yīng)該理解的是，在使用CTA的特定應(yīng)用中，CTA的線程可能會(huì)或可能不會(huì)真正互相共享數(shù)據(jù)，這取決于CTA程序，術(shù)語“CTA”和“線程陣列”在本文作為同義詞使用。圖3C為根據(jù)本發(fā)明一個(gè)實(shí)施例的圖3B的SM310的框圖。SM310包括配置為經(jīng)由L1.5高速緩存335從存儲(chǔ)器接收指令和常數(shù)的指令L1高速緩存370。線程束調(diào)度器和指令單元312從指令L1高速緩存370接收指令和常數(shù)并根據(jù)該指令和常數(shù)控制本地寄存器堆304和SM310功能單元。SM310功能單元包括N個(gè)exec（執(zhí)行或處理）單元302和P個(gè)加載-存儲(chǔ)單元（LSU）303。SM310提供具有不同級(jí)別的可訪問性的片上（內(nèi)部）數(shù)據(jù)存儲(chǔ)。特殊寄存器（未示出）對于LSU303可讀但不可寫并且用于存儲(chǔ)定義每個(gè)線程的“位置”的參數(shù)。在一個(gè)實(shí)施例中，特殊寄存器包括每線程（或SM310內(nèi)的每exec單元302）一個(gè)的存儲(chǔ)線程ID的寄存器；每個(gè)線程ID寄存器僅由各自的exec單元302可訪問。特殊寄存器還可以包括附加寄存器，其對于執(zhí)行由TMD322所代表的同一個(gè)處理任務(wù)的所有線程（或由所有LSU303）可讀，其存儲(chǔ)CTA標(biāo)識(shí)符、CTA維數(shù)、CTA所屬網(wǎng)格（grid）的維數(shù)（或隊(duì)列位置,如果TMD322編碼隊(duì)列任務(wù)而不是網(wǎng)格任務(wù)的話）、以及CTA被指派到的TMD322的標(biāo)識(shí)符。如果TMD322是網(wǎng)格TMD，則TMD322的執(zhí)行會(huì)啟動(dòng)和執(zhí)行固定數(shù)目的CTA以處理存儲(chǔ)在隊(duì)列525中的固定量的數(shù)據(jù)。將CTA的數(shù)目指定為網(wǎng)格寬度、高度和深度的乘積?？梢詫⒐潭康臄?shù)據(jù)存儲(chǔ)在TMD322中或TMD322可以存儲(chǔ)指向?qū)⒂蒀TA所處理的數(shù)據(jù)的指針。TMD322還存儲(chǔ)由CTA所執(zhí)行的程序的開始地址。如果TMD322是隊(duì)列TMD，那么使用TMD322的隊(duì)列特點(diǎn)，這意味著將要被處理的數(shù)據(jù)量不一定是固定的。隊(duì)列條目存儲(chǔ)用于由指派到TMD322的CTA所處理的數(shù)據(jù)。隊(duì)列條目還可以代表在線程執(zhí)行期間由另一個(gè)TMD322所生成的子任務(wù)，從而提供嵌套并行性。通常線程或包括線程的CTA的執(zhí)行被暫停直到子任務(wù)的執(zhí)行完成。可以將隊(duì)列存儲(chǔ)在TMD322中或與TMD322分開存儲(chǔ)，在該情況下TMD322存儲(chǔ)指向該隊(duì)列的隊(duì)列指針。有利地，當(dāng)代表子任務(wù)的TMD322正在執(zhí)行時(shí)可以將由子任務(wù)所生成的數(shù)據(jù)寫到隊(duì)列。隊(duì)列可以實(shí)現(xiàn)為循環(huán)隊(duì)列以使得數(shù)據(jù)的總量不限于隊(duì)列的大小。屬于網(wǎng)格的CTA具有指示網(wǎng)格內(nèi)各自CTA的位置的隱含網(wǎng)格寬度、高度和深度參數(shù)。在初始化期間響應(yīng)于經(jīng)由前端212從設(shè)備驅(qū)動(dòng)程序103所接收的命令來寫特殊寄存器并且在處理任務(wù)的執(zhí)行期間特殊寄存器不改變。前端212調(diào)度每個(gè)處理任務(wù)用于執(zhí)行。每個(gè)CTA與具體TMD322相關(guān)聯(lián)用于一個(gè)或多個(gè)任務(wù)的并發(fā)執(zhí)行。此外，單個(gè)GPC208可以并發(fā)執(zhí)行多個(gè)任務(wù)。參數(shù)存儲(chǔ)器（未示出）存儲(chǔ)可由同一個(gè)CTA內(nèi)的任何線程（或任何LSU303）讀取但不可由其寫入的運(yùn)行時(shí)間參數(shù)（常數(shù)）。在一個(gè)實(shí)施例中，設(shè)備驅(qū)動(dòng)程序103在引導(dǎo)SM310開始執(zhí)行使用參數(shù)的任務(wù)之前將這些參數(shù)提供給參數(shù)存儲(chǔ)器。任何CTA內(nèi)的任何線程（或SM310內(nèi)的任何exec單元302）可以通過存儲(chǔ)器接口214訪問全局存儲(chǔ)器?？梢詫⑷执鎯?chǔ)器的各部分存儲(chǔ)在L1高速緩存320中。每個(gè)線程將本地寄存器堆304用作暫存空間；每個(gè)寄存器被分配以專用于一個(gè)線程，并且在本地寄存器堆304的任何部分中的數(shù)據(jù)僅對于寄存器被分配到的線程可訪問。本地寄存器堆304可以實(shí)現(xiàn)為物理上或邏輯上分為P個(gè)通道的寄存器堆，每個(gè)通道具有一定數(shù)目的條目（其中每個(gè)條目可以存儲(chǔ)例如32位字）。將一個(gè)通道指派到N個(gè)exec單元302和P個(gè)下載-存儲(chǔ)單元LSU303的每一個(gè)，并且利用用于執(zhí)行同一個(gè)程序的不同線程的數(shù)據(jù)來填充不同通道中的相應(yīng)條目以幫助SIMD執(zhí)行?？梢詫⑼ǖ赖牟煌糠址峙涞紾個(gè)并發(fā)線程組中的不同線程組，以使得本地寄存器堆304中的給定條目僅對于特定線程可訪問。在一個(gè)實(shí)施例中，保留本地寄存器堆304內(nèi)的某些條目用于存儲(chǔ)線程標(biāo)識(shí)符，實(shí)現(xiàn)特殊寄存器之一。此外，一致L1高速緩存375存儲(chǔ)用于N個(gè)exec單元302和P個(gè)下載-存儲(chǔ)單元LSU303的每個(gè)通道的一致值或常數(shù)值。共享存儲(chǔ)器306對于單個(gè)CTA內(nèi)的線程可訪問；換言之，共享存儲(chǔ)器306中的任何位置對于同一個(gè)CTA內(nèi)的任何線程（或?qū)τ赟M310內(nèi)的任何處理引擎）可訪問。共享存儲(chǔ)器306可以實(shí)現(xiàn)為具有允許任何處理引擎對共享存儲(chǔ)器中的任何位置讀取或?qū)懭氲幕ミB的共享寄存器堆或共享片上高速緩存存儲(chǔ)器。在其他實(shí)施例中，共享狀態(tài)空間可能映射到片外存儲(chǔ)器的每CTA區(qū)上并被高速緩存在L1高速緩存320中。參數(shù)存儲(chǔ)器可以實(shí)現(xiàn)為在實(shí)現(xiàn)共享存儲(chǔ)器306的同一個(gè)共享寄存器堆或共享高速緩存存儲(chǔ)器內(nèi)的指定部分，或者實(shí)現(xiàn)為LSU303對其具有只讀訪問權(quán)限的分開的共享寄存器堆或片上高速緩存存儲(chǔ)器。在一個(gè)實(shí)施例中，實(shí)現(xiàn)參數(shù)存儲(chǔ)器的區(qū)域還用于存儲(chǔ)CTAID和任務(wù)ID，以及CTA和網(wǎng)格維數(shù)或隊(duì)列位置，實(shí)現(xiàn)特殊寄存器的各部分。SM310中的每個(gè)LSU303耦連到統(tǒng)一地址映射單元352，統(tǒng)一地址映射單元352將為在統(tǒng)一存儲(chǔ)器空間中所指定的加載和存儲(chǔ)指令所提供的地址轉(zhuǎn)換為每個(gè)各異存儲(chǔ)器空間中的地址。因此，指令可以用于通過指定統(tǒng)一存儲(chǔ)器空間中的地址來訪問本地、共享或全局存儲(chǔ)器空間中的任何一個(gè)。每個(gè)SM310中的L1高速緩存320可以用于高速緩存私有的每線程本地?cái)?shù)據(jù)還有每應(yīng)用全局?jǐn)?shù)據(jù)。在一些實(shí)施例中，可以將每CTA共享數(shù)據(jù)高速緩存在L1高速緩存320中。LSU303經(jīng)由存儲(chǔ)器和高速緩存互連380耦連到共享存儲(chǔ)器306和L1高速緩存320。嵌套執(zhí)行流圖4示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、并行處理子系統(tǒng)112上的嵌套任務(wù)執(zhí)行。如所示，CPU102在并行處理子系統(tǒng)112上發(fā)起示例性任務(wù)420的執(zhí)行。在任務(wù)420（0）完成之后，任務(wù)420（1）執(zhí)行。在任務(wù)420（1）完成之后，任務(wù)420（2）執(zhí)行。在執(zhí)行過程期間，例如任務(wù)420（1）喚起任務(wù)430（0）到430（2），以計(jì)算由任務(wù)420（1）所使用的中間結(jié)果。為了維持適當(dāng)?shù)闹噶顖?zhí)行次序，任務(wù)420（1）在繼續(xù)之前應(yīng)等待，直到任務(wù)430完成為止。為了以該方式進(jìn)行等待，任務(wù)420（1）可在任務(wù)430上的線程同步障礙（synchronizationbarrier）處阻塞。如先前所定義的，每個(gè)任務(wù)420、430可通過一個(gè)或多個(gè)線程、CTA、或網(wǎng)格來實(shí)施。雖然本文按照在圖形處理單元（GPU）的上下文內(nèi)描述并行處理子系統(tǒng)112，但本文所描述的技術(shù)可在與CPU102相關(guān)聯(lián)的任意協(xié)處理單元的上下文中實(shí)現(xiàn)。在該示例中，任務(wù)420（1）是任務(wù)430的父，因此任務(wù)430是任務(wù)420（1）的子。雖然圖4中僅示出一個(gè)級(jí)別的父子層次，但實(shí)際中可實(shí)現(xiàn)任意層次。在一個(gè)實(shí)施例中，任務(wù)420和430每個(gè)執(zhí)行為圖3B的SM310內(nèi)的至少一個(gè)CTA或至少一個(gè)線程組。為了使得具有父子關(guān)系的線程程序能夠在SM310上執(zhí)行，應(yīng)實(shí)現(xiàn)三個(gè)系統(tǒng)元件，包括用于并行處理子系統(tǒng)112的硬件功能、用于并行處理子系統(tǒng)112的軟件運(yùn)行時(shí)間功能、以及用于編程并行處理子系統(tǒng)112的語言支持構(gòu)造。支持父線程在并行處理子系統(tǒng)112內(nèi)啟動(dòng)子線程、CTA或網(wǎng)格所需的硬件功能包括啟動(dòng)來自由SM310所生成并被排隊(duì)用于執(zhí)行的對任務(wù)/工作單元207的請求的工作的新網(wǎng)格或CTA、保存用于SM310的執(zhí)行狀態(tài)、從所保存的執(zhí)行狀態(tài)繼續(xù)在SM310內(nèi)的執(zhí)行、以及促進(jìn)父和子任務(wù)之間的存儲(chǔ)器連貫性。支持父線程在并行處理子系統(tǒng)112內(nèi)啟動(dòng)子線程、CTA或網(wǎng)格所需的運(yùn)行時(shí)特征包括響應(yīng)于來自在SM310內(nèi)執(zhí)行的線程的請求而啟動(dòng)新網(wǎng)格、使得父線程能夠在子線程組上實(shí)施線程同步障礙、確保父線程和子組之間的存儲(chǔ)器連貫性、調(diào)度經(jīng)同步的線程組的工作和繼續(xù)用于所保證的向前的計(jì)算進(jìn)展、以及確保適當(dāng)?shù)膱?zhí)行語義用于父線程和子組。語言支持構(gòu)造包括用于指定來自父線程的子線程程序的啟動(dòng)、以及在子程序上執(zhí)行同步障礙的機(jī)制。使用面向線程的編程環(huán)境，諸如來自NVIDIA（tm）的CUDA（tm）編程環(huán)境來對并行處理子系統(tǒng)112進(jìn)行編程。在一個(gè)實(shí)施例中，CUDA語言規(guī)范經(jīng)擴(kuò)展以包括子啟動(dòng)構(gòu)造（“<<<>>>”）來指定用于啟動(dòng)子CUDA網(wǎng)格的細(xì)節(jié)。本文所指示為“A<<<B>>>C”的子啟動(dòng)構(gòu)造包括子程序名稱（A）、網(wǎng)格配置參數(shù)（B）、以及程序輸入?yún)?shù)（C）。CUDA運(yùn)行時(shí)環(huán)境經(jīng)擴(kuò)展以使得父線程能夠在子CUDA網(wǎng)格上實(shí)施同步障礙。雖然目前的討論在CUDA編程環(huán)境的上下文中示出本發(fā)明的實(shí)施例，但本領(lǐng)域技術(shù)人員將意識(shí)到，本文所教導(dǎo)的技術(shù)適用于任何并行編程環(huán)境和任何并行處理系統(tǒng)。同樣，對CUDA的引用僅用于例示性目的并且不旨在限制本發(fā)明的范圍或精神。下面的表1示出示例性CUDA程序中的子啟動(dòng)構(gòu)造和同步障礙的使用。表1在表1的示例中，線程程序“foo()”的實(shí)例使用具有指向由foo()所分配的存儲(chǔ)器的指針（*ptr）的線程程序“A”啟動(dòng)子網(wǎng)格。所分配的存儲(chǔ)器可由子網(wǎng)格內(nèi)的線程訪問。父線程foo()能夠在子網(wǎng)格A完成后繼續(xù)，由來自阻塞同步障礙函數(shù)調(diào)用的返回所指示，本文命名為cudaThreadSynchronize()。在GPU上所啟動(dòng)的任務(wù)通常是合格的以被立即執(zhí)行。缺少確保任務(wù)隊(duì)列內(nèi)的任務(wù)的順序執(zhí)行的機(jī)制，并行處理子系統(tǒng)112調(diào)度任何任務(wù)以開始執(zhí)行而不考慮對先前所啟動(dòng)到相同任務(wù)隊(duì)列中的任務(wù)的依賴性。順序執(zhí)行可通過如下文所描述的層次執(zhí)行圖的裝置來強(qiáng)制實(shí)行。圖5示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、包括相關(guān)聯(lián)的任務(wù)元數(shù)據(jù)描述符隊(duì)列(TMDQ)和任務(wù)的層次執(zhí)行圖。如所示，層次執(zhí)行圖包括處于嵌套深度0的線程組510、TMDQ512、任務(wù)520530540、處于嵌套深度1的執(zhí)行圖580、以及處于嵌套深度2的執(zhí)行圖590。處于嵌套深度0的線程組510包括由CPU102所創(chuàng)建和管理的線程。線程組包括任何線程集，該任何線程集包括CTA，其中所有線程存在于相同的嵌套深度。線程的嵌套深度是在線程級(jí)別之上的父網(wǎng)格的數(shù)目。例如，CPU線程具有嵌套深度0，因?yàn)樵贑PU線程之上沒有父網(wǎng)格。如果CPU線程啟動(dòng)網(wǎng)格，那么該網(wǎng)格就被稱為處于嵌套深度1。如果處于嵌套深度1的網(wǎng)格中的線程啟動(dòng)新網(wǎng)格，那么該新網(wǎng)格就被稱為處于嵌套深度2，以此類推。因?yàn)樵诰€程組510中的線程是CPU線程，所以這些線程中的每一個(gè)處于嵌套深度0。如上文結(jié)合圖2所描述的TMDQ512，包括指向被認(rèn)為是任務(wù)的數(shù)據(jù)結(jié)構(gòu)的指針，如下文所進(jìn)一步描述。每個(gè)TMDQ512指向?qū)儆谝粋€(gè)或多個(gè)流的任務(wù)。TMDQ（0）512（0）指向與第一流相關(guān)聯(lián)的任務(wù)520（0）。TMDQ（1）512（1）指向與第二流相關(guān)聯(lián)的任務(wù)530（0）和530（1）。TMDQ（2）512（2）指向與第三流相關(guān)聯(lián)的任務(wù)534（0）、540（1）、以及540（2）。在每個(gè)TMDQ512包括任意數(shù)目的任務(wù)的情況下可定義任何數(shù)目的TMDQ512。任務(wù)520530540是包括要由GPU所執(zhí)行的一個(gè)或多個(gè)命令的數(shù)據(jù)結(jié)構(gòu)。啟動(dòng)到給定的TMDQ512上的任務(wù)以順序的次序執(zhí)行。任務(wù)530（0）在任務(wù)530（1）開始執(zhí)行之前完成。同樣，任務(wù)540（0）在任務(wù)540（1）開始執(zhí)行之前完成，其依次在任務(wù)540（1）開始執(zhí)行之前完成。在TMDQ512的前面的任務(wù)一旦啟動(dòng)，該任務(wù)就開始執(zhí)行。所以，任務(wù)520（0）、530（0）、以及540（0）一旦啟動(dòng)，這些任務(wù)就執(zhí)行。在不同TMDQ512中的任務(wù)不具有順序依賴性。例如，任務(wù)530（1）可在任務(wù)540（1）之前、之后或與其并發(fā)執(zhí)行。處于嵌套深度1的執(zhí)行圖580是線程組加上相關(guān)聯(lián)的TMDQ和任務(wù)，其已經(jīng)由處于嵌套深度0的任務(wù)之一所啟動(dòng)。任何任務(wù)可啟動(dòng)一個(gè)或多個(gè)網(wǎng)格，其中這種網(wǎng)格處于比與啟動(dòng)網(wǎng)格的任務(wù)相關(guān)聯(lián)的嵌套深度大一個(gè)的嵌套深度。如所示，存在于嵌套深度0的任務(wù)540（1）在任務(wù)540（1）的執(zhí)行期間的某個(gè)時(shí)間啟動(dòng)執(zhí)行圖580。執(zhí)行圖580內(nèi)的每個(gè)任務(wù)和TMDQ基本與處于嵌套深度0的任務(wù)和TMDQ起相同的作用。當(dāng)執(zhí)行圖580內(nèi)的每個(gè)任務(wù)完成、以及任務(wù)540（1）中的所有其他命令已完成時(shí)，任務(wù)540（2）可開始執(zhí)行。處于嵌套深度2的執(zhí)行圖590是線程組加上相關(guān)聯(lián)的TMDQ和任務(wù)，其已經(jīng)由處于嵌套深度1的任務(wù)之一所啟動(dòng)。執(zhí)行圖590內(nèi)的每個(gè)任務(wù)和TMDQ基本與處于較低嵌套級(jí)別的任務(wù)和TMDQ起相同的作用。當(dāng)執(zhí)行圖590內(nèi)的每個(gè)任務(wù)完成時(shí)，之后一旦啟動(dòng)任務(wù)中的所有其他命令已完成，那么啟動(dòng)任務(wù)可完成。通過該方式，在保持流內(nèi)的任務(wù)的順序執(zhí)行的同時(shí)，在可嵌套到任意嵌套深度的任何網(wǎng)格內(nèi)保持順序執(zhí)行。依據(jù)上下文定義線程組內(nèi)的線程，其中上下文是具有對相同流和TMDQ資源的訪問權(quán)限的線程集。只要線程處于相同嵌套深度以及在相同設(shè)備上（GPU，或CPU102），那么相同上下文內(nèi)的線程就可創(chuàng)建以及共享TMDQ。對于CPU線程，上下文定義為與CUDA上下文相關(guān)聯(lián)的線程集。對于GPU線程，上下文可代表協(xié)作線程陣列（CTA）或存在于相同嵌套深度的任何線程集。當(dāng)新流由CPU線程所創(chuàng)建時(shí)，CPU102動(dòng)態(tài)地分配存儲(chǔ)器以支持流的管理。當(dāng)流在流任務(wù)的完成之后被隨后銷毀時(shí)，CPU102解放先前為流所分配的存儲(chǔ)器。GPU典型地不能動(dòng)態(tài)地分配存儲(chǔ)器。因此，GPU為可同時(shí)執(zhí)行的每個(gè)上下文預(yù)分配上下文數(shù)據(jù)。結(jié)果，與GPU網(wǎng)格相關(guān)聯(lián)的線程組具有固定數(shù)目的TMDQ，其在網(wǎng)格的執(zhí)行期間不可改變。用cudaStreamCreate()函數(shù)調(diào)用創(chuàng)建GPU網(wǎng)格內(nèi)的新流。函數(shù)調(diào)用返回指向網(wǎng)格中所預(yù)分配的TMDQ之一的整數(shù)索引。無需存儲(chǔ)器的動(dòng)態(tài)分配來創(chuàng)建流。一旦GPU流內(nèi)的所有任務(wù)已完成，就用cudaStreamDestroy()函數(shù)調(diào)用來銷毀流。因?yàn)闆]有動(dòng)態(tài)地為GPU流分配存儲(chǔ)器，所以cudaStreamDestroy()函數(shù)調(diào)用沒有存儲(chǔ)器來放回到空閑池并且因此簡單地返回到調(diào)用程序。一旦流已被創(chuàng)建，新任務(wù)就由相關(guān)聯(lián)的上下文中的一個(gè)或多個(gè)線程啟動(dòng)到流中。如果線程將新任務(wù)啟動(dòng)到當(dāng)前沒有任務(wù)的TMDQ中，那么新任務(wù)在任務(wù)啟動(dòng)之后立即開始執(zhí)行。同樣，如果TMDQ中的所有先前任務(wù)已完成執(zhí)行，那么啟動(dòng)到TMDQ中的新任務(wù)在任務(wù)啟動(dòng)之后立即開始執(zhí)行?？商娲兀绻€程將新任務(wù)啟動(dòng)到具有尚未完成執(zhí)行的一個(gè)或多個(gè)掛起的任務(wù)的TMDQ中，那么新任務(wù)啟動(dòng)到TMDQ中，但任務(wù)并不開始執(zhí)行直到掛起的在先任務(wù)完成執(zhí)行為止。無論哪種情況，新任務(wù)均經(jīng)由不要求CPU102的干預(yù)的非鎖定操作而啟動(dòng)到TMDQ中。圖6示出根據(jù)本發(fā)明的另一個(gè)實(shí)施例的、包括相關(guān)聯(lián)的TMDQ和任務(wù)的層次執(zhí)行圖。如所示，層次執(zhí)行圖包括處于嵌套深度1的線程組610、TMDQ612、任務(wù)620630640650660、處于嵌套深度2的執(zhí)行圖680、以及處于嵌套深度3的執(zhí)行圖690。層次執(zhí)行圖的部件除了下文所詳述的以外，大致與以上結(jié)合圖5所描述的起相同的作用。如所示，線程組610的每個(gè)TMDQ612具有一個(gè)或多個(gè)掛起的任務(wù)。在一個(gè)示例中，與流670相關(guān)聯(lián)的任務(wù)620（0）可能已啟動(dòng)到TMDQ612（0），但與流675相關(guān)聯(lián)的任務(wù)660（0）尚未啟動(dòng)。與一流相關(guān)聯(lián)的任務(wù)630可能已啟動(dòng)到TMDQ（1）612（1）中。同樣，與第二流相關(guān)聯(lián)的任務(wù)640可能已啟動(dòng)到TMDQ（2）612（2）中，與第三流相關(guān)聯(lián)的任務(wù)650可能已啟動(dòng)到TMDQ（N）612（N）中，并且所有干預(yù)的TMDQ612還可具有一個(gè)或多個(gè)相關(guān)聯(lián)的任務(wù)。在這時(shí)，線程組610內(nèi)的線程可嘗試創(chuàng)建新流675。然而，線程組610具有嵌套深度1，并且與GPU相關(guān)聯(lián)。因?yàn)镚PU不能動(dòng)態(tài)地分配存儲(chǔ)器，所以可能無法創(chuàng)建新TMDQ以容納新流675。在這種情況下，可將與新流675相關(guān)聯(lián)的任務(wù)660啟動(dòng)到當(dāng)前正由流670所使用的TMDQ（0）中。流675可將任務(wù)660（0）和660（1）啟動(dòng)到TMDQ（0）612（0）中。然后流670可將任務(wù)620（1）啟動(dòng)到TMDQ（0）612（0）中。然后流675可將任務(wù)660（2）啟動(dòng)到TMDQ（0）612（0）中。注意，該方法導(dǎo)致不需要的依賴性。即使流670和675相互獨(dú)立，但是TMDQ的順序性質(zhì)導(dǎo)致任務(wù)660（0）依賴于任務(wù)620（0）的完成、任務(wù)620（1）依賴于任務(wù)660（1）的完成，以此類推。雖然結(jié)果是性能可能降低，但流670中的任務(wù)620和流675中的任務(wù)660的順序排序被適當(dāng)?shù)乇Ａ簟D7示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、包括與線程組相關(guān)聯(lián)的參數(shù)和上下文信息的線程組上下文720數(shù)據(jù)結(jié)構(gòu)。如所示，線程組上下文720包括用于線程組中的每個(gè)TMDQ的最后任務(wù)指針740以及工作計(jì)數(shù)器750。最后任務(wù)指針740是指向相關(guān)聯(lián)的TMDQ中的最后任務(wù)的指針。當(dāng)新任務(wù)啟動(dòng)到TMDQ中時(shí)，最后任務(wù)指針740經(jīng)由原子操作而更新以指示新任務(wù)現(xiàn)在是TMDQ中的最后任務(wù)。下面的表2示出在示例性的CUDA程序中啟動(dòng)TMDQ中的新任務(wù)。表2在表2的示例中，用指向NewTask的指針來覆寫在地址StreamEnd處的最后任務(wù)指針740，并且最后任務(wù)指針740中的在先值作為FormerStreamEnd返回。如果FormerStreamEnd非零（即FormerStreamEnd是指向任務(wù)的指針），那么與任務(wù)相關(guān)聯(lián)的StreamNext值被更新以指向新啟動(dòng)的任務(wù)。如果FormerStreamEnd為零，那么沒有任務(wù)在TMDQ中正在掛起，并且新任務(wù)可立即開始執(zhí)行。表2的示例在操作的臨界段內(nèi)執(zhí)行以便避免在線程已將任務(wù)發(fā)布到流中、但隨后線程在啟動(dòng)新任務(wù)之前已被換出（swapout）的情況下的死鎖。在這種情況下，如果直到新任務(wù)已完成為止才允許所換出的線程被換回，那么死鎖可能發(fā)生。然而，新任務(wù)可能不開始執(zhí)行，因?yàn)樾氯蝿?wù)尚未啟動(dòng)。當(dāng)任務(wù)完成時(shí)，在并行處理子系統(tǒng)112上執(zhí)行的調(diào)度器讀取對應(yīng)于與已完成任務(wù)相關(guān)聯(lián)的TMDQ的最后流指針。如果相關(guān)聯(lián)的TMDQ的最后任務(wù)指針740不指向已完成任務(wù)，那么已完成任務(wù)不是TMDQ中的最后任務(wù)。在這種情況下，調(diào)度器使TMDQ中的下一個(gè)任務(wù)開始執(zhí)行，如下文結(jié)合圖8所描述的。如果相關(guān)聯(lián)的TMDQ的最后任務(wù)指針740指向已完成任務(wù)，那么已完成任務(wù)是TMDQ中的最后任務(wù)。在這種情況下，調(diào)度器實(shí)施原子比較和交換以將最后任務(wù)指針740設(shè)置為空（null）指針并讀取當(dāng)前存儲(chǔ)在最后任務(wù)指針740中的值。調(diào)度器以“currentEnd=atomicCAS(&StreamEnd,finishedTask,NULL),”的形式實(shí)施函數(shù)調(diào)用，其中“StreamEnd”是相關(guān)聯(lián)的TMDQ的最后任務(wù)指針740、“finishedTask”是指向已完成任務(wù)的指針、以及“NULL”是空指針。函數(shù)原子地返回存儲(chǔ)在最后任務(wù)指針740中的值，如由函數(shù)調(diào)用中的“currentEnd”所代表的。如果“currentEnd”的值是指向已完成任務(wù)的指針，那么TMDQ中的所有任務(wù)已完成，并且尚未啟動(dòng)新任務(wù)。調(diào)度器知道流中的所有任務(wù)已完成。如果“currentEnd”的值不是指向已完成任務(wù)的指針，那么新任務(wù)已啟動(dòng)，并且線程組上下文720已被更新以反映新任務(wù)的存在。在這種情況下，調(diào)度器讀取與已完成任務(wù)相關(guān)聯(lián)的StreamNext指針（下文所述）。如果與已完成任務(wù)相關(guān)聯(lián)的StreamNext指針非零，那么調(diào)度器使在地址StreamNext處的任務(wù)開始執(zhí)行。如果StreamNext的值是空指針，那么新任務(wù)已啟動(dòng)，但任務(wù)狀況尚未被更新以反映新任務(wù)的存在。在這種情況下，調(diào)度器監(jiān)視StreamNext直到值從空指針改變?yōu)橹赶蛐氯蝿?wù)的指針為止。然后調(diào)度器使由StreamNext所指向的新任務(wù)開始執(zhí)行。圖8示出根據(jù)本發(fā)明的一個(gè)實(shí)施例的、包括與計(jì)算任務(wù)相關(guān)聯(lián)的參數(shù)的任務(wù)狀況820數(shù)據(jù)結(jié)構(gòu)。如所示，任務(wù)狀況820包括任務(wù)標(biāo)識(shí)符（任務(wù)ID）840、下一個(gè)流指針842、線程組上下文標(biāo)識(shí)符（線程組上下文ID）844、以及與任務(wù)相關(guān)聯(lián)的其他參數(shù)（未示出）。任務(wù)ID840是指向與任務(wù)狀況820相關(guān)聯(lián)的任務(wù)的唯一標(biāo)識(shí)符。隨著在TMDQ上創(chuàng)建和啟動(dòng)任務(wù)，為每個(gè)新任務(wù)創(chuàng)建任務(wù)狀況820。任務(wù)ID使得調(diào)度器能夠查找與給定任務(wù)狀況820相關(guān)聯(lián)的任務(wù)。下一個(gè)流指針842是指向TMDQ中的下一個(gè)任務(wù)的指針。當(dāng)任務(wù)完成時(shí)，調(diào)度器讀取下一個(gè)流指針以確定在何處查找TMDQ中的可開始執(zhí)行的下一個(gè)任務(wù)。然后調(diào)度器使位于由下一個(gè)流指針842所指向的地址處的任務(wù)開始執(zhí)行。如果已完成任務(wù)是TMDQ中的最后任務(wù)，那么下一個(gè)流指針842設(shè)置為空指針。線程組上下文ID820是指向與任務(wù)狀況820相關(guān)聯(lián)的線程組上下文720的唯一標(biāo)識(shí)符。當(dāng)任務(wù)完成時(shí)，調(diào)度器讀取線程組上下文ID820以查找線程組上下文720。然后調(diào)度器可實(shí)施相關(guān)聯(lián)的任務(wù)完成步驟，諸如更新關(guān)閉TMDQ的工作計(jì)數(shù)器以及關(guān)閉上下文，如上文結(jié)合圖7所述。本領(lǐng)域的技術(shù)人員將意識(shí)到，本文所描述的技術(shù)僅是示例性的，可能進(jìn)行變化和修改。例如，所描述的技術(shù)足夠靈活以在任何并行編程環(huán)境和任何并行處理系統(tǒng)中所采用，無論與這種環(huán)境或系統(tǒng)相關(guān)聯(lián)的GPU或其他協(xié)處理器是否可動(dòng)態(tài)地分配存儲(chǔ)器。同樣，不論GPU是否預(yù)分配與TMDQ相關(guān)聯(lián)的存儲(chǔ)器或是否按照需要?jiǎng)討B(tài)地分配存儲(chǔ)器到TMDQ，均可采用所描述的技術(shù)。圖9是根據(jù)本發(fā)明的一個(gè)實(shí)施例的、用于處理已完成計(jì)算任務(wù)的方法步驟的流程圖。盡管結(jié)合圖1-8的系統(tǒng)來描述方法步驟，但本領(lǐng)域的普通技術(shù)人員應(yīng)予以理解的是配置為以任何次序?qū)嵤┓椒ú襟E的任何系統(tǒng)均在本發(fā)明的范圍內(nèi)。方法900在步驟902開始，其中調(diào)度器接收計(jì)算任務(wù)已完成的通知。在步驟902，調(diào)度器遞減與已完成任務(wù)所屬的線程組相關(guān)聯(lián)的工作計(jì)數(shù)器。在步驟906，調(diào)度器確定已完成任務(wù)是否是相關(guān)聯(lián)的TMDQ中的最后任務(wù)，也就是說，相關(guān)聯(lián)的TMDQ中的所有任務(wù)已完成。例如，如果與TMDQ相關(guān)聯(lián)的流結(jié)束指針指向已完成任務(wù)，那么調(diào)度器可確定已完成任務(wù)是TMDQ中的最后任務(wù)。如果已完成任務(wù)不是TMDQ中的最后任務(wù)，那么方法900前進(jìn)到步驟908，其中調(diào)度器使TMDQ中的下一個(gè)任務(wù)開始執(zhí)行。然后方法900終止。在步驟906，如果已完成任務(wù)是TMDQ中的最后任務(wù)，那么方法900前進(jìn)到步驟910，其中調(diào)度器使用原子操作將與TMDQ相關(guān)聯(lián)的流結(jié)束指針更新為空指針，這反映TMDQ現(xiàn)在為空。在步驟912，調(diào)度器確定CPU102或線程組是否已剛剛將新任務(wù)啟動(dòng)到隊(duì)列。例如，當(dāng)原子地將流結(jié)束指針更新為空指針時(shí)，調(diào)度器可確定指針已被改變以不再指向已完成任務(wù)。如果CPU102或線程組尚未啟動(dòng)新任務(wù)，那么方法900終止。在步驟912，如果CPU102或線程組已啟動(dòng)新任務(wù)，那么方法900前進(jìn)到步驟914，其中調(diào)度器等待新任務(wù)在TMDQ中排隊(duì)。例如，調(diào)度器可等待與TMDQ相關(guān)聯(lián)的流結(jié)束指針被更新以指向新任務(wù)。在步驟916，調(diào)度器使新任務(wù)開始執(zhí)行。然后方法900終止。總而言之，所公開的技術(shù)提供用于GPU將新計(jì)算任務(wù)排隊(duì)到TMDQ中的增強(qiáng)的方式。具體地，當(dāng)創(chuàng)建新TMDQ時(shí)用于上下文數(shù)據(jù)的存儲(chǔ)器被預(yù)分配，其中存儲(chǔ)器包括用于可在GPU上分開執(zhí)行的每個(gè)上下文的數(shù)據(jù)空間。當(dāng)新TMDQ被創(chuàng)建、并且CTA上下文沒有用于新TMDQ的可用條目時(shí)，新TMDQ可與現(xiàn)有的TMDQ整合，其中該TMDQ內(nèi)的計(jì)算任務(wù)包括來自初始TMDQ中的每一個(gè)的任務(wù)。排隊(duì)到TMDQ中的新計(jì)算任務(wù)可立即執(zhí)行或可等待TMDQ內(nèi)的在先任務(wù)完成。為了不使用原子鎖定操作而保留任務(wù)的順序執(zhí)行，對每個(gè)計(jì)算任務(wù)的完成執(zhí)行調(diào)度操作。有利地，所公開的技術(shù)使得GPU能夠?qū)⒂?jì)算任務(wù)在任務(wù)隊(duì)列內(nèi)排隊(duì)，還將任意數(shù)目的新任務(wù)隊(duì)列創(chuàng)建到任何任意嵌套級(jí)別，而無需CPU干預(yù)。在當(dāng)CPU創(chuàng)建任務(wù)和對任務(wù)進(jìn)行排隊(duì)的同時(shí)GPU不進(jìn)行等待的情況下，處理效率提高。因?yàn)镚PU不需要來自CPU的干預(yù)，所以避免了死鎖，甚至是在CPU消耗所有到GPU的通信信道的情況下。計(jì)算任務(wù)的順序執(zhí)行被保留用于由CPU和GPU均啟動(dòng)的任務(wù)。雖然上述內(nèi)容針對本發(fā)明的實(shí)施例，但可對本發(fā)明的其他和進(jìn)一步的實(shí)施例進(jìn)行設(shè)計(jì)而不脫離其基本范圍。例如，可以以硬件或軟件或硬件和軟件的組合來實(shí)現(xiàn)本發(fā)明的各方面。本發(fā)明的一個(gè)實(shí)施例可實(shí)現(xiàn)為與計(jì)算機(jī)系統(tǒng)一起使用的程序產(chǎn)品。該程序產(chǎn)品的程序定義實(shí)施例的各功能（包括本文中描述的方法）并且可以包含在各種計(jì)算機(jī)可讀存儲(chǔ)介質(zhì)上。示例性計(jì)算機(jī)可讀存儲(chǔ)介質(zhì)包括但不限于：（i）不可寫的存儲(chǔ)介質(zhì)（例如，計(jì)算機(jī)內(nèi)的只讀存儲(chǔ)器設(shè)備，諸如可由CD-ROM驅(qū)動(dòng)器讀取的CD-ROM盤、閃存、ROM芯片或任何類型的固態(tài)非易失性半導(dǎo)體存儲(chǔ)器），在其上永久性地存儲(chǔ)信息；和（ii）可寫的存儲(chǔ)介質(zhì)（例如，軟盤驅(qū)動(dòng)器內(nèi)的軟盤或硬盤驅(qū)動(dòng)器或者任何類型的固態(tài)隨機(jī)存取半導(dǎo)體存儲(chǔ)器），在其上存儲(chǔ)可更改的信息。當(dāng)承載針對本發(fā)明的功能的計(jì)算機(jī)可讀指令時(shí)，這樣的計(jì)算機(jī)可讀存儲(chǔ)介質(zhì)是本發(fā)明的實(shí)施例。因此，本發(fā)明的范圍由接下來的權(quán)利要求所確定。