__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 /全局內存上運行。