全域性工作佇列
考慮一系列工作項。每個工作項完成所需的時間差別很大。為了平衡塊之間的工作分配,每個塊只有在前一個塊完成時才能獲取下一個專案。這與先驗地將專案分配給塊相反。
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 程式碼完成,應該非常簡單。