0

以下全局屏障适用于 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 中有哪些新功能有助于实现全球同步?

4

1 回答 1

1

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

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

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

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

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

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

于 2013-01-09T22:55:40.367 回答