CUDA： 共享内存与同步

　　CUDA C支持共享内存， 将CUDA C关键字__shared__添加到变量声明中，将使这个变量驻留在共享内存中。对在GPU上启动的每个线程块，CUDA C编译器都将创建该变量的一个副本。线程块中的每个线程都共享这块内存，但线程无法看到也不能修改其他线程块的变量副本。共享内存缓冲区驻留在物理GPU上，而不是GPU之外的系统内存中。因此访问共享内存时的延迟远远低于访问普通缓冲区的延迟，使得共享内存像每个线程块的高速缓存或者中间结果暂存器那样高效。

const int N = 33*1024;
const int threadsPerBlock = 256;

__global__ void dot(float *a, float *b, float *c)
{
    __shared__  float  cache[threadsPerBlock];
    int tid = threadIdx.x + blockId.x*blockDim.x;
    int cacheIndex = threadIdx.x;
    float temp = 0;
    while(tid<N){
        temp += a[tid]*b[tid];
        tid += blockDim.x * gridDim.x; 
    }
    cache[cacheIndex] = temp;  
    __syncthreads();
    int i = blockDim.x/2;
    while(i != 0){
        if(cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
        i /= 2;
    }
    if(cacheIndex == 0)
        c[blockIdex.x] = cache[0];
}            

__syncthreads();

这个函数调用将确保线程块中的每个线程都执行完__syncthreads（）前面的语句后，才会执行下一条语句。