रजिस्टर दबाव और स्पिल

रजिस्टर किसी थ्रेड के पास उपलब्ध सबसे तेज़ संग्रहण होते हैं, लेकिन हर SM में इनकी संख्या सीमित होती है। कोई कर्नेल इनका कितनी अधिक मात्रा में उपयोग करता है, इसे रजिस्टर दबाव कहा जाता है।

हर सक्रिय थ्रेड प्रत्येक SM में मौजूद एक रजिस्टर फ़ाइल से रजिस्टर लेता है। हर थ्रेड को जितने अधिक रजिस्टर चाहिए, उतने ही कम थ्रेड एक साथ सक्रिय रह सकते हैं।

रजिस्टरों का अधिक उपयोग सीधे यह सीमित करता है कि किसी SM पर कितने वार्प समा सकते हैं। ऑक्यूपेंसी में यह कमी हार्डवेयर के पास विलंबता छिपाने के लिए पर्याप्त काम नहीं छोड़ सकती है।

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

स्पिल किए गए मान स्थानीय मेमोरी में जाते हैं, जो ऑफ-चिप DRAM में स्थित होती है। हर स्पिल एक मुक्त रजिस्टर अभिगम की जगह धीमी आवाजाही कर देता है।

कंपाइलर से उपयोग की रिपोर्ट देने को कहें। nvcc में --ptxas-options=-v जोड़ने पर हर कर्नेल के लिए प्रति थ्रेड रजिस्टर और स्पिल बाइट की संख्या दिखाई जाती है।

रिपोर्ट स्पिल स्टोर और लोड सूचीबद्ध करती है। स्पिल की कोई भी गैर-शून्य संख्या इस बात का चेतावनी संकेत है कि आपका कर्नेल धीमी स्थानीय मेमोरी आवाजाही की कीमत चुका रहा है।

आप कंपाइल समय पर एक ऊपरी सीमा तय कर सकते हैं। --maxrregcount फ़्लैग प्रति थ्रेड रजिस्टरों को सीमित करता है और थोड़ी गति के बदले अधिक ऑक्यूपेंसी देता है।

प्रति-कर्नेल संकेत अक्सर बेहतर होता है। __launch_bounds__ कंपाइलर को आपका ब्लॉक आकार बताता है, ताकि वह आपकी इच्छित ऑक्यूपेंसी के अनुसार रजिस्टरों का बजट तय कर सके।

अक्सर आप एक बार में कम मान रख सकते हैं। किसी सस्ते परिणाम की दोबारा गणना करना या चर का दायरा छोटा करना सक्रिय रजिस्टरों की संख्या घटाता है।

ILP और अनरोलिंग दबाव बढ़ाते हैं, जबकि सीमाएँ ऑक्यूपेंसी बढ़ाती हैं। कौशल यह है कि सबसे तेज़ चलने वाला संतुलन खोजा जाए, और यह केवल प्रोफ़ाइलर ही बता सकता है।

आपने सीखा कि अधिक रजिस्टर दबाव ऑक्यूपेंसी घटाता है और धीमे DRAM में स्पिल करा सकता है। उपयोग मापें, रजिस्टर सीमित करें और प्रोफ़ाइलर को मार्गदर्शन करने दें। 🎯