CUDA之Cooperative Groups操作，细粒度并行操作。

问题

CUDA的线程执行单元是以warp来划分的，一个warp内部包含32个线程，这32个线程存在一个隐式的线程同步。而不同warp之间是不存在隐式同步的。在一个block中，往往存在多个warp，倘若在程序中使用同步机制"__syncthreads()"令线程同步，此时可能会发生条件竞争的问题。导致指令延迟过高，性能变低，如下所示。

__global__ void test()
{
	const int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < 64)
	{
		printf("idx = %d,顺利到达。\n", idx);
		__syncthreads();
	}
}

需求

如果想只让一个块中的某些线程进行同步的话，就需要更细粒度的并行等待操作。

方法

使用CUDA的Cooperative Groups操作。头文件如下所示：

#define __CUDACC__ //该定义在开发阶段使用，方便代码自动补全
#include <cooperative_groups.h>

使用说明

对于#define __CUDACC__ 这条语句特殊说明下：由于主机端代码和设备端代码在编译的起始，是分开编译的。在设备端编译时，编译器会自定预定义__CUDACC__。但是开发阶段是在主机端，因此，可以在开发时先自定义一个，在开发完毕后再注释掉。不过不注释掉也没关系，只会警告你 重复定义。

下面说明如何使用cooperative_groups

namespace cg = cooperative_groups;  //工作区名字太长了 ，这样方便。
__global__ void binomial()
{
	cg::thread_block cta = cg::this_thread_block(); //这个块中的所有线程。特指块。这个块根据包含了线程组织分布的所有信息。
	cg::thread_group tg = cg::this_thread();  //这个线程(组)，特指当前线程(组)。这个后面说
	
	if (threadIdx.x ==0)
	{
		printf("cta.size() = %d\n", cta.size());
	
		printf("