在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：

[cpp]view plain copy
#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，获取的结果可能是不一样的，生成的图像如下，有散乱的杂点：

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

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

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

共享内存“__share__”

线程同步机制“__syncthreads()”

相关推荐

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

共享内存“share”