CN116775266A - 用于处理器中的线程组的可扩展负载均衡的技术 - Google Patents

Abstract

本公开涉及用于处理器中的线程组的可扩展负载均衡的技术。处理器通过集中工作分配来支持新的线程组层次结构，以通过跨处理核心的推测性启动和负载均衡来提供线程组阵列中的线程组的硬件保证并发地执行。通过在处理核心之间分配网格光栅化来实现效率。

Description

用于处理器中的线程组的可扩展负载均衡的技术

相关申请的交叉引用

本申请涉及以下共同转让的共同未决的美国专利申请，将这些专利申请中的每一个的全部内容通过引用合并于此：

2022年3月10日提交的题目为“用于高效访问多维数据结构和/或其他大数据块的方法和装置(Method And Apparatus For Efficient Access To Multidimensional DataStructures and/or other Large Data Blocks)”的美国申请No.17/691,276；

2022年3月10日提交的题目为“协作组阵列(Cooperative Group Arrays)”的美国申请No.17/691,621；

2022年3月10日提交的题目为“分布式共享存储器(Distributed SharedMemory)”的美国申请No.17/691,690；

2022年3月10日提交的题目为“虚拟化处理器中的硬件处理资源 (VirtualizingHardware Processing Resources in a Processor)”的美国申请 No.17/691,759；

2022年3月10日提交的题目为“跨多个计算引擎的程序控制的数据多播(Programmatically Controlled Data Multicasting Across Multiple ComputeEngines)”的美国申请No.17/691,288；

2022年3月10日提交的题目为“具有异步事务支持的硬件加速的同步(HardwareAccelerated Synchronization with Asynchronous Transaction Support)”的美国申请No.17/691,296；

2022年3月10日提交的题目为“处理器和存储器中的快速数据同步 (Fast DataSynchronization in Processors And Memory)”的美国申请 No.17/691,303；

2022年3月10日提交的题目为“高效矩阵乘法和与一组线程束相加 (EfficientMatrix Multiply and Add with a Group of warps)”的美国申请 No.17/691,406；

2022年3月10日提交的题目为“不需要硬件复位的处理组件之间的执行软件的灵活迁移(Flexible Migration of Executing Software Between Processing ComponentsWithout Need For Hardware Reset)”的美国申请 No.17/691,808；以及

2022年3月10日提交的题目为“用于高效访问多维数据结构和/或其他大数据块的方法和装置(Method And Apparatus For Efficient Access To Multidimensional DataStructures and/or other Large Data Blocks)”的美国专利申请No.17/691,422。

背景技术

用户希望随着图形处理单元(GPU)技术的改进以及处理核心单元的数量随着每代而每芯片增加，深度学习和高性能计算(HPC)计算程序继续扩展。所期望的是单个应用程序的解决方案的更快时间，不仅仅通过运行N个独立的应用程序来扩展。

现代GPU是多处理器系统。每个GPU的处理器数量(我们称它们为“流式多处理器”或“SM”，但其他供应商可能会用其他名称来称呼它们，例如“执行单元”或“核心”)逐代增加，并且变得更具挑战性以与GPU 的大小成比例地扩展应用程序的性能。图1A示出了包括顺序相关的计算密集层的长链的示例深度学习(DL)网络。使用诸如例如将输入激活与权重矩阵相乘以产生输出激活的运算来计算每个层。所述层通常通过将工作划分为输出激活图块(tile)而跨GPU或GPU集群并行化，每个输出激活图块都表示一个SM或处理核心将处理的工作。

由于潜在大量的计算深度学习需求，通常目标更快。并且与连续地执行所有这些计算相比，并行执行许多计算将加速处理具有直观意义。事实上，应用通过在给定GPU实现上运行而实现的性能益处的量通常完全取决于它可以并行化的程度。但是存在不同的并行化方法。

在概念上，为了加速进程，可以使每个并行处理器执行多个工作，或者可以替代地保持每个并行处理器上的工作量恒定并且添加更多处理器。考虑重铺几英里长的高速路的工作量。你作为项目经理想要在最短的时间量内完成重铺工作以最小化业务中断。显然，如果你有几名工作人员在道路的不同部分并行工作，道路重铺工程将更快地完成。但是，哪种方法将使工作更快速地完成呢——要求每个道路工作人员做更多的工作，还是增加更多的工作人员，每个工作人员做相同的工作量？结果是，答案取决于工作的性质和用于支持该工作的资源。

图1B的弱扩展示例示出了每个处理核心运行的激活图块的大小增长，从而表明每个处理核心做更多工作。图1C的强扩展示例同时保持每个处理核心执行的工作量恒定(固定大小的网络由固定图块大小指示)，而增加并行操作的处理核心的数量(如由省略号指示)。表现出线性强扩展的应用具有等于所使用的处理器的数量的加速。参见 https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html； https://hpc-wiki.info/hpc/Scaling_tests。对于一些应用，诸如DL训练，问题大小通常将保持恒定，因此仅强扩展适用。

因此，此类应用的用户通常想要强扩展，这意味着单个应用可以在不必改变其工作负载(例如，通过增加其批大小来创建更固有的并行性)的情况下实现更高的性能。当在提供更多并行处理器的新的、容量更大的 GPU平台上运行现有的(例如，重新编译的)应用时，用户还期望提高的速度性能。如以下描述的，GPU开发已经满足或甚至超过市场在更多并行处理器和在那些并行处理器上运行的增加数量的并行执行线程之间的更多协调/协作方面的期望——但是仍需要进一步的性能改进以实现强扩展。

增加的GPU计算并行性和复杂性

多年来，GPU硬件已经变得越来越复杂并且能够实现增加的并行性。例如，图2A示出了较旧的GPU架构，该较旧的GPU架构提供具有16个 SM的流式执行模型，这16个流式多处理器在每个具有四个SM的集群 (GPC)中，其中每个SM表示GPU实际面积(real estate)的大部分。在此上下文中，SM或“流式多处理器”是指按照诺德奎斯特(Nordquist) 的USP7,447,873中所述架构的处理器，包括对其的改进和改进，并且例如在多代NVIDIA GPU中实现。在许多情况下，此类SM已被构建为提供快速的本地共享存储器，从而在SM上执行的所有线程之间实现数据共享/ 重用和同步。

相比之下，更多最近的GPU的半导体基板布局的图2B图示示出并行计算能力的显著增加，包括非常大量的(例如，128个或更多个)SM，每个SM仅表示GPU半导体基板实际面积的一小部分——其中数学计算硬件和每个SM内的并行处理核心的数量也随时间增长。

图2C示出了包括高级计算硬件能力的现代SM的示例架构图，该现代SM包括许多并行数学核心，这些并行数学核心除了纹理处理单元之外还包括多个张量核心。例如，2017NVIDIA Volta GV100 SM被划分成四个处理块，每个处理块具有16个FP32核心、8个FP64核心、16个INT32 核心、用于深度学习矩阵算术的两个混合精度张量核心、L0指令高速缓存、一个线程束调度器、一个分派单元和64KB寄存器文件——并且将来的 GPU设计可能继续该趋势。这种增加的计算并行性使得能够显著减少计算处理时间。

同时，图3和图4示出了现代GPU可以提供各种不同的硬件分区和层次结构。在这些示例中，GPU内的SM本身可分组成较大功能单元。例如，GPU的图形处理集群(GPC)可包括多个纹理处理集群(TPC)和流式多处理器(SM)的附加阵列(例如，用于计算能力)以及其他支持硬件，例如用于实时光线追踪加速的光线追踪单元。每个SM又可被分区成多个独立的处理块，每个处理块具有一个或若干个不同种类的核心(例如，FP32、 INT32、张量等)、线程束调度器、分派单元和本地寄存器文件。

图5和图5A示出了一些GPU实现方式(例如，NVIDIA Ampere)可以如何实现作为微GPU运行的多个分区，诸如μGPU0和μGPU1，其中，每个微GPU包括整个GPU的处理资源的一部分。当GPU被分区为两个或更多个单独的更小的μGPU以供不同客户端访问时，资源(包括物理存储器设备165，诸如本地L2高速缓存存储器)通常也被分区。例如，在一种设计中，被耦合到uGPU0的物理存储器设备165的第一半部可对应于第一组存储器分区位置，且被耦合到uGPU1的物理存储器设备165的第二半部可对应于第二组存储器分区位置。GPU内的性能资源还根据两个或更多单独的较小处理器分区来分区。资源可包括二级高速缓存(L2)资源170和处理资源160。这种多实例GPU(“MIG”)特征的一个实施例允许将GPU 安全地分区成用于CUDA(“计算统一设备架构”)应用的许多单独的GPU 实例，从而为多个用户提供单独的GPU资源以加速其相应的应用。

关于此类先前GPU硬件及其如何进展的更多信息，参见例如USP8，112，614；USP7，506，134；USP7，836，118；USP7，788，468；US10909033； US20140122809；林德霍姆(Lindholm)等人，“NVIDIA Tesla：统一图形和计算架构(A Unified Graphics andComputing Architecture)，”IEEE Micro(2008)；https://docs.nvidia.com/cuda/parallel-thread-execution/index.html (2021年检索到的)；肖凯特(Choquette)等人，“volta：性能和可编程性(Performance and Programmability)”，IEEE Micro(第38卷，发布于2018年3月/4月2日)，DOI:10.1109/MM.2018.022071134。

协作组API软件实现

为了利用由现代GPU提供的增加的并行性，CUDA版本9中的 NVIDIA引入了基于软件的“协作组”API，用于定义和同步CUDA程序中的线程组以允许内核动态地组织线程组。参见例如 https://developer.nvidia.com/blog/cooperative-groups/(2021年检索的)；https://developer.nvidia.com/blog/cuda-9-features-revealed/(2021年检索的)；鲍勃·克罗维拉(Bob Crovella)等人，“协作组(Cooperative Groups)” (09/17/2020)，https://vimeo.com/461821629；US 2020/0043123。

在协作组API之前，执行控制(即，线程同步)和线程间通信两者通常限于在一个SM上执行的线程块(也称为“协作线程阵列”或“CTA”) 的级别。协作组API扩展CUDA编程模型以描述在网格内和跨网格或跨多个网格并因此潜在地(取决于硬件平台)跨设备或多个设备的同步模式。

协作组API提供CUDA设备代码API，用于定义、分区和同步线程组——其中，“组(group)”是可编程的并且可以跨线程块延伸。协作组 API还提供主机侧API，用于启动其线程全部由基于软件的调度调度以并发地启动的网格。这些协作组API基元实现CUDA内的协作并行性的附加模式，包括跨整个线程网格或甚至跨多个GPU的生产者-消费者并行性和全局同步，而不需要底层GPU平台的硬件改变。

例如，协作组API提供全网格(grid-wide)(并且因此通常是全设备 (device-wide))同步屏障(“grid.sync()”)，该同步屏障可以用于防止网格组内的线程行进超过屏障，直到定义的网格组内的所有线程已经到达该屏障。这种全设备同步基于网格组(“grid_group”)的概念，该网格组在同一网格内定义一组线程，由软件调度为驻留在设备上并且可在该设备上调度，其方式为使得网格组中的每个线程可以向前进展。线程组的大小范围可以从几个线程(小于线程束)到整个线程块、到网格启动中的所有线程块、到跨越多个GPU的网格。较新的GPU平台(诸如NVIDIA Pascal 和Volta GPU)实现全网格和多GPU同步组，并且Volta的独立线程调度实现以任意交叉线程束和子线程束粒度对线程组的显著更灵活的选择和分区。

因此提供协作组API用于跨或甚至超出网格的协作/合作线程，但具有某些限制。例如，协作组API使用软件而不是硬件来提供并发地执行。在没有硬件级的并发保证的情况下，附加的API调用对于评估GPU占用以便预测网格组是否可以启动通常是必要的——并且因此在许多情况下确定 SM占用被留给软件应用。另外，虽然在一些平台上提供对全系统同步/存储器屏障的某些硬件支持，但是缺乏用于跨在不同SM上运行的线程块并因此跨设备或跨多个设备高效地共享数据带宽的高性能机制。作为一个重要的示例，不能跨多个SM高效地利用数据读取通常将导致冗余数据检索——其产生数据带宽不能跟上计算带宽的性能瓶颈。因为协作组API是基于软件的，所以它不能在硬件级上解决这些挑战。参见例如张(Zhang)等人，NVIDIA GPU中的单设备和多设备同步方法的研究(A Study of Single andMulti-device Synchronization Methods in NVIDIA GPU)， (arXiv:2004.05371v1[cs.DC]2020年4月11日)；卢斯蒂格(Lustig)等人，“NVIDIA PTX存储器一致性模型的正式分析(A Formal Analysis of the NVIDIA PTX Memory Consistency Model)”，编程语言和操作系统的架构支持的第二十四届国际会议的论文集，第257-270页(2019年4月)https://doi.org/10.1145/3297858.3304043；韦伯(Weber)等人，“朝向模块化整数GCD算法扩展摘要的多GPU实现方式(Toward a Multi-GPU Implementation of the ModularInteger GCD Algorithm Extended Abstract)” ICPP 2018年8月13-16日，尤金(Eugene)，俄勒冈州(OR USA)(ACM 2018)；乔格(Jog)等人，“OWL：提高GPGPU性能的协作线程阵列感知调度技术(Cooperative Thread Array Aware Scheduling Techniques for ImprovingGPGPU Performance))”(ASPLOS’13，2013年3月16-20 日，休斯顿，德克萨斯州，美国)。

数据带宽跟不上处理带宽

虽然已经有可能增加每代新的GPU硬件的数学吞吐量，但是越来越难以向新的GPU硬件中的SM或处理核心(例如，张量核心)的其他集合馈送足够的数据来维持强扩展。图6对不同类型的数学计算(例如，张量浮点32位精度、浮点16精度、“脑”浮点16位精度、整数8位精度、整数4位精度以及二进制)针对各种不同GPU生成并且还针对不同数据呈现 (稀疏和密集)比较数学带宽(每时钟每SM的乘积-加法计算的数量)。图6的左手侧示出了理论数学计算带宽如何随着GPU计算硬件能力的增加而呈指数增加(例如，通过向GPU添加具有张量核心或其他核心的大规模并行SM)。同时，然而，图6的右手侧示出了保持GPU计算硬件被提供有数据的相应的数据带宽要求还没有保持步调。

经验已经显示，存储器带宽和互连带宽(例如，从存储器系统到SM 中)以及处理带宽不扩展。用于支持张量核心和其他数学计算的GPU系统内的基本数据流(即，从互连到系统DRAM存储器，到L2高速缓存，到 L1高速缓存中的共享存储器，到SM内的数学计算处理器)的图7流程图表明，为了实现强扩展，有必要提高跨计算和存储器层次结构的所有级别 (端到端)的速度以及馈送和效率。

已经尝试并实现了各种技术(诸如存储器管理改进、高速缓存改进等) 以增加数据带宽。然而，经由导线增加更多的数据带宽会耗费面积和功率。增加更多的高速缓存成本面积和功率。所需要的是在不需要彻底检修和使存储器访问/管理层次结构复杂化的情况下，利用一个或更多个算法中固有的更多并行性，同时更高效地使用现今和将来可用的处理核心和高速缓存/ 互连层次结构的方式。同时，一旦提供了解决此类数据带宽问题的新硬件平台，就非常需要跨硬件平台的并发并行处理资源提供高效的工作负载均衡。

附图说明

图1A示出了在GPU上运行的示例应用。

图1B示出了弱扩展深度学习场景。

图1C示出了强扩展深度学习场景。

图2A和图2B示出了增加的GPU硬件并行性。

图2C是GPU内的最近流式多处理器的框架构图。

图3示出了示例现有技术GPU硬件分区。

图4示出了具有图形处理集群的示例现有技术GPU硬件。

图5示出了示例现有技术μGPU分区。

图5A是最近GPU架构的框架构图，该最近GPU架构包括被分区为不同μGPU分区的流式多处理器和关联的互连。

图6示出了针对不同GPU生成的示例增加的数学吞吐量和关联的数据带宽。

图7示出了提高跨所有级别的计算和存储器层次结构的速度和馈送以及效率以实现强扩展的示例需要。

图8示出了CTA的示例现有技术网格。

图9示出了CTA的现有技术网格如何映射到GPU硬件分区上的示例。

图10A示出了CTA的示例现有技术网格。

图10B示出了示例新的CGA层次结构。

图11A示出了CTA的示例现有技术网格。

图11B示出了将CTA分组为CGA的示例新的CGA层次结构。

图12、图12A示出了示例CGA网格布置。

图13示出了包括与其他相关电路交互的计算工作分配器电路的框图。

图14示出了包括改进的CGA负载均衡器的示例计算工作分配器电路框图。

图15示出了包括用于将CTA分配到不同硬件分区级别内的SM的硬件工作分配器电路的分层计算工作分配器电路。

图16A示出了示例流程图，该示例流程图示出了使用负载均衡的推测性或“影子状态”启动。

图16B示出了示例CGA启动操作，该示例CGA启动操作包括对基于硬件的查询模型的推测性或影子状态启动，如果推测性或影子状态启动成功，则保存保留信息，然后使用保存的预留信息实际启动CGA(在真实硬件上，而不是查询模型)。

图16C、图16D示出了包括CGA的推测性或“影子状态(shadow state)”启动以提供基于硬件的并发保证的硬件实现的操作步骤的示例非限制性流程图。

图17示出了源和目标SM之间的示例通信。

图18A-18D示出了SM间通信。

图19是SM间通信的另一个视图。

图20是非程序计数器误差检查的流程图。

图21示出了示例非限制性负载均衡场景。

图21A-21Z、图21AA-21II是一起示出计算工作分配器如何将CGA 内的CTA分配给SM的浏览图表动画(I)。

图22和图23是浏览图表动画的延续，该浏览图表动画示出了示例网格的最后CTA的负载均衡分布。

图24、图25A-25G、图26和图27一起是另一个浏览图表动画(II)，该浏览图表动画(II)示出了计算工作分配器如何将CGA内的CTA分配给示例实施例中的SM。

图28A-28C一起是示出了计算工作分配器如何在CGA内跨GPU硬件的多个分区或层次结构对CTA的分配进行负载均衡的浏览图表动画(III)。

具体实施方式

实现或旨在实现扩展的一种方法是通过负载分配和负载均衡。一般而言，负载均衡是个过程，其中将固定数量的(例如，处理)资源分配给任意数量的传入任务。负载均衡的简单示例是我们许多人每次为汽车加油时所经历的。我们开车进入加油站，然后拉起燃油泵，该泵目前旁边没有汽车。当我们将燃料泵入汽车后，我们会离开，另一辆车可以代替我们。如果加油站很忙，一些汽车可能需要排队等候加油。汽车有时会从一条线路换到另一条线路，具体取决于不同驾驶者完成向各自汽车中加油、完成信用卡或其他支付交易等所需的时间。在繁忙的加油站，我们的想法是让每辆加油机保持近100％占用泵送燃料——尽管由于各种原因这并不总是可能的(例如，不同的车辆有不同大小的油箱，因此可能需要更长或更短的时间来加油，一些司机用信用卡在外面快速支付，而其他人在里面支付现金，一些司机想检查他们的油或洗他们的挡风玻璃等)。

类似的负载均衡考虑适用于在现代GPU的高度并行系统中加载SM 或其他处理资源——尽管用于在现代GPU中工作的负载均衡技术通常比在加油站管理汽车更复杂。特别是，希望在SM或其他处理资源之间平均分配计算工作，以便所有此类资源保持忙于执行指令，并且所有资源几乎同时释放用于更多工作。然而，由于各种原因，一个执行线程或一组线程可能比另一个执行线程或一组线程花费更长的时间来处理。GPU通过利用工作特性并在容量可用时动态分配工作来最大限度地利用处理资源的能力会对全局效率和处理时间产生巨大影响。参见例如陈(Chen)等人，“单个和多个GPU系统上的动态负载均衡(Dynamicload balancing on single- and multi-GPU systems)”，2010年IEEE并行和分布式处理国际研讨会 (IPDPS)，2010，第1-12页，doi:10.1109/IPDPS.2010.5470413；林(Lin) 等人，“使用混合整数非线性编程在异构GPU上实现高效工作负载均衡 (Efficient WorkloadBalancing on Heterogeneous GPUs using Mixed Integer Non-Linear Programming)”，应用研究与技术杂志(Journal of Applied Research and Technology)第12卷，第6期，第1176-1186页(2014年12 月)，https://doi.org/10.1016/S1665-6423(14)71676-1；USP9,715,413；USP9,524,138；USP9,069,609；USP8,427,474；USP8,106,913；USP8,087,029；USP8,077,181；USP7,868,891；US20050041031。

对于初学者来说，为了扩展性能，跨SM的统一利用率通常非常有帮助。因此，示例GPU实施例中的负载均衡理想地集中在单个单元，例如 CWD(＝计算工作分配器)。由于CWD是中心化的，它可以一次比较所有SM的工作量，以达到理想的统一性。然而，在示例实施例中，与先前的GPU架构相比，工作的性质和并发并行的数量发生了显著变化。具体来说，“协作组阵列”的概念在负载均衡中引入了大量附加的复杂性，但也为有效负载均衡带来了大量附加的机会，以显著提高GPU的全局处理效率和吞吐量。

先前的CTA网格层次结构

在一个示例上下文中要进行负载均衡的“工作”粒度可以是“协作线程阵列”或CTA(＝协作线程阵列)。先前的CUDA编程模型使用CTA 作为GPU软件(SW)的并行性的基本构建块。在一个这样的模型中，CTA 可以具有多达1024个线程，并且所有线程被保证在同一SM上同时启动和执行。在这种模型中，因为一个SM运行CTA中的所有线程，所以线程可以利用SM内和/或连接到SM的共享存储器资源，以在线程之间共享数据、同步、通信等——这确保在并发地执行的线程上的数据局部性和数据重用。在这种模型中，CTA的每个线程都可以分配到自身的资源，例如私有存储器，并且同步通常发生在线程级别。

在这种基于SM的编程模型中，CTA声明CTA在其上运行的SM本地的一些量的共享存储器。该共享存储器存在于CTA的有效期期间，并且对于CTA中的所有线程可见。CTA内的线程可以通过该共享存储器彼此通信以用于数据共享和同步两者。存在着色器指令(例如，“__syncthreads()”) 以跨CTA中的所有线程进行屏障同步。例如，为了协调CTA内线程的执行，可以使用屏障指令来指定线程等待，直到CTA中的所有其他线程已经到达的同步点。参见例如USP10977037。

同时，SM的线程束调度器调度CTA中的所有线程束在SM上并发地执行，以保证CTA中的所有线程同时执行。在现代GPU中，SM具有满足或超过CTA中最大线程数的并行计算执行能力——这意味着整个CTA (或在某些情况下，多个CTA)可以在同一个SM上同时执行。

由于许多应用程序需要超过1024个线程，因此用于计算应用程序的原始CUDA编程模型基于“网格”(CTA的集合称为“网格”，因为它可以或可以由多维数组表示)。图8和11A示出了示例层次结构，其中网格包括多个CTA。每个SM调度多个(例如，32个或其他可编程值)线程的并发地执行，这些线程被分组为“线程束”(使用纺织品(textile)类比)。通常，线程束在SM上以SIMD方式执行，即，线程束中的所有线程共享相同的指令流并以锁步方式一起执行(这有时称为单指令多线程，或SIMT)。

由于在单个SM上执行的单个CTA是现有模型中软件的并行性的基本单元，因此GPU硬件在现有模型中不保证跨CTA在较高级别(例如，网格级别)的任何协作。如图9所示，网格中的所有CTA在同一GPU上运行，共享同一内核并且可经由全局存储器进行通信。但是，在现有模型中网格的不同CTA可以同时全部在GPU硬件上执行，或者它们可以顺序地运行，例如取决于GPU的大小和由此网格或其他网格引起的负载。现有模型通过潜在地在不同时间在不同SM上独立地执行这些CTA，不可能在它们之间高效地共享操作(例如，存储器数据检索、同步等)。并且，即使它们确实并发地执行(诸如在协作组API下)，它们可能无法高效地共享存储器或数据带宽以提供跨该组的紧密协作耦合。例如，如果网格被拆分为多个CTA，则从硬件观点来看，机器非并发地运行那些CTA将是合法的——如果算法需要两个或所有CTA并发地运行并且来回传递信息，则导致死锁。

CTA编程模型如上描述已经很好地为开发者服务，其提供了多年和多代GPU在SM级别的数据局部性和数据重用。然而，如以上所讨论的，随着时间的推移，GPU已经变得更大，例如每GPU包括超过100个SM，并且到L2高速缓存和存储器系统的互连不再是平坦的交叉开关，而是分层的并且反映分层硬件域级别(例如，GPU、μGPU、GPC等)。在此类更高级的GPU中，将SM定义为数据局部性的基本单元的机制通常粒度太小。为了最大化性能和可扩展性，所需要的是新的编程/执行模型，其允许软件以比单个SM(其现在小于GPU的1％)大得多的单位控制局部性和并发性，同时仍维持跨所有线程(如CTA)共享数据和同步的能力。应用应该能够控制数据局部性和数据重用以最小化延迟。这对于想要通过创建跨GPU硬件的较大部分的协作线程集来进行强扩展(见上文)的深度学习和 HPC应用尤其如此。

协作组阵列

本文的示例非限制性实施例提供新级别的层次结构——“协作组阵列”(CGA)——以及关联的新编程/执行模型和支持硬件实现方式，其提供了分配所有并发工作的高效负载均衡。例如跨可用处理和其他硬件资源的 CGA。参见上文已标明的2022年3月10日提交的题为“协作组阵列 (Cooperative Group Arrays)”的美国申请No.17/691,621。

在一个实施例中，CGA是CTA的集合，其中硬件保证CGA的所有 CTA被启动至CGA指定或与其相关联的同一硬件组织级别。该硬件被配置为确保在目标硬件级中有足够的处理资源来启动CGA的所有CTA，然后才启动任何其他，并且跨这些处理资源对CTA进行负载均衡，以最大限度地提高全局执行速度和吞吐量。

如图11B所示，网格是CGA阵列，并且每个CGA是CTA阵列。在这种情况下，“CGA”和“集群(cluster)”是同义词，“CTA”是一种“线程块”。这样的CGA提供相对于应用所需的存储器和相对于彼此的共同调度，例如，对CTA在GPU中被放置/执行的位置的控制。这使得应用能够在紧密协作的CTA集群或阵列中的所有线程之间看到更多的数据局部性、减少的等待时间和更好的同步。

例如，CGA让应用利用现代GPU中的互连和高速缓存子系统的层次结构性质，并且使其在未来随着芯片增长而更容易扩展。通过利用空间局部性，CGA允许更高效的通信和更低的延时数据移动。GPU硬件改进通过允许CGA控制并发CTA线程在机器上的何处相对于彼此运行来保证新 CGA层次结构级别定义的多个CTA的线程将针对期望的空间定位并发运行。

如以上讨论的，在一个实施例中，CGA由硬件保证将同时地/并发地启动和执行的CTA组成。CGA中的CTA可以(并且在一般情况下将)在 GPU内的不同SM上执行。即使CTA在不同SM上执行，GPU硬件/系统仍然提供跨SM保证：CGA中的CTA将被调度为并发地执行。如下文详细解释的，CGA的所有CTA的这种高性能并行启动能力可以得到负载均衡算法的支持，以进一步提高性能和吞吐量。GPU硬件/系统还提供高效机制，通过该机制并发地执行的CTA可以彼此通信。这允许应用在CGA中的CTA之间明确地共享数据并且还实现CGA中的CTA的各个线程之间的同步。

在示例实施例中，CGA内的各个线程可以从公共共享存储器读取/写入——这使得CGA中的任何线程能够与CGA中的任何其他线程共享数据。 CGA中的CTA之间的共享数据节省了互连和存储器带宽，这通常是应用的性能限制器。CGA因此增加GPU性能。如上所述，在先前编程模型中，通常不可能在两个CTA之间直接共享数据，因为不能保证两个CTA将在相同的相关硬件域中同时运行。在没有CGA的情况下，如果需要两个CTA 共享相同的数据，则它们通常将都必须使用两倍的带宽来从存储器中提取该数据。就像两个父母各自去商店买牛奶一样。相反，已知有效利用数据局部性对GPU性能是重要的。参见例如拉尔(Lal)等人，“GPU高速缓存中的局部性的定量研究(A Quantitative Study of Locality in GPUCaches)”，奥尔奥格鲁(Orailoglu)等人(编)，嵌入式计算机系统：架构、建模和模拟(Embedded Computer Systems:Architectures，Modeling，and Simulation)， (SAMOS2020)，计算机科学讲座笔记(Lecture Notes in Computer Science)，第12471卷。Springer，Cham.https://doi.org/10.1007/978-3-030-60939-9_16。

现在，使用并发地执行和由硬件支持的附加共享存储器，有可能直接在一个CTA的线程与另一个CTA的线程之间共享数据——这实现了跨桥接硬件(例如，跨SM)分区的CTA的依赖性。

因为CGA保证其所有CTA与已知的空间关系并发地执行，因此其他硬件优化是可能的，诸如：

·跨GPU或GPU的分区的负载均衡

·通过向并发地执行CGA的CTA的所有SM广播标识符来高效地传达某些信息，例如网格标识符

·与另一个(或多个其他)SM共享一个SM内的存储器

·将从存储器返回的数据多播至多个SM(CTA)以节省互连带宽；

·直接SM到SM通信用于CGA中的生产者线程与消费者线程之间的较低延时数据共享和改进的同步

·用于在CGA中的所有(或任何)线程上同步执行的硬件屏障

·以及更多(参见以上列出的共同未决的共同转让的专利申请)。

这些特征通过更高效地使用处理资源、放大存储器和互连带宽、减少存储器延时、以及减少线程到线程通信和同步的开销来提供更高的性能。因此，所有这些特征最终导致应用的强扩展。

层次结构的新级别-CGA

在示例实施例中，CGA由多个CTA构成——即，被构造成协作地执行的多个线程集合或束。每个这样的线程集合或束提供早已由先前CTA提供的所有优点和结构——诸如例如在相同的SM上运行。然而，CGA提供的附加覆盖定义CTA将在何处以及何时运行，并且特别地，保证CGA的所有CTA将在公共硬件域内并发地运行，该公共硬件域提供对CTA之间的数据、消息传递和同步的动态共享，以及跨可能跨越任意数量的SM的处理资源的集合对CTA进行负载均衡的可能性。

示例实施例支持针对不同GPU硬件域、分区或其他组织级别的不同类型/级别的CGA。具体地，CGA可以定义或指定该CGA中的所有CTA 都应该运行在其上的硬件域。以此类推，正如当地高中体育团队可能在当地划分、区域或州际比赛一样，CGA可能需要其所引用的CTA全部在GPU 的同一部分(GPC和/或μGPU)上、在同一GPU上、在GPU的同一集群上等上运行。同时，可以使用集中式工作分配器形式的负载均衡硬件来均衡该硬件集合中所有SM的CGA内的CTA的加载，例如GPC和/或μGPU、 GPU和/或GPU集群。

示例分层硬件集合/分区

在示例实施例中，CGA定义/指定的层次结构被绑定到或以其他方式反映GPU硬件分区(其反映存储器访问和/或通信能力)，以便提供期望的资源和数据重用和数据局部性。例如，正如GPU可以包括如图3和图4 所示的多个GPC，GPU_CGA可以由多个GPC_CGA构成。图10B示出了提供反映与现有技术图10A情况相关的不同硬件域的附加嵌套层次结构级别的示例CGA层次结构，例如：

·GPU_CGAs

·μGPU-CGAs

·GPC_CGAs。

在示例非限制性实施例中，硬件保证某个CGA内的所有CTA并发地启动到SM上，所述SM是由与某个CGA相关联的硬件域指定符所指定的硬件域的一部分，例如：

·用于GPU_CGA的所有CTA被启动到作为同一GPU的一部分的 SM上；

·用于μGPU_CGA的所有CTA被启动到作为同一μGPU的一部分的SM上；

·用于GPC_CGA的所有CTA被启动到作为同一GPC的一部分的 SM上。

此外，相同的启动硬件可以跨GPU、μGPU和/或GPC的处理资源对CGA 的CTA进行负载均衡，以实现更高的处理效率和吞吐量。

更详细地，CGA的一些实施例还支持μGPU分区(诸如图5、图5A 中所示)并且提供若干新的能力和硬件保证，诸如：

·CGA为网格(内核)级别和CTA级别之间的线程提供新的层次结构级别

·GPC_CGA将所有CTA放置在同一GPC内

·μGPU_CGA将所有CTA放置在同一μGPU的SM内，在一些实现方式中，该同一μGPU与大型GPU内的存储器互连层次结构相匹配

·GPU_CGA将所有CTA放置在同一GPU内

·ABC_CGA将所有CTA放置在ABC硬件域内，其中“ABC”是在 GPU架构内或跨GPU架构的任何GPU硬件域、组织或层次结构。

这些示例级别(网格、GPU_CGA、μGPU_CGA、GPC_CGA和CTA ——参见图10B)可以被嵌套以进一步控制SM资源在每个级别处的放置。例如，GPU_CGA可以由μGPU_CGA构成，μGPU_CGA由GPC_CGA构成，GPC_CGA由CTA构成。此嵌套可支持层次结构的每一级别及所有级别的常规动态并行性。参见例如 https://developer.nvidia.com/blog/cuda-dynamic-parallelism-api-principles/

示例协作组阵列网格

随着CGA的添加，现在有许多更可能的网格类型示例在图12、图12A 中示出，其中每个网格类型指定或以其他方式与特定硬件域相关联：

图12(1)：CTA网格——这是传统网格。在实施例中，此网格表示 CTA的三维网格(XxYxZ)。网格的示例尺寸可以是例如18x12x1 CTA。

图12(2)：CTA的GPC_CGA网格——每个GPC_CGA的CTA的三维网格一起启动并且总是放置在同一GPC上。因此，这种类型的网格指定的硬件域是“GPC”。网格中的共同图案化的相邻正方形构成将被调度以同时运行的CGA。因此，标记为“G P C C G A”的六个CTA都将在同一GPC 上一起启动，并且它们都不会启动，直到它们都可以一起启动。示例网格尺寸是6x6x1 GPC CGA，其中每个GPC CGA具有3x2x1 CTA的尺寸。在一个示例中，GPC_CGA支持SM到SM通信、全局存储器中的CGA线性存储器和基于硬件的CGA屏障。

图12(3)：CTA的GPU_CGA的网格——这是其中每个GPU_CGA 的CTA被一起启动并且总是被放置在同一GPU上的网格。因此，由这种类型的网格指定的硬件域是“GPU”。在一些环境中，多个GPU可以被配置为GPU集群；在这种情况下，图12(3)网格迫使GPU_CGA的所有 CTA在公共GPU上运行。这个网格意味着作为CUDA的协作组API特征的替换，但是现在由硬件来保证共同调度。示例网格尺寸是2x2x1的GPU CGA，每个GPU CGA子网格包括9x6x1的CTA。

图12(4)：CTA的GPC_CGA的GPU_CGA的网格。这是具有CGA 层次结构的两个级别的网格。GPU_CGA具有在图12(3)中描述的能力，并且GPC_CGA具有在图12(2)中描述的能力。因此，这种类型的网格指定两个嵌套的硬件级别：GPU和GPC。该示例例如允许开发者调度GPUCGA在单个GPU上运行，并且该GPU CGA内的每个GPC CGA在该GPU 内的同一GPC上运行。示例网格尺寸是2x2x1的GPU CGA，每个GPU CGA子网格包括3x3x1的GPC CGA，每个GPC CGA子子网格包括3x 2x1的CTA。

基于硬件的CGA启动保证&负载均衡

在示例实施例中，上述每种CGA类型中的所有CTA是共同调度的。这意味着GPU硬件将不允许CGA中的任何CTA启动，除非/直到相关GPU 硬件域上有空间供CGA中的所有CTA启动。这种硬件保证允许软件依赖于CGA中的所有线程将同时执行的事实，因此诸如屏障同步和跨所有线程的数据共享之类的事情是可能的。CGA中的任何单个CTA都不能无限期地等待启动——在一个实施例中，要么启动整个CGA，要么不启动。此外，通过使用我们称为“CWD”的集中式分层CTA启动设施或电路，可以在GPU或其他相关硬件域的SM之间对CTA进行负载均衡。

在一个示例布置中，硬件维护CGA中正在运行的CTA(即尚未退出的CTA)的数量的计数，并且软件可以跨CGA中所有正在运行的CTA中的所有线程执行屏障同步。在示例实施例中，每个CGA具有分配给它的 (至少一个)硬件屏障，并且CGA中的所有CTA可以引用该CGA硬件屏障。参见上文标明的2022年3月10日提交的题为“具有异步事务支持的硬件加速同步(Hardware Accelerated Synchronization With Asynchronous Transaction Support)”的美国申请No.17/691,296。例如，此硬件屏障可用于引导所有CTA并确认它们已全部启动。

改进的集中式工作分配器

在图13、图14、图15所示的一个或更多个实施例中，使用主计算工作分配器(“CWD”)硬件电路420在GPU上启动CGA，同时提供CGA 的所有CTA可以同时启动的基于硬件的保证。关于用于调度工作的现有技术示例GPU CWD和MPC的更多信息，参见示例20200043123，特别是图 7和该在先专利公开的相关联的描述，并且还参见例如USP10,817,338和 USP9,069,609。这种传统的CWD 420在过去成功使用时存在三个扩展问题：

1)负载均衡：每个CTA都必须启动到利用率最低的SM。它通常一次完成一个CTA。CWD 420比较所有SM之间的利用率度量以选择一个用于CTA的SM，然后CWD重新开始下一个CTA的过程。要向许多SM启动CTA，此过程可能会变得非常低效。此外，先前的CWD没有很好的方法来处理对不同范围的分层硬件域的启动。

2)光栅化吞吐量：一旦解决了负载均衡问题#1，生成许多多维ID的数学运算(见下文讨论)将成为瓶颈。

3)CTA传输带宽：即使解决了上述问题#1和#2，也很难将许多多维 ID传输到物理上远离CWD 420的SM。例如，在像CUDA这样的计算API 中，CTA ID是3D并表示为32b-16b-16b元组。将64b数据传输到许多(可能数百个)SM将需要大量带宽。

值得一提的是，就网格启动开销而言，更高的CTA启动吞吐量也是可取的。例如，如果SM有足够的资源来容纳8个CTA，则现有技术的CWD 将需要超过1000个周期来填充所有128个SM。这需要很长时间，会引入延迟并降低效率。

另一个需要解决的问题是上述新型工作单元的负载均衡，称为CGA (＝协作组阵列)。如上所述，CGA可以被视为CTA的集合，这些CTA被共同安排并限制在“GPC”中(作为一个示例)。如上所述，GPC是GPU 内的一组SM。例如，一个示例GPU架构由8个GPC组成，每个GPC由18个SM组成。如上所述，CGA的概念可以扩展到多个层次结构。例如， GPC可以由多个“CPC”组成，每个“CPC”都包括一定数量的SM。本文的示例性非限制性技术向多级CGA嵌套和相关的分层硬件域和范围提供负载均衡算法。

具有负载均衡和推测性/影子状态启动的增强型CWD结构和功能

在本文的实施例中，图13、图14和图15示出的改进的CWD 420是集中式电路，其被扩展/增强以提供负载均衡的推测性或基于影子状态CGA 的硬件启动能力，以确认资源可用于启动跨相关硬件域的CGA中的所有 CTA。如果CGA的所有CTA不能同时启动，则CWD 420不启动CGA的任何CTA，而是等待，直到相关GPU硬件域的足够资源变得可用，以使得CGA的所有CTA可被启动，从而它们并发运行。在示例性实施例中， CWD 420使用提供负载均衡的多级工作分布架构支持多级CGA的嵌套 (例如，GPU-CGA内的多个GPC-CGA)。因为CWD 420是集中式的，它可以查看所有SM的状态并根据这些状态做出良好的负载均衡决策。

更详细地，图13&图14中示出的CWD 420在使用推测性或影子状态启动技术确定CGA的所有CTA可适合在指定硬件域中可用的硬件资源上之后启动CGA中的CTA。以这种方式，一个示例模式中的CWD 420确保在启动CGA的任何CTA之前，跨所有GPC或其他相关硬件域有足够的资源用于CGA的所有CTA。在一个实施例中，启动CGA的CTA的算法可从传统(非CGA)网格启动借用一些技术，同时首先确认CGA的所有 CTA可以确保它们将同时运行的方式启动。另参见图16C，其示出了CPU 执行操作550-560以生成网格启动命令并将其发送到GPU以供CWD 420 处理的示例流程图。

图14示出了CWD 420的基本架构，CWD 420包括负载均衡器422、资源跟踪器(TRT)425(0)、425(1)、…425(N-1)、TPC启用表430、本地存储器(LMEM)块索引表432、信用计数器434、任务表436和优先级排序的任务表438。每个TRT 425(0)、425(1)、…425(N-1)与对应的TPC 340(0)、340(1)、…340(N-1)通信。关于这些各种结构的传统操作的更多细节，参见例如USP10,817,338；US20200043123； US20150178879；USP10,217,183；以及USP9,921,873。在示例实施例中，这些和其他结构的功能沿着以下线在示例实施例中被增强以适应CGA：

负载均衡示例

本文的示例技术通过一次对所有SM进行负载均衡来解决负载均衡问题(例如，与一次负载均衡一个SM的现有方法相反)。CWD 420首先查询所有SM的空闲槽。SM的空闲槽数定义为“在SM的剩余资源(例如共享存储器、寄存器和线程束ID)中可以容纳多少CTA”。参见USP9,921,873。每个SM都可以报告它在网格光栅化开始时拥有的空闲槽的数量。可以在图14所示的任务表436中维护这样的空闲槽的计数(例如，将多达预定数量的任务分配给相关联的任务ID的多位值)。空闲槽可以抽象出CGA或网格的CTA需要在GPU处理硬件上运行的资源需求(例如，每个CTA需要运行的资源可以由网格定义为对网格中的所有CTA是统一的，例如如执行线程数、寄存器数、存储器分配等)。空闲槽越多，SM被利用/占用的越少，这表明未来向SM发起的CTA对数学逻辑和存储器带宽的竞争将更少，这通常会导致更快的执行。因此，空闲槽是衡量即将推出的CTA的预期“速度”的指标。

在一个实施例中，GPC内部有(至少)两种负载均衡分布模式：

A)负载均衡模式：它是在下面讨论的图21中使用的模式。它相当于将具有与CGA相同数量的CTA的网格启动到具有与GPC相同数量的 SM的GPU时的分布。

B)多播模式：这是一种负载均衡模式，但有一个限制：只有一个 CTA可以从这个CGA发送到同一个SM(可以向一个SM发送多个CTA，但不能从同一个CGA)。这可以用来确保CGA可以享受多播存储器加载机制，因为在某些提供程序多播的实施例中，在同一SM上运行的同一CGA 中只有一个CTA能够访问来自存储器系统的多播数据。在其他实施例中，可能不存在这样的约束，但是将CGA的仅一个CTA分配给每个SM而不是双倍、三倍等可能仍然存在优势。

下面将更详细地讨论这两种负载均衡模式。以下算法描述适用于“负载均衡模式”。

示例非限制性负载均衡算法

大量CTA的负载均衡看起来很像将水(CTA)倒入干涸的湖中，其深度与空闲槽成正比。参见图21、21A-21Z、21AA-21II、22、23，它们是浏览图表动画，显示了CWD 420将CTA分配给空闲SM槽的示例。在所示示例中，相关的有32个SM硬件分区，但任意数量的SM可能存在于给定的GPU或分区中。

在这个特定示例中，一些SM，例如SM3、SM4、SM6、SM8、SM18 和SM31，已经几乎或完全被先前分配的工作占用。这些SM不能接受附加的工作。他们没有可用的空闲槽(slot)来承担新的工作(在某些实现中可能希望不要完全加载任何SM，因为SM尝试同时执行的工作越多，执行每项工作的速度就越慢)。

一些SM，如SM10、SM11、SM12、SM17、SM18等，已经在执行一些先前分配的工作，但有开放的空闲槽，这意味着更多的处理资源可以潜在地用于执行新工作。他们有一些空闲位置，但并非空置。

一些SM，如SM2、SM5、SM7等，没有执行任何先前分配的工作，因此完全可以用于新工作。他们有许多(所有)空闲槽开放。

相关硬件分区的示例使用状态如图21所示，其中一些SM有很多空闲槽，而一些SM没有空闲槽，例如由于动态GPU操作，可能会发生这种情况。当GPU在重置或其他主要初始化事件后首次启动时，SM不做任何工作，并且所有SM都有最大数量的空闲槽。当CWD 420将第一个CGA 的工作分配给这个完全未被占用的SM集合时，它可以在所有SM之间自由分配工作，而无需任何特殊的负载均衡算法来实现高效的负载均衡。但是，随着GPU继续运行，CGA中的某些CTA可能会先于其他CTA完成，从而释放出可供新启动的CGA重用的槽。就像上面描述的加油站类比一样，加油站不需要完全空置来接受需要加油的新车——而是可以在加油槽可用时动态地为汽车分配加油槽。然而，CGA比CTA提出了更困难的调度挑战，因为CGA需要并行分配多个资源(即CGA中的每个CTA一个 SM空闲槽)。一个类比可以是跨多个工人团队(CGA)调度项目与跨一组个人工人(CTA)调度任务。因此，CWD 420在示例实施例中执行的负载均衡比加油站的简单“先到先得”负载均衡甚至CTA的负载均衡更复杂且不同。

图21示出了示例CGA(由2D网格表示)，其中包括137个CTA，将分布在包括32个SM或其他核心处理器的GPU上执行。由于在一个实施例中构造CWD 420的分层方式(详情如下)，32个SM的这个抽象集合可以分布在GPU内的一个或更多个GPU硬件域中，例如在一个GPC或更多个GPC内、在一个CPC或更多个CPC等。在一个实施例中，它们甚至可以分布在多个GPU上。在此示例中，CWD 420的工作是(a)保证(在硬件中)所有137个CTA将同时启动和运行，以及(b)以最大化速度的方式跨SM集合对CTA进行负载均衡和吞吐量。CDW 420通过使用推测性或影子状态启动机制(见下文)预先安排所有CTA的所有启动，然后(假设并发地执行保证适用于当前加载的SM硬件集合)，“同时”启动所有CTA。这里，“同时”不限于字面上在单个处理周期中向所有SM启动所有CTA，因为实际上，CWD 420可能需要一些时间来向许多SM中的每一个传送适当的指针信息和指令可能分布在半导体晶片的表面上(参见图2B)。

负载均衡浏览图表动画I

图21及以下是浏览图表动画(I)。要在您的计算机或其他设备上使用本专利的电子副本查看此浏览图表动画，请将图21的大小调整为刚好填满屏幕并以横向而不是纵向查看(例如，使用顺时针“旋转”)，然后反复按下计算机上的“下页(Page Down)”按钮，依次进入图21A-21Z、图 21AA-21II。然后，您将看到CDW 420如何填充SM0-SM31的空闲槽的动画。当然，您也可以翻阅这些数字的纸质副本。此动画显示CWD 420似乎一次处理每个CTA以用于推测性(speculative)/影子状态启动目的。然而，即使在那里，当将CTA分配给SM以用于推测性/影子状态启动时， CWD 420也以集中方式运行以考虑SM集合中所有SM的当前加载。

从这些图中可以看出，CTA“填充”了GPU的SM(意味着每个CTA 被调度在一个SM上执行)，CTA被调度在空闲槽数最多的SM上执行，然后逐渐填充“湖”向上(图16A框202、204)。填充“湖”的可用“水”量(填充处理资源的工作)对应于网格中CTA的数量。请注意，随着浏览图表动画的进行，CWD 420首先将工作分配给具有最多空闲槽位的SM。

在一个实施例中，不允许CWD 420启动比SM可以接受的更多的CTA (图16A框2006)，因此如果“湖”在网格用完CTA之前变满，则暂停过程(图16框2008))并等待已启动的CTA完成以释放更多空闲槽。如果“湖”容量(空闲槽的总数)多于CTA，则“水”位上升到一定水平并停止。

有几个选项可以在具体逻辑中实现这一点。一种是二分查找。如果最大空闲槽位(最大湖深)为N，CWD 420会试探性地尝试N/2的最终水位，并查看水量(已启动CTA的数量)是否会大于网格中剩余的CTA。如果是，由于CWD 420没有足够的CTA来满足这个水位，CWD接下来尝试 N/4。如果否，CWD 420尝试3N/4。它重复这个过程，直到找到均衡。

上述方法对于较大的N值是有效的，但对于较小的值，CWD 420可以简单地将SM从最大空闲时隙扫描到较低值(在图21中向上)。在每次迭代中，CWD 420推测性地(影子状态)向具有当前空闲槽的SM启动 CTA。如果使用这种方法，则要特别注意最终的空闲时隙级别，如图22所示。这是流体类比失效的地方，因为CTA工作单元是离散的，因此在一个实施例中的CTA不会被传播跨多个SM。

在一个实施例中，CWD 420通过对图22所示的该最终空闲时隙级别应用附加标准来同时选择SM的子集以分配工作。在一个实施例中，该标准是SM之间的预定优先级或其他优先级。在图22示例中，SM0具有最高优先级，并且SM优先级随着SM#向右增加而降低。CWD420同时计算每个SM的“位1计数(popcount)”(population count)，从左到右的最终空闲槽级别，然后选择位1计数(popcount)小于剩余CTA数量的SM。在一个实施例中，“位1计数”通过自下而上的减少和自上而下的累积在O (logN)时间内获得(图16A框2004)。图23示出了一个示例结果，其在一个实施例中同时选择多个SM以将CTA启动到而不是一次选择一个这样的SM。SM优先级可以按任何顺序分配，“填湖”不需要从左到右执行(它可以从右到左或以任何其他预定或非预定顺序执行)。

浏览图表动画II

图24、25A-25G、26和27一起是另一个浏览图表动画(II)，示出了计算工作分配器可以将CGA内的CTA分配给示例实施例中的SM的另一种方式。在该实施例中，CWD 420不是一次分发一个CTA，而是一次分发一行或一层CTA。如上所述，对于较小的N值，CWD 420可以简单地将SM从最大空闲时隙扫描到较低值(图24中向上)。在每次迭代中，CWD 420 推测性地向具有当前空闲槽的SM发起CTA。基本上，这个实施例不是一一显示CTA的下降，而是一次将“湖”填满一整行。推测性启动成功后，可以将CTA并行加载到SM(TPC)中(见下文)。上面描述的二分搜索算法也有效。对于CGA，也会同时启动多个CGA。

浏览图表动画I和II的替代结局——推测性启动失败

另一方面，如果推测性或影子状态启动应确定CGA太大而无法在可用处理资源上启动(图16A，“是”退出到框2006)，CWD 420将确定没有足够的SM空闲槽可同时启动所有CTA，因此此时会为此CGA发出“失败”指示(图16A，框2008)。没有实际启动该CGA的CTA，系统必须稍后尝试启动此CGA(例如，当更多SM槽可用时)但可能会尝试启动较小的CGA，再次使用上面讨论的负载均衡算法(图16A框2012、2000)。当然，如果网格由正常(即非CGA)CTA组成，如果CWD 420的CTA比总空闲槽位多，那么简单地所有空闲槽位应立即由CTA填充，并且随着更多 CTA完成，CWD将启动剩余的CTA。

负载均衡的GPC组织

如上所述，许多GPU具有分区或其他细分，例如GPC。在一个实施例中，CWD 420在负载均衡时将这些分区或其他细分考虑在内。为了启动 CGA，CWD 420选择GPC，然后为CGA中的CTA选择GPC中的SM。选择的目标是最大化性能：哪些GPC和SM最有可能让CGA运行最快并最快完成？这里的一个假设是CGA中的CTA经常同步并进行锁步进度，因此最慢的CTA决定了CGA的执行速度。

附加浏览图表动画III

图28A、28B、28C是浏览图表动画(III)，它示出了硬件域层次结构的CGA调度算法——在这种情况下，多个GPC每个都包括SM，其中 CGA-GPC必须在单个GPC上运行并且不能跨越GPC之间的边界。在示例实施例中，该算法还以分层方式操作。在顶层(图15框420a)，选择“最佳”GPC，并行启动CGA。CWD的下一个管线阶段(图15框420b)为 CGA中的CTA选择SM。因此，CTA是并行启动的，并且问题#1再次得到妥善解决。

现在，如何选择最好的GPC？在一个实施例中，它是通过从每个GPC 查询“速度”来完成的，如图28A所示。如上所述，我们希望将CGA发送到“最快”的GPC。但在一个实施例中，“GPC”并不是真正的GPC。从CWD到GPC的往返延迟很大，应避免查询该距离。相反，在一个实施例中，CWD 420在内部实例化一组模仿GPC，并将查询发送给它们。这样的模仿模拟了SM向CWD 420报告的空闲槽(参见上面的讨论)，因此CWD 考虑了每个GPC中每个SM的当前加载。在图28A所示的示例中，GPC0 包括8个SM，其中5个(SM3-SM7)满载，其中两个(SM1，SM2)部分加载，其中一个(SM0)完全未占用。等等。

GPC的这种模仿或模型允许CWD 420提供为上述CGA网格定制的推测性或影子状态启动能力。每个查询都以样本CGA的形式出现。收到样本CGA(在本图中为CGA1)后，模仿GPC会尽力将CGA的CTA发送给SM。模仿GPC有两种跨SM负载均衡CTA的模式——一种是简单的负载均衡，另一种是“多播”模式。在模仿GPC尝试CTA启动后，如果它没有找到足够的空闲槽，它会向顶层回复“失败”。如果它确实找到了足够的空闲槽，使用负载均衡模式之一，它分配CTA，然后计算收到CTA 的SM之间的最小剩余空闲槽。如上所述，剩余的空闲时段是预期CTA的“速度”的指标。而CGA中最慢的CTA成为瓶颈。因此，剩余的最小空闲槽位就是CGA的“速度”。

顶层重复这个查询启动过程，直到网格中的CGA或空闲槽用完为止。

该算法可以扩展到多级CGA层次结构。图28A、28B、28C显示了当我们具有三个层次结构时的情况：一个网格由GPC CGA组成，这些GPC CGA由预定义数量的CPC CGA组成，每个CPC CGA由CTA组成。在顶层，该算法向每个模仿GPC查询GPC CGA的速度(图15)。然后，每个模仿GPC重复向每个CPC(在GPC中)查询CPC CGA的速度，并“推测性地”启动CPC CGA，直到启动GPC CGA的所有CPC CGA(图15)。 GPC CGA的速度就是上次查询的速度，因为GPC CGA的速度会是GPC CGA中最慢的CTA的速度。一旦顶层决定接受GPC CGA速度并向模仿 GPC发送启动消息，则为推测性GPC CGA保留的空闲槽被提交，并且GPC CGA在实际硬件上真正启动。否则，预订将被取消。更多层次的层次可以通过附加的CWD 420硬件层和/或通过递归地重用每个层次的可用硬件 (例如，以节省芯片空间和相关的功率)来适应。

在此图中，由于GPC0已被占用，因此无法在GPC0上启动任何 GPC-CGA。同时，GPC1确实有足够的资源来启动GPC-CGA，但这样的启动会耗尽GPC的所有可用处理资源。但同时，GPC2和GPC3都可以启动一个有空闲槽的GPC-CGA，这意味着通过选择GPC2和/或GPC3的负载均衡算法能够最大限度地提高启动的GPC-CGA的运行速度。在所示示例中，即使GPC2比GPC3负载更轻，CWD 420不一定更喜欢GPC2而不是GPC3，但是其他示例实施例可以根据例如有多少(以及哪些特性) GPC-CGA等待运行来考虑这种差异。图28C之后的场景是CWD 420将启动GCA2到GPC2，以便两个CGA将在同一个GPC上运行。

也就是说，在一个实施例中，具有相同最大速度的所有GPC同时接收CGA。在图25C中，GPC2和GPC3同时接收CGA。最重要的是，每次 CGA发布都包括同时进行的多个CTA发布。因此可以看出，在示例实施例中，一次(同时)调度多个CTA，而不是一次一个。

基于多维ID和网格光栅化的扩展

如上所述，出于各种原因，示例实施例中的每个CTA都被分配了多维ID，CTA和GPU可以使用该多维ID来识别CTA并将CTA与数据和其他资源相关联。例如，如上所述，对于数据并行编程模型，每个CTA都与支持CTA执行的数据相关联。例如，CTA可以查询多维ID以查找分配给它的数据。例如，如果将图像过滤器网格应用于3D图像，则每个CTA可能会查询XYZ元组来确定它应该处理图像的哪个部分。如果硬件仅提供整数值作为ID，则每个CTA必须执行昂贵的除法以将其分解为XYZ。因此，当CTA被发送到SM时，硬件生成多维(例如，XYZ)ID是有帮助的。这种多维ID生成可以称为“网格光栅化”，因为它类似于计算机图形算法中从屏幕上的三角形中扫描出像素的行为。

当CTA是CGA的一部分时，这种网格光栅化变得更加复杂，因为分配给CTA的多维ID现在也应该编码或识别或以其他方式与CTA所属的 CGA相关联。

一旦使用如上所述的集中式CWD 420解决了负载均衡问题，就会出现光栅化吞吐量问题；在启动时生成许多多维ID的数学运算可能会成为瓶颈。一个相关的问题是CTA传输带宽；即使解决了上述问题，也很难将许多ID传输到物理上远离集中式CWD 420的SM。例如，在像CUDA这样的计算API中，CTA ID是3D并表示为32b-16b-16b元组。向许多SM传输64b数据需要大量带宽。

将光栅化分布到SM而不需要CWD 420或其他集中式硬件来执行光栅化来解决这两个问题。特别地，当CWD 420通知SM它们被分配了新的 CTA工作时，SM内的MPC 404可以通过来自其他SM的广播信息被通知所有SM的所有CTA的XYZ ID分配。在一个实施例中，这些XYZ坐标由SM生成是确定性的；例如，可以先增加X坐标值，一旦增加了Y坐标值，再从初始值开始，同样可以在增加Z坐标值后从初始值重新开始Y坐标值。每个SM都被编程为根据它接收到的广播启动分组独立地执行此过程。这种分布式ID坐标分配过程使CWD 420不必向任何SM发送任何XYZ ID分配。因此，XYZ分配分布在所有SM上，并且可以由同时执行的硬件同时执行。示例实施例进一步提供了下文描述的用于软件查询硬件以获知 XYZ ID分配的机制。该解决方案提供了两全其美：用于实施负载均衡算法的集中式决策电路，以及用于在CWD确实安排CTA执行时执行详细记录保存(CTA ID生成)的分散(分布式)电路。

在一个实施例中，ID生成的负担从CWD 420转移到SM。CWD 420 执行负载均衡并将消息发送到负载均衡算法选择的SM以触发CTA启动。在一个实施例中，CWD 420向所有SM广播所有CTA启动消息，使得每个SM都可以观察到整个CTA启动序列。因此，每个SM都可以告诉网格内的当前光栅化位置。该方案解决了传输带宽问题，因为CWD到SM的流量不再需要携带多维ID。

这个想法的幼稚实现仍然会留下光栅化吞吐量问题，因为即使所有 SM并行执行光栅化，每个SM将消耗与CWD执行相同计算的情况相同的时间。为了解决这个问题，每个SM只偶尔计算一次多维ID。大多数情况下，它只是计算自上次计算多维ID以来的CTA数量。可称为计数delta(Δ)。而当多维ID是绝对需要的时候，即SM发起CTA或者delta(Δ)溢出时，SM将delta(Δ)分解为多维版本，然后将分解后的delta(Δ)添加到最后计算的ID。该技术解决了光栅化吞吐量问题。这是一个3D示例。

a)将delta分解为XYZ的步骤。在代码中，gridDim.X和Y分别是网格的X和Y维度。“％”是类似于C++的模运算符。请注意，由于我们可以将delta的精度保持在较低水平，因此这些除法不需要昂贵的逻辑。

Zdelta＝delta/(gridDim.X*gridDim.Y)

delta_tmp＝delta％(gridDim.X*gridDim.Y)

Ydelta＝delta_tmp/gridDim.X

delta_tmp＝delta_tmp％gridDim.X

Xdelta＝delta_tmp

b)更新当前X、Y和Z的步骤

Xcurrent＝Xcurrent+Xdelta

If Xcurrent>＝gridDim.X then Xcurrent-＝gridDim.X and C＝1 else C＝0

Ycurrent＝Ycurrent+Ydelta+C

If Ycurrent>＝gridDim.Y then Ycurrent-＝gridDim.Y and C＝1 else C＝0

Zcurrent＝Zcurrent+Zdelta+C

该算法能够实现典型案例的峰值CTA启动吞吐量，与以前以集中方式将多维ID分配给CTA的方法相比，该算法提高了36倍。

CWD 420的详细实施

在一个实施例中，CWD 420从前端212接收用于在与GPU协作的 CPU上执行的各种进程的任务。在示例实施例中，每个任务可以对应于一个CGA。在CPU上执行的每个进程都可以发出此类任务。

在示例实施例中，调度器410从前端212接收任务并将它们发送到 CWD 420(图16，框502、504)。CWD 420从多个CGA查询和启动CTA。在一个实施例中，它一次在一个CGA上工作，但在其他实施例中，协调并行是可能的。在一个实施例中，对于每个CGA，CWD 420针对由查询模型定义的影子状态推测性地启动CGA中的所有CTA，如上所述(使用 GPC的基于硬件的模型而不是实际的GPC，如图16B)，递增“启动”寄存器以存储推测性/影子状态启动(参见图16B)。如果在针对查询模型定义的影子状态推测性地启动CGA的所有CTA之前SM或硬件域中的其他处理器中的所有空闲槽用尽，则CWD 420终止该启动，并可稍后再次尝试。相反，如果对于CGA中的所有CTA存在足够的空闲槽，则CWD 420 从在推测性/影子状态启动过程中累积的“启动”寄存器生成sm_mask(该 sm_mask数据结构存储用于在CGA启动的相关硬件域中的每个SM上运行的每个CTA的预留信息(图16A，框2010；参见图16B))，并且一直移动到下一CGA。硬件分配CGA序列号，并将其附加到每个sm_mask，该 CGA序列号指定哪个SM在启动时获得哪个CTA。其还将CGA位的end_ 附加到最后一个，以防止来自不同CGA的sm_mask的交错。在一个实施例中，发送到SM的启动分组可以包括针对每个sm_mask的多个CTA启动指令——通过不再需要启动的CTA和启动分组之间的一一对应关系来提高性能。应当注意，在一个实施例中，CWD 420维护整个GPU的两个模型——一个是不断更新的，另一个是不断更新的模型的快照，它定义了影子状态并用于推测性启动。只要推测性启动负载均衡会话发生得非常快，即在统计上可能快照变得足够过时以导致严重的低效率(CWD 420控制所有启动，因此唯一的不准确之处是某些CTA可能会完成)，这种安排就可以工作并且正在执行它们的SM相应地变得可用，并且在使用现已过时的模型的正在进行的负载均衡会话中没有考虑到这一点)。

在一个实施例中，GPU CGA序列号被附加到启动命令，并且被预加到为每个GPCCGA生成的sm_mask。这个GPU CGA序列号用于将每个GPC CGA的sm_mask映射到GPU CGA，并且在将掩码发送到各个SM中的M管道控制器(MPC)之前由任何重排序单元使用。硬件因此可以提供掩码的多个迭代波，以确定何时将CGA中的所有CTA映射到SM，从而使得CGA可以启动。一旦SM掩码准备好，它们被广播(与关联的CGA ID) 到GPU的所有SM工作调度器(参见图16B)。广播也是lmem_blk_idx分组，其将lmem_blk_idx(参见图14的LMEM块索引表432)从CWD 420 运载到SM。这些操作完成了在真实硬件上的实际启动。

图15示出，在一个实施例中，CWD 420包括用于分配由GPC CGA 构成的GPU CGA的两个级别的工作分配器(WD)：

·GPU2GPC工作分配器420a

·多个GPC2SM工作分配器420b(0)、420b(1)、420b(2)、...。如上所述，第一级420a跨GPC分配GPC CGA。第二级(GPC到SM工作分配器420b)将CTA分配到GPC内的SM。在GPU到GPC级别之前或高于其的另一个级别可以用于将μGPU CGA分配给μGPU(在一个实施例中，当存在μGPU时，GPU由μGPU构成，μGPU由GPC构成，并且 GPC由TPC或SM构成)。

GPU2GPC WD 420a将GPU CGA的可能多个(1个或更多个)成分 (constituent)GPCCGA分配到相应的GPC2SM工作分配器(图16，框 506)。GPC2SM工作分配器420b各自将GPCCGA的CTA分配到GPC内的SM(使用例如负载均衡模式或多播模式，如下所述)。图15的统一工作分配器(UWD)420a/420b保证GPU CGA中的所有GPC CGA可以一起启动，并且每个GPC CGA中的所有CTA可以一起启动。在支持CGA的更深嵌套的其他实施例中，这个UWD可以被扩展到所需要的任何数量的级别(例如，用于增加GPC到CPC的层次结构)和任何层次结构的任何数量的同时调度的硬件范围。

在一个实施例中，UWD 420a、420b执行以下进程：

I.CGA的推测性启动(图16，框508)

阶段1：

第一步骤是状态快照：从任务表436(图14)读取剩余数量的GPU CGA，并且基于remaining_GPU_CGA将其箝位。在实施例中，负载均衡会话可以一次限于GPU_GPC_CGA。

阶段2：

对于GPC CGA，CWD 420执行查询+启动进程，直到不再有剩余的 GPC CGA，其中“查询”构成针对查询模型的“推测性”(“影子状态”) 启动，并且“启动”构成实际启动。因此，在一个实施例中，“推测性”或“查询”启动是针对查询模型在影子状态上执行的，但尚未执行实际启动 (在其他实施例中，例如，具有重复的硬件，“推测性”可以指以下操作：实际执行，但在知道它是否有效之前执行，例如在“推测性执行”中)。在一个实施例中，在启动任何CTA之前，对CGA结构中的所有CTA完成“查询”。例如，在具有多个GPC CGA的GPU CGA的情况下，CWD 420将仅在其GPC CGA的所有成分都被保证接收GPU上的空闲槽时才启动GPU CGA。为了确定，(GPU CGA的)每个成分GPC CGA在任何CTA被启动之前被推测性地启动并且检查(但是实际上没有被启动到SM)。

在一个实施例中，每个GPU CGA可以在两次传递(pass)中被处理：

传递I：推测性启动以“检查所有成分GPC CGA是否将找到家”

假设GPU CGA中的GPC CGA的数量是“N”。为了确定以上情况， CWD 420推测性地启动N个GPC CGA。

参见图15，GPU2GPC WD 420a向所有GPC2SM WD 420b发送查询命令。每个个体GPC2SM对分配给它的GPC CGA的所有CTA执行推测性调度，并为该查询产生快速和有效的响应。在示例实施例中，由于推测性启动测试将针对GPU CGA内的每个GPC CGA重复，所以每个GPC2SM 包括空闲槽寄存器和每SM的启动槽寄存器，用于存储其先前的响应。在每SM具有单个空闲槽和启动槽寄存器的实现方式中，在GPC CGA的第一推测性调度之后的迭代中使用的每SM的空闲槽值可以是“空闲槽值”——“当前启动槽值”，以考虑已经推测性地调度的CGA。

GPU2GPC WD收集来自GPC2SM WD的响应，对“有效”的数量进行计数并且累积至计数器。这完成了第一查询迭代。GPU2GPC WD继续再次查询所有GPC2SM WD，直到计数器达到每GPU CGA的GPC CGA 的数量。如果GPU2GPC WD未能收集足够的“有效”，则GPU2GPC WD 将终止会话，因为没有足够的空闲槽来保证GPU CGA中的所有GPC CGA 中的所有CTA可以被一起启动(图16D，“否”退出到判定框510)。

在一些实施例中，不同的GPC可以具有不同数量的SM。在一个实施例中，CWD 420还可实现每GPC的计数器，用于追踪在给定GPC上可同时执行的GPC CGA的数量。基于对应GPC中的SM的数目(例如，对于给定芯片数目)初始化每个计数器。每当启动新的GPC CGA时，CWD420递减适当的GPC计数器，并且每当从给定的GPC到达cga_complete 分组时递增适当的计数器。

在示例实施例中，CWD 420可以使用不同的基于硬件的模式来将 GPC_CGA中的CTA分配到GPC内的SM/核心，所述基于上文描述的硬件的模式：

·LOAD_BALANCING——(上文描述的)CTA被发送至GPC或其他硬件域内的最小负载的SM/核心。该模式允许CWD 420将 CTA放置在GPC或其他相关硬件域内的任何位置。例如，这可导致来自同一CGA的多于一个CTA(或者甚至用于小CTA的所有CTA)在同一SM上运行。

·MULTI_CAST——CWD 420跨GPC或其他相关硬件域内的 SM/核心分配CTA，每SM至多来自同一CGA的一个CTA。该模式保证每个CTA将在不同的SM上运行——这意味着由那些多个SM提供的所有互连和资源可被使得用于承载执行CGA。在一个实施例中，CTA被首先调度到其中两者(所有)SM/核心都可获取CTA的分区上，然后调度到其中仅一个(少于所有)SM可用的分区上。

MULTI_CAST模式保证CTA跨SM/核心很好地分布(而不是允许在同一SM上的多个CTA)，这为CGA提供最大互连资源。例如MULTI_CAST 模式可以用在GPC_CGA上，GPC_CGA想要利用SM和通用网络接口控制器(GNIC)中的新多播硬件和软件，例如张量存储器访问单元(TMA)，如2022年3月10日提交的题为“用于高效访问多维数据结构和/或其他大数据块的方法和装置(Method And Apparatus For Efficient Access To Multidimensional DataStructures And/Or Other Large Data Blocks)”的上述美国专利申请No.17/691,276中所描述的。关于MULTI_CAST方法的更多信息可以在2022年3月10日提交的题为“跨多个计算引擎的程序控制的数据多播(Programmatically Controlled Data Multicasting AcrossMultiple Compute Engines)”的上述美国专利申请No.17/691,288中找到。

传递II：“重置。然后，查询+启动”CGA的实际启动(图16D，框 512)

如果传递1(推测性启动)成功，则为整个GPU CGA保证足够的空闲资源，CWD 420开始传递2＝＞这是实际启动。这涉及：

·重置所有GPC2SM WD的启动槽寄存器；

·分配GPU CGA序列号(用于重新排序)；

·逐一启动(GPU CGA的)成分GPC CGA；以及

·对每个GPC2SM WD重复“查询”和“启动”，以在SM上的每个 GPC CGA中启动CTA。

在示例实施例中，CWD 420还负责在线性存储器池中分配CGA存储器槽以及刷新和回收槽。假设CWD 420确定存在足够的资源并且以上阶段2完成或正在进行，则CWD 420将信息传递到驻留在GPC内的GPM 功能电路块。每个GPM分配屏障槽，当GPC CGA中的所有CTA完成时，还分配CGA_id和轨道。每个SM中的MPC(M-管道控制器)电路404同时跟踪每CTA的槽，并且参与启动CTA到其关联的SM上以实际进行工作。在此上下文中，在一个实施例中，CDW420就像纸牌游戏中的发牌者一样，并且SM是纸牌游戏的参与者。CWD 420的工作是将“启动”卡分发给游戏中的每个SM“玩家”。同时，所有的SM都在看着卡被分发给彼此的SM，这样当一个SM最终收到自己的启动卡时，它可以根据之前收到的广播将自己得出的XYZ ID值写在卡上在比赛期间。任何新的SM都可以参与游戏，只要被告知另一个SM上次分配的当前XYZID值是什么。在一个实施例中，还有一种状态同步，其中CWD 420向下发送束或网格的某些维度，然后在启动时，CWD将当前XYZ偏移传送给任何新的SM 及其相关的内部MPC。在其他实施例中，每个新到达的网格可以提示重新初始化XYZ ID坐标。

当工作完成时，SM向GPM报告CTA完成状态。当GPM电路接收到CGA中的所有CTA已完成的状态信息(图16D，框514)以及对CGA 的所有存储器分配已被刷新的状态信息(图16D，框516)时，GPM电路可发信号通知CWD 420以释放池中的CGA存储器槽，使得它可被分配给另一CGA(图16D，框518)。

使用上述技术，应用程序可以在GPC或其他硬件分区中启动许多小型CGA，但随着CGA大小的增长，数量会减少。在某个点(取决于硬件平台)，没有CGA可以再适合GPC或其他硬件分区，这可能会损害代码的可移植性。如果假设每个平台至少有一个GPC和4个TPC，那么保证跨未来架构兼容性的最大CGA大小是8个CTA。给定的应用程序可以根据查询平台动态调整CGA大小，以确定可在GPU中同时运行的CGA数量，作为1)CTA资源需求和2)每个CGA的CTA数量的函数。

CTA分配和跟踪

示例硬件实现在每个SM中提供了一个新的S2R寄存器，有助于跟踪CGA内的CTA(即，允许CTA确定它是CGA内的哪个CTA)。在一个实施例中，SM执行S2R(特殊寄存器到寄存器)操作以返回CGA内的线性CTA ID。特别是，称为gpc_local_cga_id的附加基于硬件的多位标识符 (使用的位数可能取决于支持的同时活动CGA的数量)用于识别GPC命名空间内的CGA并跟踪数量该CGA的活跃CTA。这个相同的值 gpc_local_cga_id可以例如用于索引分布式共享本地存储器、引用屏障和其他CTA间通信机制。

S2R寄存器使着色器软件能够读取此线程的gpc_local_cga_id。 Gpc_local_cga_id在每次GPC_CGA启动时分配给本地GPC，并在CGA启动时跨相关硬件域广播。它在CGA的生命周期内被跟踪，并在CGA中的最后一个线程组完成时被释放。在一个实施例中，只要硬件看到GPC CGA 启动的第一个数据包(见下文)，它就会分配唯一的gpc_local_cga_id，然后跟踪其本地GPC内的所有活动GPC CGA。每当硬件接收到GPC CGA 中所有CTA的共享存储器刷新指示时，它都会回收gpc_local_cga_id。硬件维护可用gpc_local_cga_id的空闲列表或空闲向量，如果gpc_local_cga_id 用完则停止CGA启动。

在上面讨论并显示在图12A中的示例网格中，标记为“C”的CTA 需要能够告诉(学习)它在六CTA CGA中的哪个CTA，一旦由硬件分配 (即，每个协作线程阵列现在有序合作组数组的一部分)。知道整个网格的尺寸和上面讨论的各种局部网格层次结构的尺寸，就可以将整个网格内的 CTA的坐标转换为CTA在其CGA内的坐标。在示例实施例中，每个网格或CGA是根据使用3维坐标的层次结构中的下一级来定义的。如上所述，每个CTA通过硬件寄存器将其CTA id(X，Y，Z)暴露给着色器中的软件。对于GPC_CGA，新的S2R硬件寄存器可用于确定或发现GPC_CGA 中由启动程序预设的一维CGA_CTA_id。在一个实施例中，该CGA_CTA_id 可以直接用于共享存储器索引(这在寻址共享存储器时很有用，因为可以使用其对应的CGA_CTA_id来引用共享存储器的每个段)。

图12A的示例是针对包含GPC_CGA的网格内的CTA#C，整个网格内CTA 的坐标为(7，3，0)，但GPC_CGA内的CGA_CTA_id是一维坐标CgaCtald ＝4。这样就维持了基于整个网格内的3D坐标的编程模型，同时在其CGA 内为CTA提供了一个附加的坐标。

CGA跟踪

在如图13所示的一个实施例中，GPM跟踪CGA中活动CTA的总数。例如，当CGA启动时，GPM将计数设置为CGA中已经启动的CTA的总数。当MPC指示CTA已经退出时，GPM递减该计数。当计数已经递减到零(意味着CGA中不再有CTA是活动的)时，GPM确定CGA已经完成。但在示例实施例中，GPM尚未将CGA ID释放到池中以供重新使用。这是因为即使CGA中的所有CTA都已经完成，也可能存在一些未完成的 (outstanding)DSMEM(分布式共享存储器)访问请求。因而，示例实施例提供协议以确保CGA中的CTA在释放与那些CTA相关联的CGA ID之前已经完成其所有DSMEM存储器访问(和其他访问)。在一个实施例中， GPC不释放CGA ID，直到CGA中的每个CTA已经退出并且其所有存储器指令/访问已经完成。

这样做是为了防止新的CGA读取或写入(或从其接收读取或写入) 先前使用相同的CGA ID的失效的(defunct)CGA。在一个实施例中， gpc_local_cga_id提供针对此的保护，因为当新的CGA启动时，可能不存在来自CGA ID的非当前用户的进行中(in flight)的DSMEM访问。

如以上所讨论的，当CGA完成执行时，基于硬件的调度器(GPM) 释放之前由CGA使用的资源(例如，共享存储器、在SM上运行所需的线程束槽等)，从而使得CWD 420可以重新分配资源以启动新的CGA。类似地，当CTA完成执行时，基于硬件的调度器(GPM)释放CTA先前使用的资源(例如，共享存储器、在SM上运行所需的线程束槽等)。一旦CTA 结束，就使用协议来挑剔(fault)对该CTA的共享存储器的任何DSMEM 存储器访问。在一个实施例中，当CGA中的所有CTA都完成执行时，基于硬件的调度器保留CGA ID并向正在该CGA中运行CTA的SM中的每一个发送DSMEM存储器刷新(图16D，框516)，然后等待响应。一旦运行CGA中的CTA的所有SM确认先前分配给该CGA的共享存储器的存储器刷新，则GPM最终可将CGA ID释放至重用池。

在启动侧，CGA中的每个CTA需要知道CGA中的所有其他CTA在哪里执行，因此CTA可以将事务发送给那些其他CTA。该映射信息在启动期间被编程。

SM内的映射表

Gpc_local_cga_id

图18A-18B和19示出了用于允许SM与其他SM通信的示例实施例架构的不同视图。SM可以与另一个SM通信的消息之一是SM正在执行的CTA的local_cga_id。在一个实施例中，这种SM到SM消息的分组格式包括U008字段“gpc_local_cga_id”。每个GPC都有自己的CGAID池， GPM在CGA启动时将其中一个号码分配给CGA。然后，该分配的编号用作例如指向由CGA中的各种CTA使用的DSMEM分布式内存段的指针。在一个实施例中，“gpc_local_cga_id”还用作用于跟踪每个GPC_CGA的屏障状态的id。

图17示出了由每个SM维持的示例映射表布置。在一个实施例中， SM基于分段的地址确定目标，然后选择正确的分组类型以使互连知道这是SM2SM事务，并且基于如图17所示的路由表中的查找来提供物理SM id。在一个实施例中，SM将GPC_CGA内的逻辑CTA ID映射到CTA在其上运行的物理SM，以及在SM上的该CTA的物理共享存储器。每次CTA 启动时，GPC上的所有SM可能需要知道它，因为这些SM中的任何一个都可能正在执行作为相同CGA的一部分的CTA。在一个实施例中，MPC 404在每次启动新的CTA时通知(广播消息到)所有SM。作为响应，每个SM更新其维护的映射表。在一个实施例中，CAM结构被用于该映射以允许来自远程(其他)SM的DSMEM寻址。如图17所示，CAM结构被存储在RAM中，作为由SMCGAslot值索引的SM到SM映射表5004。映射表5004向SM标识CGA中的其他CTA在哪些其他SM上执行。定义5004示例表的伪码如下所示：

在本示例中，gpc_local_cga_id因此被用作该CGA中的所有SM可以引用的本地CGAID。该表允许每个SM查找tpc_id和sm_id_in_tpc，其有效地为另一个SM的地址。该结构的索引是CGA中的(逻辑)CTA ID(该 ID对于每个CGA是本地的)。因此，给定指示(可能正在运行的所有CGA 中的)哪个CGA和逻辑CTA ID的槽ID，SM可查找正在运行该CTA的该其他SM的SM_id，使得它可通过互连与该其他SM通信，以用于涉及例如分配给该其他SM上的该CTA的DSMEM段的事务。

随着附加CTA被启动和完成，表5004继续被更新，其中，每个SM 随时间推移维护其自身的映射表5004。同时，硬件(MPC和GPM与SM 协作)阻止CGA同步屏障是有效的，直到CGA中的所有CTA已经启动并且所有SM已经接收到广播信息以构建其映射表5004，以便防止CGA中的任何CTA脱离屏障同步体制。

在一个实施例中，图17中所示的第二表5002由每个SM维护以将线程束映射至CGA槽。特别是，SM自身的内部线程束调度器依据线程束调度执行槽(例如，某一数目(例如64个线程束)可同时在任何给定SM上运行)。SM维护映射信息，以将线程束数目映射到CGA_slot信息。因此，例如，一个SM上的线程束可发出LD指令，该LD指令被映射到执行同一CGA中的其他线程束(CTA)的另一SM的DSMEM中。它首先使用表格5002识别CGA_slot(槽)，然后使用表格5004来确定将指令传递到哪个SM。综上所述，在源SM中，当CTA(SM的物理线程束ID＝X)访问另一CTA的共享存储器(由同一CGA中的逻辑cta_id＝A寻址)时，CTA 首先查找bl_table，以获得sm_cga_slot，然后查找cga_cga2Sm_dir，以获得每以下伪代码的目的地SM(tpc_id和sm_id_in_tpc的元组)的 gpc_local_cga_id和sm_id：

gpc_local_cga_id＝cga_cta2sm_dir[bl_table[X].sm_cga_slot].gpc_local_cga_id；

destination_sm_id＝cga_cta2sm_dir[bl_table[X].sm_cga_slot].sm_id_of_cta_in_cga[A]；

源SM然后使用根据以上指令格式的gpc_local_cga_id和sm_id将指令在互连5008上引导至目标SM的DSMEM内的位置。

图17还示出了目标SM在互连5008上接收的请求。当目标SM接收到该请求时，它可以使用如在以下伪代码中描述的表5010执行查找以找到 DSMEM基数和大小：

·传入SM,CAM匹配[CTA中的gpc-local CGA_ID和CTA_ID]以找到共享存储器基数和大小

目标SM匹配gpc_local_cga_id和cta_id_in_cga(注意：cta_id_in_cga 被包括，因为可能存在运行在给定SM上的CGA的多于一个CTA)。如果存在匹配，则生成有效查找标签(如果不存在匹配，则这可能意味着CTA 不再在SM上运行并且接收SM相应地生成错误通知，该接收SM将该错误通知发送给始发SM)。假设有效的查找标签，所述表然后用于在保持共享存储器的物理存储中查找DSMEM基数和大小(DSMEM分配是可重定位的，并且因此可以在物理存储中的任何地方)。如上所述，表5010(在一些实施例中可以是内容可寻址存储器或CAM)可以在硬件中复制以提供多个并发查找。目标SM然后将检查伴随指令出现的偏移，确保其在范围内，且接着对指定的DSMEM存储器偏移执行读、写、原子操作或其他请求动作。如果指令指定超出范围的偏移，则检测到错误并且向源SM通知该错误。

CGA/CTA退出和错误处理协议

在一个实施例中，某些种类的错误不归因于程序计数器(PC)。通常，实施例将保留过去的PC的FIFO，并且可以将任何存储器错误与给定的线程束、线程和PC相关联。当确定不存在可归因于该PC的错误时，PC可从FIFO的末尾掉落。然而，某些类型的错误在目标SM处被检测或可检测，但不由源SM检测或可检测，因此不能与源SM的PC相关联。这种错误例如可以包括在目标处“未发现CGA/CTA”或者具体地目标SM检测到gpc_local_cga_id和cta_id_in_cga不在shmem_base CAM中(通常因为 CTA已经退出)，或者目标SM检测到绑定之外的地址，诸如Address Offset ＞shmem_base+shmem_size(例如，由于目标SM将其DSMEM共享存储器分配的一部分提前释放至CGA)。为了处理这样的错误，一个实施例不向目标或目的地SM报告错误，而是使目标SM负责使用类似于确认消息传递的错误消息传递向源SM报告这样的错误。当接收到错误分组时，源 SM发布错误并将其归到CGA，但不一定将其归到特定的线程束和/或PC，因为该信息可能不再可用。在源SM处，陷阱处理程序可以使用SR寄存器读取坏线程束的gpc_local_cga_id和cta_id_in cga。如果CGA已经退出 (对于存储和原子来说是可能的)，则错误可以被忽略/丢弃，因为它现在是有争议的。

在源SM侧上可检测的其他类型的错误可以提供有效的线程束ID (warpID)和PC，例如：

Cta_id_in_cga＞CGA中的CTA的最大数目

Cta_id_in_cga在SM2SM表中具有无效的SM_id

地址偏移＞可能的最大共享存储器大小

CGA退出

在一个实施例中，CGA退出是多步骤过程。首先，运行CTA的SM 检测到线程束已经发送了Warp_exit命令。这意味着CTA想要退出，但如上所述，DSMEM SM到SM写入和对L2线性存储器的CGA写入可能仍在进行中。因此，CTA实际上不退出，而是MPC被通知并且CTA等待MPC授予退出许可。当CTA中的所有线程束完成时，MPC向SM发送 inval_cta_entry，以使图17中所示的CGA共享存储器sm_cga_cta_slot CAM 条目无效。MPC然后将cta_complete发送到GPM和CWD，并将CTA标记为需要存储器刷新。当CGA中的所有CTA都完成时，MPC将包括 sm_cga_slot的CGA资源解除分配，并向SM发布DSMEM刷新。在接收到刷新完成的确认之后，MPC发送dsmem_flush_done。作为响应，在接收到来自CGA中所有CTA的dsmem_flush_done后，GPM回收 gpc_local_cga_id，并向CWD发送cga_complete。

由此，虽然CGA线程块分组构造对于保证跨SM的并发性和负载均衡是有用的，但用于保证并发性的其他技术可代替地或组合地使用。例如，一些实施例可以使用软件布置(如协作组API)来布置以实现跨线程块集合的并发和负载均衡，或者还可使用其他技术来提供或保证GPU硬件的同一相关硬件域或分区内的并发性和跨GPU硬件的同一相关硬件域或分区的负载均衡(例如，利用分布式共享存储器的所有线程不仅并发地运行，而且可以在SM上启动和发现，所有这些SM都在特定硬件域(诸如被称为GPC的GPU的子部分)内，例如，因为各个线程可以通过查询已经在哪个GPC上启动了线程来测试)。虽然这样的其他技术是可能的，但CGA 层次结构在效率和确定性方面提供了某些优势。

出于所有目的，本文引用的所有专利、专利申请以及公开物均通过引用结合在此，如同明确阐明一样。

虽然已经结合目前被认为是最实际和优选的实施例描述了本发明，但应理解的是，本发明并不限于所公开的实施例，而是相反，意图覆盖包括在所附权利要求的精神和范围内的各种修改和等效布置。

Claims (20)

1.一种处理系统，包括：

一组处理器，以及

工作分配器，其将线程块分配给所述一组处理器以供执行，所述工作分配器被配置为：

(a)跨所述一组处理器均衡所述线程块的加载，以及

(b)保证所述一组处理器能够并发地执行所述线程块，

其中相应线程块被指派标识符坐标以在所述处理器上执行。

2.根据权利要求1所述的处理系统，其中所述线程块由网格表示，并且所述处理器中的每一个处理器被配置为光栅化所述网格的相应部分。

3.根据权利要求2所述的处理系统，其中所述网格包括三维网格。

4.根据权利要求1所述的处理系统，其中所述处理器包括流式多处理器，并且所述工作分配器包括硬件电路。

5.根据权利要求1所述的处理系统，其中所述工作分配器包括被配置为跨处理器集合分配工作的第一工作分配器，以及被构造为将工作指派给各个处理器的多个第二工作分配器。

6.根据权利要求1所述的处理系统，其中所述工作分配器包括所述一组处理器的查询模型，并且使用所述查询模型针对所述一组处理器的影子状态来启动所述线程块，以测试所述线程块是否能够并发地启动。

7.根据权利要求6所述的处理系统，其中所述工作分配器维护持续更新的实时查询模型，以及存储影子状态的进一步的查询模型。

8.根据权利要求6所述的处理系统，其中所述工作分配器以迭代或递归方式使用所述查询模型来测试线程块组的多个层次级别的启动。

9.根据权利要求1所述的处理系统，其中所述工作分配器通过同时选择多于一个处理器以启动其上线程块，而跨所述一组处理器对所述线程块进行负载均衡。

10.根据权利要求1所述的处理系统，其中所述相应线程块是协作组阵列CGA的一部分。

11.根据权利要求10所述的处理系统，其中所述工作分配器选择性地不启动多于一个线程阵列，所述线程阵列是所述处理器中的任何一个处理器上的公共阵列的一部分。

12.根据权利要求1所述的处理系统，其中所述一组处理器各自包括独立地导出或计算唯一线程块标识符的硬件。

13.根据权利要求1所述的处理系统，其中所述工作分配器被配置为基于所述处理器的相应加载水平来确定哪些处理器可能最快地执行新工作。

14.一种处理方法，包括：

接收表示线程块的协作组阵列的网格；

推测性启动所述线程块，其包括基于占用水平跨一组处理器对所述线程块进行负载均衡；以及

如果所述推测性启动揭示所述网格将在所述一组处理器上并发地执行，则在所述一组处理器上启动所述线程块。

15.根据权利要求14所述的处理方法，还包括处理器中的每一个处理器以分布式方式光栅化所述网格。

16.根据权利要求15所述的处理方法，还包括向处理器广播线程块指派，每个处理器通过响应于所述广播确定多维标识符的全局进程并在所述全局进程中为其自身的线程块指派生成多维标识符来光栅化所述网格。

17.根据权利要求14所述的处理方法，其中所述推测性启动由硬件电路执行。

18.根据权利要求14所述的处理方法，包括跨所述处理器执行所述负载均衡。

19.根据权利要求14所述的处理方法，其中网格是三维的并且所述标识符是三维的。

20.一种处理系统，包括：

启动测试电路，其被连接以接收启动线程组阵列的指令，所述启动测试电路被配置为确定所述线程组阵列中的所有线程组是否能够在一组处理器上并发地执行，其中所述一组处理器中的至少一些处理器已经在执行其他任务；以及

启动电路，所述启动电路以所述启动测试电路确定所述线程组阵列中的所有线程组能够并发地执行为条件，在均衡跨所述组的所述处理器的加载的同时，并发地启动所述线程组阵列中的所有线程组。