1

我正在研究我的游戏项目（塔防），我试图使用共享内存计算所有小动物和带有 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 时，我都会得到不同的输出。 .我相当确定它与共享内存和我处理它的方式有关，因为无论何时执行代码，重写代码以使用全局内存都会给出相同的结果。有任何想法吗？

4

1 回答 1

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 中内核的执行正确性。

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

于 2011-04-16T16:29:37.933 回答