__shared__ सारणियाँ घोषित करें

प्रत्येक ब्लॉक को ऑन-चिप स्मृति का एक छोटा, तेज़ पूल मिलता है, जिसे साझी स्मृति कहते हैं। इसे ऐसा स्क्रैचपैड समझें जिसे पूरा ब्लॉक मिलकर लिख और पढ़ सकता है।

साझी स्मृति streaming multiprocessor पर ही स्थित होती है, इसलिए यह वैश्विक स्मृति की तुलना में लगभग 100x तेज़ होती है। धीमी ऑफ-चिप DRAM पर बार-बार निर्भरता से बचने के लिए इसका उपयोग करें। 🚀

आप इसे kern के भीतर __shared__ qualifier से घोषित करते हैं। इसके बाद यह एक ऐरे ब्लॉक के प्रत्येक थ्रेड द्वारा साझा की जाती है।

जब आप कोष्ठकों में आकार देते हैं, तो वह स्थिर साझा स्मृति होती है। कम्पाइलर को आकार पता होना चाहिए, इसलिए वह runtime मान नहीं, बल्कि एक स्थिरांक होना चाहिए।

मुख्य विचार यह है: __shared__ ऐरे प्रति ब्लॉक केवल एक बार बनती है, प्रत्येक थ्रेड के लिए नहीं। सभी थ्रेड ठीक उसी ऐरे को देखते हैं।

क्योंकि प्रत्येक थ्रेड समान डेटा देखता है, साझा स्मृति थ्रेडों को सहयोग करने देती है। एक थ्रेड कोई मान रख सकता है और पड़ोसी थ्रेड उसे उठा सकता है।

इसका जीवनकाल ब्लॉक के समान होता है। ऐरे ब्लॉक के शुरू होने पर बनती है और उसके समाप्त होने पर मिट जाती है, इसलिए इसका केवल ब्लॉक क्षेत्र होता है। 🧱

एक सामान्य पैटर्न है कि प्रत्येक थ्रेड के लिए एक स्लॉट हो और उसे threadIdx.x से इंडेक्स किया जाए। प्रत्येक थ्रेड अपना तत्व समानांतर रूप से साझा स्मृति में लोड करता है।

साझी स्मृति छोटी होती है, अक्सर प्रत्येक SM के लिए केवल 48 से 100 KB। प्रत्येक ब्लॉक के लिए बहुत अधिक माँग करने पर एक साथ चल सकने वाले ब्लॉकों की संख्या घट जाती है।

सबसे सामान्य कार्य वैश्विक डेटा की एक टाइल तैयार करना है, ताकि ब्लॉक धीमी DRAM पर लौटे बिना उसका कई बार पुनः उपयोग कर सके।

सीमा याद रखें: साझा स्मृति अपने ब्लॉक तक सीमित होती है। दो अलग-अलग ब्लॉकों को अपनी अलग-अलग प्रति मिलती है और वे एक-दूसरे में झाँक नहीं सकते।

आपने सीखा कि __shared__ प्रत्येक ब्लॉक को एक तेज़ ऑन-चिप स्क्रैचपैड देता है, जो प्रति ब्लॉक एक बार बनता है और पुनः उपयोग योग्य डेटा तैयार करने के लिए आदर्श है। अगला विषय: थ्रेडों को तालमेल में रखना। 🎯