CUDA C/C++：相同的可执行文件给出了第一次运行的不同结果

Question

以我之前的回答问题：My Previous Question，顺便说一句，正确回答了Robert Crovella。

我想出了另一个计算随机步骤到一个点的内核（通过使用与我以前的问题相同的RNG）并计算该点相对于其先前位置（坐标）的能量差异。这是内核：

__global__ 
void DeltaE(float *X, float *Y,float *Xn, float *Yn, float *step,float *DELTA, curandState *state, const int N,const int n){ 

    int tIdx = blockIdx.x*blockDim.x + threadIdx.x; 
    int bIdx = blockIdx.x; 
    //int sIdx = blockIdx.x*blockDim.x; 

    float x , y; 
    float rdmn1, rdmn2; 

    float dIntE = 0.0e0f, dConfE = 0.0e0f, dTotE = 0.0e0f; 
    if(tIdx < N){ 

     if(tIdx == n){ 
      step[tIdx] = 0.2; 
      rdmn1 = curand_uniform(&state[tIdx]); 
      rdmn2 = curand_uniform(&state[tIdx]); 

      Xn[tIdx] = X[tIdx] + step[tIdx]*(2.0e0f*rdmn1 - 1.0e0f); 
      Yn[tIdx] = Y[tIdx] + step[tIdx]*(2.0e0f*rdmn2 - 1.0e0f); 
      dConfE = - (X[tIdx]*X[tIdx] + Y[tIdx]*Y[tIdx]); 
      dConfE += Xn[tIdx]*Xn[tIdx] + Yn[tIdx]*Yn[tIdx]; 

     } 
     else{ 
      x = X[tIdx] - X[n]; 
      y = Y[tIdx] - Y[n]; 

      dIntE += -1.0e0f/sqrt(x*x + y*y); 
     } 
     __syncthreads(); 
     if(tIdx != n){ 
      x = X[tIdx] - Xn[n]; 
      y = Y[tIdx] - Yn[n]; 

      dIntE += 1.0e0f/sqrt(x*x + y*y); 
     }  
     dTotE = dConfE + dIntE; 
     dTotE = ReduceSum2(dTotE); 
     if(threadIdx.x == 0)DELTA[bIdx] = dTotE; 


    } 
} 

然后我做CPU上最后一笔：

cudaMemcpy(&delta[0],&d_delta[0],blocks.x*sizeof(float),cudaMemcpyDeviceToHost); 
float dE = 0; 
for(int i = 0; i < blocks.x; i++){ 
    dE += delta[i]; 
} 

我的内核具有以下配置推出：

dim3 threads(BlockSize,BlockSize); 
dim3 blocks(ceil(Np/threads.x),ceil(Np/threads.y)); 
DeltaE<<<blocks.x,threads.x,threads.x*sizeof(float)>>>(d_rx,d_ry,d_rxn,d_ryn,d_step,d_delta,d_state,Np,nn); 

中，Np是多少点（我用1k - 4k）。我有一个GeForce 9500 GT，它不支持双倍。我编译时不使用flag/no选项。例如，取Np = 1k。当我编译然后运行时，结果是dE = 6.557993。当我跑第二，第三，第四，无论什么时候，它都是dE = -0.3515406。有谁知道这是从哪里来的？

P.S .:我忘了提及，在DeltaE之前就可以找到My Previous Question的相同内核AvgDistance。我不知道这是否有什么可做的事，但我认为值得一提。 P.2：nn是任何选择的点（粒子）。

Answer 1

正如通过评论指出，通过上述Robert Crovella，什么可能发生的事情是，虽然tIdx = n被计算Xn[n]和Yn[n]，其他线程使用这个值，它可能还没有被计算。在这种情况下，其他运行（除第一个之外）唯一的原因是获得相同（正确）值的方式是由Xn和Yn指向的内存已被正确的值占用，即使使用该同步问题应用程序返回正确的值。

在任何情况下，我通过拆分内核分为二避免了同步的问题，就像我被Robert Crovella通过评论建议：

__global__ 
void DeltaE1(float *X, float *Y,float *Xn, float *Yn, float *step,float *DELTA, curandState *state, const int N,const int n){ 

    int tIdx = blockIdx.x*blockDim.x + threadIdx.x; 
    float x , y; 
    float rdmn1, rdmn2; 

    if(tIdx < N){ 
     DELTA[tIdx] = 0.0e0f; 
     if(tIdx == n){ 
      step[tIdx] = 0.2e0f; 
      rdmn1 = curand_uniform(&state[tIdx]); 
      rdmn2 = curand_uniform(&state[tIdx]); 

      Xn[tIdx] = X[tIdx] + step[tIdx]*(2.0e0f*rdmn1 - 1.0e0f); 
      Yn[tIdx] = Y[tIdx] + step[tIdx]*(2.0e0f*rdmn2 - 1.0e0f); 
      DELTA[tIdx] = - (X[tIdx]*X[tIdx] + Y[tIdx]*Y[tIdx]); 
      DELTA[tIdx] += Xn[tIdx]*Xn[tIdx] + Yn[tIdx]*Yn[tIdx]; 

     } 
     else{ 
      x = X[tIdx] - X[n]; 
      y = Y[tIdx] - Y[n]; 

      DELTA[tIdx] += -1.0e0f/sqrt(x*x + y*y); 
     }   
    } 
} 

__global__ 
void DeltaE2(float *X, float *Y,float *Xn, float *Yn,float *DELTA,const int N,const int n){ 

    int tIdx = blockIdx.x*blockDim.x + threadIdx.x; 
    int bIdx = blockIdx.x; 

    float x , y; 
    float dTotE = 0.0e0f; 
    if(tIdx < N){ 
     if(tIdx != n){ 
      x = X[tIdx] - Xn[n]; 
      y = Y[tIdx] - Yn[n]; 

      DELTA[tIdx] += 1.0e0f/sqrt(x*x + y*y); 

     } 
     dTotE = DELTA[tIdx]; 
     dTotE = ReduceSum2(dTotE); 
     if(threadIdx.x == 0)DELTA[bIdx] = dTotE; 

    } 

}