c - 尽管 atomicAdd 函数 (CUDA) 存在竞争条件？

Question

我有一个在两个层面上并行的问题:我有大量的 (x0, x1, y0, y1) 坐标对，它们被转换成变量 vdx， vdy, vyy 并且对于这些集合中的每一个，我都试图计算由它们组成的所有“单项式”的值，直到 n 度(即它们的不同幂的所有可能组合，例如 vdx^3*vdy*vyy^2 或 vdx*1*vyy^4)。然后将这些值加到所有集合上。

我的策略(现在我只是想让它工作，它不必用多个内核或复杂的减少来优化，除非它真的必须)是让每个线程处理一组坐标对并计算所有它们对应的单项式的值。每个 block 的共享内存保存所有单项式总和，当 block 完成时， block 中的第一个线程将结果添加到全局总和。由于每个 block 的共享内存都被所有地方的所有线程访问，所以我使用的是 atomicAdd；与 block 和全局内存相同。

不幸的是，某处似乎仍然存在竞争条件，因为每次运行内核时我的结果都不同。

如果有帮助，我目前正在使用 degree = 3 并省略其中一个变量，这意味着在下面的代码中，最内层的 for 循环(在 evbl) 什么都不做，只是重复 4 次。 Indeed, the output of the kernel looks like this: 51502,55043.1,55043.1,51502,47868.5,47868.5,48440.5,48440.6,46284.7,46284.7,46284.7,46284.7,46034.3,46034.3,46034.3,46034.3,44972.8,44972.8, 44972.8,44972.8,43607.6,43607.6,43607.6,43607.6,43011,43011,43011,43011,42747.8,42747.8,42747.8,42747.8,45937.8,45937.8,46509.9,46509.9,... and it's noticable that there is a 4 元组的(粗略)模式。但每次我运行它时，值都非常不同。

一切都是 float 的，但我使用的是 2.1 GPU，所以这应该不是问题。 cuda-memcheck 也没有报告错误。

有更多 CUDA 经验的人可以给我一些指导，告诉我如何在这里追踪竞争条件吗？

__global__ void kernel(...) {

  extern __shared__ float s_data[];

  // just use global memory for now
  // get threadID:
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if(idx >= nPairs) return;

  // ... do some calculations to get x/y...

  // calculate vdx, vdy and vyy
  float vdx = (x1 - x0)/(float)xheight;
  float vdy = (y1 - y0)/(float)xheight;
  float vyy =  0.5*(y0 + y1)/(float)xheight;


  const int offs1 = degree + 1;
  const int offs2 = offs1 * offs1;
  const int offs3 = offs2 * offs1;
  float sol = 1.0;

  // now calculate monomial results and store in shared memory

  for(int evdx = 0; evdx <= degree; evdx++) {
    for(int evdy = 0; evdy <= degree; evdy++) {
      for(int evyy = 0; evyy <= degree; evyy++) {
        for(int evbl = 0; evbl <= degree; evbl++) {
          s = powf(vdx, evdx) + powf(vdy, evdy) + powf(vyy, evyy);
          atomicAdd(&(s_data[evbl + offs1*evyy + offs2*evdy +
                offs3*evdx]), sol/1000.0 ); 

        }
      }
    }
  }

  // now copy shared memory to global
  __syncthreads();
  if(threadIdx.x == 0) {
    for(int i = 0; i < nMonomials; i++) {
      atomicAdd(&outmD[i], s_data[i]);
    }
  }
}

Answer 1

您正在使用共享内存，但您从未对其进行初始化。