कोएलस्ड बनाम स्ट्राइडेड रीड

कोएलैसिंग इस बात पर निर्भर करती है कि प्रत्येक थ्रेड किस पते को छूता है। थ्रेड से पते का मानचित्र तय करता है कि क्या पूरा वार्प एक ही लेन-देन से काम चला सकता है।

जब पड़ोसी थ्रेड पड़ोसी पतों को पढ़ते हैं, तो वार्प एक सन्निहित अनुक्रम को कवर करता है। यह सुव्यवस्थित विन्यास कोएलैस्ड अभिगम कहलाता है।

थ्रेड 0 इंडेक्स 0 को, थ्रेड 1 इंडेक्स 1 को और इसी तरह आगे पढ़ता है। सभी 32 पते एक संरेखित पंक्ति में आते हैं, इसलिए GPU को केवल एक लेन-देन की आवश्यकता होती है।

स्ट्राइड उन पतों के बीच का निश्चित अंतर है जिन्हें क्रमिक थ्रेड छूते हैं। जैसे ही यह अंतर एक अवयव से अधिक होता है, कोएलैसिंग टूटने लगती है।

2 के स्ट्राइड में, थ्रेड 0, 0 पढ़ता है; थ्रेड 1, 2 पढ़ता है; और थ्रेड 2, 4 पढ़ता है। अब वार्प दोगुनी स्मृति में फैल जाता है, इसलिए उसे अधिक लेन-देन चाहिए।

स्ट्राइड को दोगुना करने पर छुई जाने वाली पंक्तियों की संख्या भी लगभग दोगुनी हो जाती है। बड़े स्ट्राइड वार्प को बहुत-से लेन-देन की ओर धकेल सकते हैं, जबकि प्रत्येक पंक्ति के केवल छोटे-से हिस्से का उपयोग होता है।

पंक्ति-प्रमुख मैट्रिक्स का पूरा कॉलम प्रत्येक थ्रेड को देने से बहुत बड़ा स्ट्राइड बनता है। कोड में यह सुव्यवस्थित दिखता है, लेकिन चुपचाप हर वार्प को बिखेर देता है।

अक्सर समाधान केवल यह बदलना होता है कि कौन-सा इंडेक्स थ्रेड के साथ सबसे तेज़ी से बदलता है। क्रमिक थ्रेडों को क्रमिक स्मृति पर चलने दें और पठन फिर से कोएलैस्ड हो जाएगा।

यदि कोएलैस्ड पैटर्न पंक्ति के बीच से शुरू हो, तो उसे भी नुकसान होता है। गलत-संरेखित ऑफ़सेट वार्प को सीमा के आर-पार ले जाता है और एक पठन को दो भागों में बाँट देता है।

किसी पैटर्न का मूल्यांकन करते समय केवल एक थ्रेड की कल्पना न करें। सभी 32 लेनों को एक साथ देखें और पूछें कि उनके पते मिलकर कितनी पंक्तियों को कवर करते हैं। 🔍

सबसे तेज़ी से बदलने वाले इंडेक्स को threadIdx.x के अनुसार चलाएँ। यह एक आदत आपकी अधिकांश वैश्विक रीड को बिना अतिरिक्त प्रयास के कोएलैस्ड बनाए रखती है।

आपने देखा कि कोएलैस्ड रीड पड़ोसी तत्वों को साथ रखती हैं, जबकि स्ट्राइड उन्हें पंक्तियों में बिखेर देता है। सबसे तेज़ इंडेक्स को चलाने का काम threadIdx.x को दें। 🎉