CUDA中__syncthreads()和__threadfence_block()和__trheadfence()

        作为cuda小白，我看完书上对这三个函数的解释，仍然不懂，于是做了以下几个实验来理解这三个函数的使用区别。

        我们先来看看__threadfence_block()是在干啥，这个实验非常简单，A数组的长度为1024（一个block的大小），先往A[0:512]里面写2，再往A[512:1024]里面写1，最后按照倒序把A数组复制到B数组里面，简单预测一下B里面正确的结果应该是11111...22222，即前512个元素全是1，后512个元素全是2。

#include<cuda_runtime.h>
#include<iostream>

template<typename scalar_t>
__global__ void swap(scalar_t* A, scalar_t* B)
{
	unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;

	// 往A的前512个元素写2，后512个元素写1，然后再把A的所有元素倒着写到B内
	if (idx < 512)
		A[idx] = 2;
	else
		A[idx] = 1;
	//__threadfence_block();	// warp内写完了再往下走
	//__threadfence();		// block内写完了再往下走
	//__syncthreads();		// block内执行到这里再往下走
	B[idx] = A[1023 - idx];		// 正确的B结果应该是111....222
}

int main()
{
	unsigned int len = 1024;
	unsigned int size = len * sizeof(int);
	int* A = new int[len];

	int* A_cuda, * B_cuda;
	cudaMalloc((void**)&A_cuda, size);
	cudaMalloc((void**)&B_cuda, size);
	cudaMemset(A_cuda, 0, size);
	cudaMemset(B_cuda, 0, size);

	swap<<<1, 1024>>>(A_cuda, B_cuda);

	int* B = new int[len];
	cudaMemcpy(B, B_cuda, size, cudaMemcpyDeviceToHost);
	cudaFree(A_cuda);
	cudaFree(B_cuda);

	for (int i = 0; i < len; ++i)
		printf("%d, ", B[i]);
	return 0;
}

不同步的情况下，即不使用代码中14,15,16行的三种同步指令，结果如图，得到了错误的结果，这很容易理解，因为不同warp间执行的顺序并没有做同步，所以A的写操作还没有完成，B就对A进行了复制，所以得到了很多0元素。我还仔细数了一下，相邻的一组0元素正好是32个，等于warp_size。

 使用__threadfence_block之后，结果如图，仍然得到了错误的结果，和不使用__threadfence_block效果差不多，我猜测__threadfence_block的功能是阻塞warp的内存延迟隐藏使之重新暴露，并没有同步不同的warp，书上也说它不会同步任何线程，所以和没使用的效果差不多。

使用__syncthreads之后，结果如图，终于得到了正确的结果，这很容易理解，它同步了一个block内的线程，所以B的复制操作在所有的A写操作之后。

 使用__threadfence之后，结果如图，结果也是正确的，因为__threadfence阻塞了所有块的线程， 所有块都

我们终于弄清楚__threadfence_block和其他两个函数的差别了，但是__syncthreads和__threadfence的差距似乎还没看出来，简单改动代码之后（如下），我把原来一个block拆分为了两个block，并且写了一个循环来拖慢block2的速度（这个循环不会影响正确结果），核函数的任务仍然不变。

#include<cuda_runtime.h>
#include<iostream>

template<typename scalar_t>
__global__ void swap(scalar_t* A, scalar_t* B)
{
	unsigned int idx = threadIdx.x + blockDim.x * blockIdx.x;

	// 拖慢block2的速度
	if (blockIdx.x == 1)
		for (int i = 0; i < 1024; ++i)
			A[idx] = -1;
	
	// 往A的前512个元素写2，后512个元素写1，然后再把A的所有元素倒着写到B内
	if (idx < 512)
		A[idx] = 2;
	else
		A[idx] = 1;
	//__threadfence_block();	// warp内写完了再往下走
	//__threadfence();		// block内写完了再往下走
	//__syncthreads();		// block内执行到这里再往下走
	B[idx] = A[1023 - idx];		// 正确的B结果应该是111....222
}

int main()
{
	unsigned int len = 1024;
	unsigned int size = len * sizeof(int);
	int* A = new int[len];

	int* A_cuda, * B_cuda;
	cudaMalloc((void**)&A_cuda, size);
	cudaMalloc((void**)&B_cuda, size);
	cudaMemset(A_cuda, 0, size);
	cudaMemset(B_cuda, 0, size);

	swap<<<2, 512>>>(A_cuda, B_cuda);

	int* B = new int[len];
	cudaMemcpy(B, B_cuda, size, cudaMemcpyDeviceToHost);
	cudaFree(A_cuda);
	cudaFree(B_cuda);

	for (int i = 0; i < len; ++i)
		printf("%d, ", B[i]);
	return 0;
}

用__syncthreads函数的结果如图，结果是错误的，因为__syncthreads只能同步块内线程，而B的复制是块间进行的，所以结果是错的。

 用__threadfence函数的结果如图，结果是正确的，因为__threadfence能够同步不同块之间的线程，B的复制操作在所有块都完成了A的写操作之后。

至此我们就理清了这三个函数的区别和作用了：

        1.__threadfence_block是阻塞warp直至warp发出的写操作完成，但由于warp本身就是单指令多线程，这个操作就比较多余，一般没什么用。但在分支语句中不能使用__syncthreads时就能派上用场了。

        2.__syncthreads是阻塞block直至block内的线程全都执行到这一行，但不能对块间进行同步。

        3.__threadfence是阻塞grid直至grid内的线程发出的读写操作完成，可以实现块间同步。

PS: 以上均是我的个人理解，如有错误，感谢指出