CUDA Warps के बारे में जानना क्यों परेशान है?

Question

मेरे पास GeForce GTX460 एसई है, इसलिए यह है: 6 एसएम एक्स 48 CUDA कोर = 288 CUDA कोर। यह ज्ञात है कि एक वार में 32 धागे होते हैं, और एक ब्लॉक में एक साथ (एक समय में) केवल एक वार को निष्पादित किया जा सकता है। कि, एक मल्टीप्रोसेसर (एसएम) में एक साथ केवल एक ब्लॉक, एक वार और केवल 32 धागे निष्पादित कर सकते हैं, भले ही 48 कोर उपलब्ध हों?

और इसके अलावा, कंक्रीट थ्रेड और ब्लॉक वितरित करने के लिए एक उदाहरण threadIdx.x और blockIdx.x का उपयोग किया जा सकता है। उन्हें आवंटित करने के लिए कर्नेल < < < ब्लॉक, थ्रेड >>>() का उपयोग करें। लेकिन कैसे एक विशिष्ट संख्या को वारप-एस आवंटित करने और उन्हें वितरित करने के लिए, और यदि यह संभव नहीं है तो Warps के बारे में जानना क्यों परेशान है?

Answer 1

स्थिति तुम क्या वर्णन से काफ़ी अधिक जटिल है।

एएलयू (कोर), लोड/स्टोर (एलडी/एसटी) इकाइयों और विशेष समारोह इकाइयों (एसएफयू) (छवि में हरा) पाइपलाइन इकाइयां हैं। वे पूरा होने के विभिन्न चरणों में, एक ही समय में कई कंप्यूटेशंस या संचालन के परिणाम रखते हैं। तो, एक चक्र में वे एक नए ऑपरेशन को स्वीकार कर सकते हैं और एक अन्य ऑपरेशन के परिणाम प्रदान कर सकते हैं जो बहुत समय पहले शुरू हुआ था (अगर मुझे सही याद है तो एएलयू के लिए लगभग 20 चक्र)। इसलिए, सिद्धांत में एक एसएम में 48 * 20 चक्र = 960 एएलयू संचालन एक ही समय में संसाधनों के संसाधन हैं, जो 960/32 धागे प्रति वार = 30 युद्ध हैं। इसके अलावा, यह एलडी/एसटी संचालन और एसएफयू संचालन को उनके विलंबता और थ्रूपुट पर भी संसाधित कर सकता है।

वार्प शेड्यूलर (छवि में पीला) प्रति चक्र पाइपलाइनों के लिए 2 * 32 थ्रेड प्रति वार = 64 धागे शेड्यूल कर सकते हैं। तो यह परिणाम की संख्या है जो प्रति घड़ी प्राप्त की जा सकती है। इसलिए, यह देखते हुए कि कंप्यूटिंग संसाधनों का मिश्रण है, 48 कोर, 16 एलडी/एसटी, 8 एसएफयू, जिनमें प्रत्येक की अलग-अलग विलंबताएं हैं, एक ही समय में युद्धपोतों का मिश्रण संसाधित किया जा रहा है। किसी दिए गए चक्र पर, वार शेड्यूलर एसएम के उपयोग को अधिकतम करने के लिए शेड्यूल करने के लिए दो युगल "युग्मित" करने का प्रयास करते हैं।

वार्प शेड्यूलर अलग-अलग ब्लॉक से या उसी ब्लॉक में विभिन्न स्थानों से युद्ध जारी कर सकते हैं, यदि निर्देश स्वतंत्र हैं। तो, एक ही समय में कई ब्लॉक से warps संसाधित किया जा सकता है।

जटिलता को जोड़ना, उन निर्देशों को निष्पादित करने वाले युद्ध जो 32 से कम संसाधन हैं, सभी थ्रेडों के लिए कई बार जारी किए जाने चाहिए। उदाहरण के लिए, 8 एसएफयू हैं, इसलिए इसका मतलब है कि एक निर्देश जिसमें एक निर्देश है जिसमें एसएफयू की आवश्यकता होती है उसे 4 बार निर्धारित किया जाना चाहिए।

यह विवरण सरलीकृत है। ऐसे अन्य प्रतिबंध भी हैं जो खेल में आते हैं और यह निर्धारित करते हैं कि जीपीयू कैसे काम करता है। आप "फर्मि आर्किटेक्चर" के लिए वेब खोजकर अधिक जानकारी प्राप्त कर सकते हैं।

तो, अपने वास्तविक प्रश्न के लिए आ रहा,

क्यों Warps बारे में पता करने के लिए परेशान?

एक वार्प में धागे की संख्या को जानना और इसे ध्यान में रखना महत्वपूर्ण हो जाता है जब आप अपने एल्गोरिदम के प्रदर्शन को अधिकतम करने का प्रयास करते हैं।यदि आप इन नियमों का पालन नहीं करते हैं, तो आप प्रदर्शन खोना:

• गिरी मंगलाचरण में, <<<Blocks, Threads>>>, कि एक ताना में धागे की संख्या के साथ समान रूप से विभाजित करता धागे के एक नंबर चुना है की कोशिश करो। यदि आप नहीं करते हैं, तो आप निष्क्रिय थ्रेड वाले ब्लॉक को लॉन्च करने के साथ समाप्त होते हैं।

• अपने कर्नेल में, प्रत्येक थ्रेड को एक वार में एक ही कोड पथ का पालन करने का प्रयास करें। यदि आप नहीं करते हैं, तो आपको वार्प विचलन कहा जाता है। ऐसा इसलिए होता है क्योंकि जीपीयू को अलग-अलग कोड पथों के माध्यम से पूरे वार को चलाने की ज़रूरत होती है।

• अपने कर्नेल में, प्रत्येक थ्रेड को एक वार्प लोड में रखने और विशिष्ट पैटर्न में डेटा स्टोर करने का प्रयास करें। उदाहरण के लिए, ग्लोबल मेमोरी में लगातार 32-बिट शब्दों तक पहुंचने के धागे हैं।

Answer 2

सूत्र, अनिवार्य रूप से क्रम में Warps में बांटा 1 - 32, 33 - 64 ...?

हां, प्रोग्रामिंग मॉडल गारंटी देता है कि थ्रेड को उस विशिष्ट क्रम में वारस में समूहीकृत किया जाता है।

अलग-अलग कोड पथों को अनुकूलित करने का एक सरल उदाहरण 32 धागे के समूहों में ब्लॉक में सभी धागे को अलग करने के लिए उपयोग किया जा सकता है? उदाहरण के लिए: स्विच (threadIdx.s/32) {केस 0:/* 1 वार */ब्रेक; मामला 1:/* 2 वार */ब्रेक;/* आदि * /}

बिल्कुल :)

कितने बाइट्स एकल वार्प के लिए एक समय में पढ़ा जाना चाहिए: 4 बाइट्स 32 धागे, 8 बाइट्स * * 32 धागे या 16 बाइट्स * 32 धागे? जहां तक ​​मुझे पता है, एक समय में वैश्विक स्मृति के लिए एक लेनदेन 128 बाइट प्राप्त करता है।

हां, वैश्विक मेमोरी के लेनदेन 128 बाइट हैं। इसलिए, यदि प्रत्येक थ्रेड लगातार पतों से 32-बिट शब्द पढ़ता है (उन्हें शायद 128-बाइट भी गठबंधन करने की आवश्यकता है), तो वार्प के सभी थ्रेडों को एक ही लेनदेन के साथ सर्विस किया जा सकता है (4 बाइट्स * 32 थ्रेड = 128 बाइट्स)। यदि प्रत्येक थ्रेड अधिक बाइट्स पढ़ता है, या यदि पते लगातार नहीं हैं, तो अधिक लेनदेन जारी किए जाने की आवश्यकता है (प्रत्येक अलग 128-बाइट लाइन को छूने के लिए अलग-अलग लेनदेन के साथ)।

यह सीयूडीए प्रोग्रामिंग मैनुअल 4.2, सेक्शन एफ.4.2, "ग्लोबल मेमोरी" में वर्णित है। वहां एक अस्पष्टता भी कह रही है कि स्थिति डेटा के साथ अलग है जो केवल एल 2 में कैश की जाती है, क्योंकि एल 2 कैश में 32-बाइट कैश लाइनें होती हैं। मुझे नहीं पता कि डेटा को केवल एल 2 में कैश किया जा सकता है या कितने लेनदेन समाप्त होते हैं।