CN103425533B - 用于管理嵌套执行流的方法和系统 - Google Patents

Abstract

本公开的一个实施例阐述用于GPU将新计算任务排队到任务源数据描述符队列(TMDQ)中的增强的方式。具体地，当创建新TMDQ时用于上下文数据的存储器被预分配。新TMDQ可与现有的TMDQ整合，其中该TMDQ内的计算任务包括来自初始TMDQ中的每一个的任务。为了不使用原子锁定操作而保留任务的顺序执行，对每个计算任务的完成执行调度操作。所公开的技术的一个优点是使得GPU能够将计算任务在TMDQ内排队，并且还将任意数目的新TMDQ创建到任何任意嵌套级别，而不用CPU干预。在当CPU创建任务和对任务进行排队的同时GPU不进行等待的情况下，处理效率提高。

Description

用于管理嵌套执行流的方法和系统

技术领域

本发明总地涉及计算机架构，并且，更具体地，涉及用于管理嵌套（nested）执行流的方法和系统。

背景技术

在具有中央处理单元(CPU)和图形处理单元(GPU)二者的常规计算系统中，CPU确定由GPU实施哪些具体计算任务以及以什么次序实施。GPU计算任务典型地包括跨并行数据集的高度并行、高度类似的操作，该并行数据集诸如图像或图像集。在常规GPU执行模型中，CPU通过选择相应的线程程序并且指导GPU执行线程程序的并行实例集来发起特定计算任务。在常规GPU执行模型中，CPU经常是可在GPU上发起线程程序的执行的仅有的实体。在所有线程实例完成执行后，GPU必须通知CPU并等待将由CPU所发出的另一个计算任务。通知CPU并等待下一个计算任务是使GPU内的某些资源暂时闲置的典型的阻塞型、序列化操作，从而降低整体系统性能。

在某些场景中可通过在入栈缓冲区中对顺序的计算任务进行排队，GPU可从该入栈缓冲区中拉取工作用于执行而不用等待CPU，从而改进性能。当CPU能够足够快地生成用于GPU的工作，使每当GPU能够开始新任务时工作均在入栈缓冲区内挂起（pending）时，包括固定数据流处理管线的计算任务从该入栈缓冲区模型中受益。然而，依赖于数据的计算任务仍在GPU结果、CPU任务管理、以及后续的必须由CPU来启动的GPU任务执行之间存在顺序依赖性。解决该问题的一个办法是提供用于GPU线程程序的机制来对附加的计算任务进行排队而不要求来自CPU的干预，并等待那些计算任务的完成。然而，这种方法有几个缺点。首先，常规地，CPU具有动态地分配存储器的装置，但GPU没有。当新计算任务由GPU所启动时，计算任务被分配到存储器以存储在任务执行期间所访问的上下文和参数信息。在这种情况下，GPU使CPU分配存储器用于新计算任务。然后，在对新任务进行排队之前，GPU等待CPU给计算任务分配存储器，从而降低性能。

其次，在CPU和GPU均能够启动新计算任务进入入栈缓冲区之处可能发生死锁的情况。CPU可能出于对新计算任务进行排队的目的而占据对GPU的所有通信信道。然后GPU可能对为完成而访问CPU的新计算任务进行排队。在这种情况下，CPU在释放任意通信信道前等待GPU任务以完成，同时GPU任务不能完成直到任务被许可经由所阻塞的通信信道之一访问CPU为止，这导致死锁。

最后，对新计算任务进行排队和从入栈缓冲区拉取任务用于执行典型地利用锁定操作以确保任务顺序地执行以及确保正确地保存和管理入栈缓冲区中的信息。尽管GPU实施类似的锁定操作，但锁定操作本身是缓慢的。如果GPU在对新任务进行排队的同时采用锁定操作，那么系统性能将受到负面的影响。

如前所述，本领域中所需要的是允许GPU更有效率地对工作进行排队用于执行的技术。

发明内容

本发明的一个实施例阐述用于处理正由第一组线程所执行并存储在多个任务元数据描述符队列(TMDQ)内的多个任务的计算机实现的方法。方法包括接收包括在多个任务中的第一任务已完成的通知，以及在协处理单元内确定是否包括在多个任务的子集中并与第一TMDQ相关联的所有任务已经执行。如果并非包括在多个任务的子集中的所有任务已经执行，那么方法进一步包括启动包括在多个任务中的第二任务。如果包括在多个任务的子集中的所有任务已经执行，那么方法进一步包括更新与第一TMDQ相关联的第一数据结构中的指针、确定将在第一TMDQ中对包括在多个任务中的第三任务进行排队、以及启动第三任务。

所公开的技术的一个优点是GPU使能以在任务队列内对计算任务进行排队，还将任意数目的新任务队列创建到任何任意嵌套级别，而不用CPU干预。在当CPU创建任务并对任务进行排队的同时GPU不进行等待的情况下，提高了处理效率。

附图说明

因此，可以详细地理解本发明的上述特征，并且可以参考实施例得到对如上面所简要概括的本发明更具体的描述，其中一些实施例在附图中示出。然而，应当注意的是，附图仅示出了本发明的典型实施例，因此不应被认为是对其范围的限制，本发明可以具有其他等效的实施例。

图1是示出配置为实现本发明的一个或多个方面的计算机系统的框图；

图2是根据本发明的一个实施例的、用于图1的计算机系统的并行处理子系统的框图；

图3A是根据本发明的一个实施例的、图2的前端的框图；

图3B是根据本发明的一个实施例的、图2的并行处理单元之一内的通用处理集群的框图；

图3C是根据本发明的一个实施例的、图3B的流多处理器的一部分的框图；

图4示出根据本发明的一个实施例的、并行处理子系统上的嵌套任务执行；

图5示出根据本发明的一个实施例的、包括相关联的任务元数据描述符队列(TMDQ)和任务的层次执行图；

图6示出根据本发明的另一个实施例的、包括相关联的TMDQ和任务的层次执行图；

图7示出根据本发明的一个实施例的、包括与线程组相关联的参数和上下文信息的线程组上下文数据结构；

图8示出根据本发明的一个实施例的、包括与计算任务相关联的参数的任务状况数据结构；以及

图9是根据本发明的一个实施例的、用于处理已完成计算任务的方法步骤的流程图。

具体实施方式

在下面的描述中，将阐述大量的具体细节以提供对本发明更透彻的理解。然而，本领域的技术人员应该清楚，本发明可以在没有一个或多个这些具体细节的情况下得以实施。

系统概述

图1为示出了配置为实现本发明的一个或多个方面的计算机系统100的框图。计算机系统100包括经由可以包括存储器桥105的互连路径通信的中央处理单元（CPU）102和系统存储器104。存储器桥105可以是例如北桥芯片，经由总线或其他通信路径106（例如超传输（HyperTransport）链路）连接到I/O（输入/输出）桥107。I/O桥107，其可以是例如南桥芯片，从一个或多个用户输入设备108（例如键盘、鼠标）接收用户输入并且经由通信路径106和存储器桥105将该输入转发到CPU102。并行处理子系统112经由总线或第二通信路径113（例如外围部件互连（PCI）Express、加速图形端口或超传输链路）耦连到存储器桥105；在一个实施例中，并行处理子系统112是将像素传递到显示设备110（例如传统的基于阴极射线管或液晶显示器的监视器）的图形子系统。系统盘114也连接到I/O桥107。交换器116提供I/O桥107与诸如网络适配器118以及各种插卡120和121的其他部件之间的连接。其他部件（未明确示出），包括通用串行总线（USB）或其他端口连接、压缩磁盘（CD）驱动器、数字视频光盘（DVD）驱动器、胶片录制设备及类似部件，也可以连接到I/O桥107。图1所示的各种通信路径包括具体命名的通信路径106和113可以使用任何适合的协议实现，诸如PCI-Express、AGP（加速图形端口）、超传输或者任何其他总线或点到点通信协议，并且如本领域已知的，不同设备间的连接可使用不同协议。

在一个实施例中，并行处理子系统112包含经优化用于图形和视频处理的电路，包括例如视频输出电路，并且构成图形处理单元（GPU）。在另一个实施例中，并行处理子系统112包含经优化用于通用处理的电路，同时保留底层（underlying）的计算架构，本文将更详细地进行描述。在又一个实施例中，可以将并行处理子系统112与一个或多个其他系统元件集成在单个子系统中，诸如结合存储器桥105、CPU102以及I/O桥107，以形成片上系统（SoC）。

应该理解，本文所示系统是示例性的，并且变化和修改都是可能的。连接拓扑，包括桥的数目和布置、CPU102的数目以及并行处理子系统112的数目，可根据需要修改。例如，在一些实施例中，系统存储器104直接连接到CPU102而不是通过桥，并且其他设备经由存储器桥105和CPU102与系统存储器104通信。在其他替代性拓扑中，并行处理子系统112连接到I/O桥107或直接连接到CPU102，而不是连接到存储器桥105。而在其他实施例中，I/O桥107和存储器桥105可能被集成到单个芯片上而不是作为一个或多个分立设备存在。大型实施例可以包括两个或两个以上的CPU102以及两个或两个以上的并行处理子系统112。本文所示的特定部件是可选的；例如，任何数目的插卡或外围设备都可能得到支持。在一些实施例中，交换器116被去掉，网络适配器118和插卡120、121直接连接到I/O桥107。

图2示出了根据本发明一个实施例的并行处理子系统112。如所示的，并行处理子系统112包括一个或多个并行处理单元（PPU）202，每个并行处理单元202都耦连到本地并行处理（PP）存储器204。通常，并行处理子系统包括U个PPU，其中U≥1。（本文中，类似对象的多个实例需要时以标识对象的参考数字和标识实例的括号中的数字来表示。）PPU202和并行处理存储器204可使用一个或多个集成电路设备来实现，诸如可编程处理器、专用集成电路（ASIC）或存储器设备，或者以任何其他技术可行的方式来实现。

再参考图1以及图2，在一些实施例中，并行处理子系统112中的一些或所有PPU202是具有渲染管线的图形处理器，其可以配置为实施与下述相关的各种操作：经由存储器桥105和第二通信路径113从CPU102和/或系统存储器104所供应的图形数据生成像素数据，与本地并行处理存储器204（可被用作图形存储器，包括例如常规帧缓冲区（buffer））交互以存储和更新像素数据，传递像素数据到显示设备110等等。在一些实施例中，并行处理子系统112可包括一个或多个作为图形处理器而操作的PPU202以及一个或多个用于通用计算的其他PPU202。这些PPU可以是同样的或不同的，并且每个PPU可具有专用并行处理存储器设备或不具有专用并行处理存储器设备。并行处理子系统112中的一个或多个PPU202可输出数据到显示设备110，或者并行处理子系统112中的每个PPU202可输出数据到一个或多个显示设备110。

在操作中，CPU102是计算机系统100的主处理器，控制和协调其他系统部件的操作。具体地，CPU102发出控制PPU202的操作的命令。在一些实施例中，CPU102写入用于每个PPU202的命令流到数据结构中（在图1或图2中未明确示出），该数据结构可位于系统存储器104、并行处理存储器204、或CPU102和PPU202都可访问的其他存储位置中。将指向每个数据结构的指针写到入栈缓冲区（pushbuffer）以发起对数据结构中的命令流的处理。PPU202从一个或多个入栈缓冲区读取命令流，然后相对于CPU102的操作异步地执行命令。可以经由设备驱动程序103由应用程序为每个入栈缓冲区指定执行优先级以控制对不同入栈缓冲区的调度。

现在返回参考图2和图1，每个PPU202包括经由连接到存储器桥105（或者，在一个替代性实施例中，直接连接到CPU102）的通信路径113与计算机系统100的其余部分通信的I/O（输入/输出）单元205。PPU202到计算机系统100的其余部分的连接也可以变化。在一些实施例中，并行处理子系统112可实现为可插入到计算机系统100的扩展槽中的插卡。在其他实施例中，PPU202可以和诸如存储器桥105或I/O桥107的总线桥集成在单个芯片上。而在其他实施例中，PPU202的一些或所有元件可以和CPU102集成在单个芯片上。

在一个实施例中，通信路径113是PCI Express链路，如本领域所知的，其中专用通道被分配到每个PPU202。也可以使用其他通信路径。I/O单元205生成用于在通信路径113上传送的包（或其他信号），并且还从通信路径113接收所有传入的包（或其他信号），将传入的包引导到PPU202的适当部件。例如，可将与处理任务相关的命令引导到主机接口206，而将与存储器操作相关的命令（例如，对并行处理存储器204的读取或写入）引导到存储器交叉开关单元210。主机接口206读取每个入栈缓冲区，并且将存储在入栈缓冲区中的命令流输出到前端212。

有利地，每个PPU202都实现高度并行处理架构。如详细示出的，PPU202（0）包括处理集群阵列230，该阵列230包括C个通用处理集群（GPC）208，其中C≥1。每个GPC208能够并发执行大量的（例如，几百或几千）线程，其中每个线程是程序的实例（instance）。在各种应用中，可分配不同的GPC208用于处理不同类型的程序或用于执行不同类型的计算。GPC208的分配可以取决于因每种类型的程序或计算所产生的工作量而变化。

GPC208从任务/工作单元207内的工作分布单元接收所要执行的处理任务。工作分布单元接收指向编码为任务元数据（TMD）并存储在存储器中的处理任务的指针。指向TMD的指针包括在存储为入栈缓冲区并由前端单元212从主机接口206接收的命令流中。可以编码为TMD的处理任务包括所要处理的数据的索引，以及定义数据将被如何处理（例如，什么程序将被执行）的状态参数和命令。任务/工作单元207从前端212接收任务并确保在每一个TMD所指定的处理发起前，将GPC208配置为有效状态。可以为每个TMD指定用来调度处理任务的执行的优先级。还可从处理集群阵列230接收处理任务。可选地，TMD可包括控制将TMD添加到处理任务列表（或指向处理任务的指针的列表）的头部还是尾部的参数，从而提供除优先级以外的另一级别的控制。

存储器接口214包括D个分区单元215，每个分区单元215直接耦连到并行处理存储器204的一部分，其中D≥1。如所示的，分区单元215的数目一般等于动态随机存取存储器（DRAM）220的数目。在其他实施例中，分区单元215的数目也可以不等于存储器设备的数目。本领域的技术人员应该理解DRAM220可以用其他合适的存储设备来替代并且可以是一般常规的设计。因此省略了详细描述。诸如帧缓冲区或纹理映射图的渲染目标可以跨DRAM220加以存储，这允许分区单元215并行写入每个渲染目标的各部分以有效地使用并行处理存储器204的可用带宽。

任何一个GPC208都可以处理要被写到并行处理存储器204内的任何DRAM220的数据。交叉开关单元210配置为路由每个GPC208的输出到任何分区单元215的输入或到另一个GPC208用于进一步处理。GPC208通过交叉开关单元210与存储器接口214通信，以对各种外部存储器设备进行读取或写入。在一个实施例中，交叉开关单元210具有到存储器接口214的连接以和I/O单元205通信，以及到本地并行处理存储器204的连接，从而使得在不同GPC208内的处理内核能够与系统存储器104或对于PPU202而言非本地的其他存储器通信。在图2所示的实施例中，交叉开关单元210直接与I/O单元205连接。交叉开关单元210可使用虚拟信道来分开GPC208与分区单元215之间的业务流。

另外，GPC208可被编程以执行与种类繁多的应用相关的处理任务，包括但不限于，线性和非线性数据变换、视频和/或音频数据过滤、建模操作（例如，应用物理定律以确定对象的位置、速率和其他属性）、图像渲染操作（例如，曲面细分（tessellation）着色器、顶点着色器、几何着色器、和/或像素着色器程序）等等。PPU202可将数据从系统存储器104和/或本地并行处理存储器204转移到内部（片上）存储器中，处理该数据，并且将结果数据写回到系统存储器104和/或本地并行处理存储器204，其中这样的数据可以由其他系统部件访问，所述其他系统部件包括CPU102或另一个并行处理子系统112。

PPU202可配备有任何容量（amount）的本地并行处理存储器204，包括没有本地存储器，并且可以以任何组合方式使用本地存储器和系统存储器。例如，在统一存储器架构（UMA）实施例中，PPU202可以是图形处理器。在这样的实施例中，将不提供或几乎不提供专用的图形（并行处理）存储器，并且PPU202会以排他或几乎排他的方式使用系统存储器。在UMA实施例中，PPU202可集成到桥式芯片中或处理器芯片中，或作为具有高速链路（例如，PCI Express）的分立芯片提供，所述高速链路经由桥式芯片或其他通信手段将PPU202连接到系统存储器。

如上所示，在并行处理子系统112中可以包括任何数目的PPU202。例如，可在单个插卡上提供多个PPU202、或可将多个插卡连接到通信路径113、或可将一个或多个PPU202集成到桥式芯片中。在多PPU系统中的PPU202可以彼此同样或不同。例如，不同的PPU202可能具有不同数目的处理内核、不同容量的本地并行处理存储器等等。在存在多个PPU202的情况下，可并行操作那些PPU从而以高于单个PPU202所可能达到的吞吐量来处理数据。包含一个或多个PPU202的系统可以以各种配置和形式因素来实现，包括台式电脑、笔记本电脑或手持式个人计算机、服务器、工作站、游戏控制台、嵌入式系统等等。

多个并发任务调度

可以在GPC208上并发执行多个处理任务并且处理任务在执行期间可以生成一个或多个“子”处理任务。任务/工作单元207接收任务并动态调度处理任务和子处理任务用于由GPC208执行。

图3A为根据本发明一个实施例的图2的任务/工作单元207的框图。任务/工作单元207包括任务管理单元300和工作分布单元340。任务管理单元300基于执行优先级级别来组织所要调度的任务。对于每个优先级级别，任务管理单元300将指向与任务相对应的TMD322的指针的列表存储在调度器表321中，其中所述列表可以实现为链表。可以将TMD322存储在PP存储器204或系统存储器104中。任务管理单元300接受任务并将任务存储在调度器表321中的速度与任务管理单元300调度任务用于执行的速度是解耦的。因此，任务管理单元300可以在调度任务之前收集数个任务。之后可以基于优先级信息或使用其他技术诸如轮叫调度来调度所收集的任务。

工作分布单元340包括具有槽的任务表345，每个槽可以被用于正在执行的任务的TMD322所占用。当任务表345中有空闲槽时，任务管理单元300可以调度任务用于执行。当没有空闲槽时，未占用槽的较高优先级任务可以驱逐占用槽的较低优先级任务。当任务被驱逐时，该任务被停止，并且如果该任务的执行没有完成，则将指向该任务的指针添加到所要调度的任务指针的列表以使得任务的执行稍后恢复。当生成子处理任务时，在任务的执行期间，将指向该子任务的指针添加到所要调度的任务指针的列表。可以由在处理集群阵列230中执行的TMD322生成子任务。

不同于由任务/工作单元207从前端212接收的任务，子任务从处理集群阵列230接收。子任务不被插入入栈缓冲区或传送到前端。当生成子任务或将用于子任务的数据存储在存储器中时不通知CPU102。通过入栈缓冲区提供的任务与子任务之间的另一个区别是通过入栈缓冲区提供的任务由应用程序来定义而子任务是在任务执行期间动态生成的。

任务处理概述

图3B为根据本发明一个实施例的在图2的PPU202之一内的GPC208的框图。每个GPC208可配置为并行执行大量线程，其中术语“线程”是指在特定输入数据集上执行的特定程序的实例。在一些实施例中，单指令、多数据（SIMD）指令发出技术用于在不提供多个独立指令单元的情况下支持大量线程的并行执行。在其他实施例中，单指令、多线程（SIMT）技术用于使用配置为向GPC208中的每一个内的处理引擎集发出指令的公共指令单元来支持大量一般来说同步的线程的并行执行。不同于所有处理引擎通常都执行同样指令的SIMD执行机制，SIMT执行通过给定线程程序允许不同线程更容易跟随分散执行路径。本领域普通技术人员应该理解SIMD处理机制代表SIMT处理机制的功能子集。

经由将处理任务分布到流多处理器（SM）310的管线管理器305来有利地控制GPC208的操作。管线管理器305还可配置为通过为由SM310所输出的处理数据指定目的地来控制工作分布交叉开关330。

在一个实施例中，每个GPC208包括M个SM310，其中M≥1，每个SM310配置为处理一个或多个线程组。另外，如本领域已知的，每个SM310有利地包括可以管线化的同样功能执行单元集（例如执行单元和加载-存储单元—在图3C中示出为Exec单元302和LSU303），其允许在前一个指令完成之前发出新指令。可提供功能执行单元的任何组合。在一个实施例中，功能单元支持各种各样的操作，包括整数和浮点运算（例如加法和乘法）、比较操作、布尔操作（AND、OR、XOR）、移位和各种代数函数的计算（例如平面插值、三角函数、指数函数和对数函数等等）；以及相同功能单元硬件可均衡地用来实施不同的操作。

如本文之前所定义的，传送到特定GPC208的一系列指令构成线程，并且跨SM310内的并行处理引擎（未示出）的某一数目的并发执行线程的集合在本文中称为“线程束（warp）”或“线程组”。如本文所使用的，“线程组”是指对不同输入数据并发执行相同程序的一组线程，所述组的一个线程被指派到SM310内的不同处理引擎。线程组可以包括比SM310内的处理引擎数目少的线程，在这种情况下一些处理引擎在该线程组正在被处理的周期期间处于闲置状态。线程组还可以包括比SM310内的处理引擎数目多的线程，在这种情况下处理在连续的时钟周期内发生。因为每个SM310可以并发支持多达G个线程组，结果是在任何给定时间在GPC208中可以执行多达G*M个线程组。

此外，多个相关线程组可以在SM310内同时活动（在执行的不同阶段）。该线程组集合在本文中称为“协作线程阵列”（“CTA”）或“线程阵列”。特定CTA的大小等于m*k，其中k是线程组中并发执行线程的数目并且通常是SM310内的并行处理引擎数目的整数倍，以及m是SM310内同时活动的线程组的数目。CTA的大小一般由编程者以及可用于CTA的硬件资源诸如存储器或寄存器的容量来确定。

每个SM310包括一级（L1）高速缓存（图3C所示）或使用用于实施加载和存储操作的SM310外部的相应L1高速缓存中的空间。每个SM310都还有权访问在所有GPC208之间共享并且可用于在线程之间转移数据的二级（L2）高速缓存。最后，SM310还有权访问片外“全局”存储器，所述“全局”存储器可以包括例如并行处理存储器204和/或系统存储器104。应该理解，PPU202外部的任何存储器可用作全局存储器。此外，一点五级（L1.5）高速缓存335可以包括在GPC208内，其配置为接收并保持由SM310所请求的经由存储器接口214从存储器获取的数据，包括指令、一致（uniform）数据和常数数据，并将所请求的数据提供给SM310。在GPC208中具有多个SM310的实施例有利地共享了高速缓存在L1.5高速缓存335中的公共指令和数据。

每个GPC208可以包括配置为将虚拟地址映射到物理地址中的存储器管理单元（MMU）328。在其他实施例中，MMU328可以驻留在存储器接口214内。MMU328包括用于将虚拟地址映射到像素块（tile）的物理地址的页表条目（PTE）集和可选地包括高速缓存行索引。MMU328可以包括地址转换后备缓冲区（TLB）或可以驻留在多处理器SM310或L1高速缓存或GPC208内的高速缓存。物理地址经处理以分布表面数据访问位置来允许高效请求在分区单元215之间交错。高速缓存行索引可用于确定用于高速缓存行的请求是命中还是未命中。

在图形和计算应用中，GPC208可配置为使得每个SM310耦连到用于实施纹理映射操作例如确定纹理样本位置、读出纹理数据以及过滤该纹理数据的纹理单元315。从内部纹理L1高速缓存（未示出）或者在一些实施例中从SM310内的L1高速缓存读出纹理数据并根据需要从在所有GPC208之间共享的L2高速缓存、并行处理存储器204或系统存储器104中获取纹理数据。为了将所处理的任务提供给另一个GPC208用于进一步处理或为了经由交叉开关单元210将所处理的任务存储在L2高速缓存、并行处理存储器204或系统存储器104中，每个SM310将所处理的任务输出到工作分布交叉开关330。preROP（预光栅操作）325配置为从SM310接收数据、将数据引导到分区单元215内的ROP单元以及针对颜色混合实施优化、组织像素颜色数据和实施地址转译。

应该理解本文所述的内核架构是示例性的并且变化和修改都是可能的。任何数目的处理单元例如SM310或纹理单元315、preROP325可以包括在GPC208内。进一步地，如图2所示，PPU202可以包括任何数目的GPC208，所述GPC208有利地在功能上彼此相似以使得执行行为不取决于哪个GPC208接收特定处理任务。进一步地，每个GPC208有利地使用分开且各异的处理单元、L1高速缓存来独立于其他GPC208操作以为一个或多个应用程序执行任务。

本领域普通技术人员应该理解图1、2、3A和3B所描述的架构决不限制本发明的范围并且在不脱离本发明范围的情况下本文所教导的技术可以在任何经适当配置的处理单元上实现，所述处理单元包括但不限于一个或多个CPU、一个或多个多核CPU、一个或多个PPU202、一个或多个GPC208、一个或多个图形或专用处理单元等等。

在本发明的实施例中，使用计算系统的PPU202或其他处理器来使用线程阵列执行通用计算是可取的。为线程阵列中的每个线程指派在线程的执行期间对于线程可访问的唯一的线程标识符（“线程ID”）。可被定义为一维或多维数值的线程ID控制线程处理行为的各方面。例如，线程ID可用于确定线程将要处理输入数据集的哪部分和/或确定线程将要产生或写输出数据集的哪部分。

每线程指令序列可包括定义线程阵列的代表性线程和一个或多个其他线程之间的协作行为的至少一个指令。例如，每线程指令序列可能包括在序列中的特定点处暂停用于代表性线程的操作执行直到诸如其他线程的一个或多个到达该特定点的时间为止的指令、用于代表性线程将数据存储在其他线程的一个或多个有权访问的共享存储器中的指令、用于代表性线程原子地读出和更新存储在其他线程的一个或多个基于它们的线程ID有权访问的共享存储器中的数据的指令等等。CTA程序还可以包括计算数据将从其读出的共享存储器中的地址的指令，该地址是线程ID的函数。通过定义合适的函数并提供同步技术，可以以可预测的方式由CTA的一个线程将数据写入共享存储器中的给定位置并由同一个CTA的不同线程从该位置读出数据。因此，数据在线程之间共享的任何期望模式可以得到支持，以及CTA中的任何线程可以与同一个CTA中的任何其他线程共享数据。如果存在数据在CTA的线程之间的共享，则其范围由CTA程序确定；因此，应该理解的是，在使用CTA的特定应用中，CTA的线程可能会或可能不会真正互相共享数据，这取决于CTA程序，术语“CTA”和“线程阵列”在本文作为同义词使用。

图3C为根据本发明一个实施例的图3B的SM310的框图。SM310包括配置为经由L1.5高速缓存335从存储器接收指令和常数的指令L1高速缓存370。线程束调度器和指令单元312从指令L1高速缓存370接收指令和常数并根据该指令和常数控制本地寄存器堆304和SM310功能单元。SM310功能单元包括N个exec（执行或处理）单元302和P个加载-存储单元（LSU）303。

SM310提供具有不同级别的可访问性的片上（内部）数据存储。特殊寄存器（未示出）对于LSU303可读但不可写并且用于存储定义每个线程的“位置”的参数。在一个实施例中，特殊寄存器包括每线程（或SM310内的每exec单元302）一个的存储线程ID的寄存器；每个线程ID寄存器仅由各自的exec单元302可访问。特殊寄存器还可以包括附加寄存器，其对于执行由TMD322所代表的同一个处理任务的所有线程（或由所有LSU303）可读，其存储CTA标识符、CTA维数、CTA所属网格（grid）的维数（或队列位置,如果TMD322编码队列任务而不是网格任务的话）、以及CTA被指派到的TMD322的标识符。

如果TMD322是网格TMD，则TMD322的执行会启动和执行固定数目的CTA以处理存储在队列525中的固定量的数据。将CTA的数目指定为网格宽度、高度和深度的乘积。可以将固定量的数据存储在TMD322中或TMD322可以存储指向将由CTA所处理的数据的指针。TMD322还存储由CTA所执行的程序的开始地址。

如果TMD322是队列TMD，那么使用TMD322的队列特点，这意味着将要被处理的数据量不一定是固定的。队列条目存储用于由指派到TMD322的CTA所处理的数据。队列条目还可以代表在线程执行期间由另一个TMD322所生成的子任务，从而提供嵌套并行性。通常线程或包括线程的CTA的执行被暂停直到子任务的执行完成。可以将队列存储在TMD322中或与TMD322分开存储，在该情况下TMD322存储指向该队列的队列指针。有利地，当代表子任务的TMD322正在执行时可以将由子任务所生成的数据写到队列。队列可以实现为循环队列以使得数据的总量不限于队列的大小。

属于网格的CTA具有指示网格内各自CTA的位置的隐含网格宽度、高度和深度参数。在初始化期间响应于经由前端212从设备驱动程序103所接收的命令来写特殊寄存器并且在处理任务的执行期间特殊寄存器不改变。前端212调度每个处理任务用于执行。每个CTA与具体TMD322相关联用于一个或多个任务的并发执行。此外，单个GPC208可以并发执行多个任务。

参数存储器（未示出）存储可由同一个CTA内的任何线程（或任何LSU303）读取但不可由其写入的运行时间参数（常数）。在一个实施例中，设备驱动程序103在引导SM310开始执行使用参数的任务之前将这些参数提供给参数存储器。任何CTA内的任何线程（或SM310内的任何exec单元302）可以通过存储器接口214访问全局存储器。可以将全局存储器的各部分存储在L1高速缓存320中。

每个线程将本地寄存器堆304用作暂存空间；每个寄存器被分配以专用于一个线程，并且在本地寄存器堆304的任何部分中的数据仅对于寄存器被分配到的线程可访问。本地寄存器堆304可以实现为物理上或逻辑上分为P个通道的寄存器堆，每个通道具有一定数目的条目（其中每个条目可以存储例如32位字）。将一个通道指派到N个exec单元302和P个下载-存储单元LSU303的每一个，并且利用用于执行同一个程序的不同线程的数据来填充不同通道中的相应条目以帮助SIMD执行。可以将通道的不同部分分配到G个并发线程组中的不同线程组，以使得本地寄存器堆304中的给定条目仅对于特定线程可访问。在一个实施例中，保留本地寄存器堆304内的某些条目用于存储线程标识符，实现特殊寄存器之一。此外，一致L1高速缓存375存储用于N个exec单元302和P个下载-存储单元LSU303的每个通道的一致值或常数值。

共享存储器306对于单个CTA内的线程可访问；换言之，共享存储器306中的任何位置对于同一个CTA内的任何线程（或对于SM310内的任何处理引擎）可访问。共享存储器306可以实现为具有允许任何处理引擎对共享存储器中的任何位置读取或写入的互连的共享寄存器堆或共享片上高速缓存存储器。在其他实施例中，共享状态空间可能映射到片外存储器的每CTA区上并被高速缓存在L1高速缓存320中。参数存储器可以实现为在实现共享存储器306的同一个共享寄存器堆或共享高速缓存存储器内的指定部分，或者实现为LSU303对其具有只读访问权限的分开的共享寄存器堆或片上高速缓存存储器。在一个实施例中，实现参数存储器的区域还用于存储CTA ID和任务ID，以及CTA和网格维数或队列位置，实现特殊寄存器的各部分。SM310中的每个LSU303耦连到统一地址映射单元352，统一地址映射单元352将为在统一存储器空间中所指定的加载和存储指令所提供的地址转换为每个各异存储器空间中的地址。因此，指令可以用于通过指定统一存储器空间中的地址来访问本地、共享或全局存储器空间中的任何一个。

每个SM310中的L1高速缓存320可以用于高速缓存私有的每线程本地数据还有每应用全局数据。在一些实施例中，可以将每CTA共享数据高速缓存在L1高速缓存320中。LSU303经由存储器和高速缓存互连380耦连到共享存储器306和L1高速缓存320。

嵌套执行流

图4示出根据本发明的一个实施例的、并行处理子系统112上的嵌套任务执行。如所示，CPU102在并行处理子系统112上发起示例性任务420的执行。在任务420（0）完成之后，任务420（1）执行。在任务420（1）完成之后，任务420（2）执行。在执行过程期间，例如任务420（1）唤起任务430（0）到430（2），以计算由任务420（1）所使用的中间结果。为了维持适当的指令执行次序，任务420（1）在继续之前应等待，直到任务430完成为止。为了以该方式进行等待，任务420（1）可在任务430上的线程同步障碍（synchronization barrier）处阻塞。如先前所定义的，每个任务420、430可通过一个或多个线程、CTA、或网格来实施。虽然本文按照在图形处理单元（GPU）的上下文内描述并行处理子系统112，但本文所描述的技术可在与CPU102相关联的任意协处理单元的上下文中实现。

在该示例中，任务420（1）是任务430的父，因此任务430是任务420（1）的子。虽然图4中仅示出一个级别的父子层次，但实际中可实现任意层次。在一个实施例中，任务420和430每个执行为图3B的SM310内的至少一个CTA或至少一个线程组。为了使得具有父子关系的线程程序能够在SM310上执行，应实现三个系统元件，包括用于并行处理子系统112的硬件功能、用于并行处理子系统112的软件运行时间功能、以及用于编程并行处理子系统112的语言支持构造。

支持父线程在并行处理子系统112内启动子线程、CTA或网格所需的硬件功能包括启动来自由SM310所生成并被排队用于执行的对任务/工作单元207的请求的工作的新网格或CTA、保存用于SM310的执行状态、从所保存的执行状态继续在SM310内的执行、以及促进父和子任务之间的存储器连贯性。支持父线程在并行处理子系统112内启动子线程、CTA或网格所需的运行时特征包括响应于来自在SM310内执行的线程的请求而启动新网格、使得父线程能够在子线程组上实施线程同步障碍、确保父线程和子组之间的存储器连贯性、调度经同步的线程组的工作和继续用于所保证的向前的计算进展、以及确保适当的执行语义用于父线程和子组。语言支持构造包括用于指定来自父线程的子线程程序的启动、以及在子程序上执行同步障碍的机制。

使用面向线程的编程环境，诸如来自NVIDIA（tm）的CUDA（tm）编程环境来对并行处理子系统112进行编程。在一个实施例中，CUDA语言规范经扩展以包括子启动构造（“<<<>>>”）来指定用于启动子CUDA网格的细节。本文所指示为“A<<<B>>>C”的子启动构造包括子程序名称（A）、网格配置参数（B）、以及程序输入参数（C）。CUDA运行时环境经扩展以使得父线程能够在子CUDA网格上实施同步障碍。虽然目前的讨论在CUDA编程环境的上下文中示出本发明的实施例，但本领域技术人员将意识到，本文所教导的技术适用于任何并行编程环境和任何并行处理系统。同样，对CUDA的引用仅用于例示性目的并且不旨在限制本发明的范围或精神。

下面的表1示出示例性CUDA程序中的子启动构造和同步障碍的使用。

表1

在表1的示例中，线程程序“foo()”的实例使用具有指向由foo()所分配的存储器的指针（*ptr）的线程程序“A”启动子网格。所分配的存储器可由子网格内的线程访问。父线程foo()能够在子网格A完成后继续，由来自阻塞同步障碍函数调用的返回所指示，本文命名为cudaThreadSynchronize()。

在GPU上所启动的任务通常是合格的以被立即执行。缺少确保任务队列内的任务的顺序执行的机制，并行处理子系统112调度任何任务以开始执行而不考虑对先前所启动到相同任务队列中的任务的依赖性。顺序执行可通过如下文所描述的层次执行图的装置来强制实行。

图5示出根据本发明的一个实施例的、包括相关联的任务元数据描述符队列(TMDQ)和任务的层次执行图。如所示，层次执行图包括处于嵌套深度0的线程组510、TMDQ512、任务520530540、处于嵌套深度1的执行图580、以及处于嵌套深度2的执行图590。

处于嵌套深度0的线程组510包括由CPU102所创建和管理的线程。线程组包括任何线程集，该任何线程集包括CTA，其中所有线程存在于相同的嵌套深度。线程的嵌套深度是在线程级别之上的父网格的数目。例如，CPU线程具有嵌套深度0，因为在CPU线程之上没有父网格。如果CPU线程启动网格，那么该网格就被称为处于嵌套深度1。如果处于嵌套深度1的网格中的线程启动新网格，那么该新网格就被称为处于嵌套深度2，以此类推。因为在线程组510中的线程是CPU线程，所以这些线程中的每一个处于嵌套深度0。

如上文结合图2所描述的TMDQ512，包括指向被认为是任务的数据结构的指针，如下文所进一步描述。每个TMDQ512指向属于一个或多个流的任务。TMDQ（0）512（0）指向与第一流相关联的任务520（0）。TMDQ（1）512（1）指向与第二流相关联的任务530（0）和530（1）。TMDQ（2）512（2）指向与第三流相关联的任务534（0）、540（1）、以及540（2）。在每个TMDQ512包括任意数目的任务的情况下可定义任何数目的TMDQ512。

任务520530540是包括要由GPU所执行的一个或多个命令的数据结构。启动到给定的TMDQ512上的任务以顺序的次序执行。任务530（0）在任务530（1）开始执行之前完成。同样，任务540（0）在任务540（1）开始执行之前完成，其依次在任务540（1）开始执行之前完成。在TMDQ512的前面的任务一旦启动，该任务就开始执行。所以，任务520（0）、530（0）、以及540（0）一旦启动，这些任务就执行。在不同TMDQ512中的任务不具有顺序依赖性。例如，任务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资源的访问权限的线程集。只要线程处于相同嵌套深度以及在相同设备上（GPU，或CPU102），那么相同上下文内的线程就可创建以及共享TMDQ。对于CPU线程，上下文定义为与CUDA上下文相关联的线程集。对于GPU线程，上下文可代表协作线程阵列（CTA）或存在于相同嵌套深度的任何线程集。

当新流由CPU线程所创建时，CPU102动态地分配存储器以支持流的管理。当流在流任务的完成之后被随后销毁时，CPU102解放先前为流所分配的存储器。GPU典型地不能动态地分配存储器。因此，GPU为可同时执行的每个上下文预分配上下文数据。结果，与GPU网格相关联的线程组具有固定数目的TMDQ，其在网格的执行期间不可改变。用cudaStreamCreate()函数调用创建GPU网格内的新流。函数调用返回指向网格中所预分配的TMDQ之一的整数索引。无需存储器的动态分配来创建流。一旦GPU流内的所有任务已完成，就用cudaStreamDestroy()函数调用来销毁流。因为没有动态地为GPU流分配存储器，所以cudaStreamDestroy()函数调用没有存储器来放回到空闲池并且因此简单地返回到调用程序。

一旦流已被创建，新任务就由相关联的上下文中的一个或多个线程启动到流中。如果线程将新任务启动到当前没有任务的TMDQ中，那么新任务在任务启动之后立即开始执行。同样，如果TMDQ中的所有先前任务已完成执行，那么启动到TMDQ中的新任务在任务启动之后立即开始执行。可替代地，如果线程将新任务启动到具有尚未完成执行的一个或多个挂起的任务的TMDQ中，那么新任务启动到TMDQ中，但任务并不开始执行直到挂起的在先任务完成执行为止。无论哪种情况，新任务均经由不要求CPU102的干预的非锁定操作而启动到TMDQ中。

图6示出根据本发明的另一个实施例的、包括相关联的TMDQ和任务的层次执行图。如所示，层次执行图包括处于嵌套深度1的线程组610、TMDQ612、任务620630640650660、处于嵌套深度2的执行图680、以及处于嵌套深度3的执行图690。层次执行图的部件除了下文所详述的以外，大致与以上结合图5所描述的起相同的作用。

如所示，线程组610的每个TMDQ612具有一个或多个挂起的任务。在一个示例中，与流670相关联的任务620（0）可能已启动到TMDQ612（0），但与流675相关联的任务660（0）尚未启动。与一流相关联的任务630可能已启动到TMDQ（1）612（1）中。同样，与第二流相关联的任务640可能已启动到TMDQ（2）612（2）中，与第三流相关联的任务650可能已启动到TMDQ（N）612（N）中，并且所有干预的TMDQ612还可具有一个或多个相关联的任务。在这时，线程组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的顺序排序被适当地保留。

图7示出根据本发明的一个实施例的、包括与线程组相关联的参数和上下文信息的线程组上下文720数据结构。如所示，线程组上下文720包括用于线程组中的每个TMDQ的最后任务指针740以及工作计数器750。

最后任务指针740是指向相关联的TMDQ中的最后任务的指针。当新任务启动到TMDQ中时，最后任务指针740经由原子操作而更新以指示新任务现在是TMDQ中的最后任务。下面的表2示出在示例性的CUDA程序中启动TMDQ中的新任务。

表2

在表2的示例中，用指向NewTask的指针来覆写在地址StreamEnd处的最后任务指针740，并且最后任务指针740中的在先值作为FormerStreamEnd返回。如果FormerStreamEnd非零（即FormerStreamEnd是指向任务的指针），那么与任务相关联的StreamNext值被更新以指向新启动的任务。如果FormerStreamEnd为零，那么没有任务在TMDQ中正在挂起，并且新任务可立即开始执行。

表2的示例在操作的临界段内执行以便避免在线程已将任务发布到流中、但随后线程在启动新任务之前已被换出（swap out）的情况下的死锁。在这种情况下，如果直到新任务已完成为止才允许所换出的线程被换回，那么死锁可能发生。然而，新任务可能不开始执行，因为新任务尚未启动。

当任务完成时，在并行处理子系统112上执行的调度器读取对应于与已完成任务相关联的TMDQ的最后流指针。如果相关联的TMDQ的最后任务指针740不指向已完成任务，那么已完成任务不是TMDQ中的最后任务。在这种情况下，调度器使TMDQ中的下一个任务开始执行，如下文结合图8所描述的。如果相关联的TMDQ的最后任务指针740指向已完成任务，那么已完成任务是TMDQ中的最后任务。在这种情况下，调度器实施原子比较和交换以将最后任务指针740设置为空（null）指针并读取当前存储在最后任务指针740中的值。调度器以“currentEnd=atomicCAS(&StreamEnd,finishedTask,NULL),”的形式实施函数调用，其中“StreamEnd”是相关联的TMDQ的最后任务指针740、“finishedTask”是指向已完成任务的指针、以及“NULL”是空指针。函数原子地返回存储在最后任务指针740中的值，如由函数调用中的“currentEnd”所代表的。

如果“currentEnd”的值是指向已完成任务的指针，那么TMDQ中的所有任务已完成，并且尚未启动新任务。调度器知道流中的所有任务已完成。如果“currentEnd”的值不是指向已完成任务的指针，那么新任务已启动，并且线程组上下文720已被更新以反映新任务的存在。在这种情况下，调度器读取与已完成任务相关联的StreamNext指针（下文所述）。如果与已完成任务相关联的StreamNext指针非零，那么调度器使在地址StreamNext处的任务开始执行。如果StreamNext的值是空指针，那么新任务已启动，但任务状况尚未被更新以反映新任务的存在。在这种情况下，调度器监视StreamNext直到值从空指针改变为指向新任务的指针为止。然后调度器使由StreamNext所指向的新任务开始执行。

图8示出根据本发明的一个实施例的、包括与计算任务相关联的参数的任务状况820数据结构。如所示，任务状况820包括任务标识符（任务ID）840、下一个流指针842、线程组上下文标识符（线程组上下文ID）844、以及与任务相关联的其他参数（未示出）。

任务ID840是指向与任务状况820相关联的任务的唯一标识符。随着在TMDQ上创建和启动任务，为每个新任务创建任务状况820。任务ID使得调度器能够查找与给定任务状况820相关联的任务。

下一个流指针842是指向TMDQ中的下一个任务的指针。当任务完成时，调度器读取下一个流指针以确定在何处查找TMDQ中的可开始执行的下一个任务。然后调度器使位于由下一个流指针842所指向的地址处的任务开始执行。如果已完成任务是TMDQ中的最后任务，那么下一个流指针842设置为空指针。

线程组上下文ID820是指向与任务状况820相关联的线程组上下文720的唯一标识符。当任务完成时，调度器读取线程组上下文ID820以查找线程组上下文720。然后调度器可实施相关联的任务完成步骤，诸如更新关闭TMDQ的工作计数器以及关闭上下文，如上文结合图7所述。

本领域的技术人员将意识到，本文所描述的技术仅是示例性的，可能进行变化和修改。例如，所描述的技术足够灵活以在任何并行编程环境和任何并行处理系统中所采用，无论与这种环境或系统相关联的GPU或其他协处理器是否可动态地分配存储器。同样，不论GPU是否预分配与TMDQ相关联的存储器或是否按照需要动态地分配存储器到TMDQ，均可采用所描述的技术。

图9是根据本发明的一个实施例的、用于处理已完成计算任务的方法步骤的流程图。尽管结合图1-8的系统来描述方法步骤，但本领域的普通技术人员应予以理解的是配置为以任何次序实施方法步骤的任何系统均在本发明的范围内。

方法900在步骤902开始，其中调度器接收计算任务已完成的通知。

在步骤902，调度器递减与已完成任务所属的线程组相关联的工作计数器。在步骤906，调度器确定已完成任务是否是相关联的TMDQ中的最后任务，也就是说，相关联的TMDQ中的所有任务已完成。例如，如果与TMDQ相关联的流结束指针指向已完成任务，那么调度器可确定已完成任务是TMDQ中的最后任务。如果已完成任务不是TMDQ中的最后任务，那么方法900前进到步骤908，其中调度器使TMDQ中的下一个任务开始执行。然后方法900终止。

在步骤906，如果已完成任务是TMDQ中的最后任务，那么方法900前进到步骤910，其中调度器使用原子操作将与TMDQ相关联的流结束指针更新为空指针，这反映TMDQ现在为空。在步骤912，调度器确定CPU102或线程组是否已刚刚将新任务启动到队列。例如，当原子地将流结束指针更新为空指针时，调度器可确定指针已被改变以不再指向已完成任务。如果CPU102或线程组尚未启动新任务，那么方法900终止。

在步骤912，如果CPU102或线程组已启动新任务，那么方法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）可写的存储介质（例如，软盘驱动器内的软盘或硬盘驱动器或者任何类型的固态随机存取半导体存储器），在其上存储可更改的信息。当承载针对本发明的功能的计算机可读指令时，这样的计算机可读存储介质是本发明的实施例。

因此，本发明的范围由接下来的权利要求所确定。

Claims (10)

1.一种用于处理正由第一组线程所执行并存储在多个任务元数据描述符队列TMDQ内的多个任务的计算机实现的方法，所述方法包括：

接收包括在所述多个任务中的第一任务已完成的通知；

在协处理单元内确定是否包括在所述多个任务的子集中并且与第一TMDQ相关联的所有任务已执行；

如果并非包括在所述多个任务的所述子集中的所有任务已执行，那么启动包括在所述多个任务中的第二任务；并且

如果包括在所述多个任务的所述子集中的所有任务已执行，那么：

更新与所述第一TMDQ相关联的第一数据结构中的指针；

确定包括在所述多个任务中的第三任务即将在所述第一TMDQ中排队；以及

执行所述第三任务。

2.一种用于处理正由第一组线程所执行并存储在多个任务元数据描述符队列TMDQ内的多个任务的子系统，包括：

任务管理单元，其配置为实施以下步骤：

接收包括在所述多个任务中的第一任务已完成的通知；

在协处理单元内确定是否包括在所述多个任务的子集中并且与第一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中的最后任务的指针指向所述第一任务。