全局工作队列

考虑一系列工作项。每个工作项完成所需的时间差别很大。为了平衡块之间的工作分配，每个块只有在前一个块完成时才能获取下一个项目。这与先验地将项目分配给块相反。

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 代码完成，应该非常简单。