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

【问题标题】：CUDA (JCUDA) shared memory (?) problems / undefined behaviourCUDA (JCUDA) 共享内存 (?) 问题/未定义行为
【发布时间】：2016-02-15 05:37:16
【问题描述】：

我正在开发我的游戏项目（塔防），并尝试使用 JCuda 使用共享内存计算所有小动物与塔之间的距离。对于每座塔，我用 N threds 运行 1 个街区，其中 N 等于地图上小动物的数量。我正在计算给定块的所有生物与塔之间的距离，并将迄今为止发现的最小距离存储在块的共享内存中。我当前的代码如下所示：

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 ...我相当确定它与共享内存以及我处理它的方式有关，因为无论何时执行代码，重写代码以使用全局内存都会给出相同的结果。有什么想法吗？

【问题讨论】：

标签： cuda jcuda

【解决方案1】：

这里的内核中存在内存竞争问题（因此读写正确）：

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

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

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

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

写入之前的屏障确保写入到最小值将在“最终”（尽管不一致）值存储在最小值之前完成，但随后块中的每个线程都将执行写入。

如果您的意图是让每个线程计算一个距离，然后将块上的最小距离值写入全局内存，您将不得不使用原子内存操作（对于共享内存，这是仅在计算 1.2/1.3 和 2.x 设备上受支持），或编写显式共享内存缩减。之后，只有一个线程应该执行回写到全局内存。

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

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

如果块中的某些线程可以绕过屏障分支并退出，而其他线程在屏障处等待，则会导致内核挂起。 __syncthreads() 调用周围不应该有任何分支分歧，以确保 CUDA 中内核的执行正确性。

所以，总而言之，在当前代码中至少三个问题回到绘图板上。

【讨论】：