c++ - CUDA 内核中的竞争条件

标签 c++ cuda race-condition

我有一个 CUDA 内核，它似乎有竞争条件，我试图查明这种竞争条件的来源。我知道 cuda-memcheck 的“racecheck”工具，但是 racecheck 告诉我使用小输入时没有危险，这实际上也与我自己的调查一致。对于大输入，虽然 racecheck 似乎永远(字面上)，所以我不能使用它。 简单解释一下，定义为 __device__ 变量的一维 vector d_mat_3d 填充为 0 并加载到全局内存中。作为内核输入的两个大数组(d_A 和d_v)也在main 中定义并传递给内核。数组 d_mat_3d 的一段，称为 mat_2d 被剪切，加载到共享内存中，并将对其进行一些处理。然后，mat_2d 将被写回全局内存中的 d_mat_3d。

如此处所示，使用原子操作是因为如果不使用原子操作 mat_2d 会遇到不同线程的竞争条件。

我猜我仍然存在某种竞争条件的原因是 mat_3d 的结果每次都不同。

有没有想过这种竞争条件可能来自哪里？我可以采取任何步骤来清除它(工具 racecheck 除外)？如果您认为，没有竞争条件的证据，您能解释一下为什么每次执行内核时都会为 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;
}

最佳答案

是的，您的代码中至少有 2 个不同的竞争条件。

  1. 由于您在循环中加载整个共享内存(即在循环中一遍又一遍地加载所有内容)，因此有必要使用 __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，这可能会加剧这个问题.)

  2. 由于每个线程 block 都在写入整个 d_mat_3d ，因此当每个线程 block 尝试写入各种值时，您会遇到竞争条件。 threadblock 执行的顺序(CUDA 未定义)将主要决定在那里结束的内容，并且这很容易随运行而变化。我知道在没有完全重写内核的情况下解决这个问题的唯一简单方法是简单地启动 1 个线程 block (它仍然会填充 d_mat_3d 的相同区域)。这种竞争条件是全局内存竞争，目前 cuda-memcheck 无法发现这种竞争。我犹豫是否要对此读太多，但这段代码并没有任何意义，要么表明缺乏对合理代码的关注，要么表明缺乏对 CUDA 执行模型的理解(尤其是与下面的第 2 项相结合。 )

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

  1. 您在最后一个线程 block 中使用 __syncthreads() 可能是非法的。这个结构:

      if(n >= N_ELEM)
    return;

    将允许(最后一个)线程 block 中的一些线程提前退出，这意味着它们将不会参与后续的 __syncthreads() 语句。这在 CUDA 中是非法的，限制包含在 the programming guide 中。这可以通过删除早期返回并使用 if (n < N_ELEM) 或类似代码保护内核循环的各个部分(__syncthreads() 语句除外)来解决。

  2. 您的内核代码通常很奇怪，正如您已经在评论中指出的那样。这方面的一个例子是，您让 block 中的每个线程都对共享内存执行完全相同的加载和存储操作。从几个方面来说，这在性能方面是浪费的。

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

$ 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

关于c++ - CUDA 内核中的竞争条件，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/48086305/

上一篇：c++ - 将 QTablewidget 的更改值转换为 QStringList

下一篇：c++ - C++程序打印到哪里？

相关文章：

c++ - 如何在不改变原始颜色的情况下使用着色器

c++ - 模拟3D游戏中的鼠标移动？

c++ - 有没有一种方法可以使用 Thrust 根据索引 vector 设置标志

mysql - 如何确保在递增字段时 MySQL 数据库中没有竞争条件？

c++ - 分配给结构映射的语法

c++ - 如何在 C 中获得到核心映射的进程？

cuda - CUDA 事件会为 cudaMalloc 和 cudaMemcpy 执行计时吗？

c++ - 为什么 CUDA 8.0(有时)内存访问错误，而 7.5 却没有？

c - 当 parent 和 child 都写入标准输出时，输出没有交错

php - 如何保护PHP中的临界区？

©2024 IT工具网  联系我们