CUDA学习（五十五）

线程组和线程块：
任何CUDA程序员都已经熟悉某个线程组：线程块。 协作组扩展引入了一个新的数据类型thread_block，以在内核中明确表示这个概念。 该组可以初始化如下：

thread_block g = this_thread_block();

thread_block数据类型是从更通用的thread_group数据类型派生的，可用于表示更广泛的组类。 thread_group提供以下功能：

void sync(); // Synchronize the threads in the group
unsigned size(); // Total number of threads in the group
unsigned thread_rank(); // Rank of the calling thread within [0, size]
bool is_valid(); // Whether the group violated any APIconstraints

而thread_block提供以下额外的特定于块的功能：

dim3 group_index(); // 3-dimensional block index within the grid
dim3 thread_index(); // 3-dimensional thread index within the block

例如，如果组g如上所述被初始化，则：

g.sync();

将同步块中的所有线程（即相当于__syncthreads（））。 请注意，组中的所有线程都必须参与集体操作，否则行为未定义。
平铺分区:
tiled_partition（）函数可用于将线程块分解为多个较小的协作线程组。 例如，如果我们首先创建一个包含块中所有线程的组:

thread_block wholeBlock = this_thread_block();

那么我们可以将它分成更小的组，每个大小为32个线程：

thread_group tile32 = tiled_partition(wholeBlock, 32);

此外，我们可以将这些组中的每一个划分为更小的组，每个组都有4个线程：

thread_group tile4 = tiled_partition(tile32, 4);

例如，如果我们要包含以下代码行：

if (tile4.thread_rank()==0) printf(“Hello from tile4 rank 0\n”);

那么语句将被块中的每个第四个线程打印：每个tile4组中的等级为0的线程，这些线程对应于全部组中等级为0,4,8,12 ...的那些线程。
请注意，目前仅支持2次幂且不大于32的大小。
Thread Block Tiles：
tiled_partition函数的替代模板版本是可用的，其中使用模板参数来指定瓦片的大小：在编译时已知这有助于更优化执行的可能性。 类似于上一节中的内容，以下代码将创建两组分别为32和4的平铺组：

thread_block_tile<32> tile32 = tiled_partition<32>(this_thread_block());
thread_block_tile<4> tile4 = tiled_partition<4>(this_thread_block());

请注意，这里使用了thread_block_tile模板数据结构，并且组的大小作为模板参数而不是参数传递给tiled_partition调用。
线程块Tiles还公开了其他功能，如下所示：

.shfl()
.shfl_down()
.shfl_up()
.shfl_xor()
.any()
.all()
.ballot()
.match_any()
.match_all()

其中这些协作同步操作类似于Warp Shuffle函数和Warp Vote 函数中描述的操作。 然而，在这些用户定义的合作组的背景下，它们的使用提供了更高的灵活性和生产力。 该功能将在本附录的后面进行介绍。
合并组：
在CUDA的SIMT体系结构中，在硬件级别，多处理器以32个称为warp的组执行线程。 如果在应用程序代码中存在依赖于数据的条件分支，使得warp内的线程发散，则warp将按顺序执行禁用不在该路径上的线程的每个分支。 在路径上保持活动的线程称为合并。 合作组具有发现和创建包含所有合并线程的组的功能，如下所示：

coalesced_group active = coalesced_threads();

例如，考虑在代码中存在分支的情况，其中每个warp中仅有第2,4,8线程处于活动状态。 放置在该分支中的上述呼叫将创建（对于每个）一个活动的组，其具有三个线程（等级0-2）