CUDA学习（五十四）

COOPERATIVE GROUPS：
介绍：
COOPERATIVE GROUPS：是对CUDA 9中引入的CUDA编程模型的扩展，用于组织通信线程组。 合作组允许开发人员表达线程通信的粒度，帮助他们表达更丰富，更高效的并行分解。
历史上，CUDA编程模型为同步协作线程提供了一个简单的构造：跨越线程块的所有线程的屏障，如使用__syncthreads（）内部函数实现的。但是，程序员希望以其他粒度定义和同步线程组，以“集体”组范围功能接口的形式实现更高的性能，设计灵活性和软件重用。为了表达更广泛的并行交互模式，许多面向性能的程序员采用编写自己的临时和不安全基元来同步单个warp内的线程，或者跨单个GPU上运行的线程块集。虽然实现的性能改进常常是有价值的，但这导致了越来越多的脆弱代码集合，随着时间的推移以及GPU代之间的编写，调优和维护成本很高。合作组通过提供安全且面向未来的机制来实现高性能代码来解决这个问题。
协作组编程模型扩展描述了CUDA线程块内部和跨CUDA线程块的同步模式。 它提供了应用程序定义自己的线程组以及同步它们的接口的方法。 它还提供了新的启动API，可以强制实施某些限制，因此可以保证同步工作。 这些原语支持CUDA内部的新合作并行模式，包括生产者 - 消费者并行(producer-consumer parallelism)，机会并行(opportunistic parallelism)以及跨整个网格的全局同步。
由于集合函数可以接收表示参与线程组的显式对象，因此将组表示为一等程序对象可以改进软件组合。 这个对象也使得程序员的意图是明确的，这消除了不合理的架构假设，导致代码变得脆弱，编译器优化的不受限制的限制以及与新一代GPU的更好的兼容性。
协作组编程模型由以下元素组成：

• 用于表示合作线程组的数据类型;
• 获取由CUDA启动API定义的内在组的操作（例如，线程块）;
• 将现有组划分为新组的操作;
• 同步给定组的屏障操作
• 以及检查团队属性以及团队特定集体的操作。

块内组:
在本节中，我们将介绍可用于在可同步和协作的线程块内创建线程组的功能。 请注意，使用协作组来跨线程块或设备进行同步需要一些其他注意事项，如本附录后面所述。
合作组织需要CUDA 9.0或更高版本。 要使用协作组，请包括头文件：

#include <cooperative_groups.h>

并使用命名空间：

using namespace cooperative_groups;

然后可以使用nvcc以正常方式编译包含任何块内协作组功能的代码。