CUDA 全球屏障——适用于开普勒而不是费米

Question

以下全局屏障适用于 Kepler K10 而不是 Fermi GTX580：

__global__ void cudaKernel (float* ref1, float* ref2, int* lock, int time, int dim) {
  int gid  = blockIdx.x * blockDim.x + threadIdx.x;
  int lid  = threadIdx.x;                          
  int numT = blockDim.x * gridDim.x;               
  int numP = int (dim / numT);                     
  int numB = gridDim.x;

  for (int t = 0; t < time; ++t) {
    // compute @ time t
    for (int i = 0; i < numP; ++i) {
      int idx  = gid + i * numT;
      if (idx > 0 && idx < dim - 1)
        ref2 [idx]  = 0.333f * ((ref1 [idx - 1] + ref1 [idx]) + ref1 [idx + 1]);
    }

    // global sync
    if (lid == 0){
      atomicSub (lock, 1);
      while (atomicCAS(lock, 0, 0) != 0);
    }
    __syncthreads();

    // copy-back @ time t
    for (int i = 0; i < numP; ++i) {
      int idx  = gid + i * numT;
      if (idx > 0 && idx < dim - 1)
        ref1 [idx]  = ref2 [idx];
    }

    // global sync
    if (lid == 0){
      atomicAdd (lock, 1);
      while (atomicCAS(lock, numB, numB) != numB);
    }
    __syncthreads();
  }
}

因此，通过查看发送回 CPU 的输出，我注意到一个线程（第一个或最后一个线程）逃脱了屏障并比其他线程更早地恢复执行。我正在使用 CUDA 5.0。块数也总是小于 SM 数（在我的一组运行中）。

知道为什么相同的代码不能在两种架构上运行吗？ Kepler 中有哪些新功能有助于实现全球同步？

Answer 1

所以我怀疑屏障代码本身可能以相同的方式工作。这似乎是在与屏障功能本身无关的其他数据结构上发生的事情。

Niether Kepler 和 Fermi 都有彼此一致的 L1 缓存。您发现（尽管它与您的屏障代码本身无关）是 Kepler 和 Fermi 之间的 L1 缓存行为不同。

特别是，Kepler L1 缓存在上述链接中描述的全局负载上不起作用，因此缓存行为在设备范围的 L2 级别处理，因此是连贯的。当 Kepler SMX 读取其全局数据时，它会从 L2 获取一致的值。

另一方面，Fermi 具有也参与全局加载的 L1 缓存（默认情况下 - 尽管可以关闭此行为），并且上面链接中描述的 L1 缓存对于每个 Fermi SM 都是唯一的，并且是非- 与其他 SM 中的 L1 缓存一致。当 Fermi SM 读取其全局数据时，它从 L1 获取值，这可能与其他 SM 中的其他 L1 缓存不一致。

这是您所看到的“一致性”的差异，即您在障碍之前和之后操纵的数据。

正如我所提到的，我相信屏障代码本身在两种设备上的工作方式可能相同。