ब्लॉक, धागे, warpSize

Question

#blocks & ब्लॉक आकार का चयन करने के बारे में बहुत सी चर्चा हुई है, लेकिन मुझे अभी भी कुछ याद आ रहा है। मेरी चिंताओं में से कई लोग इस सवाल का पता: How CUDA Blocks/Warps/Threads map onto CUDA Cores? (चर्चा आसान बनाने के लिए, वहाँ है पर्याप्त perThread & perBlock स्मृति मेमोरी सीमा एक मुद्दा यहां नहीं हैं।।)

kernelA<<<nBlocks, nThreads>>>(varA,constB, nThreadsTotal); 

1) एस एम संभव के रूप में व्यस्त रखने के लिए, मुझे warpSize के एकाधिक में सेट करना चाहिए। सच?

2) एक एसएम केवल एक समय में एक कर्नेल निष्पादित कर सकता है। एसएम के सभी HWcores केवल कर्नेल ए निष्पादित कर रहे हैं। (कुछ HWcores कर्नेलए नहीं चल रहे हैं, जबकि अन्य कर्नेलबी चलाते हैं।) तो अगर मेरे पास केवल एक धागा चलाने के लिए है, तो मैं अन्य HWcores को "बर्बाद कर रहा हूं"। सच?

3) यदि वार्प-शेड्यूलर warpSize (32 धागे) की इकाइयों में काम करता है, और प्रत्येक एसएम में 32 एचडब्ल्यूकोर्स होते हैं, तो एसएम का पूर्ण उपयोग किया जाएगा। क्या होता है जब एसएम के पास 48 एचडब्ल्यूओआरओआर होते हैं? जब शेड्यूलर 32 के हिस्सों में काम जारी कर रहा है तो मैं सभी 48 कोर का पूर्ण उपयोग कैसे कर सकता हूं? (यदि पिछला अनुच्छेद सत्य है, तो क्या बेहतर होगा यदि शेड्यूलर HWcore आकार की इकाइयों में काम जारी करता है?)

4) ऐसा लगता है कि एक समय में वार्प-शेड्यूलर कतार 2 कार्य करता है। इसलिए जब वर्तमान में निष्पादित कर्नेल स्टाल या ब्लॉक करता है, तो दूसरा कर्नेल बदल जाता है। (यह स्पष्ट नहीं है, लेकिन मुझे लगता है कि यहां कतार 2 कर्नेल से अधिक गहरी है।) क्या यह सही है?

5) यदि मेरे एचडब्ल्यू की 512 धागे-प्रति-ब्लॉक (एनटीएचड्समैक्स) की ऊपरी सीमा है, तो इसका मतलब यह नहीं है कि 512 धागे वाले कर्नेल एक ब्लॉक पर सबसे तेज़ी से चलेंगे। (दोबारा, कोई मुद्दा याद नहीं है।) एक अच्छा मौका है यदि मैं 512-थ्रेड कर्नेल को कई ब्लॉक में फैलाता हूं, तो केवल एक ही नहीं। ब्लॉक एक या कई एसएम पर निष्पादित किया जाता है। सच?

5 ए) मैं छोटे से बेहतर सोच रहा हूं, लेकिन क्या इससे कोई फर्क पड़ता है कि मैं nBlocks कितना छोटा करता हूं? सवाल यह है कि, nBlocks का मान कैसे चुनें, यह सभ्य है? (जरूरी नहीं कि इष्टतम।) nBlocks चुनने के लिए कोई गणितीय दृष्टिकोण है, या यह केवल परीक्षण-एन-गलती है।

Answer 1

मुझे अपने प्रश्नों का उत्तर एक-एक करके करने का प्रयास करें।

1. यह सही है।
2. आपका मतलब वास्तव में "HWcores" से क्या है? आपके कथन का पहला भाग सही है।

3. NVIDIA Fermi Compute Architecture Whitepaper के अनुसार: "32 समानांतर धागे warps बुलाया के समूह में एसएम कार्यक्रम धागे प्रत्येक एसएम दो ताना शेड्यूलर और दो अनुदेश प्रेषण इकाइयों, की अनुमति देता है दो warps जारी किए गए और समवर्ती निष्पादित फर्मी के दोहरे ताना अनुसूचक चयन किए जाने की सुविधा है।। दो warps, और प्रत्येक warp से सोलह कोर, सोलह लोड/स्टोर इकाइयों, या चार एसएफयू के एक समूह के लिए एक निर्देश जारी करता है। क्योंकि युद्ध स्वतंत्र रूप से निष्पादित करते हैं, फर्मि के शेड्यूलर को निर्देश धारा के भीतर निर्भरताओं की जांच करने की आवश्यकता नहीं है "।

   इसके अलावा, NVIDIA Keppler Architecture Whitepaper कहता है: "केप्लर क्वाड वार शेड्यूलर चार युद्धों का चयन करता है, और प्रति चक्र दो स्वतंत्र निर्देश प्रत्येक चक्र को प्रेषित किया जा सकता है।"

   इसलिए "अतिरिक्त" कोर का उपयोग एक समय में एक से अधिक युद्ध को शेड्यूल करके किया जाता है।

4. वार्प शेड्यूलर शेड्यूल उसी कर्नेल के warps, अलग कर्नेल नहीं।

5. बिल्कुल सही नहीं: प्रत्येक ब्लॉक को एक एकल एसएम में बंद कर दिया जाता है, क्योंकि वह कहां साझा की गई स्मृति रहता है।
6. यह एक मुश्किल मुद्दा है और यह निर्भर करता है कि आपका कर्नेल कैसे कार्यान्वित किया जाता है। आप वसीली वोल्कोव द्वारा एनवीडिया वेबिनार Better Performance at Lower Occupancy पर एक नज़र डालना चाहते हैं जो कुछ और महत्वपूर्ण मुद्दों को बताता है। मुख्य रूप से, हालांकि, मैं सुझाव दूंगा कि आप CUDA Occupancy Calculator का उपयोग करके अधिग्रहण में सुधार के लिए अपनी थ्रेड गिनती चुनें।

Answer 2

1) हां।

2) सीसी 2.0 - 3.0 डिवाइस एक साथ 16 ग्रिड तक निष्पादित कर सकते हैं। प्रत्येक एसएम 8 ब्लॉक तक सीमित है ताकि पूर्ण समेकन तक पहुंचने के लिए डिवाइस में कम से कम 2 एसएम हो।

3) हां वार्प शेड्यूलर समय पर युद्ध का चयन करते हैं और जारी करते हैं। सीयूडीए कोर की अवधारणा को भूल जाओ वे अप्रासंगिक हैं। विलंबता को छिपाने के लिए आपको उच्च निर्देश स्तर समांतरता या उच्च अधिभोग होना चाहिए। सीसी 1.x के लिए 25% और सीसी> = 2.0 के लिए> 50% होने की अनुशंसा की जाती है। सामान्य सीसी 3.0 में शेड्यूलर की दोगुना होने के कारण 2.0 उपकरणों की तुलना में अधिक अधिभोग की आवश्यकता होती है लेकिन प्रति एसएम में केवल 33% की वृद्धि होती है। नाइट वीएसई इश्यु दक्षता प्रयोग यह निर्धारित करने का सबसे अच्छा तरीका है कि क्या आपके पास निर्देश और स्मृति विलंबता को छिपाने के लिए पर्याप्त युद्ध हैं या नहीं। दुर्भाग्यवश, विजुअल प्रोफाइलर में यह मीट्रिक नहीं है।

4) वार्प शेड्यूलर एल्गोरिदम दस्तावेज नहीं है; हालांकि, यह नहीं मानता कि थ्रेड ब्लॉक किस ग्रिड की उत्पत्ति हुई थी। सीसी 2.x और 3.0 उपकरणों के लिए सीयूडीए कार्य वितरक अगले ग्रिड से ब्लॉक वितरित करने से पहले सभी ब्लॉक को ग्रिड से वितरित करेगा; हालांकि, यह प्रोग्रामिंग मॉडल द्वारा गारंटी नहीं है।

5) एसएम व्यस्त रखने के लिए आपको डिवाइस को भरने के लिए पर्याप्त ब्लॉक होना चाहिए। इसके बाद आप यह सुनिश्चित करना चाहते हैं कि उचित अधिग्रहण तक पहुंचने के लिए आपके पास पर्याप्त युद्ध हो। बड़े थ्रेड ब्लॉक का उपयोग करने के लिए पेशेवर और विपक्ष दोनों हैं। सामान्य रूप से बड़े थ्रेड ब्लॉक कम निर्देश कैश का उपयोग करते हैं और कैश पर छोटे पैरों के निशान होते हैं; हालांकि, बड़े थ्रेड ब्लॉक syncthreads पर stall (एसएम कम कुशल हो सकता है क्योंकि से चुनने के लिए कम warps हैं) और समान निष्पादन इकाइयों पर निष्पादन निर्देश रखने के लिए जाते हैं। मैं शुरू करने के लिए प्रति थ्रेड ब्लॉक 128 या 256 धागे की कोशिश करने की सलाह देते हैं। बड़े और छोटे धागे दोनों ब्लॉक के लिए अच्छे कारण हैं। 5 ए) अधिभोग कैलकुलेटर का उपयोग करें। थ्रेड ब्लॉक आकार के बहुत बड़े पिकिंग से आपको अक्सर रजिस्टरों द्वारा सीमित किया जा सकता है। थ्रेड ब्लॉक आकार के बहुत छोटे पिकिंग को आप साझा मेमोरी या एसएम सीमा के 8 ब्लॉक द्वारा सीमित कर सकते हैं।