下载掌阅APP，畅读海量书库
立即打开
线程束是GPU执行的基本单元，同一线程束中的所有线程同时执行相同的指令。可以从逻辑视角（Logic View）和硬件视角（Hardware View）来理解线程束和线程块的执行过程，如图5-7所示。在核函数启动时，程序员可以指定线程块中包含的线程数量，当指定的线程数量超过waveSize时（waveSize=64），曦云架构GPU会将这些线程拆分成多个线程束，并在同一个AP上执行。不同的线程束可能会并行执行，也可能串行执行，这取决于GPU的调度策略和GPU的负载情况。
图5-7 线程束和线程块执行过程的逻辑视角和硬件视角
例如，一个包含200个线程的线程块在曦云架构GPU中会被分配到4个线程束上执行，每个线程束实际执行的线程数量分别是64、64、64、8。由此可见，为了提高GPU的计算效率，线程数量应该尽量设置为waveSize的整数倍。
曦云架构GPU使用SIMT架构。在同一线程束中的所有线程同时执行相同的指令，但是不同的线程可能有不同的指令执行路径。为了使线程束中的每个线程从同一个指令地址开始执行且可以有不同的指令执行路径，GPU引入了线程束分化机制。
示例代码5-2给出了一段包含分支的核函数代码，线程束中的线程有两种指令执行路径：其中有32个线程的条件判断结果为true，剩余32个线程的条件判断结果为false，此时GPU会串行地执行每条指令分支，禁用不执行此条指令的线程。通过这种方式，曦云架构GPU能够执行带有分支、循环的核函数，只是由于执行过程中部分线程被禁用了，线程束分化会降低程序的性能。编程时应该避免出现核函数存在大量分支的情况。
当发生线程束分化时，每个线程的执行情况如图5-8所示，一个线程束中64个线程在线程束中分化时可能存在的四种不同状态，分别是
● coherent code：64个线程都在执行相同的代码。
● if clause：64个线程中满足if条件的在执行相同的代码。
● then clause：64个线程中不满足if条件的在执行相同的代码。
● stall execution：闲置的线程，它们在等待线程束中其他线程的执行代码结束，进而导致GPU运算资源的浪费。
图5-8 线程束分化示意图
线程束主要使用的资源包括指令地址计数器、寄存器（Register）、私有内存（Private Memory）（可选）、工作组共享内存（Work-group Shared Memory）（可选）、全局内存/常量内存（Global Memory/Constant Memory）等，如图5-9所示。
图5-9 线程束可用资源示意图
每个线程都有自己的指令地址计数器，同时由自己的寄存器来存放局部变量、函数入参等。由于一个AP上的寄存器总数是有限的，因此当寄存器的数量无法满足线程束执行的需要时，会在HBM上划分一块内存来存储局部变量等数据，这块内存被称为私有内存（Private Memory）。在一个线程块内部，可以通过工作组共享内存（Work-group Shared Memory）在线程间共享数据，工作组共享内存位于GPU片上，其访问速度要比私有内存快。
线程束中的所有线程从同一个指令地址开始执行。当一个线程网格执行完成时，寄存器、私有内存、工作组共享内存会被自动回收，而程序员申请的全局内存需要程序员自己回收。
和私有内存一样，常量内存（Constant Memory）也从HBM上划分一块内存用于MXMACA程序的常量内存。除常量内存和私有内存占用的部分，HBM剩余的部分全部用作MXMACA程序的全局内存（Global Memory），也就是GPU全局内存。
一条指令从开始执行到完成执行的时钟周期（Clock Cycle）被称为延迟。为了尽可能地提升计算吞吐量，GPU尽可能地使每个时钟周期都有线程束在执行，这样就隐藏了指令的延迟。具体来说，当一个线程束由于访问内存而被阻塞时，AP会马上转为执行其他准备好的线程束，直到其被内存访问阻塞，随后返回上一个线程束。
曦云架构GPU中程序执行的基本单位是线程束，单个AP的占用率按式（5-1）来计算。
MXMACA运行时库提供了函数mcGetDeviceProperties供程序员查询当前硬件下可以并行执行的线程束的最大数量，其定义如下。
其中，*prop是一个指向mcDeviceProp结构体的指针，该结构体被用于存储设备属性的信息，可以通过该信息获得waveSize和maxThreadsPerMultiProcessor这两个设备属性的值（设备属性的描述参见表5-1）。可以并行执行的线程束的最大数量等于maxThreadsPerMultiProcessor/waveSize。
MXMACA运行时库也提供了几个占用率计算函数，这些函数特别有用，它们可以帮助程序员了解并行度级别和线程块在GPU上的调度情况，或者启发式地计算实现最大AP级占用率的执行配置（优化核函数参数和调整线程块大小）。这有助于程序员优化代码，提高程序的执行效率和GPU利用率。相关的函数有以下几个。
● 函数mcOccupancyMaxActiveBlocksPerMultiprocessor用于预测给定的核函数在每个AP上的活动线程块（Active Block）数量。这个预测是基于核函数的线程块大小和共享内存的使用情况来实现的。该函数的返回值乘以每个块的线程束数可以得出单个AP上的并行线程束数，将并行线程束数除以一个AP支持的最大线程束数，可以计算出单个AP的占用率。
● 函数mcOccupancyMaxPotentialBlockSize用于计算在固定共享内存配置下的最大潜在线程块大小。这个函数可以帮助程序员了解线程块的并行度潜力，并指导他们如何调整线程块大小以获得更好的性能。
● 函数mcOccupancyMaxPotentialBlockSizeVariableSMem则考虑了可变共享内存的使用情况。它提供了更准确的预测，适用于那些共享内存使用量随线程块大小变化的核函数。这个函数可以帮助程序员了解在共享内存使用量可变的情况下，如何调整线程块大小以获得最大AP级占用率。
示例代码5-3提供了一种计算核函数MyKernel占用率的方法。首先，用函数mcGetDeviceProperties获取核函数MyKernel所在的GPU硬件单元的最大线程束数（见示例代码5-3中maxWaves的计算）。然后，用函数mcOccupancyMaxActiveBlocksPerMultiprocessor计算核函数MyKernel实际并行执行的线程束数目（见示例代码5-3中activeWaves的计算）。最后，用activeWaves除以maxWaves并乘以100，即可获得核函数MyKernel的占用率。
MXMACA提供了一些同步原语，用于完成GPU与CPU间、GPU内部的同步，以确保GPU编程中不同组件间的协调执行。MXMACA中的同步机制主要有以下三种。
（1）系统级同步。这种同步机制用于完成CPU和GPU之间的同步，是确保数据在CPU和GPU间一致性的重要手段。在MXMACA编程中，应用程序既可以通过调用函数mcStreamSynchronize来等待特定流中所有先前排队的工作完成，也可以通过调用函数mcDeviceSynchronize来等待所有先前排队的设备工作完成（包括该设备所有已创建的流中所有的工作）。此外，MXMACA还提供了函数mcEventSynchronize，应用程序可以通过该函数来阻塞调用的线程，直到特定事件完成。这意味着所有在该事件记录之前启动的工作（如内核执行、内存复制等）都将完成。
（2）线程块级/线程束级同步。在MXMACA编程中，线程块是一组同时执行的线程集合，而线程束是一组同时执行的64个线程。线程块级同步可以通过调用函数__syncthreads来实现，这个函数在核函数中由每个线程调用，直到所有线程块内的线程都到达该同步点。线程束内同步函数__syncwave支持通过函数入参来传递一个掩码（mask），从而指定哪些线程需要被同步。在默认情况下，函数__syncwave会同步整个线程束（mask为0xffffffff）。这种同步确保了同一个线程块或线程束内的所有线程在继续执行下一条指令之前完成当前指令。
（3）用户自定义的线程协作组级同步。在MXMACA编程中，应用程序可以通过原子操作、锁机制或使用共享内存等经典方法来实现用户自定义的同步。此外，MXMACA还提供了协作组编程API，以方便程序员根据特定的业务需求和程序设计需要，自定义线程协作组的粒度。
一些高效的并行算法往往需要线程协作（Threads Cooperate）和共享数据（Share Data）来完成复杂业务的集合计算（Collective Computation）。要共享数据，线程间必然涉及同步，而共享的粒度因算法而异，因此，线程间的同步应尽量灵活。例如，程序员可以显式地指定线程间同步，以确保程序的安全性、可维护性和模块化设计。基于该思想，MXMACA编程模型支持协作组（Cooperative Group）的概念，以允许程序员开发核函数时动态地组织线程组来满足这些需求，如图5-10所示。
图5-10 协作组支持灵活线程组的显式同步
协作组编程模型提供了MXMACA线程块内和跨线程块的同步模式，并提供了一套函数用于定义、划分和同步线程组，这些函数使程序员能更方便地管理线程的执行和同步。程序员用这些函数可以在MXMACA中启用新的协作并行模式，例如生产者-消费者并行，甚至跨整个线程网格的全局同步（包括一个或多个GPU）。这种模式允许线程之间协同工作，通常被用于数据流的处理和计算任务的划分。通过线程间的同步，可以确保数据的正确传输和处理，从而提高了程序的效率和可靠性。
将分组表示为一级程序对象可以改进软件的组合：集合函数可以采用显式参数表示参与线程的组。考虑一个库函数，它对调用者提出了要求。显式分组将这些需求显式化，从而降低了误用库函数的可能性。显式分组和同步有助于使代码不那么脆弱，减少对编译器优化的限制，并提高向前兼容性。
协作组编程模型在MXMACA中是一个重要的概念，它由以下元素组成。
● 表示协作线程组的数据类型。协作组编程模型提供了一种数据类型，用于表示一组协作执行的线程。这种数据类型通常被用于定义线程组，以便管理和同步。
● 与协作组配套的启动核函数管理机制。MXMACA提供了一套启动核函数的机制，用于创建和管理线程组。这套机制允许程序员获取与隐式定义的线程组相关的信息，例如线程组的属性、成员等。
● 用于将现有组划分为新组的集合操作。协作组编程模型提供了一组集合操作，允许程序员将现有的线程组划分为更小的子组。这种划分有助于更好地组织和管理线程，以便进行更细粒度的同步和数据传输。
● 用于数据移动和修改的集合算法。这些算法用于在线程组之间移动和修改数据，如memcpy_async、reduce、scan。它们通常由硬件加速实现，并提供了一套API供程序员使用。这些算法对于实现高效的并行计算至关重要，特别是处理大规模数据集的情况。
● 同步组内所有线程的操作。协作组编程模型提供了同步机制，用于确保线程组内的所有线程在继续执行之前达到某个同步点。这种同步机制有助于避免竞态条件和数据不一致的问题，并确保线程之间的正确协作。
● 检查线程组属性的操作。程序员可以用这些操作来检查线程组的属性（如线程组的成员、状态等），这些信息对于调试和优化并行程序非常重要。
● 程序员可见的群组集合操作。这些操作通常是由硬件加速的集合操作，并对程序员可见。它们提供了一种高效的方式来执行复杂的并行计算任务，同时隐藏了底层硬件的细节。通过这些操作，程序员可以更专注于编写并行逻辑，而不需要关心底层的实现细节。
接下来将深入探讨MXMACA为协作组编程所提供的各种数据类型和API。进一步了解如何利用线程组进行集合操作，并展示线程组应用的实例，以便更好地理解和应用这些功能。
1.协作组API
协作组API被用于定义和同步MXMACA程序中的线程组。要使用协作组API，请包含头文件cooperative_groups.h。
协作组类型和API是在cooperative_groups C++命名空间中定义的，所以，可以用cooperative_groups：：作为所有名称和函数的前缀，或者用using指令加载命名空间或其类型。
通常会给命名空间定义一个别名，下面的示例中会用到别名是“cg”的命名空间。
包含任何块内协作组功能的代码都可以使用mxcc以正常方式进行编译。
协作组中的基本数据类型是thread_group，它是一组线程的句柄。该句柄只能由它所代表的组的成员访问。线程组公开一个简单的API。可以使用函数size获取线程组的大小（线程总数）。
要在组中查找和调用线程（介于0和size()-1之间）的索引，请使用函数thread_rank。
最后，可以使用函数is_valid检查分组的有效性。
2.线程组集合操作
线程组提供了在组中的所有线程之间执行集合操作（Collective Operation）的能力。集合操作是需要在一组指定的线程之间进行同步或通信的操作。由于需要同步，每个被标识为参与集合操作的线程都必须对该集合操作进行匹配调用。最简单的集合操作就是一个屏障，不传输任何数据，只是同步组中的线程。
MXMACA程序员可以通过调用函数sync或cooperative_groups：：sync来同步线程组，此时线程组内的所有线程之间将会进行一个同步，如图5-10所示。这就类似于函数__synthreads以一个线程块为单位进行同步，函数__syncwave以一个线程束为单位进行同步，而函数g.sync以指定的线程组为单位进行同步。
接下来讨论如何在MXMACA程序中创建线程组。协作组引入了一种新的数据类型thread_block，按照如下的方法初始化得到的thread_block实例是MXMACA线程块中线程组的句柄。
与其他MXMACA程序一样，执行该代码行程序的每个线程都有自己的变量块实例。MXMACA内置变量blockIdx值相同的线程属于同一个线程块组。函数__syncthreads的作用就是同步一个线程块组（Thread_block Group）。下面的几行代码都是同步操作，只是同步的颗粒度不一样。
数据类型thread_block扩展了thread_group API，提供了以下两种thread_block的线程索引方法。
这里的坐标变量group_index是线程网格内的 N 维线程块组索引，可类比第3.4节中线程块在线程网格内的索引（坐标变量blockIdx）。坐标变量thread_index是线程块组内的 N 维线程索引，可类比第3.4节中线程在线程块内的索引（坐标变量threadIdx）。
3.线程组示例
下面介绍一个简单的核函数sum_kernel_block。首先，核函数sum_kernel_block调用核函数thread_sum并行计算许多的部分和，其中的每个线程在输入向量数组里跨过数组下标间距大小为blockDim.x×gridDim.x的两个数据先计算一部分和（这里使用向量化加载以获得更高的内存访问效率）。然后，核函数sum_kernel_block使用数据类型thread_block的线程组执行协作求和，使用核函数reduce_sum来计算输入数组中所有值的总和（每个线程块组负责本组内输入数组的求和）。最后，调用函数atomicAdd来完成各个线程块组计算结果的求和。
完整的线程组示例代码见示例代码5-4。核函数sum_kernel_block实现了一个并行计算部分和的核函数，并使用协作组编程模型来执行线程块内的求和操作。
在示例代码5-4中，核函数sum_kernel_block利用了GPU的并行处理能力，通过向量化加载、并行计算部分和以及协作求和，实现了高效的数组求和操作。
● 并行计算部分和。sum_kernel_block首先调用另一个核函数thread_sum来并行计算部分和。每个线程负责计算数组的一个固定部分，以实现向量化加载并提高内存访问效率。通过使用一定大小（blockDim.x×gridDim.x）的输入数组下标跨度，每个线程能够计算数组的一个连续片段的部分和，这可以更高效地利用GPU的并行处理能力，同时减少线程间的数据传输开销。
● 协作求和。完成部分和的计算后，sum_kernel_block使用一个线程组来执行协作求和。这是通过调用核函数reduce_sum来完成的，该核函数负责在每个线程块组内计算所有值的总和。协作求和的目的是将每个线程块组内的所有部分和合并为一个总和。这样可以减少所需的内存带宽，并利用GPU的并行处理能力来加速求和操作。
● 完成各个线程块组间的求和。sum_kernel_block使用内置的函数atomicAdd来完成各个线程块组间的求和操作。函数atomicAdd是一个原子操作，通常用于在多线程环境中安全地更新单个值，在这里，它被用于将线程块组计算出的总和添加到全局总和中。通过原子操作，可以确保在多线程环境中对全局总和的更新是线程安全的，从而避免了竞态条件和数据不一致的问题。
将示例代码5-4保存到文件syncWithCooperativeGroups.cpp中，然后用mxcc编译它。
随着曦云架构GPU的不断发展，让面向早期GPU编写的GPU程序能够充分利用最新GPU的计算资源是非常重要的。在当前的MXMACA执行模型中，一个线程网格可以包含若干个线程块。一个线程块的所有线程必须在同一个AP上执行，而不同线程块的线程可以在不同的AP上执行。如果在设计GPU程序时设定合适的线程块数量和线程数量，程序可以更好地兼容未来更高性能的GPU。
例如，如果核函数启动时配置的线程块数量为2，在有4个AP的GPU上，这个线程网格最多使用GPU中的2个AP，其余2个AP一直空闲。如果将线程块数量配置为4或者更大的值，就可以使用GPU的全部AP资源。不过，更多的线程块意味着会使用更多的GPU共享内存，如果线程块中的线程数量小于waveSize，计算资源就会被浪费。因此，程序员在配置线程块数量和线程块中的线程数量时应充分评估后进行平衡。
从扩展性方面来说，MXMACA程序员可以更多地采用宏而不是常数来定义所使用的资源，并更多地使用系统资源API来获取可用资源数目，通过这样规划程序的资源以适配未来的GPU。
CPU线程通常执行复杂的业务功能，程序相对复杂，可以使用大量的操作系统API，有完整的同步机制。GPU线程的主要任务是完成计算，代码逻辑相对简单，主要支持线程块内和线程束内的同步。一个GPU线程占用的内存远少于CPU线程。
另外，CPU线程由操作系统完成调度，线程是调度的基本单位，线程间切换的开销较大。GPU线程主要由硬件完成调度，线程束是调度的基本单位，线程切换的开销很小。同时，由于GPU的核心数量远远超过CPU，在GPU上能够同时执行的线程数量也远多于CPU。
为了让你继续提高技能，这里有几件事你可以尝试。
● 在核函数中使用函数printf进行实验。尝试打印出部分或所有线程的threadIdx.x和blockIdx.x的值。它们是按顺序打印的吗？为什么呢？
● 在核函数中打印threadIdx.y、threadIdx.z或blockIdx.y的值（同样适用于blockDim和gridDim）。这些为什么存在？如何让它们采用0以外的值？