全域性工作佇列

考慮一系列工作項。每個工作項完成所需的時間差別很大。為了平衡塊之間的工作分配,每個塊只有在前一個塊完成時才能獲取下一個專案。這與先驗地將專案分配給塊相反。

class WorkQueue {
private:
  WorkItem* gItems;
  size_t totalSize;
  size_t current;
public:
  __device__ WorkItem& fetch() {
    __shared__ WorkItem item;
    if (threadIdx.x == 0) {
      size_t itemIdx = atomicAdd(current,1);
      if (itemIdx<totalSize)
        item = gItems[itemIdx];
      else
        item = WorkItem::none();
    }
    __syncthreads();
    return item; //returning reference to smem - ok
  }
}

假設:

  • WorkQueue 物件以及 gItem 陣列駐留在全域性記憶體中
  • 沒有新的工作項新增到從中獲取的核心中的 WorkQueue 物件
  • WorkItem 是工作分配的小表示,例如指向另一個物件的指標
  • WorkItem::none() 靜態成員函式建立一個完全不代表任何工作的 WorkItem 物件
  • WorkQueue::fetch() 必須由塊中的所有執行緒統一呼叫
  • 在沒有另一個 __syncthreads() 之間沒有 2 個 WorkQueue::fetch() 的呼叫。否則會出現競爭條件!

該示例不包括如何初始化 WorkQueue 或填充它。它由另一個核心或 CPU 程式碼完成,應該非常簡單。