CUDA内核中的竞争条件

Question

我有一个CUDA内核似乎有竞争条件，并试图找出这种竞争条件的来源。我知道cuda-memcheck的'racecheck'工具，但是racecheck告诉我使用小输入时没有危险，这实际上与我自己的调查一致。对于大输入，虽然racecheck似乎需要永远（字面意思），所以我不能使用它。简要说明，定义为d_mat_3d变量的1D向量__device__用0填充并加载到全局存储器中。作为内核输入的两个大数组（d_A和d_v）也在main中定义并传递给核心。切割一个名为d_mat_3d的数组mat_2d，加载到共享内存中，并对其进行一些处理。然后，mat_2d将在全球记忆中写回d_mat_3d。

如此处所示，原子操作被用作不使用原子操作，mat_2d将遇到不同线程的竞争条件。

我猜我还有某种竞争条件的原因是mat_3d的结果每次都不同。

有没有想过这种竞争条件可能来自哪里？我可以采取任何步骤来清除它（除了工具竞赛检查）？如果你认为，没有竞争条件的证据，你能解释为什么每次执行内核时为d_mat_3d分配不同的值吗？

CUDA 9.0 / NVidia Titan Black / Ubuntu 16.04

#include <cstdlib>
#include <sstream>
#include <cstdio>
#include <cuda.h>
#include <cuda_runtime_api.h>

#define W 7              // fix limit for loops in kernel
#define SIZE 100         // defining matrix dimension
#define N_ELEM 10000     // no of elements in each vector
#define NTPB 1024        // no of threads per block

using namespace std;

__device__ float d_mat_3d[SIZE*SIZE*SIZE]; 

__global__ void cuda_kernel(float *d_A, float *d_v){

  __shared__ float mat_2d[SIZE*SIZE]; // a 2D slice of 3D matrix d_mat_3d

  unsigned int n = blockDim.x*blockIdx.x+threadIdx.x;

  if(n >= N_ELEM)
    return;

  int x, y, z, i;
  float r;
  float A = d_A[n];
  float v = d_v[n];

  #pragma unroll
  for(x=0; x<SIZE; x++){

    // load mat_2d (on shared memory) using d_mat_3d (on global memory)
    for(i=0; i<SIZE*SIZE; i++){
      mat_2d[i] = d_mat_3d[i+x*SIZE*SIZE];
    }

    // sync threads as mat_2d is on shared memory
    __syncthreads();

    for(y=SIZE/2; y<SIZE/2+W; y++){ 
      for(z=SIZE/2; z<SIZE/2+W; z++){
        r = sqrt( pow(A,2) / v );  // no need to be in these loops. I know, but for my real case, it must be.
        atomicAdd(&mat_2d[z+y*SIZE], r); // atomically add r 
      }
    }

    __syncthreads();
    // write mat_2d (shared memory) back to mat_3d (global memory)
    for(i=0; i<SIZE*SIZE; i++){
      d_mat_3d[i+x*SIZE*SIZE] = mat_2d[i];
    }
  }
}

// this function writes h_mat_3d to disk. 
void write_image(float *h_mat_3d){
  ostringstream o_addToFile;
  o_addToFile << "mat3d.bin";
  FILE *pFile; 
  pFile = fopen(o_addToFile.str().c_str(), "wb");
  for(int i=0; i<SIZE*SIZE*SIZE; i++){ 
    fwrite(&h_mat_3d[i], sizeof(float), 1, pFile);
  }
  fclose (pFile);
}

int main(){

  int i;
  float *h_A = new float[N_ELEM]; // some large vector
  float *h_v = new float[N_ELEM]; // some other large vector
  float h_mat_3d[SIZE*SIZE*SIZE]; // will be filled w/ 0
  float *d_A; // device variables
  float *d_v;

  for(i=0; i<N_ELEM; i++){
    h_A[i] = 0.2f+(float)i/N_ELEM; // fill out with some calculations
    h_v[i] = 0.5f+2.f*i/N_ELEM;
  }
  for(i=0; i<SIZE*SIZE*SIZE; i++){
    h_mat_3d[i] = 0.f; // fill h_mat_3d with 0 
  }

  cudaMalloc((void **)&d_A, sizeof(float)*N_ELEM); // allocate variables on device
  cudaMalloc((void **)&d_v, sizeof(float)*N_ELEM);

  cudaMemcpy(d_A, h_A, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice); // copy from host to device
  cudaMemcpy(d_v, h_v, sizeof(float)*N_ELEM, cudaMemcpyHostToDevice);
  cudaMemcpyToSymbol(d_mat_3d, &h_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // copy h_mat_3d to device

  cuda_kernel<<<(N_ELEM+NTPB-1)/NTPB,NTPB>>>(d_A, d_v); // execute kernel

  cudaMemcpyFromSymbol(h_mat_3d, d_mat_3d, sizeof(float)*SIZE*SIZE*SIZE); // write it back to h_mat_3d

  write_image(h_mat_3d); // write h_mat_3d to disk for checking

  cudaFree(d_A); // free memory
  cudaFree(d_v);
  delete [] h_A;
  delete [] h_v;

  return 0;
}