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

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

本文关键字：内存 加载 CUDA 共享 数据安全 全局      更新时间：2023-10-16

My kernel：

__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（）？

假设线程块的一个经线完成第一次迭代并开始下一次迭代，而其他经线仍在第一次迭代中工作。如果你在标签sync2没有__syncthreads，你最终会把这种扭曲写入共享内存，而其他人正在从该共享内存读取，这是竞争条件。

为了

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

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

相关文章：