CUDA (JCUDA)共享内存问题/未定义行为

Question

我正在制作我的游戏项目(塔防)，我试图使用共享内存计算所有critters和JCuda塔之间的距离。对于每个塔，我运行一个带有N个线程的块，其中N等于地图上critters的数量。我正在计算给定块中所有critters与该塔之间的距离，并将迄今为止找到的最小距离存储在块的共享内存中。我当前的代码是这样的:

extern "C"
__global__ void calcDistance(int** globalInputData, int size, int
critters, int** globalQueryData, int* globalOutputData) {
  //shared memory
  __shared__ float minimum[2];
  int x = threadIdx.x  + blockIdx.x * blockDim.x;
  int y = blockIdx.y;
  if (x < critters) {
    int distance = 0;
    //Calculate the distance between tower and criter
    for (int i = 0; i < size; i++) {
      int d = globalInputData[x][i] - globalQueryData[y][i];
      distance += d * d;
    }
    if (x == 0) {        
      minimum[0] = distance;
      minimum[1] = x;
    }
    __syncthreads();
    if (distance < minimum[0]) {
      minimum[0] = distance;
      minimum[1] = x;
    }
   
    __syncthreads();
    globalOutputData[y * 2]     = minimum[0];
    globalOutputData[y] = minimum[1];
  }
}

问题是，如果我使用相同的输入多次重新运行代码(每次运行后我释放主机和设备上的所有内存)，我每次执行代码时都会得到不同的输出块(塔)号>27日……我很确定它与共享内存和我处理它的方式有关，因为重写代码以使用全局内存在代码执行时给出相同的结果。什么好主意吗?

Answer 1

内核中存在内存争用问题(即写后读正确性):

   if (distance < minimum[0]) {
     minimum[0] = distance;
     minimum[1] = x;
   }

执行时，块中的每个线程都会尝试同时读取和写入最小值。不能保证当一个warp中的多个线程尝试写入相同的共享内存位置时会发生什么，也不能保证同一块中的其他warp在从正在写入的内存位置加载时将读取什么值。内存访问不是原子的，并且没有锁定或序列化，这将确保代码执行你似乎想要做的缩减操作类型。

同样问题的一个比较温和的版本适用于内核末尾回写全局内存:

   __syncthreads();
   globalOutputData[y * 2]     = minimum[0];
   globalOutputData[y] = minimum[1];

写操作之前的屏障确保在"最终"(尽管不一致)值被存储在最小值之前完成对最小值的写操作，但是块中的每个线程都将执行写操作。

如果您的意图是让每个线程计算一个距离，然后将块上的最小距离值写入全局内存，则必须使用原子内存操作(对于共享内存，这在compute 1.2/1.3和2上是支持的)。(仅限X设备)，或者写一个显式的共享内存缩减。在此之后，只有一个线程应该执行写回全局内存。

最后，还有一个可能导致内核挂起的潜在同步正确性问题。__syncthreads()(它映射到PTX条指令)要求块中的每个线程在内核继续之前到达并执行该指令。有这样的控制流:

 if (x < critters) {
 ....
   __syncthreads();
 ....
 }

如果块中的一些线程可以绕过屏障进行分支并退出，而其他线程在屏障处等待，则

将导致内核挂起。在__syncthreads()调用周围不应该有任何分支分歧，以确保CUDA内核的执行正确性。

因此，总结一下，回到绘图板上，至少解决了当前代码中的三个问题。