CN116774914A

CN116774914A - 分布式共享存储器

Info

Publication number: CN116774914A
Application number: CN202210915158.6A
Authority: CN
Inventors: P·班加罗尔·普拉巴卡尔; 广田源太郎; R·克拉辛斯基; 龙泽; B·帕里斯; R·达什; J·塔基; J·F·小杜鲁克; L·姗; L·杜兰特; J·肖凯特; E·韦尔奈斯; N·戈维尔; M·帕特尔; S·德布; S·纳瓦达; J·埃德蒙森; G·帕尔默; W·甘地; R·马尼亚姆
Original assignee: Nvidia Corp
Current assignee: Nvidia Corp
Priority date: 2022-03-10
Filing date: 2022-08-01
Publication date: 2023-09-19
Also published as: US20230289189A1; DE102023105575A1

Abstract

公开了分布式共享存储器，其包括跨处理器(诸如GPU)分布或者分散的存储器块。在一个存储器块本地的处理核心上执行的线程能够访问不同处理核心本地的存储器块。在一个实施例中，通过处理核心之间的通信来实现对跨处理核心的集合分布的这些DSMEM分配的共享访问。这样的分布式共享存储器为位于存储器块附近的处理核心提供了非常低延时的存储器访问，并且还为更远的处理核心提供了也以这样的方式使用不干扰处理核心对诸如由L2高速缓存支持的主存储器或全局存储器的访问的互连访问存储器块的方式。这样的分布式共享存储器通过允许先前仅在同一处理核心内可能的数据共享和通信来支持跨多个处理核心的协作并行性和强扩展。

Description

分布式共享存储器

相关申请的交叉引用

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

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

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日提交的题目为“虚拟化处理器中的硬件处理资源(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日提交的题目为“用于处理器中的线程组的可扩展负载均衡的技术(Techniques for Scalable Load Balancing of Thread Groups in a Processor)”的美国专利申请No.17/691,872；

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 MultidimensionalData Structures and/or other Large Data Blocks)”的美国专利申请No.17/691,422。

背景技术

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

图1A示出了包括顺序相关的计算密集层的长链的示例深度学习(DL)网络。使用诸如例如将输入激活与权重矩阵相乘以产生输出激活的运算来计算每个层。所述层通常通过将工作划分为输出激活块(tile)而跨GPU或GPU集群并行化，每个输出激活块都表示一个处理核心将处理的工作。

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

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

计算机科学家将第一种方法称为“弱扩展(weak scaling)”，将第二种方法称为“强扩展(strong scaling)”。图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)的大部分。相比之下，更多最近的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、张量等)、线程束调度器、分派单元和本地寄存器文件，如在图2C中所反映的。

图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 and ComputingArchitecture)，”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年检索的)；BobCrovella等人，“协作组(Cooperative Groups)”(09/17/2020)，https：//vimeo.com/461821629；US 2020/0043123。

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

协作组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)(ACM2018)；乔格(Jog)等人，“OWL：提高GPGPU性能的协作线程阵列感知调度技术(Cooperative Thread Array Aware Scheduling Techniques for ImprovingGPGPU Performance))”(ASPLOS’13，2013年3月16-20日，休斯顿，德克萨斯州，美国)。

先前的“共享存储器”

在某些先前GPU架构中，每个SM包括本地的片上读/写存储器，该存储器对于SM是私有的并且被视为该SM的一部分，并且可以在SM上执行的多个线程之间共享。例如，SM的共享存储器可以包含寄存器文件和/或高速暂存(scratchpad)存储器和/或可以分配为“共享存储器”或L1高速缓存存储器的存储器。在SM上执行的任何计算线程或线程束可以使用诸如LDS(从共享存储器加载)、STS(存储到共享存储器)、或ATOMS(执行对共享存储器的原子存储器访问)的指令来访问此“共享存储器”。参见例如US 20210124582。

使用这些指令，在SM上执行的计算线程可以与在同一SM的同一或不同处理核心上执行的另一计算线程交换数据。然而，不存在用于在SM上执行的计算线程与在不同SM上执行的另一计算线程直接交换数据的机制。相反，如果在一个SM上执行的线程想要与在不同SM上执行的线程交换数据，则两个线程都需要利用全局存储器(诸如L2高速缓存存储器)。例如，在一个SM上执行的线程可将数据写入主存储器以供在另一个SM上执行的线程读取。因此，在不同的SM上运行的CTA之间共享数据是可能的，但是只能通过也用于访问主片外存储器的存储器层次结构的下一级(即，L2高速缓存存储器)。这样的全局存储器访问在资源方面是昂贵的，并且还占用带宽，否则该带宽可以用于从主(片外)存储器检索数据。然而，由于所有核心和SM的执行线程能够对SM私有的本地存储器进行访问，所以只要大部分数据共享被限制于同一CTA的线程，该同一CTA的线程毫无疑问地都在同一SM上执行，该先前布置就工作良好。这种情形就好像允许同一家庭的所有成员在他们自己之间共享图书馆图书，但是然后要求图书被返回到图书馆以便与邻居共享图书。

数据带宽跟不上处理带宽

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

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

已经尝试并实现了各种技术(诸如存储器管理改进、高速缓存改进等)以增加数据带宽。然而，经由导线增加更多的数据带宽会耗费面积和功率。增加更多的高速缓存成本面积和功率。所需要的是在不需要彻底检修和使存储器访问/管理层次结构复杂化的情况下，利用一个或更多个算法中固有的更多并行性，同时更高效地使用现今和将来可用的处理核心和高速缓存/互连层次结构的方式。还需要一种新的硬件模型，其利用低延时本地存储器来实现存储器共享并在处理核心或SM的组之间直接通信。

附图说明

图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示出了用于向/从CGA分配和解分配线性共享存储器池的线性存储器槽的示例电路布置。

图14A示出了具有槽的示例CGA共享线性存储器池。

图14B示出了软件可以如何划分线性共享存储器池以用于分配给CGA。

图15A示出了示例计算工作分配器电路框图。

图15B示出了用于将CTA分配给指定硬件分区内的SM的硬件工作分配器。

图15C-1示出了由CPU执行以生成CGA启动命令并将该启动命令发送至GPU的操作步骤的示例非限制性流程图。

图15C-2示出了包括CGA的推测性启动以提供基于硬件的并发保证的GPU硬件实现的操作步骤的示例非限制性流程图。

图16A示出了示例分布式共享存储器分配。

图16B示出了示例分布式共享存储器访问。

图17A、图17B、图17C示出了在CGA的CTA之间共享的示例存储器。

图18示出了包括分布式共享存储器的示例全局存储器映射。

图19示出了示例分布式共享存储器路由表。

图20示出了示例非限制性MEMBAR共享存储器屏障。

图21A示出了包括分布式共享存储器互连的总体示例系统的框图。

图21A-1示出了在图21A的框图的上下文内的示例非限制性流式多处理器架构。

图21B示出了分布式共享存储器互连的示例非限制性附加视图。

图21C示出了SM到SM通信的SM处理的另一视图。

图21D示出了SM到SM请求和响应处理的进一步视图。

图22示出了跨示例GPU架构的共享存储器请求和响应流。

图23A和图23B示出了图22架构如何执行示例操作。

具体实施方式

需要一种新的编程和硬件执行模型，该模型将允许在比单个SM或一个或更多个处理核心的其他集合或集群更大的级别处的更紧密的协调和协作。需要新的基于硬件的机制和硬件执行控制来控制数据和执行资源的放置，以确保执行并发性并且允许在比SM更大的级别处的高效数据共享和同步——同时避免检修基本就位存储器互连和高速缓存技术的需要。

本文的技术提供了一种新的编程/执行模型和关联的硬件支持，其实现具有关联的数据共享和同步的跨多个SM的并发应用的分布，以便更灵活且高效地利用跨GPU硬件组织和分区的不同层次结构级别的数据局部性(data locality)和数据重用。不再是并发协作执行和关联的高效数据带宽支持被限制到硬件处理器的有限集合(诸如在单个SM内)，新技术跨任何期望的硬件域扩展这样的并发协作执行并且支持具有基于硬件的并发保证、数据共享、同步和其他机制的这样的并发协作执行。

与这些目标一致并且支持这些目标，本技术引入新类型的共享存储器：分布式共享存储器(DSMEM)。这种DSMEM包括跨多个SM分布的存储器块，并且其使在一个SM上执行的线程能够访问不同SM内或与不同SM相关联的存储器块。参见图16A和图16B。在一个实施例中，DSMEM被实现为各个SM本地的并且可操作地与各个SM相关联的存储器分配。通过SM之间的通信实现对跨SM的集合或集群分布的这些DSMEM分配的共享访问。因此，通过允许先前仅可能跨同一SM的数据共享和通信，此类共享存储器支持跨多个SM的协作并行性和强扩展。

如上所述的这种更高性能编程模型允许线程组阵列(诸如“协作组阵列”(CGA)(参见下文))内的CTA直接访问在CTA创建时在每个SM内分配的共享存储器。这被称为分布式共享存储器(DSMEM)，因为逻辑共享的存储器地址空间的区段物理上位于每个SM“内”。DSMEM可以用于整个CGA中的线程之间的数据共享或者同步。直接SM到SM通信被支持用于读、写和原子(atomic)。与对L2高速缓存或帧缓冲区(FB)存储器的全局存储器访问相比，SM-SM通信的较低延迟允许应用更容易地跨CGA中的N个SM强扩展。

作为示例，DSMEM可以用于加速深度学习应用中的拆分-K并行减少。在没有DSMEM的情况下，必须在L2高速缓存中执行拆分-K减少。使用DSMEM，由于一个SM能够直接写入另一SM的共享存储器，拆分-K减少更快。

作为第二示例，SM在进行GEMM内核的内部循环时的数学吞吐量已经变得如此之高，使得跨CGA中的所有CTA每64-256个时钟周期需要同步。为DSMEM通信建立的低延时SM2SM网络允许这些同步不再是用于计算的限制器。

本文所描述的CGA保证了CGA中的所有CTA共存于机器上。在一个实施例中，硬件推测性地启动CGA以确保在实际启动CGA之前它将适合可用的处理资源。该基于硬件的并发保证意味着可以共享SM内的资源，示例是分布式共享存储器(DSMEM)。利用基于软件的协作组API，仅有全局存储器可用于共享。需要硬件改变来暴露来自CGA中的每个SM的DSMEM，并且CUDA编程模型被扩展至匹配。共享DSMEM允许更快的同步、多播以减少来自全局存储器的数据业务(traffic)，并且具有更接近每个SM中的处理元件的一些深度学习(DL)参数数据，而不是必须从存储器重新获取它们。DSMEM暴露于程序员，程序员可根据需要将数据移动到此较近的存储器中以优化算法性能。使DSMEM可从作为同一CGA的一部分的任何SM访问，是CGA的硬件改进，其允许强扩展。

进一步的改进包括若干新的硬件能力，诸如：

·将SM到SM通信引导到作为同一CGA的一部分的所有CTA的共享存储器

οDSMEM支持读/写/原子指令

ο基于CTA_id_within_CGA的分段寻址

οCGA_id分配和回收协议

οCAM结构，用于允许从远程SM的DSMEM寻址

ο在GPCRAB中的新的低延时SM2SM通信网络

οDSMEM合并的写确认和优化的存储器屏障(Membar)处理

·用于在CGA创建时以及在CGA上下文切换过程中同步DSMEM使用的硬件屏障

ο使用硬件屏障确保在任何CTA引用(reference)DSMEM之前，DSMEM在CGA的所有CTA中可用。

ο使用硬件屏障来确保CGA中的所有CTA在针对上下文切换保存SM状态之前已经完成了所有DSMEM访问，并且在允许新的DSMEM访问之前已经通过上下文恢复恢复了所有DSMEM。

·使用DSMEM的CGA/CTA退出和错误处理协议。

在GPC_CGA中从任何CTA访问DSMEM的能力实现了若干其他特征，诸如：

·DSMEM支持程序性多播和TMA单元。

·DSMEM支持SM之间的快速数据交换。

背景：CUDA协作线程阵列(CTA)编程模型

CUDA中的编程器将它们的计算描述为被称为“协作线程阵列”或CTA的并行线程块的网格。在此上下文中，CTA是并发地或并行地执行内核的线程阵列。先前的CUDA编程模型使用CTA作为GPU软件(SW)的并行性的基本构建块。在一个这样的模型中，CTA可以具有多达1024个线程，并且所有线程被保证在同一SM上同时启动和执行。在这种模型中，因为一个SM运行CTA中的所有线程，所以线程可以利用SM内和/或连接到SM的共享存储器资源，以在线程之间共享数据、同步、通信等——这确保在并发执行的线程上的数据局部性和数据重用。

先前的CTA网格层次结构

因为许多应用需要多于1024个线程(即，多个CTA)，所以用于计算应用的原始CUDA编程模型基于“网格”——CTA阵列，其中每个CTA通常启动到GPU中的最小加载SM上。存在CTA可以含有的最大数量的线程。然而，执行同一内核的CTA可以被分批在一起成为CTA网格，使得在单个内核调用中可以启动的线程的总数非常大。这以降低的线程通信和同步为代价，因为不同CTA中的线程不能彼此通信和同步。取决于平台和平台加载，多个CTA可以并发地和并行地执行，或者它们可以顺序地执行。

每个CTA在CTA网格内具有唯一的CTA标识符(“ctaid”)。CTA的每个网格具有由称为nctaid的参数指定的1D、2D或3D形状。每个网格还具有唯一的时间网格标识符(“gridid”)。线程能够通过预定义的只读专用寄存器(诸如“％tid”、“％ntid”、“％ctaid”、“％nctaid”和“％gridid”)读取和使用这些值。参见例如并行线程执行ISA：应用指南(NVidia v5.02017年6月)。

图8和图11A示出了示例层次结构，其中网格包括多个CTA。在这种模型中，CTA的每个线程被分配它自己的资源，例如私有存储器，并且同步通常发生在线程级别。每个SM调度了分组在一起作为“线程束”(使用织物类比)的若干个(例如，32个或其他可编程/平台相关值)线程的并发执行。“线程束”是来自单个CTA的线程的最大子集，使得线程在同一SM上同时地执行相同指令。一般来说，在SM上以SIMD方式执行的线程束，即，在线程束中的所有线程共享相同指令流且在锁定步骤中一起执行(这有时被称为单指令多线程或SIMT)。同时，如图9所示，SM的线程束调度器调度CTA中的所有线程束，用于在SM上并发地启动，以保证CTA中的所有线程并发地执行。在现代GPU中，SM具有满足或超过CTA中的线程的最大数量的并行计算执行能力——这意味着整个CTA(或在一些GPU中，多个CTA)总是可以在同一SM上并发地启动。

在图9和/或图2C的上下文中，SM或“流式多处理器”意指如在USP7,447,873中所描述的Nordquist设计的(包括对其的改进及其进步)并且例如在许多代NVIDIA GPU中实现的处理器。例如，SM可包括多个处理引擎或核心，所述多个处理引擎或核心被配置为并发地执行布置在多个单指令、多数据(SIMD)组(例如，线程束)或单指令、多线程(SIMT)中的多个线程，其中在SIMD或SIMT组中的同一组中的每个线程执行同一数据处理程序，所述数据处理程序包括关于不同输入对象的指令序列，且使用处理引擎或核心中的不同处理引擎或核心来执行在SIMD或SIMT组中的同一组中的不同线程。SM通常还可提供(a)具有多个通道的本地寄存器文件，其中每个处理引擎或核心被配置为访问通道的不同子集；以及指令发布逻辑，其被配置为选择SIMD或SIMT组中的一个并且将相同数据处理程序的指令中的一个指令并行地发布到多个处理引擎中的每一个，其中每个处理引擎使用其可访问的本地寄存器文件通道的子集与每个其他处理引擎并行地执行相同指令。SM通常还包括核心接口逻辑，该核心接口逻辑被配置成启动一个或更多个SIMD或SIMT组的执行。如图2C所示，这样的SM已经被构造成提供快速的本地共享存储器，这使得能够在SM上执行的CTA的所有线程之间进行数据共享/重用和同步。

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

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

迫切需要新方法增加GPU并行性/复杂性

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

协作组阵列

本文的示例非限制性技术引入新级别的层次结构——“协作组阵列”(CGA)——以及关联的新编程/执行模型和支持硬件实现方式。本实施例还提供了基于新的CGA层次结构的网格的新的编程模型。

在一个实施例中，CGA是CTA的集合，其中硬件保证CGA的所有CTA被启动至CGA指定或与其相关联的同一硬件组织级别。CGA与硬件域或分区之间的这种空间亲和性提供了某些优点，如下面详细讨论的。该硬件被配置为确保在目标硬件级中有足够的处理资源来启动CGA的所有CTA，然后才启动任何其他。

如图11B所示，“网格”的构造被扩展为CGA阵列。每个CGA是CTA阵列(参见下文图12、图12A的讨论)。这样的CGA提供相对于应用所需的存储器和相对于彼此的共同调度，例如，对CTA在GPU中被放置/执行的位置的控制。CGA中的CTA不需要在同一SM上运行，也不需要它们在作为彼此的邻居或物理上彼此靠近或接近的SM上运行。但是，在一个实施例中，CGA中的CTA可在多个SM上并发地启动，所述多个SM处于相同硬件域中并因此连接到处理系统以能够以某种方式交互和/或通信(例如，连接到系统存储器层次结构的同一部分、能够彼此通信、以其他方式彼此耦合等)。这使得应用能够在紧密协作的CTA中的所有线程之间看到更多的数据局部性、减少的等待时间和更好的同步。

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

在一个实施例中，每个CGA由硬件保证将同时地/并发地启动和执行的CTA组成。CGA中的CTA可以(并且在一般情况下将)在GPU内的不同(以及可能许多不同的)SM上执行。即使CTA在不同SM上执行，GPU硬件/系统仍然提供跨SM保证：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与已知的空间关系并发地执行，因此其他硬件优化是可能的，诸如：

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

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

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

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

如以上所讨论的，示例实施例提供分布式共享存储器(DSMEM)。DSMEM由新的硬件支持实现和提供，该新的硬件支持移除仅SM上执行的线程块可读取、写入或以其他方式访问SM本地的存储器的约束。在一个示例性实施例中，DSMEM是在同一GPC上执行的CTA之间共享的——也就是说，它们是同一GPC-CGA的一部分并且因此由硬件保证在同一GPC上并发地运行。

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

层次结构的新级别-CGA

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

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

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

更详细地，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/

硬件保证CTA的并发执行

在示例实施例中，CGA中的所有CTA被共同调度。这意味着GPU硬件将不允许CGA中的任何CTA启动，除非/直到GPU(或大于CGA具有硬件亲和性的SM的其他硬件分区)上存在空间用于CGA中的所有CTA启动。该硬件保证允许软件对CGA中的所有线程将同时执行的事实进行计数，使得像跨所有线程的屏障同步、数据共享等事情是可能的。CGA中没有单个CTA可以无限期地被卡住等待启动——在一个实施例中，或者整个CGA被启动或者没有CGA被启动。

在示例实施例中，每个CGA具有分配给它的(至少一个)硬件屏障，并且CGA中的所有CTA可以引用(reference)该CGA硬件屏障。参见上面标识的、在2022年3月10日提交的名称为“具有异步交易支持的硬件加速同步机制(Hardware Accelerated SynchronizationMechanism With Asynchronous Transaction Support)”的美国专利申请No.17/691,296。

在一个示例部署中，硬件维护CGA中运行的CTA(即，尚未退出的CTA)的数量的计数，并且软件可跨CGA中所有正在运行的CTA中的所有线程执行屏障同步。该硬件屏障例如对引导所有CTA并确认它们都已经被启动是有用的。

示例协作组阵列网格

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

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

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

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

CTA分配和跟踪

示例硬件实现方式在每个SM中提供有助于跟踪CGA内的CTA(即，以允许CTA确定其是CGA内的哪个CTA)的新的S2R寄存器。例如，使用新的S2R来指定GPC_CGA内的1维CGA_CTA_id(CTA的X、Y、Z坐标仍可以按照整个网格编号，忽略任何CGA)。在一个实施例中，SM实现S2R(特殊寄存器到寄存器)操作以返回CGA内的线性CTA ID。具体而言，称为gpc_local_cga_id(所使用的位数可取决于所支持的同时活动的CGA的数量)的附加的基于硬件的多位标识符被用来标识GPC的命名空间内的CGA并跟踪该CGA的活动CTA的数量。如下面将说明的，该相同值gpc_local_cga_id用于索引分布式共享本地存储器。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的空闲列表或空闲向量，并且如果CGA用完gpc_local_cga_id，则停止CGA启动。

在图12A所示的示例中，标记为“C”的CTA需要能够告诉(学习)一旦由硬件分配，其是6-CTA CGA内的哪个CTA(即，每个协作线程阵列现在是有序协作组阵列的一部分)。已知整个网格的尺寸和上述各个局部网格层次结构的尺寸，可以将CTA在整个网格内的坐标转换为该CTA在其CGA内的坐标。在示例实施例中，每个网格或CGA是根据层次结构中的下一级使用3维坐标来定义的。每个CTA经由如上所讨论的硬件寄存器将其CTA id(X，Y，Z)暴露给着色器中的软件。对于GPC_CGA，可以使用新的S2R硬件寄存器来确定或发现启动过程所预设的GPC_CGA内的1维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坐标的编程模型，同时为CTA在其CGA内提供附加的坐标。

基于硬件CGA启动保证

在图15A、图15B、图15C-1和图15C-2所示的一个或更多个实施例中，使用主计算工作分配器(“CWD”)硬件电路420在GPU上启动CGA，同时提供CGA的所有CTA可以同时启动的基于硬件的保证。关于用于调度工作的示例GPU CWD和MPC的更多信息，参见示例20200043123，特别是图7和相关联的描述。在本文的实施例中，CWD 420包括寄存器、组合逻辑和硬件语句机器。扩充/增强其功能性来提供推测性CGA启动能力，以确认资源可用于启动CGA中的所有CTA。如果CGA的所有CTA不能同时启动，则CWD 420不启动CGA的任何CTA，而是等待，直到相关GPU硬件域的足够资源变得可用，以使得CGA的所有CTA可被启动，从而它们并发运行。在示例性实施例中，CWD 420使用多级工作分布架构支持多级CGA的嵌套(例如，GPU-CGA内的多个GPC-CGA)，以在相关联的硬件亲和性/域上提供CGA启动。

更详细地，图15A中示出的CWD 420在使用推测性执行技术确定CGA的所有CTA可适合在指定硬件域中可用的硬件资源上之后启动CGA中的CTA。以这种方式，一个示例模式中的CWD 420确保在启动CGA的任何CTA之前，跨所有GPC或其他相关硬件域有足够的资源用于CGA的所有CTA。在一个实施例中，启动CGA的CTA的算法可从传统(非CGA)网格启动借用一些技术，同时首先确认CGA的所有CTA可以确保它们将同时运行的方式启动。

图15A示出了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。在示例实施例中，这些和其他结构的功能沿着以下线在示例实施例中被增强：

在一个实施例中，CWD 420从与GPU协作的CPU 212接收用于在CPU 212上执行的各种进程的任务。在示例实施例中，由GPU执行的每个计算任务可以对应于CGA(尽管非CGA任务也可以被容纳)。在CPU 212上执行的每个进程或应用可以发出这样的任务。例如，CPU212可以执行存储在诸如全局存储器之类的非暂时性存储器中的一个或更多个程序，以生成命令GPU启动CGA网格的CGA启动命令。

在操作中，CPU 212执行为GPU生成“网格启动”(和其他)命令的驱动程序(参见图15C-2)。网格启动命令具有定义由GPU执行的CGA网格的关联的状态参数。在一个实施例中，状态参数包括指定网格中的CGA的数量、每个CGA中的CTA的数量和每个CTA中的线程的数量的大小参数(参见图15C-2，框552、554)。

如果线程标识符是多维的(例如，2-D或3-D)，则指定CTA在每个维度上的大小；因此，对于具有1-D线程ID的CTA，大小可以被指定为n0，或者对于具有3-D线程ID的CTA，大小可以被指定为n0＝d0*d1*d2。类似地，如果CTA或CGA标识符是多维的，则指定网格在每个维度上的大小。状态参数还标识要由每个线程执行的CTA程序、用于网格的输入数据的全局存储器(见图15)中的源位置(例如，阵列)以及用于由网格产生的输出数据的全局存储器中的目的地位置(例如，阵列)。例如参见USP7,937,567；USP9,513,975；和USP9,928,109，以了解CPU如何使用例如面向线程的编程环境(例如来自NVIDIATM的CUDATM编程环境)来启动网格的背景。CPU 212还布置要由SM执行的线程被存储在例如全局存储器中，使得GPU的直接存储器访问硬件可通过系统的存储器管理单元(MMU)检索线程以供SM执行(见图15)。

示例CGA启动命令

在示例实施例中，从CPU 212到CWD 420的启动命令可指定CGA网格，该CGA网格包括复合线程块和CGA的各个维度的枚举。作为一个示例，CGA网格启动命令可以指定运行10240个CGA，其中每个CGA是8个CTA，其中每个CTA具有256个线程，其中每个线程具有(需要)64个寄存器，并且其中每个CTA分配128KB的共享存储器，等等。这些数字被编码成如{10240，8，256，64，128}的启动命令，并且其是硬件工作分配器CWD 420在启动SM上的线程或CTA时处理的信息。CPU 212将这样的启动命令发送至GPU内的调度器410(图15C-2，框558)。在另一实施例中，SM可向CWD 420发出这些命令，即描述为由CPU 212执行的任务也可由SM完成。

使用以上技术，应用程序可以在GPC或其他硬件分区中启动许多小的CGA，但数量随着CGA大小的增长而减小。在某一点(取决于硬件平台)，CGA不再能够适合GPC或其他硬件分区，这可能损害代码便携性。如果假设每个平台都具有具有4个TPC的至少一个GPC，则保证在未来架构之间的兼容性的最大CGA大小是8个CTA。给定的应用程序可以基于查询平台来动态调整CGA大小，以根据1)CTA资源需求和2)每CGA的CTA数量来确定能够在GPU中并发运行的CGA的数量。

GPU CGA调度和启动

在示例实施例中，GPU内的调度器410从CPU 212接收任务并将它们发送到CWD 420(图15C-1，框502、504)。CWD 420从多个CGA查询和启动CTA。在一个实施例中，一次对一个CGA进行操作。对于每个CGA，CWD 420推测性地启动CGA中的所有CTA，递增“启动”寄存器来存储推测性启动。如果在推测性地启动CGA的所有CTA之前SM或硬件域中的其他处理器中的所有空闲槽用尽，则CWD 420终止该启动，并可稍后再次尝试。相反，如果对于CGA中的所有CTA存在足够的空闲槽，则CWD 420从在推测性启动过程中累积的“启动”寄存器生成sm_mask(该sm_mask数据结构存储用于在CGA启动的相关硬件域中的每个SM上运行的CTA的数量的预留信息)，并且一直移动到下一CGA。硬件分配CGA序列号，并将其附加到每个sm_mask。其还将CGA位的end_附加到最后一个，以防止来自不同CGA的sm_mask的交错。

示例CGA启动分组

基于成功的推测性启动，CWD 420将诸如以下的启动分组发送至GPC(SM)。这样的启动分组可以例如包括以下字段：

/>

CWD 420可以提供sm_mask的多个迭代波，以将CGA中的所有CTA映射到SM，从而使得CGA可以启动。一旦SM掩码准备好，上述启动分组被用于将它们(与关联的CGA ID)广播到GPU的所有SM工作调度器。在一个实施例中，CPU 212将GPU CGA顺序号附加到其发送到GPU的启动命令。该序列号被预加到为每个GPC CGA生成的sm_mask，并且用于将每个GPC CGA的sm_mask映射到GPU CGA(其还可以在将掩码发送到各个SM中的M管道控制器(MPC)之前由任何重排序单元使用)。

向所有SM广播启动分组允许SM内的所有MPC观察CGA/CTA启动的整个序列。通过观察CGA和CTA的流，每个SM的MPC(网格当前被指派给该MPC)能够冗余地并且独立地执行光栅化。广播也是lmem_blk_idx分组，其将lmem_blk_idx(参见图15A的LMEM块索引表432)从CWD420运载到SM。

多级统一工作分配器

图15B示出，在一个实施例中，CWD 420包括用于分配CGA工作的若干级别的工作分配器(WD)。例如，在GPU CGA由GPC CGA构成的情况下，CWD 420可以实现两级工作分配器：

·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工作分配器(图15C-2，框506)。GPC2SM工作分配器420b各自将GPCCGA的CTA分配到GPC内的SM(使用例如负载均衡模式或多播模式，如下所述)。图15B的统一工作分配器(UWD)420a/420b保证GPU CGA中的所有GPC CGA可以一起启动，并且每个GPCCGA中的所有CTA可以一起启动。

在支持CGA的更深嵌套的其他实施例中，这个UWD可以被扩展到所需要的任何数量的级别。更详细地，在一个实施例中，CWD 420可包括或激活需要这种功能的CGA网格的分层三级统一工作分配器(UWD)：

·GPU2SM工作分配器(GPU2SM WD)处理CTA和由CTA组成的GPU CGA。

·GPU2GPC工作分配器(GPU2GPC WD)协调GPC CGA和由GPC CGA构成的GPU CGA的负载平衡。它会谈到最低级别的工作分配器——GPC2SM WD

·GPC2SM工作分配器(GPC2SM WD)处理GPC CGA的实际负载平衡。在UWD中存在N个GPC2SM WD，GPU中的每个虚拟GPC对应一个。

在一个实施例中，UWD因此知道GPC/TPC层次结构，以便促进CGA的空间亲和性(例如，来自GPC CGA的所有CTA将在同一GPC上启动)。然而，在一个实施例中，UWD不执行CTA网格坐标的光栅化；相反，该功能(在一些GPU的先前实现方式中其由CWD执行)被移动到SM(并且具体地，被移动到SM内的MPC)，如通过状态同步操作所同步的。

示例CWD 420推测性启动过程

在实施例中，UWD 420a、420b响应于从CPU 212接收到CGA启动命令来执行以下进程：

I.CGA的推测性启动(图15C-2，框508)

阶段1：

第一步骤是状态快照：从任务表436(图15A)读取剩余数量的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)中被处理：推测性启动和实际启动。

推测性CGA启动

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

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

参见图15B，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可以被一起启动(图15C-2，“否”退出到判定框510)。

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

示例CGA负载均衡

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

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

·MULTI_CAST——基于硬件的调度器跨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中所描述的。

关于这些LOAD_BALANCING和MULTI_CAST方法的更多信息可以在2022年3月10日提交的名称为“跨多个计算引擎的程序控制的数据多播(Programmatically ControlledData Multicasting Across Multiple Compute Engines)”的上述美国专利申请No.17/691,288中找到；以及2022年3月10日提交的名称为“用于处理器中的线程组的可扩展负载均衡的技术(Techniques for Scalable Load Balancing of Thread Groups in aProcessor)”的上述美国专利申请No.17/691,872。

实际CGA启动

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

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

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

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

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

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

在线性存储器中分配CGA存储器槽

在示例实施例中，CWD 420还负责在线性存储器池(见下文)中分配CGA存储器槽以及刷新和回收槽。假设CWD确定存在足够的资源并且以上阶段2完成或正在进行，则CWD 420将信息传递到驻留在GPC内的GPM功能电路块。每个GPM分配屏障槽，当GPC CGA中的所有CTA完成时，还分配CGA_id和轨道。每个SM中的MPC(M-管道控制器)电路同时跟踪每CTA的槽，并且参与启动CTA到其关联的SM上以实际进行工作。当工作完成时，SM向GPM报告CTA完成状态。当GPM电路接收到CGA中的所有CTA已完成的状态信息(图15C-2，框514)以及对CGA的所有存储器分配已被刷新的状态信息(图15C-2，框516)时，GPM电路可发信号通知CWD 420以释放池中的CGA存储器槽，使得它可被分配给另一CGA(图15C-2，框518)。

CGA共享存储器

在示例实施例中，做出新的支持和规定，以实现CGA的CTA之间的存储器分配的共享。现在，CGA的CTA可以在不同的SM上运行，示例实施例为跨SM并且因此跨CGA共享存储器提供硬件支持。

实施例提供不同类型的这种共享存储器，诸如：

·分布式共享存储器(DSMEM)

·线性共享全局存储器。

这种共享存储器是SM内的和/或SM本地的存储器的补充和/或改进，所述SM在运行于SM上的CTA的线程之间共享。

例如，在示例性实施例中，GPC_CGA中的每个CTA可以引用作为其CGA的一部分的所有其他CTA(甚至在不同SM上执行的那些CTA)的共享存储器(例如，在L1高速缓存中)。硬件还支持整个GPC_CGA的共享存储器内的负载/存储/原子学的直接SM到SM通信。

如以下详细描述的并且如图16A和图16B中所示，分布式共享存储器DSMEM由每个SM本地的/每个SM内的存储器的至少一部分、以及允许SM访问彼此的本地存储器的硬件互连和其他功能组成。该DSMEM使得每个CTA能够暴露其存储器的部分作为连续(contiguous)段，该连续段在每个CTA的段之间的地址空间中具有空洞。这种共享存储器提供了利用数据定位和数据重用的增加的能力。为了减少数据交换延时，本文的示例技术还通过CGA共享存储器提供更快的数据交换。基于改进的基于硬件的同步，这样的数据交换比经由全局存储器的数据交换快许多倍。参见例如上述的2011年3月10日提交的名称为“硬件加速同步与异步事务支持(Hardware Accelerated Synchronization With Asynchronous TransactionSupport)”的美国专利申请No.17/691,296。

CGA内的CTA还可以基于常规协议通过全局存储器进行交换(例如，生产者将数据存储在全局存储器中，然后执行MEMBAR.GPU命令以使得存储器位置跨GPU或其他硬件分区可见，并且设置写入标志；消费者轮询该标志，使其自身的L1高速缓存值无效，并且从全局存储器加载数据)。在示例实施例中，硬件改进支持CGA内的所有CTA也可共享的线性共享全局存储器。

图17A和图17B示出了示例CGA场景，其中CGA在任意数量(在这种情况下为四个或更多个)的SM上执行。在示例实施例中，CGA内的所有CTA线程可以引用两种不同类型的共同可访问的共享存储器：分布式共享存储器(DSMEM)和CGA线性共享存储器。对于DSMEM，GPU中的硬件支持允许CGA中的不同CTA读取和写入彼此的共享存储器。因此，第一CTA的加载、存储和原子存储器访问可以以第二CTA的共享存储器为目标，其中第一CTA和第二CTA在同一CGA内。

CTA线性共享存储器

在示例实施例中，GPC_CGA中的CTA还可以或替代地从全局存储器中的公共数据池中分配存储器。在一些实施例中，该数据池完全处于具有某些策略性硬件支持(例如，具有节流的存储器槽分配)的软件控制下。该池的大小可以被确定成使得由所有执行的GPC_CGA请求的存储器总是适合近存储器(如L2高速缓存)以降低延时，或者它的大小可以被确定成提供比可以永远适合L1或L2高速缓存中的大得多的共享存储器结构。如图17B所示，这种共同的“CGA线性共享存储器”数据池可用于不具有附加层次结构或者不适合其他类型的CGA共享存储器的数据。这种CGA线性共享存储器被分配给CGA，由CGA中CTA的所有线程可同等访问，以提供统一的访问共享存储器，从而简化了编程模型。

在一个实施例中，该CGA线性共享存储器被线性地寻址。因为CGA线性共享存储器在全局存储器中，它可以基于CGA指定的特定硬件域中的(高速缓存)DRAM，其中硬件帮助将这样的线性共享全局存储器的“CGA存储器槽”分配给在CGA启动时提供的每个CGA。CGA线性共享存储器通过提供跨CGA存储器来提供易于编程，跨CGA存储器在上下文中类似于由先前GPU生成提供的CTA内的所有线程可访问的CTA共享存储器。可以为上下文(不同的上下文可以使用不同的池)分配存储器池，并且软件可以控制诸如组织、总大小、每CGA的大小、跨CGA的共享等之类的参数。

图17C示出了提供线性共享全局存储器的示例新的GPC_CGA能力，该线性共享全局存储器是基于由CGA指定的硬件域分区的并且还提供跨CGA共享的CTA DSMEM。在这种情况下，所示的两个GPC_CGA中的每一个由在两个SM上执行的CTA构成。因此，第一GPC_CGA(0)在包括SM0、SM1的GPC上执行；并且第二GPC_CGA(1)在包括SM2、SM3的另一GPC上执行。在每个GPC内，所有CTA线程还可以引用CGA线性共享全局存储器和DSMEM。然而，第一GPC_CGA无法看到第二GPC_CGA的CGA线性共享全局存储器或DSMEM，反之亦然。这些资源每次仅对一个GPC_CGA可见。因此，该示例示出了在一些非限制性示例中，硬件资源可以跨CGA指定的特定硬件域提供CGA线性共享全局存储器和DSMEM，但是不准许在该特定硬件域之外共享。由此，硬件支持跨CGA的所有线程组共享数据但不向不是CGA的一部分的附加线程组共享数据的共享存储器。此类共享布置和相关联的限制可例如在其中不同用户或不同应用在不同GPU硬件分区上运行的多用户上下文中是有用的。

线性共享全局存储器分配和控制

示例线性共享全局存储器实现基于具有由硬件分配和回收的全局CGA_linear_memory_slot索引。参见图13和图14A。然后，大多数存储器管理可由软件基于提供给每个运行CGA的唯一的CGA_linear_memory_slot来完成，该CGA需要CGA线性共享全局存储器。线性共享全局存储器可以是常规全局存储器，其可以高速缓存在L2高速缓存中并且由DRAM中的物理地址支持，使得在存储器管理单元(MMU)或L2高速缓存中不需要特殊逻辑。

在一个示例实施例中，硬件提供每GPC大小的CGA的唯一的全局CGA线性存储器槽索引，其标识该CGA正在使用池中的哪些缓存区，并且使用这个槽索引来防止CGA启动，直到在网格指定的范围中存储器槽是可用的。在这种实现方式中，硬件提供的CGA_linear_memory_slot索引在所有运行的CGA上都是唯一的。如图14B所示，这允许来自不同虚拟引擎的不同网格(其可以竞争资源)同时在硬件上运行。在一个实施例中，由于槽索引的数量受到限制，因此指定与其他CGA相同范围的CGA可以防止那些其他CGA启动。

此外，API可以被扩展为包括以下相关参数：

·CgaLinearSlotLowerLimit

·CgaLinearSlotUpperLimit

·CgaLinearSlotLimitEnable(标志)。

CGA_linear_memory_slot索引在CgaLinearMemorySlotLimitLower至CgaLinearMemorySlotLimitUpper的范围内，其中，这两个值均在网格启动时由软件进行配置。因此，硬件提供了确保CGA_linear_memory_slot在CgaLinearMemorySlotLimitLower至CgaLinearMemorySlotLimitUpper的范围内的分配器。如上所述，硬件防止GPC CGA的启动，GPC CGA不能在所需范围内分配CGA_linear_memory_slot。如果在所需范围内没有可用的槽，则这样的运行时节流阻止CGA启动。在示例实施例中，这样的硬件节流被用于最小化总存储器占用。

如以上所讨论的，硬件可以经由S2R寄存器将CGA槽暴露给着色器代码，从而使得软件可以设置对开发者期望的任何槽进行分区。在一个实施例中，硬件仅跟踪CGA存储器槽号，并将其他分配控制留给软件，从而提供灵活性。

在一个实施例中，CGA线性共享存储器缓冲区的虚拟地址被称为CGA_linear_mem_base。以字节(S)为单位的缓冲区大小被称为CGA_linear_mem_size。这些值均由着色器代码使用以在执行时计算其共享存储器区的虚拟地址。这些值可使用常数传递到着色器，或直接编译到着色器代码中。这里是着色器代码可以实现的示例等式：

Linear_mem_base_for_this_CGA＝CGA_linear_mem_base+(CGA_linear_memory_size*CGA_slot)。

软件被期望为每任意组的网格分配视频(全局)存储器中的缓冲区，以用作给定上下文的CGA线性共享存储器区域。从概念上讲，如图14A所示，这个缓冲区被分解成N个大小相等的S个字节段。执行的每个CGA基于其CGA_linear_memory_slot被给予对N个段中的一个的访问。S可以基于GPC大小的CGA需要的共享存储器的大小进行设置，N可以基于允许在硬件域上并发地执行的GPC大小的CGA的数量进行设置。总缓冲区大小将随后为N*S个字节。如上所述，软件可以具有每上下文的多个池。

假设要运行，在一个示例中，每个CGA要求CGA线性存储器中的1-MB缓冲区。如果网格具有10,000个CGA，则运行整个网格将需要10,000个1-MB缓冲区。然而，大多数平台将不能同时运行10,000个CGA。示例实施例通过提供CGA线性共享存储器的池，并且使CGA声明多少可以同时运行来减少存储器占用。例如，如果网格内的CGA声明最多能同时在平台上运行N个CGA(N＜10,000)，则最多只需要分配N个1-MB的缓冲区(不是10,000)。硬件节流跟踪已分配多少个缓冲区，并阻止N+1个CGA启动，直到先前启动的CGA完成并释放其分配的缓冲区。这样，软件可以限制能够并发执行的CGA的总数，从而限制整个共享存储器缓冲区的大小。

分布式共享存储器

本文的技术的多个方面还涉及基于硬件的支持机制，其使得在SM上并发地运行的CTA能够从在同一相关硬件域内(诸如在同一GPC内)的其他SM上运行的其他CTA读取、写入以及进行对分配给该其他CTA的存储器的原子访问。如2022年3月10日提交的名称为“跨多个计算引擎的程序控制的数据多播(Programmatically Controlled Data MulticastingAcross Multiple Compute Engines)”的上述美国专利申请No.17/691,288中所描述的程序多播技术，和如2022年3月10日提交的名称为“处理器和存储器中的快速数据同步(FastData Synchronization In Processors And Memory)”的美国专利申请No.17/691,303中所描述的对SM之间的通信的改进可以由相同或类似的基于硬件的支持机制支持，但DSMEM能力是单独的、不同的、且高度有利的。

此外，由于不必使所有CTA在同一周期上启动，本技术的一个实施例使用由GPM(GPC的组件)实现的改进的屏障技术，如在2022年3月10日提交的名称为“硬件加速的同步与异步事务支持(Hardware Accelerated Synchronization With AsynchronousTransaction Support)”的美国专利申请No.17/691,296中所描述的。具体地，CGA层次结构和关联硬件支持保证CGA的所有CTA将并发运行，但不一定保证它们将全部在精确地同一时刻启动或在精确地同一时刻完成。CGA BAR能力在一些实施例中被用来确保CGA中的所有CTA准备好用于SM间通信。这种硬件屏障或其他机制可由需要与其他CTA交互以确保那些其他CTA事实上全部存在和/或在完成时和/或在抢占时存在的CTA使用。

示例实施例进一步提供允许运行在一个SM上的线程访问不同SM内的或与其相关联的共享分布式存储器的ISA改进。这样的ISA改进可以例如包括被设计成访问通用存储器并且在一个实施例中使用改进的硬件来访问DSMEM共享存储器地址的加载(LD)、存储(ST)和“原子”(ATOM)。先前的CUDA版本通过使用诸如LDS、STS和ATOMS之类的特殊共享存储器指令来支持共享存储器操作，该特殊共享存储器指令允许在SM上执行的线程访问该SM的被分配为在SM上执行的线程之间共享的存储器。在一个实施例中，这样的先前指令仍然有效，但是仍然限于访问SMEM而不是DSMEM(尽管如果需要它们可以被增强以访问DSMEM)。同时，在一个实施例中，加载、存储和原子存储器访问指令(以及诸如TMA指令和SYNC指令之类的其他指令)用于访问被映射到多个SM中的每一个本地的存储器的通用存储器空间——从而使能分布式共享本地存储器，其中多个SM中的每一个可在原子级上访问其他SM的本地存储器(基本上就好像其他SM的存储器在执行共享存储器访问的SM本地一样)。

包括私有、共享和全局存储器的“通用”存储器的示例整体存储器映射

图18是一个示例实施例的“通用”存储器的存储器映射。在一个实施例中，该存储器映射定义SM的存储器地址空间，并且包括虚拟地址和物理地址两者。在该示例中，通用存储器映射因此限定“通用”存储器地址空间，该“通用”存储器地址空间包括全局存储器(包括上述CGA共享线性存储器)并且具有本地(DSMEM)存储器被映射到其中的窗口。该本地存储器中的一些对SM或在SM上运行的任务是私有的，但是本文的实施例中的其他本地存储器与在其他SM上运行的其他任务(即在同一CGA内的其他CTA)共享或可共享。

在一个实施例中，不是共享存储器的通用地址空间的大部分被映射到全局存储器。因此，在一个实施例中，这种通用地址空间由全局存储器和DSMEM组成。其他实施例可以包括用于其他存储器空间的其他特殊区域，诸如例如线程私有栈存储器。在该示例中，全局存储器是由L2高速缓存支持的DRAM。因此，该全局存储器是SM可以通过GPU的主存储器管理单元(MMU)访问以读取和写入数据的系统的“主存储器”。全局存储器可包括例如用于显示图像的帧缓冲区存储器；程序存储；数据存储；纹理存储；光线追踪BVH存储；以及包括CGA线性共享存储在内的许多其他类型的数据。

图18存储器映射进一步示出了DSMEM地址存储器块，其在右手侧被分解以示出共享分布式存储器的不连续块。在现有架构中，此DSMEM存储器块小得多(例如，16MB)并且被映射到在逻辑上是每个SM的一部分的“共享存储器”，并且可以在SM上执行的所有线程之间共享，但是不能被运行在不同SM上的其他进程共享。换言之，此“共享存储器”对SM是私有的并且使得SM能够访问其自身的本地存储器并且被称为“共享存储器”，因为在SM上执行的不同线程或线程组能够共享存储器并且使用它来交换数据。没有提供使一个SM能够共享或访问另一SM的共享存储器的能力。不同的SM各自将看到相同的(例如，16KB的“共享存储器”)，但该映射使得特定SM能够仅访问在线程组或运行在该特定SM上的CTA之间共享的其自身的本地共享存储器。

在本文的示例实施例中，此“共享存储器”窗口现在已经扩展为包括针对GPC-CGA中的其他(并且在一个实施例中，每个)CTA的映射。换言之，“共享存储器”的本地存储器窗口已经被扩展以允许对运行(或在一些实施例中，可以运行)线程组或CGA内的CTA的所有其他SM的本地存储器的部分进行访问。在图18所示的一个非限制性示例中，可以激活256个这样的映射以适应GPC-CTA中的高达32、64、128、256或任何其他数量的CTA。当然，特定硬件平台可以根据需要或期望支持每GPC-CTA更少、更多或不同数量的CTA。

在一个实施例中，硬件在任何时间分配的此类区域的数量取决于CGA中的CTA的实际数量。因此，如果CGA包括32个CTA，则硬件将分配并启用32个共享存储器窗口——一个窗口用于CGA中的每个激活的CTA。类似地，如果CGA仅包括23个CTA，则硬件将分配和启用23个这样的共享存储器窗口。硬件可以动态地分配/解除分配这样的共享存储器窗口，例如与上述并发执行保证一致的附加CTA启动/完成。

SM执行的加载、存储和原子指令可以由CTA索引为地址中的用于选择共享存储器区域的某些位，以及特定于该特定CTA的共享存储器区域内的特定位置的附加(例如，较低阶)地址位。因此，对于访问共享存储器的LDS和STS指令，示例共享存储器地址可以看起来像以下：

CGA内的CTA ID

SMEM偏移

通过将“CGA内的CTA ID”设置为零(CTA可以读取S至R硬件寄存器以确定哪个CTAID被分配给它)，这种寻址布置可以为CGA-不知道的代码提供后向兼容性，从而允许在SM上运行的CTA寻址该特定SM本地的该CTA自己的共享存储器。传统使用因此被支持并且被扩展成允许CTA访问它们全部被分组在其内的CGA内的其他CTA的分布式共享存储器。

同时，以下格式可以用于允许SM发布LD、ST和原子指令以访问共享存储器：

0

DSMEM/SMEM孔径

CGA内的CTA ID

SMEM偏移

CTA可以通过读取指定图18中所示的共享存储器窗口的孔径(aperture)地址的硬件寄存器来确定孔径值。以上地址格式用于访问现在跨越CTA共享存储器阵列的通用存储器中的共享存储器窗口，每个CTA共享存储器具有一定大小。

在一个实施例中，可以由软件以编程方式调整要分配的每CTA的DSMEM分布式共享存储器区域的大小。例如，分配可以是每CTA 8KB、每CTA 16KB、每CTA 32KB等，直到可以取决于硬件平台和可用资源的预设限制。如上所述，在一个实施例中，图18所示的通用存储器地址空间的该可访问区被限于实际分配的共享存储器。换言之，虽然新的硬件支持通常将允许任何SM经由SM间通信访问任何其他SM的本地存储器，但是提供附加支持以检查访问是否合法/被允许-一个示例约束是目标是用于访问的特定位置被分配给访问线程被分组在其中的同一CGA内的CTA。

在一个实施例中，由通用存储器支持的一些原子操作可以或可以不由DSMEM支持。具体来说，这些原子操作中的一些可指示需要硬件支持的某些类型的读取-修改-写入操作来执行某些计算。在一些实施例中，DSMEM支持可为一些类型的原子指令提供比例如由全局存储器提供的硬件计算支持更有限的硬件计算支持。在其他实施例中，可为对DSMEM的原子访问提供相同或更大的硬件计算支持。由此，能够访问DSMEM的原子指令集可包含比能够访问通用存储器的其他部分的原子指令集多、少或不同的指令。

网络上的SM2SM分组

在一个实施例中，在SM之间交换的分组包括读/写/原子/写ACK(确认)/读数据/写数据、读响应(其包含读数据)、写数据分组(其包含写数据)、以及读错误。分组编码优化支持不同业务类型(例如，小掩蔽有效载荷大小和大未掩蔽有效载荷大小)，用于使性能和片上实际面积最大化。下表示出在SM之间通信以提供对DSMEM的访问的示例分组：

/>

注意，上表中的归约(reduction)操作和原子操作是密切相关的。原子是还返回值给请求者的归约操作。归约是“单向的”，因为数据被发送到存储器并且原子地与当前内容组合，但是除了ack之外，没有响应被发送回或至少被要求。原子和归约两者都做不能被中断的读取-操作-修改/写入(read-operate-modify/write)，因此“原子地”达到目标存储器中的结果。不同实施例中的硬件可支持所有此类直接SM到SM(SM-to-SM)通信功能或任何期望子集。

以下是对上表中的短写命令进行编码的示例分组：

struct xx2shmem_byte_wr_cmd_pd{

mem_datum_t wdat[4]；//有效负载

U016 we_；//写使能掩码

U004 mask；//指定高速缓存行的哪一部分被寻址

U001 upper_half；//指定掩码部分的哪一半被寻址

sm2sm_sm_gpc_id_fields dst_sm；//由网络用来路由到正确的目的地SM

sm2sm_sm_gpc_id_fields src_sm；//由目的地SM使用以将写ACK发送回源SM

U008 gpc_local_cga_id；//CAM查找目标SM

U005 cta_id_in_cga；//CAM查找目标SM

U011 cgamem_offset；//目的地SM的共享存储器中分配给这个CGA中的给定CTA的地址偏移量

U001 phase_id；

U015 barrier_addr；

U001 barrier_addr_valid；

}

Gpc_local_cga_id

上述分组格式包括U008字段“gpc_local_cga_id”。如上所述，每个GPC都有自己的CGA ID池，GPM在启动CGA时将这些数字中的一个分配给该CGA。然后，该指派的数字用作指向正被CGA中的各个CTA使用的DSMEM分布式存储器段的指针。在一个实施例中，“gpc_local_cga_id”还充当用于跟踪每个GPC_CGA的屏障状态的id。

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

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

DSMEM映射表

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

在一个实施例中，图19中所示的第二表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内的位置。

图19还示出了目标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通知该错误。

DSMEM合并写确认和优化的存储器屏障处理

如以上所讨论的，GPM连续地跟踪CGA中继续执行的CTA的数量。在一个实施例中，此计数还用于使GPM能够支持屏障。屏障例如用于出于任何原因而将CGA中的所有CTA同步。例如，CTA可以执行指示它们希望等待一个或更多个其他CTA产生的结果的指令。这样的数据依赖性在并发执行的线程中是常见的。然而，使用DSMEM的示例实施例为屏障提供特殊用例。

即使CGA的所有CTA被保证并发地执行，但这并不意味着所有CTA将恰好在同一时间立即启动。相反，GPM启动所有CTA需要一些时间量。正如一所大学的本科生在学期开始时需要一些时间来移入他们各自的宿舍并准备去他们的第一堂课一样，不同的CTA需要一些时间来分配给SM并为其分配DSMEM。但不像上午8点直到周一上午才开始上课的宿舍情况，首先要启动的(first-to-launch)CTA开始尽快尝试做他们的工作，这可以包括访问其他CTA的DSMEM段，这些DSMEM段尚未准备好并且没有任何DSMEM段分配给他们。因此，在该启动时间窗口期间，一个实施例在允许任何CTA访问任何其他CTA的任何DSMEM共享存储器之前要求每个CTA在屏障上门控。参见图20。这个硬件屏障实例(被示出为图20的左手铁路道口门)指示何时所有CTA都已经成功启动，从而使得CGA中的每个(每一个)CTA的DSMEM段现在可供SGA中的任何其他CTA访问。

在执行过程的另一端，CTA中的一些可以在其他CTA之前完成执行——但如果允许它们完全退出，则分配给它们的DSMEM段将被解除分配，并且因此将不再可用于仍然做工作(这可能取决于那些DSMEM段所包含的数据)的其他CTA。这和学期末的本科生一样，他们想在上次考试后马上离开，但他们对班组的贡献仍然没有改变。在这样的情况下，一个实施例使另一硬件屏障实例(图示为图20的右手铁路道口门)可用于对CTA退出进行门控，使得CTA不能实际退出(因此它们保持活动以便维持它们的DSMEM段分配并且对其服务任何附加请求)，直到所有CTA准备好退出。

一个示例实施例提供对ISA的增强，该ISA提供两个新的屏障指令，这两个新的屏障指令类似于传统同步线程指令，但是提供对新的CGA屏障的访问，该新的CGA屏障可以用于如上文所讨论的对DSMEM指令的访问和释放进行门控。这两个新指令是：

CGABAR.ARRIVE//信号到达屏障，但不阻塞

CGABAR.WAIT//现在阻塞，直到屏障清除。

在一个实施例中，这些指令不指定特定屏障，当指令执行时，硬件向CGA分配硬件屏障。

执行到达屏障指令首先使用软件将线程束中的所有线程同步，然后将线程束级到达消息触发到MPC。它还设置

Warp[WarpiD].CgaBarWait＝1。

如果CgaBarWait已经＝1，则到达屏障指令将停止发布。

CGA中的用于执行到达屏障指令的每个CTA将执行上述相同的步骤。

执行等待屏障指令是阻塞实际发生的地方，并且每个CTA将在屏障上阻塞，直到所有CTA到达并且MPC清除屏障。

因此，在示例实施例中，存储器屏障(MEMBAR)指令被用于确保存储器操作已经完成。在MEMBAR指令之前的所有读取和写入已经完成或至少具有在MEMBAR指令之后的读取和写入的执行之前可见的结果。参见例如USP9223578；USP8997103；USP9324175。例如，memoryBarrierNV()OpenGL着色语言(GLSL)操作连同“MEMBAR”汇编操作一起提供确保着色器线程内的读和写操作的适当排序的显式同步。当存储器屏障命令在执行中完成时，被调度用于在存储器屏障命令之前执行的存储器操作全部被保证已完成到一致性点。在一个实施例中，读取和写入可以不按顺序执行，并且当MEMBAR重要时(即，当后续读取或写入命令需要依赖于所有已经完成的先前读取和写入时)，MEMBAR被用来将它们按顺序放置。MEMBAR还可以用于计算抢占(CILP)以确保CGA中的所有CTA在针对上下文切换保存SM状态之前已经完成所有DSMEM访问，并且在允许新的DSMEM访问之前已经通过上下文恢复恢复了所有DSMEM。

在过去，MEMBAR是通过向存储器中的每个目的地发送消息并且从那些目的地中的每一个目的地接收到确认来实现的。这例如可以在诸如L2高速缓存的存储器域上施加遵循MEMBAR确认的任何操作将在之前出现的每个操作之后发生的规则。然而，在一个实施例中，不同的MEMBAR实现被用于DSMEM。

在一个实施例中，对DSMEM的每个未完成的请求被跟踪。当源SM向目标SM发送请求时，它保持跟踪它期望来自目标SM的响应。如果请求是LD(加载)命令，则源SM在等待某数量的读响应，即，源SM在等待它期望来自目标SM的所有数据。如果源SM发送ST(存储)命令，则目标SM将发送特别用于允许源SM跟踪未完成的ST命令的确认。在一个实施例中，SM甚至将在内部确认由一个CTA发送到运行在同一SM上的另一CTA的ST命令。由此，跟踪负载，直到接收到所有响应数据，并且跟踪存储，直到接收到确认。

为了提供这样的跟踪，每个SM保持未完成的事务计数器。在一个实施例中，未完成的事务计数器被合并，意味着每个源SM简单地对所有其他目标SM的未完成事务的总数进行计数(例如，而不是跟踪每个目标SM的未完成事务的数量)。但其他实现方式是可能的。

当源SM发送请求时，它递增其未完成事务计数器的计数。当源SM接收到确认(对于ST命令)或返回数据(对于LD命令)时，它递减计数。其他实现可使用两个不同的未完成事务计数器，一个用于ST命令，一个用于LD命令。

为了使总线带宽最小化，确认被合并而不是单独通过通信总线发送。在一个实施例中，充当ST命令的目标的每个SM跟踪它欠源SM的确认的数量。在这种情况下，目标SM可为发送它的ST命令的每个源SM维护单独的计数器。例如，当计数达到某值时，或者当总线空闲时，可以以合并的方式发送确认命令。当发送确认时，其包括到特定源SM的累积确认的计数，并且在接收到累积计数时，源SM可以将其未完成事务计数器递减在累积确认消息中接收的累积计数的量。

在一个实施例中，GPM保持expected_cta_arrival_count，其跟踪在GPM释放屏障之前多少CTA应到达屏障。在一个实施例中，GPM将expected_cta_arrival_count设置为仍然在运行的CTA的数量。

当遇到MEMBAR时，系统必须等待，直到所有SM的所有未完成事务计数器变为零，以确保所有未完成DSMEM访问请求已经完成。这涉及停止CGA的所有新的(在发出MEMBAR命令之后)DSMEM存储器访问命令，直到未完成事务计数器指示所有未完成的存储器访问请求已完成为止。然而，在一般情况下，每个SM可执行来自各种不同CGA的CTA，不仅仅是已针对其发布MEMBAR命令的CGA的CTA。不是停止所有存储器访问(例如，包括来自不同CGA中的CTA的存储器访问)，直到未完成事务计数器变为零，而是每个SM维持一组两个(多个)未完成事务计数器-例如，相位(phase)0计数器和相位1计数器。在遇到MEMBAR时，SM翻转相位以使用不同的未完成事务计数器来跟踪未完成的请求(所以如果它们先前使用相位0计数器，则它们开始使用相位1计数器，或者如果它们先前使用相位1计数器，则它们开始使用相位0计数器)。SM硬件因此维持两个未完成事务计数器状态——一个来自相位翻转之前，另一个来自相位翻转之后。这还意味着发送的每个存储器事务和每个对应的确认或数据响应识别相位(相位0或相位1)，因此对应的确认可以更新正确的相位未完成请求计数器。因此，在一个实施例中，所有未完成请求的记账(accounting)都是每相位的。

当考虑特定相位的未完成请求最终变为零时，在MEMBAR完成之前发出的所有ST命令以及在MEMBAR完成之前发出的所有LD命令也已经完成。在等待旧的相位计数变为零的同时，如果一个或更多个新的MEMBAR同时到来，则新的MEMBAR请求被停止和合并，直到旧的相位计数变为零。一旦旧的相位计数达到零，则硬件再次翻转相位以针对任何停止的MEMBAR重复该过程。这个过程可以无限地重复，随着新MEMBAR的到来，相位来回翻转。

在一个实施例中，这些未完成事务账目还在CTA退出时被使用，以便(例如，对于MPC)确定与CTA相关联的所有未完成DSMEM事务何时已经完成。

以上过程用于实现SYNC类型的访问。如2022年3月10日提交的名称为“硬件加速同步与异步事务支持(Hardware Accelerated Synchronization With AsynchronousTransaction Support)”的美国专利申请No.17/691,296中所述，复制以上内容以为ASYNC类型的访问提供附加的功能。注意，可以使用除了硬件同步屏障之外的机制来确保CTA不开始访问尚未被分配的DSMEM或者在DSMEM已经被解除分配之后不继续访问DSMEM。例如，可以使用存储器屏障或提供CTA间通信的任何其他通信机制。

使用DSMEM的CGA/CTA退出和错误处理协议

在一个实施例中，某些种类的错误不归因于程序计数器(PC)。通常，实施例将保留过去的PC的FIFO，并且可以将任何存储器错误与给定的线程束、线程和PC相关联。当确定不存在可归因于该PC的错误时，PC可从FIFO的末尾掉落。然而，在DSMEM事务的情况下，某些类型的错误在目标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，以使图19中所示的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。

示例全功能管线

图21A示出了定义包括多个SM的GPC的示例全功能管线。图21A示出各自组织/包括多个SM的多个CPC(CPC0、CPC1、CPC2、……CPCN)，所述SM中的每一者提供可由GPC中的任何其他SM访问的内部DSMEM存储器段。CPC内的SM各自连接至相应的总线(“tex2gnic”)，该总线将SM连接至GPCMMU存储器管理单元。GPCMMU将每个SM总线上的虚拟存储器地址转换成外部存储器层次结构(例如，L2高速缓存)物理存储器地址，并将具有物理地址的请求发送到GPCARB(在一个实施例中，到GPCARB中的总线对于每个SM保持分开且不同)。GPCARB仲裁来自不同SM的请求以通过总线“L2 REQ”(黑色数据路径)发送，以访问L2高速缓存存储器。这样的请求通过MXBAR传递到L2高速缓存和平台外部存储器层次结构。

在所示的示例中，GPC内部的新的层次结构GXBAR网络/交换机专用于提供替代的“SM2SM”输出路径(斑点数据路径)，该输出路径允许GPC中的任何SM与GPC中的任何其他SM通信。GXBAR交换机因此使得在一个SM中运行的CTA能够通过访问另一SM中的共享存储器与在另一SM中运行的另一CTA协作。

每个CPC包括GPCARB运算符，该运算符将CPC内的SM连接至GXBAR交换机。在一个实施例中，GPCARB块充当路由器以将L2存储器访问请求路由到MXBAR(黑色数据路径)并且将SM到SM请求路由到GXBAR(斑点数据路径)。此外，在此示例中，CPC内的所有SM与公共GPCARB通信，且GPC内的所有CPC内的所有SM与公共GXBAR交换机通信。因此，该示例中的CPC是共享公共交换机(GPCARB)和GPCMMU的TPC的集合。

在一个实施例中，SM到SM连接性局限于GPC层次结构。该定位实现了有助于最小化同步开销的低延迟。在一个实施例中，可供SM经由GXBAR彼此通信的带宽不同于可供SM经由MXBAR与L2高速缓存通信的带宽。此外，在该示例布置中，GPC内的SM到SM业务不使用可能需要与其他GPC共享的任何MXBAR带宽。因此L2高速缓存与GPC内的SM到SM通信隔离，使干扰L2高速缓存访问的SM到SM通信的可能性最小化，反之亦然(例如，一个GPC可以进行低延时SM到SM通信，而另一个GPC可以使L2高速缓存饱和)。

如图21A中所示，GXBAR(“SM2SM12”)的输出在此示例中经由GPCARB“EG”(出口)块(例如，EG0、EG1；EG2、EG3；EG4、EG5)被分配至GPC中的SM的不同输入端口(在此实施例中，三条“SM2SM4”路径中的每条路径分别通信至SM的CPC组……但是可以使用任何所希望的总线宽度)。在一个实施例中，这些SM2SM4路径(斑点的)与用于从L2高速缓存向SM的这些相同CPC组传送响应的“L2RESP”路径(以黑色示出)不同。在一个实施例中，每个SM具有用于SM到SM通信和用于L2RESP通信的相应不同端口。当然，在其他实现方式中，相同的SM端口可以用于混合的SM到SM和L2业务。

此互连网络可以包括一个或更多个交叉开关，其包括读交叉开关和/或写交叉开关。互连网络可以支持多个并行读取和写入操作，并且可以支持使用直接寻址访问存储在共享数据存储器中的数据和使用标签查找访问存储在Ll数据存储器中的数据。互连网络可以支持与共享存储器、L1高速缓存和/或寄存器中的库(bank)的数量相对应的同时读取和写入的数量。在一个示例中，互连网络可支持等于本地共享存储器中的存储器库的数量的同时读取和写入的数量。参见例如US 20210124582。

图21A-1示出了在图21A系统的上下文内的示例SM内部架构。在一个实施例中，SM的本地共享存储器与SM的处理核心紧密关联并且绑定到SM的处理核心。例如，SM可包括能够经由地址总线和数据总线访问物理上本地的或紧密连接的半导体存储器的多个处理核心。通常，作为SM的一部分的存储器仲裁电路用于仲裁(并且在一些情况下缓存或排队)本地共享存储器访问。本地存储器本身可以物理地位于半导体衬底的紧密接近SM的核心的一部分上——即，半导体衬底的区域被指定为“SM”，并且每个都包括处理核心、存储器仲裁电路和半导体存储器块。在一些示例中，本地半导体存储器块可以包括多个存储器库，并且存储器仲裁电路可以包括交叉开关电路，用于允许任何处理核心访问本地存储器的任何存储器库内的任何位置。在示例实施例中，扩展存储器仲裁电路以允许基于包含在从SM外部(即，由其他SM)提供的分组中的命令、地址和数据信息对本地共享存储器的读取、写入和读取-修改-写入(即，原子)访问。但是在其他实施例中，用于实现DSMEM的物理存储器可以位于与处理核心相同的位置以外的位置，其中提供互连(导线)以减小或最小化延时。

图21B示出了GXBAR的另一视图，该视图为SM间业务提供基于地址的路由。要注意的是，在图21B中，在左侧显示为源(“SRC”)SM的SM以及在右侧显示为目的地(“DEST”)SM的SM是相同的SM。图21B示出GXBAR由多个k x k交换机实现以避免协议死锁(例如，其中第一SM在等待来自第二SM的第一响应，并且第二SM同时在等待来自第一SM的第二响应)。在所示的示例中，k＝6以提供全连接的12x12交换机，但是任何期望的维度都是可能的。

在所示示例中，三个入口块“IG0”、“IG1”、“IG2”(IG代表“入口”)分别对应于如图21A中所示的CPC0、CPC1、CPC2的三个“GPCARB IG”块。类似地，图21B中的出口块EG0、EG1、EG2、EG3、EG4、EG5(如上所述，EG代表“出口”)对应于图21A中所示的GPCARB EG0、EG1；GPCARB EG2、EG3；以及GPCARB EG4、EG5。图21B因此示出了由源SM发起的SM到SM请求传递通过对应的GPCARB IG块，传递到GXBAR交换机，GXBAR交换机然后基于地址通过对应的EG块将请求路由到目的地或目标SM。在该基于地址的路由中，给定地址确定从源SM到目标SM的通过GXBAR的唯一路径。但是，从任何源SM到任何目标SM存在多条路径。在一个实施例中，图21A的GPCMMU块实现uTLB(微转换后备缓冲区)硬件，该uTLB硬件实现地址散列而不是虚拟到物理地址转换，以便跨GXBAR网络中的多条路径均匀地分布请求，从而最大化网络利用率并避免预占(camping)情形。

图21B架构提供鲁棒性，因为电线可以运行长距离(由于互连的SM的数量和硅上的GPC的物理尺寸)，总线编码方案用于针对电力完整性覆盖切换速率，并且如果网络上发现任何硅问题，可编程(例如，通过固件)“旋钮”使用小于全部网络带宽的某个部分(例如，通过关闭GX0或GX1的一半带宽)。还可提供网络带宽控制以将每SM的最大带宽(例如，SM可在网络上每时钟传送的数据量)限制到某一限制(例如，以提供可预测的性能)和/或允许SM机会性地在峰值可用网络带宽下操作。此架构还通过切断有缺陷的TPC或CPC来支持TPC/CPC的扫除。

图21C示出了以上讨论的路由的另一视图，详述了从源SM到GXBAR和MXBAR的路径。可见，在一个实施例中，DSMEM路径的一部分与到L2高速缓存的路径是共用的(即，路径的一部分在SM到SM事务与到全局存储器的事务之间共享)。该公共路径定义或包括以下各项：

SM→TPCARB→GPCMMU uTLB→GPCARB-IG

用于两种不同类型的业务的分支点是图底部的GPCARB IG，其将全局存储器业务路由到MXBAR并且将SM到SM通信路由到GXBAR。使用公共路径节省了面积并重用线路/逻辑。然而，使用完全不同的路径的其他实现方式也是可能的。

图21D示出了SM到SM业务的示例路由的又一个视图。该图的左侧示出了来自请求SM的请求。该请求穿过TPCARB，该TPCARB在同一TPC中的多个SM之间仲裁。如以上所讨论的，uTLB不在虚拟地址与物理地址之间转换，而是用于散列该地址以增加网络利用率并防止预占。GPCARB在同一GPC内的不同SM之间仲裁。仲裁器然后通过GXBAR发送请求，如上所述，GXBAR提供互连，源SM可通过该互连与在相同CGA内运行CTA的任何目标SM通信。如图的右侧所示，GXBAR将请求通信至正确的GPCARB EG(出口)。在该实施例中，在SM到SM响应网络和全局存储器(GNIC)响应网络之间共享GPCARB EG。图21D进一步示出从目标SM到源SM的响应，该响应采用与源SM所使用的传送原始请求到目标SM的路径非常相似的路径(请求和响应在一个实施例中可以采用通过GCBAR的不同路由路径)。

图22示出了共享存储器请求/响应处理的另一视图，并且图23A-23B是图22的示例架构执行的操作的流程图。如在这些图中可以看到的，源SM核心产生DSMEM访问请求，该DSMEM访问请求包括用于目标SM的共享存储器的地址/存储器槽(gpc_local_cga_id和sm_id)(使用如上所述的查找表确定的)并且将其发送到互连(图23A，框6002、6004)。互连将请求进行排队并传送到目标SM(图23A，框6006、6008)。目标SM接收、排队、仲裁和检查该请求以找到DSMEM基数和大小，如上所述(图23A，框6010、6012)。目标SM的共享存储器(“SHM”)处理(经受如上所述的屏障阻塞)该请求(图23，框6014)。在目标SM处的共享存储器操作完成时，目标SM硬件对共享存储器输出对应的读数据(如果读请求)进行排队，或者目标SM硬件生成确认消息(如果例如写请求)(图23A，框6016)。目标SM硬件经由Tex2genic向互连发送读响应或合并的写确认消息(图23A，框6018)。互连对读响应或合并的确认进行排队并传送到源SM(图23B，框6020、6022)。源SM接收该消息并对其进行排队(图23B，框6024)，并使用如上所述的跟踪和事务计数器来执行信用检查以将接收到的响应与请求相匹配(图23B，框6026)。如上所述，陷阱处理程序和其他机制处理不基于源SM的程序计数器的错误条件。如果适用，则源SM硬件响应于该核心的读请求将来自读响应的数据呈现给该核心(图23B，框6028)。

因此，在一个实施例中，由GPU的一起工作的许多不同部分支持/提供/启用SM到SMDSMEM共享存储器访问；例如，源SM核心/本地存储器/硬件、互连、目标SM核心/本地存储器/硬件。在一个示例中，互连、网络或路由器/交换机可在SM之间路由DSMEM访问消息，但SM内的功能本身可使一个SM能够通过互连寻址另一SM的共享存储器，并且其他硬件或软件机制可提供附加支持以实现恰当或高效的SM到SM DSMEM访问、错误处理等。例如，在一个实施例中，源SM基于分段的地址确定目标SM和存储器槽，并且基于SM维护的(CAM)路由表中的查找提供物理SM id，然后选择正确的分组类型，以使互连知道这是SM到SM事务。同时，目标SM可执行地址边界/有效性检查，对其共享存储器执行所请求的操作，和合并写确认以减少互连上的业务。CGACPC内的附加硬件或软件支持可支持CGA并发启动(以确保跨执行CGA的SM分布的DSMEM可用)、CGA_id分配和回收协议、有序的CGA分解和DSMEM刷新和其他功能。不同实施例可在互连、SM或SM的各部分、其他硬件支持机制、其他软件支持机制、系统的其他组件等之间不同地分配这些分布式功能。

尽管本文中描述的共享存储器布置由CGA层次结构和关联硬件支持所提供的并发性保证实现，但它们不限于此。具体地，任何其他布置、协议或硬件支持可以用于确保在任何类型的处理核心(不限于“流式多处理器”)的集合上执行的执行线程组或线程组的执行并发性，每个处理核心都具有关联的本地存储器，这提供可以在执行线程组或线程组之间共享的本地存储器分配的分布式阵列。关于提供并发性保证的不同方式，参见例如布雷希尔(Breshears)的并发性技术：写入并行应用的线程猴指南(The Art of Concurrency：AThread Monkey’s Guide to Writing Parallel Applications)(O’Reilly 2009)和拉杰瓦尔(Rajwar)等人的“推测性锁省略：实现高度并发的多线程执行(Speculative lockelision:enabling highly concurrent multithreaded execution)”，微结构MiCRO-34的第34届ACM/IEEE国际研讨会论文集(2001年12月1日至5日)。如文献显而易见的，处理布置(诸如图形处理单元)内的这样的处理核心组可以由不同名称(诸如执行单元、计算单元、流式多处理器或其他术语)引用。这些处理核心组可为多线程的，即，支持由例如多个内核或着色器实例组成的多个并发线程或线程组。组中的处理核心可以具有基于例如包括相同或不同精度的相同或不同计算逻辑的相同或不同ALU(算术逻辑单元)硬件的不同技术能力(例如，组中的一些处理核心可以包括用于执行基于整数的计算的ALU硬件，组中的其他处理核心可以包括用于执行浮点计算的ALU硬件，组中的其他处理核心还可以包括用于执行张量或矩阵计算的ALU硬件等)。组中的处理核心可以基于SIMT(单指令、多线程)执行模型和/或SIMD(单指令、多数据)执行模型来执行线程指令。该组中的处理核心可各自具有对相同资源中的至少一些的直接访问，诸如执行流、可直接寻址的本地存储器、高速缓存存储器、寄存器组、地址、数据流路径等。它们可一起形成共享共同指令高速缓存和共同指令解码器的并发处理的单个单元或集合，所述共同指令解码器使用共同程序计数器来检索和解码指令，同时例如具有允许处理核心使用相同或不同数据来执行/计算的相同或不同栈指针。

由此，虽然CGA构造对于保证跨SM的并发性是有用的，但用于保证并发性的其他技术可代替地或组合地用于支持DSMEM。例如，一些实施例可以使用软件布置(如协作组API)来布置以实现并发性，或者还可使用其他技术来提供或保证GPU硬件的同一相关硬件域或分区内的并发性(例如，利用分布式共享存储器的所有线程不仅并发地运行，而且可以在SM上启动和发现，所有这些SM都在特定硬件域(诸如被称为GPC的GPU的子部分)内(例如，因为各个线程可以通过查询已经在哪个GPC上启动了线程来测试))。虽然此类其他技术是可能的，使得DSMEM不限于CGA特定实现，但CGA层次结构提供良好支持DSMEM的效率和确定性方面的某些优点。此外，DSMEM可用于支持并发线程块，而不管如何提供并发性以允许并发线程块访问高速半导体存储器的块，所述块可跨GPU分布或分散以提供在任何数目的不同处理核心集合或分组之间共享的统一逻辑存储器。这样的分布式共享存储器为位于存储器块附近的处理核心提供了非常低的延时存储器访问，并且还为更远的处理核心提供了以不干扰处理核心对诸如由L2高速缓存支持的主存储器或全局存储器的访问的方式和使用不干扰处理核心对主存储器或全局存储器的访问的互连来访问存储器块的方式。

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

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

Claims

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

第一半导体存储器块，其与执行协作组阵列CGA的第一协作线程阵列CTA的至少一个关联的第一处理核心位于同一位置；

第二半导体存储器块，其与执行所述CGA的第二CTA的至少一个关联的第二处理核心位于同一位置；以及

互连，其可操作地连接所述第一处理核心与所述第二处理核心，

其中所述第一CTA被实现为能够经由所述互连访问所述第二半导体存储器块。

2.根据权利要求1所述的处理系统，其中在将所述第二半导体存储器块分配给所述第二CTA时，所述处理系统控制所述第一CTA对所述第二半导体存储器块的访问。

3.根据权利要求1所述的处理系统，其中所述处理系统被配置为保证所述第一CTA和所述第二CTA并发地运行。

4.根据权利要求1所述的处理系统，其中所述处理核心包括流式多处理器。

5.根据权利要求1所述的处理系统，其中所述处理系统被配置为使得所述第一CTA能够发出用于访问所述第二半导体存储器块的读命令、写命令和原子命令。

6.根据权利要求1所述的处理系统，其中所述互连包括交换机或网络。

7.根据权利要求1所述的处理系统，其中所述互连使得能够在处理核心与分配给作为同一CGA的一部分的CTA的所有共享存储器之间进行直接通信。

8.根据权利要求1所述的处理系统，其中所述处理系统被配置为按照执行CGA基于唯一标识符提供分段寻址。

9.根据权利要求1所述的处理系统，其中所述处理系统被配置为支持硬件驻留的CGA标识符分配和回收协议。

10.根据权利要求1所述的处理系统，其中每个所述处理器都存储用于实现来自远程处理器的DSMEM寻址的关联结构。

11.根据权利要求1所述的处理系统，其中所述互连包括低延时的处理器到处理器通信网络。

12.根据权利要求1所述的处理系统，其中所述处理系统合并写确认以用于跨互联的通信。

13.根据权利要求1所述的处理系统，其中所述处理系统包括用于在CGA创建时以及在CGA上下文切换期间同步存储器块使用的硬件屏障。

14.根据权利要求13所述的处理系统，其中所述硬件屏障确保在任何CTA通过所述互连引用所述存储器块之前，所述存储器块在CGA的所有CTA中可用。

15.根据权利要求13所述的处理系统，其中所述硬件屏障确保CGA中的所有CTA在状态保存之前已完成所有访问，并且在状态恢复期间已恢复所有DSMEM。

16.根据权利要求1所述的处理系统，其中所述处理系统进一步被配置为向DSMEM提供CGA/CTA退出和错误处理协议。

17.一种操作处理系统的方法，包括：

在第一处理核心上启动第一线程块；

向所述第一线程块分配所述第一处理核心本地的第一半导体存储器块；

在第二处理核心上启动第二线程块；

向所述第二线程块分配所述第二处理核心本地的第二半导体存储器块；

在第三处理核心上启动第三线程块；

向所述第三线程块分配所述第三处理核心本地的第三半导体存储器块；以及

实现所述第一线程块、所述第二线程块和所述第三线程块中的每一者能够访问所分配的第一半导体存储器块、所述第二半导体存储器块和所述第三半导体存储器块中的每一者。

18.根据权利要求17所述的方法，其中所述处理系统被设置在基板上，并且所述第一存储器块、所述第二存储器块和所述第三存储器块被设置在所述基板的不同部分上并彼此分离。

19.根据权利要求17所述的方法，还包括：使用硬件屏障以确保在允许由所述第一线程块、所述第二线程块和所述第三线程块中的任何线程块访问之前，分配所述第一半导体存储器块、所述第二半导体存储器块和所述第三半导体存储器块。

20.根据权利要求17所述的方法，其中所述实现包括在所述第一处理核心、所述第二处理核心和所述第三处理核心之间提供通信。

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

硬件工作分配器电路，其在多个处理器上并发地启动线程组的集合；以及

所述多个处理器之间的通信接口，其实现每个线程组能够访问被分配给所述集合中的任何其他线程组的存储器。

22.根据权利要求21所述的处理系统，其中所述通信接口包括：

所述多个处理器之间的通信链路，以及

由所述处理器执行的处理器间消息传递功能。

23.根据权利要求21所述的处理系统，其中所述通信接口禁止不在所述集合中的线程组访问被分配给所述集合中的任何其他线程组的存储器部分。

24.根据权利要求21所述的处理系统，其中所述多个处理器包括流式多处理器。

25.根据权利要求21所述的处理系统，其中所述多个处理器每个都包括一组处理核心以及所述组中的每个处理核心本地的并且能够由其直接寻址的存储器。