线程组和线程块:
任何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)