代码之家 › 专栏 › 技术社区 › If_You_Say_So jmilloy

CUDA内核中的竞争条件

race-condition cuda c++

If_You_Say_So jmilloy · 技术社区 · 7 年前

我有一个CUDA内核,它似乎有竞态条件,我正试图找出这种竞态条件的来源。我知道cuda memcheck的“racecheck”工具,但racecheck告诉我,使用小输入时没有危险,这实际上也与我自己的调查一致。对于大输入,虽然racecheck似乎需要花费很长时间(字面上),所以我不能使用它。简要说明,一维向量 d_mat_3d 定义为 __device__ 变量用0填充并加载到全局内存中。作为内核输入的两个大型数组( d_A 和 d_v )也在中定义 main 并传递给内核。数组的一段 d\U mat\U 3d ,调用 mat_2d 将被剪切、加载到共享内存中,并对其进行一些处理。然后 mat\U 2d 将写回 d\U mat\U 3d 在全局内存上。

如图所示,原子操作与未使用原子操作一样使用 mat\U 2d 将遇到带有不同线程的争用条件。

我想我仍然有某种比赛状态的原因是 mat_3d 每次都不一样。

有没有想过这种比赛状态可能来自哪里?我可以采取哪些步骤来清除该问题(工具线检查除外)?如果你认为,没有证据表明存在种族状况,你能解释一下为什么分配给不同的值吗 d\U mat\U 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;
}

1 回复 | 直到 7 年前

Robert Crovella 7 年前

是的,您的代码中至少有2种不同的比赛条件。

由于您正在循环中加载整个共享内存(即,在循环中反复加载所有共享内存),因此有必要使用 __syncthreads() . 这样做将减少从一次运行到另一次运行的可变性,直至第6位或第7位有效十进制数字,这与 ordinary float variability in floating-point operations ,其中操作顺序不重复(此处通常如此)。

添加以下行:
```
  for(x=0; x<SIZE; x++){
    __syncthreads();  // add this line
    // 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();
```
应该基本上纠正这个问题。如果不这样做,当内核循环时 x ,一些扭曲可以“抢先”开始加载共享内存,而以前的扭曲仍忙于中的前一个循环迭代 x个 (注意下面的评论2,这可能会加剧这个问题。)
自从 每个螺纹块 正在向整个 d_mat_3d ,当每个threadblock尝试写入不同的值时,您有一个竞争条件。线程块的执行顺序(CUDA未定义)将主要决定最终的结果,这很容易因运行而异。我所知道的在不进行完整内核重写的情况下解决这个问题的唯一简单方法是只启动1个threadblock(它仍然会填充 d\U mat\U 3d ). 这种竞争条件是一种全局内存竞争 cuda-memcheck 目前无法发现此类种族。我不想读太多,但这段代码实际上没有任何意义,要么表示对合理的代码缺乏关注,要么表示对CUDA执行模型缺乏了解(尤其是与下面的第2项结合在一起)

我还想指出一些其他的事情。

您对的使用 __syncthreads() 在最后一个线程块中可能非法。此构造:
```
  if(n >= N_ELEM)
    return;
```
将允许(最后一个)线程块中的某些线程提前退役,这意味着它们将不参与后续的 __syncthreads() 声明。这在CUDA中是非法的,该限制包含在 the programming guide . 这可以通过删除早期返回并使用 if (n < N_ELEM) 或类似。
正如您已经在评论中指出的那样,您的内核代码通常很奇怪。这方面的一个例子是,您让块中的每个线程执行与共享内存完全相同的加载和存储。从几个方面来看,这是浪费性能的。

我并不是说这涵盖了代码的所有问题,仅仅是我注意到的事情。这是一个相对完整的测试用例,我用它来验证我的发现。它包括对我上面提到的项目的一些更改,以及对我来说很重要的其他各种更改:

$ cat t268.cu
#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;


  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++){
  __syncthreads();
if (n < N_ELEM){
    // 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();
if (n < N_ELEM){
    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)
if (n < N_ELEM){
    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){
  for (int i = 0; i < SIZE*SIZE; i++){
    for (int j = 0; j < SIZE; j++)
      if (h_mat_3d[i*SIZE+j] > 1.0f) printf("%d:%f\n ", i*SIZE+j,  h_mat_3d[i*SIZE+j]);
    printf("\n");}
}

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 = new float[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+i/(float)N_ELEM; // fill out with some calculations
    h_v[i] = 0.5f+2.f*i/(float)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<<<1,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
  delete [] h_A;
  delete [] h_v;

  return 0;
}
$ nvcc -arch=sm_52 -o t268 t268.cu
$ ./t268 > out1.txt
$ ./t268 > out2.txt
$ diff out1.txt out2.txt |more
51,57c51,57
< 5050:330.657715
<  5051:330.657715
<  5052:330.657715
<  5053:330.657715
<  5054:330.657715
<  5055:330.657715
<  5056:330.657715
---
> 5050:330.657654
>  5051:330.657593
>  5052:330.657593
>  5053:330.657593
>  5054:330.657593
>  5055:330.657593
>  5056:330.657593
59,65c59,65
< 5150:330.657715
<  5151:330.657715
<  5152:330.657715
<  5153:330.657715
<  5154:330.657745
<  5155:330.657745
<  5156:330.657745
---
> 5150:330.657593
>  5151:330.657593
>  5152:330.657593
>  5153:330.657593
>  5154:330.657593
>  5155:330.657593
>  5156:330.657593
67,73c67,73
< 5250:330.657745
<  5251:330.657745
<  5252:330.657745
<  5253:330.657745
<  5254:330.657715
<  5255:330.657715
<  5256:330.657715
---
> 5250:330.657593
>  5251:330.657593
>  5252:330.657623
>  5253:330.657593
>  5254:330.657593
>  5255:330.657593
>  5256:330.657593
75,81c75,81
< 5350:330.657715
<  5351:330.657715
<  5352:330.657715
<  5353:330.657715
<  5354:330.657715
<  5355:330.657745
<  5356:330.657715
---
> 5350:330.657593
>  5351:330.657593
$

可以看出,剩余的变化在第7位有效十进制数字中:

51,57c51,57
< 5050:330.657715
...
---
> 5050:330.657654