如何將全局內存中的數據加載到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中的共享內存中？

回答

相關問題