如何在 CUDA 中安全地将全局内存中的数据加载到共享内存中？

Question

我的内核：

__global__ void myKernel(float * devData, float * devVec, float * devStrFac,
int Natom, int vecNo) {

extern __shared__ float sdata[];
int idx = blockIdx.x * blockDim.x + threadIdx.x;

float qx=devVec[3*idx];
float qy=devVec[3*idx+1];
float qz=devVec[3*idx+2];
__syncthreads();//sync_1

float c=0.0,s=0.0;
for (int iatom=0; iatom<Natom; iatom += blockDim.x) {
    float rtx = devData[3*(iatom + threadIdx.x)];//tag_0
    float rty = devData[3*(iatom + threadIdx.x)+1];
    float rtz = devData[3*(iatom + threadIdx.x)+2];
    __syncthreads();//sync_2
    sdata[3*threadIdx.x] = rtx;//tag_1
    sdata[3*threadIdx.x + 1] = rty;
    sdata[3*threadIdx.x + 2] = rtz;
    __syncthreads();//sync_3

    int end_offset=  min(blockDim.x, Natom - iatom);

    for (int cur_offset=0; cur_offset<end_offset; cur_offset++) {
        float rx = sdata[3*cur_offset];
        float ry = sdata[3*cur_offset + 1];
        float rz = sdata[3*cur_offset + 2];
        //sync_4  
        float theta = rx*qx + ry*qy + rz*qz;

        theta = theta - lrint  (theta);
        theta = theta * 2 * 3.1415926;//reduce theta to [-pi,pi]

        float ct,st;
        sincosf(theta,&st,&ct);

        c += ct;
        s += st;
    }

}

devStrFac[idx] += c*c + s*s;
}

为什么需要标记为 sync_2 的“__syncthreads()”？没有 sync_2，sdata[] 得到错误的数字，我得到错误的结果。 "tag_1" 行使用了 "tag_0" 行的结果，所以在我看来 sync_2 是没有必要的。我哪里错了？如果由于指令执行混乱，我应该在 "sync_4" 行放一个 __syncthreads() 吗？

Answer 1

考虑线程块的一个 warp 完成第一次迭代并开始下一次迭代，而其他 warp 仍在进行第一次迭代。如果您在标签 sync2 处没有 __syncthreads，您最终会在其他人从该共享内存读取时将此 warp 写入共享内存，这是竞争条件。

为了清楚起见，您可以将标签 sync2 处的 __syncthreads() 移动到外循环的末尾。

"cuda-memcheck --tool racecheck"应该告诉你问题出在哪里了。

如何在 CUDA 中安全地将全局内存中的数据加载到共享内存中？

How to load data in global memory into shared memory SAFELY in CUDA?

c++

cuda