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

我的内核：如何将全局内存中的数据加载到CUDA中的共享内存中？

__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; 
}

为什么 “__syncthreads（）” 需要标注sync_2？如果没有sync_2，sdata []会得到错误的数字，我会得到错误的结果。行“tag_1”使用行“tag_0”的结果，所以在我的脑海里，sync_2是没有必要的。我错在哪里？如果由于执行混乱的指令，我应该将__syncthreads（）放在“sync_4”行中？

来源

2015-06-14 DengPan

'sync_2'应该位于外部循环的底部。在任何线程开始将新数据加载到共享数组之前，您必须确保所有线程warp在给定的迭代中都使用共享内存完成 – talonmies

考虑一个线程块的变形完成第一次迭代并开始下一次迭代，而其他变形仍在第一次迭代中工作。如果你没有__syncthreads标签sync2，你将最终将这个warp写入共享内存，而其他人从共享内存中读取，这是竞争条件。

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

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

来源

2015-06-14 11:42:09

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

回答

相关问题