CUDA-死锁-例-点积

好好体会-Dot product

（CUDA by example 55页）

这个例子值得好好感悟

点积：两个长度相同的向量A和B，对应元素相乘后相加。

CUDA实现思路：每个线程分别读取A和B中对应位置的元素，紧接着执行相乘操作，最后将相乘结果存入shared memory的对应位置。

注意，当向量元素个数远远超过一个block中的threads数量时的处理。

const int N = 33*1025;
const int threadsPerBlock = 256;
const int blocksPerGrid = (N+threadsPerBlock-1)/threadsPerBlock;

__global__ void dotProduction(float* a, float* b){
/// 第一步：
    // thread计算得到a和b对应元素的乘积，
    // 存入这个thread对应的shared memory中的位置
    // 每个block对应的shared memory的大小为这个block中的threads数量
    __shared__ float cache[threadsPerBlock];

    int tid = threadIdx.x + blockDim.x*blocIdx.x;

    // 每个block的shared memory的索引就是threadIdx.x，与blockIdx.x 无关,
    // 这个block和那个block中的thread的ID是一摸一样的。
    int cacheIndex = threadsIdx.x;

    // 执行乘法操作，更新threads ID，
    // 同stride-loop
    // 看图一的过程
    float tmp = 0.0f;
    while (tid<N){
        tmp += a[tid]*b[tid];
        tid += blockDim.x*gridDim.x;
    }
    cache[cacheIndex] = tmp;

    // 同步这个block中的所有threads
    __syncthreads();
    // 确保所有threads完成工作之后，执行后续指令

/// 第二步：
    // 对于每个block对应的shared memory中的元素，进行规约求和。
    // 其中threadPerBlock必须是2的指数。
    // 同for-loop
    int i = blockDim.x/2;
    while(i != 0){
        if (cacheIndex <i )
            cache[cacheIndex] += cache[cacheIndex + i];
        // 确保上一轮所有和得到，所以要同步
        __syncthreads();
        i = i >> 2;
    }
    // 把每个shared memory中的第一个元素，也就是这个shared memory中
    // 所有元素之和，写入c中对应的位置。
    if(cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}
...

// 在主函数中，把c从device拷贝到host，以及之后：
cudaMemcpy(h_c, c, cudaMemcpyDeviceToHost);

// 在CPU上将h_c中的结果求和，于此同时GPU上可以后其他任务执行。
// CPU和GPU并行执行。隐藏延时
float sum = 0;
for (int i=0;i < blocksPerGrid; i++){
    sum += h_c[i];
}
// sum即是最终点积结果。

其中这一部分：

float tmp = 0.0f;
while (tid<N){
    tmp += a[tid]*b[tid];
    tid += blockDim.x*gridDim.x; // 更新tid，自加不是1，而是所有threads数量。
}
cache[cacheIndex] = tmp;

当所有threads的数目小于a或b中的元素个数时（不论有多少个blocks），上述保证正确，与stride更新效果相同。当threads个数等于元素个数时，也正确。所以这样写，分析见下图：

上图中只使用了一个block，所以在第二步归约计算时，就可以在第0个位置上得到最终结果。当使用多个blocks时，第二步得到每个block的第0个元素，而这些若干个第0个元素保存于c（Global）中，最终还要将c中元素求和。

而下面这种实现：

if (tid<N){
    cache[threadIdx.x] = a[tid]*b[tid];
}

只适用于thread个数等于元素个数时（如下图）。但通常元素个数会远大于threads数。所以不适用此法。

• 技能：多个blocks中的各个shared memory 缓存同时被操作。

• 为什么要将最后的结果传回host计算？

  因为，事实证明，想GPU这种大规模并行机器在执行最后的规约步骤时，通常会浪费计算资源，因为此时的数据集非常小。比如，当使用480 个threads将32 个数相加时，将难以充分使用每一个threads。

总结一下，使用shared memory优化dot-production为什么有效，因为减少了写入global memory的次数，并且复制回host的数据量减少。性能增加。

敲黑板注意threadIdx.x 与threads ID的区别，前者相对ID后者绝对ID。访存Shared memory时，一定使用threadIdx.x。

__syncthreads() 放错位置会导致死锁

规约程序中：

int i=blockDim.x/2;
while(i != 0){
    if (cacheIndex <i )
        cache[cacheIndex] += cache[cacheIndex + i];
    __syncthreads();
    i /= 2;
}

如果将__syncthreads()放入if语句，会产生死锁：

int i=blockDim.x/2;
while(i != 0){
    if (cacheIndex <i ){
        cache[cacheIndex] += cache[cacheIndex + i];
        __syncthreads();
    }
    i /= 2;
}

解释一下，当出现线程发散时，发散的分支会使得某些threads处于空闲状态，而其他threads将执行分支中的代码。而对于__syncthreads()而言，CUDA架构确保，一个block中的所有threads都执行到__syncthreads()，才能执行__syncthreads()之后的语句。这样一来，上述代码块，只要有一个threads没有执行if语句，也就不能够执行__syncthreads()，其他执行了if语句的threads会等待哪一个thread，一直等下去，造成死锁。

所以，对于__syncthreads()要谨慎使用。