समांतर कमी

Question

मैंने मार्क हैरिस द्वारा सीयूडीए में समांतर कमी को अनुकूलित करने वाले लेख को पढ़ा है, और मैंने इसे वास्तव में बहुत उपयोगी पाया है, लेकिन फिर भी मैं कभी-कभी 1 या 2 अवधारणाओं को समझने में असमर्थ हूं।

//First add during load 

// each thread loads one element from global to shared mem 

unsigned int tid = threadIdx.x; 

unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; 

sdata[tid] = g_idata[i]; 
__syncthreads(); 

अनुकूलित कोड:: यह स्नातकोत्तर 18 पर लिखा है 2 भार और कमी के 1 जोड़ने के साथ:

// perform first level of reduction, 

// reading from global memory, writing to shared memory 
unsigned int tid = threadIdx.x;         ...1 

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x;   ...2 

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x];     ...3 

__syncthreads();             ...4 

मैं लाइन 2 को समझने में असमर्थ हूँ, अगर मेरे पास 256 तत्व हैं, और यदि मैं 128 को अपने ब्लॉकइज़ के रूप में चुनता हूं, तो मैं इसे 2 से गुणा क्यों कर रहा हूं? कृपया बताएं कि ब्लॉकइज़ को कैसे निर्धारित किया जाए?

Answer 1

यह मूल रूप से प्रदर्शन कर आपरेशन नीचे चित्र में दिखाया:

इस कोड को मूल रूप से "कह" है धागे की है कि आधे होगा प्रदर्शन वैश्विक स्मृति और साझा स्मृति के लिए लिख, के रूप में से पढ़ने तस्वीर में दिखाया गया।

आप एक कर्नेल निष्पादित करते हैं, और अब आप कुछ मानों को कम करना चाहते हैं, तो आप उपरोक्त कोड को केवल थ्रेड के कुल भाग तक पहुंचने तक सीमित कर देते हैं। , कल्पना आप 4 ब्लॉक है प्रत्येक सूत्र एक 512 के साथ, आप केवल दो पहले ब्लॉक द्वारा निष्पादित किया जाना है इसके बाद के संस्करण कोड की सीमा है, और आप एक g_idate [4 * 512] है:

unsigned int i = blockIdx.x*(blockDim.x*2) + threadIdx.x; 

sdata[tid] = g_idata[i] + g_idata[i+blockDim.x]; 

तो:

thread 0 of block = 0 will copy the position 0 and 512, 
thread 1 of block = 0 position 1 and 513; 
thread 511 of block = 0 position 511 and 1023; 
thread 0 of block 1 position 1024 and 1536 
thread 511 of block = 1 position 1535 and 2047 

ब्लॉकडिम।x * 2 का उपयोग किया जाता है क्योंकि प्रत्येक थ्रेड i और i+blockDim.x पर स्थिति तक पहुंच जाएगा, इसलिए आप गारंटी देते हैं कि अगले आईडी ब्लॉक पर थ्रेड g_idata की पहले से ही गणना की गई स्थिति को ओवरलैप नहीं करते हैं।

Answer 2

अनुकूलित कोड में आप कर्नेल को गैर-अनुकूलित कार्यान्वयन के रूप में आधे से बड़े ब्लॉक के साथ चलाते हैं।

चलो गैर-अनुकूलित कोड work में ब्लॉक का आकार कॉल करें, इस आकार के आधे को unit कहा जाए, और इन आकारों को अनुकूलित कोड के लिए भी समान संख्यात्मक मान दें।

गैर-अनुकूलित कोड में आप कर्नेल को work के रूप में कई धागे के साथ चलाते हैं, यह blockDim = 2 * unit है। प्रत्येक ब्लॉक में कोड 2 * unit आकार की साझा स्मृति में एक सरणी में g_idata का हिस्सा कॉपी करता है।

अनुकूलित कोड blockDim = unit में, इसलिए अब 1/2 धागे हैं, और साझा स्मृति में सरणी 2x छोटी है। लाइन 3 में पहला सारांश भी इकाइयों से आता है, जबकि विषम इकाइयों से दूसरा। इस तरह कमी के लिए आवश्यक सभी डेटा को ध्यान में रखा जाता है।

उदाहरण: आप blockDim=256=work (एकल ब्लॉक, unit=128) के साथ गैर अनुकूलित गिरी चलाते हैं, तो अनुकूलित कोड blockDim=128=unit के एक विशाल खंड है। चूंकि यह ब्लॉक blockIdx=0 प्राप्त करता है, *2 कोई फर्क नहीं पड़ता; पहला धागा g_idata[0] + g_idata[0 + 128] करता है।

आप 512 तत्वों, था और चलाने गैर अनुकूलित आकार 256 (work=256, unit=128) के 2 ब्लॉक, तो अनुकूलित कोड, 2 ब्लॉक के साथ लेकिन आकार 128 दूसरे खंड में पहली धागा (blockIdx=1) की अब करता है g_idata[2*128] + g_idata[2*128+128]।