CUDA线程协作之共享存储器“__shared__”&&“__syncthreads()”

在GPU并行编程中，一般情况下，各个处理器都需要了解其他处理器的执行状态，在各个并行副本之间进行通信和协作，这涉及到不同线程间的通信机制和并行执行线程的同步机制。

共享内存“__share__”

CUDA中的线程协作主要是通过共享内存实现的。使用关键字“__share__”声明共享变量，将使这个变量驻留在共享内存中，该变量具有以下特征：

• 位于线程块的共享存储器空间中
• 与线程块具有相同的生命周期
• 仅可通过块内的所有线程访问

对于GPU上启动的每个线程块，CUDA C编译器都将创建该变量的一个副本。 线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。 这就使得一个线程块中的多个线程能够在计算上进行通信和协作。而且，共享内存缓冲区驻留在物理GPU上，在访问共享内存时的延迟要远远低于访问普通缓冲区的延迟，使得共享内存的访问非常高效。

线程同步机制“__syncthreads()”

关键字“__share__”只是声明了共享变量，位于同一个线程块中的不同线程都可以访问该变量，如果没有同步机制，将会发生竞态条件 (Race Condition)，导致错误的运行结果。

CUDA确保同步的方法是调用“__syncthreads()”。__syncthreads()将确保线程块中的每个线程都执行完 __syncthreads()前面的语句后，才会执行下一条语句。

以下是CUDA和OpenCV的应用中，绘制一幅图像，Grid的尺寸大小是60*60，Block的尺寸大小是10*10，在各个线程块内声明了一个共享变量sharedMem：

#include "cuda_runtime.h"    
#include <highgui.hpp>    

using namespace cv;

#define DIM 600   //图像长宽
#define PI 3.1415926535897932f  

__global__ void kernel(unsigned char *ptr)
{
	// map from blockIdx to pixel position    
	int x = threadIdx.x + blockIdx.x * blockDim.x;
	int y = threadIdx.y + blockIdx.y * blockDim.y;
	int offset = x + y * blockDim.x * gridDim.x;

	__shared__  float sharedMem[16][16];
	const float period = 128.0f;
	sharedMem[threadIdx.x][threadIdx.y] =
		255 * (sinf(x*2.0f*PI / period) + 1.0f) *
		(sinf(y*2.0f*PI / period) + 1.0f) / 4.0f;
	__syncthreads();

	ptr[offset * 3 + 0] = 0;
	ptr[offset * 3 + 1] = sharedMem[15 - threadIdx.x][15 - threadIdx.y];
	ptr[offset * 3 + 2] = 0;
}

// globals needed by the update routine    
struct DataBlock
{
	unsigned char   *dev_bitmap;
};

int main(void)
{
	DataBlock   data;
	cudaError_t error;

	Mat image = Mat(DIM, DIM, CV_8UC3, Scalar::all(0));

	data.dev_bitmap = image.data;
	unsigned char    *dev_bitmap;

	error = cudaMalloc((void**)&dev_bitmap, 3 * image.cols*image.rows);
	data.dev_bitmap = dev_bitmap;

	dim3    grid(DIM / 10, DIM / 10);
	dim3   block(10, 10);
	//DIM*DIM个线程块  
	kernel << <grid, block >> > (dev_bitmap);

	error = cudaMemcpy(image.data, dev_bitmap,
		3 * image.cols*image.rows,
		cudaMemcpyDeviceToHost);

	error = cudaFree(dev_bitmap);

	imshow("__share__ and __syncthreads()", image);
	waitKey();
}

如果线程间不加入__syncthreads()同步机制，同一线程块内不同线程访问sharedMem，获取的结果可能是不一样的，生成的图像如下，有散乱的杂点：

加入__syncthreads()同步机制，保证了同一线程块中不同的线程都执行完成__syncthreads()这个集合点之前的部分之后，才继续往下执行，所以不同的线程访问sharedMem获取的结果是一致的，图像无杂散点，是一个规律的排布：