__threadfence（）和L1高速缓存一致性

Question

这是我的理解（请参阅，例如How can I enforce CUDA global memory coherence without declaring pointer as volatile?,CUDA block synchronization differences between GTS 250 and Fermi devices和this post in the nvidia Developer Zone）__threadfence()保证在线程继续之前全局写入对其他线程可见。但是，在__threadfence()已返回之后，另一个线程仍然可以从L1高速缓存读取陈旧值，即使也是如此。

即：

线程A的一些数据写入到全局存储器，然后调用__threadfence()。然后，在之后__threadfence()已经返回，并且写入可见到全部其他线程，线程B被要求从这个内存位置读取。它发现它具有L1中的数据，因此会加载该数据。不幸的是，对于开发者来说，线程B的L1中的数据是陈旧的（即和线程A更新这个数据一样）。

首先：这是正确的吗？

假如是的话，它似乎我__threadfence()只有有用的，如果任何一个可以是某些该数据不会在L1（有点不太可能？），或者如果例如读取总是绕过L1（例如易失性或原子性）。它是否正确？

---

我问，因为我有一个比较简单的用例 - 传播的数据了一个二叉树 - 用原子级组标志和__threadfence()：第一个线程到达某个节点退出，而第二写入数据到它基于其两个孩子（例如他们的最小数据）。这适用于大多数节点，但通常至少失败一次。声明数据volatile可获得始终如一的正确结果，但会导致99％以上的未从L1中获取陈旧值的情况下的性能下降。我想确定这是该算法的唯一解决方案。下面给出了一个简化的例子。请注意，节点数组的排列宽度优先，叶子从索引start开始，并已填充数据。

__global__ void propagate_data(volatile Node *nodes, 
           const unsigned int n_nodes, 
           const unsigned int start, 
           unsigned int* flags) 
{ 
    int tid, index, left, right; 
    float data; 
    bool first_arrival; 

    tid = start + threadIdx.x + blockIdx.x*blockDim.x; 

    while (tid < n_nodes) 
    { 
     // We start at a node with a full data section; modify its flag 
     // accordingly. 
     flags[tid] = 2; 

     // Immediately move up the tree. 
     index = nodes[tid].parent; 
     first_arrival = (atomicAdd(&flags[index], 1) == 0); 

     // If we are the second thread to reach this node then process it. 
     while (!first_arrival) 
     { 
      left = nodes[index].left; 
      right = nodes[index].right; 

      // If Node* nodes is not declared volatile, this occasionally 
      // reads a stale value from L1. 
      data = min(nodes[left].data, nodes[right].data); 

      nodes[index].data = data; 

      if (index == 0) { 
       // Root node processed, so all nodes processed. 
       return; 
      } 

      // Ensure above global write is visible to all device threads 
      // before setting flag for the parent. 
      __threadfence(); 

      index = nodes[index].parent; 
      first_arrival = (atomicAdd(&flags[index], 1) == 0); 
     } 
     tid += blockDim.x*gridDim.x; 
    } 
    return; 
} 

Answer 1

首先：这是正确的？

是的，__threadfence()将数据推入L2并输出到全局内存。它对其他 SM中的L1缓存没有影响。

这是正确的吗？

是的，如果你有volatile结合__threadfence()对于全局存储器访问，你应该有信心，价值最终会以其他threadblocks可见。但请注意，线程块之间的同步在CUDA中并不是一个明确的概念。没有明确的机制可以这样做，也不能保证线程块的执行顺序，所以仅仅因为你的代码有一个__threadfence()某处在volatile项上运行，但仍然不能保证另一个线程块可能获得哪些数据。这也取决于执行的顺序。

如果您使用volatile，应该绕过L1（如果启用 - current Kepler devices don't really have L1 enabled for general global access）。如果您不使用volatile，那么当前正在执行__threadfence()操作的SM的L1应该在完成__threadfence()操作时与L2（和全局）一致/一致。

请注意，L2缓存在设备中是统一的，因此始终是“一致的”。至少从设备代码的角度来看，对于您的用例来说，无论您在使用哪个SM，L2和全局内存之间没有区别。

而且，如您所示，（全局）原子总是在L2 /全局内存上运行。