cuda
इंटर-ब्लॉक संचार

टिप्पणियों

अंतिम-ब्लॉक गार्ड

कुछ कार्यों पर काम करने वाले ग्रिड पर विचार करें, जैसे एक समानांतर कमी। प्रारंभ में, प्रत्येक ब्लॉक स्वतंत्र रूप से अपना काम कर सकता है, कुछ आंशिक परिणाम उत्पन्न कर सकता है। हालांकि, अंत में, आंशिक परिणामों को एक साथ मिलाने और विलय करने की आवश्यकता होती है। एक विशिष्ट उदाहरण एक बड़े डेटा पर कमी एल्गोरिदम है।

एक विशिष्ट दृष्टिकोण दो गुठली को लागू करना है, एक आंशिक गणना के लिए और दूसरा विलय के लिए। हालांकि, यदि विलय एकल ब्लॉक द्वारा कुशलतापूर्वक किया जा सकता है, तो केवल एक कर्नेल कॉल की आवश्यकता होती है। इसे lastBlock रूप में परिभाषित गार्ड द्वारा प्राप्त किया जाता है:

__device__ bool lastBlock(int* counter) {
  __threadfence(); //ensure that partial result is visible by all blocks
  int last = 0;
  if (threadIdx.x == 0)
    last = atomicAdd(counter, 1);
  return __syncthreads_or(last == gridDim.x-1);
}

__device__ bool lastBlock(int* counter) {
  __shared__ int last;
  __threadfence(); //ensure that partial result is visible by all blocks
  if (threadIdx.x == 0) {
    last = atomicAdd(counter, 1);
  }
  __syncthreads();
  return last == gridDim.x-1;
}

इस तरह के एक गार्ड के साथ अंतिम ब्लॉक को अन्य सभी ब्लॉकों द्वारा उत्पादित सभी परिणामों को देखने की गारंटी है और विलय कर सकते हैं।

__device__ void computePartial(T* out) { ... }
__device__ void merge(T* partialResults, T* out) { ... }

__global__ void kernel(int* counter, T* partialResults, T* finalResult) {
    computePartial(&partialResults[blockIdx.x]);
    if (lastBlock(counter)) {
      //this is executed by all threads of the last block only
      merge(partialResults,finalResult);
    }
}

मान्यताओं:

• काउंटर को एक वैश्विक मेमोरी पॉइंटर होना चाहिए, जो कर्नेल के आह्वान से पहले 0 से शुरू होता है।
• lastBlock फ़ंक्शन को सभी ब्लॉकों में सभी थ्रेड्स द्वारा समान रूप से लागू किया गया है
• कर्नेल को एक-डायमेंशनल ग्रिड (उदाहरण की सादगी के लिए) में लगाया जाता है
• T आपको किसी भी प्रकार का नाम देता है, लेकिन उदाहरण C ++ अर्थ में टेम्पलेट होने का इरादा नहीं है

वैश्विक कार्य कतार

कार्य आइटम की एक सरणी पर विचार करें। प्रत्येक कार्य आइटम को पूरा करने के लिए आवश्यक समय बहुत भिन्न होता है। ब्लॉक के बीच कार्य वितरण को संतुलित करने के लिए प्रत्येक ब्लॉक के लिए अगले आइटम को लाने के लिए यह विवेकपूर्ण हो सकता है कि पिछले एक पूरा होने पर ही। यह ब्लॉक करने के लिए आइटमों को निर्दिष्ट करने वाली प्राथमिकताओं के विपरीत है।

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 सरणी वैश्विक मेमोरी में रहती है
• कर्नेल में वर्कक्यू ऑब्जेक्ट में कोई नया कार्य आइटम नहीं जोड़ा गया है जो इससे प्राप्त हो रहा है
• WorkItem कार्य असाइनमेंट का एक छोटा प्रतिनिधित्व है, उदाहरण के लिए किसी अन्य ऑब्जेक्ट के लिए एक सूचक
• WorkItem::none() स्थिर सदस्य फ़ंक्शन एक WorkItem ऑब्जेक्ट बनाता है जो WorkItem भी काम का प्रतिनिधित्व नहीं करता है
• ब्लॉक में सभी थ्रेड्स द्वारा WorkQueue::fetch() को समान रूप से कहा जाना चाहिए
• बीच में एक और __syncthreads() बिना WorkQueue::fetch() 2 इनवोकेशन नहीं हैं। अन्यथा एक दौड़ की स्थिति दिखाई देगी!

उदाहरण में शामिल नहीं है कि कैसे WorkQueue को इनिशियलाइज़ करें या इसे पॉप्युलेट करें। यह एक और कर्नेल या सीपीयू कोड द्वारा किया जाता है और इसे बहुत सीधा-आगे होना चाहिए।