हमारी पिछली पोस्ट में वेक्टर-ऐड कर्नेल को nvcc से वार्प्स तक देखा गया था। इस बार, हम एक महत्वपूर्ण SASS निर्देश (ग्लोबल लोड) को RTX4090 के हार्डवेयर से होते हुए देखते हैं। जांचे गए CUDA कर्नेल में दो पंक्तियाँ हैं: __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}। संकलित SASS दिखाता है: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]। ये निर्देश वेक्टर b के तत्वों को ग्लोबल मेमोरी से रजिस्टरों में लोड करते हैं।
एक LDG.E 32 लेन में से प्रत्येक में चार बाइट मांगता है। इसे पूरा करने के लिए चार 32-बाइट सेक्टर, एक कैश लाइन, एक पता अनुवाद, एक क्रॉसबार पार, छत्तीस L2 स्लाइस में से एक, और जब हर जगह मिस होता है, तो DRAM चिप पर एक एक्टिवेट और चार कॉलम रीड की आवश्यकता होती है। हमारा वार्प SM के चार सब-पार्टीशन में से एक पर ग्यारह अन्य निवासी वार्प्स के साथ रहता है। निर्देश लोड/स्टोर यूनिट (LSU) को जारी किया जाता है, जो ओपकोड, सक्रिय लेन मास्क, गणना किए गए पते और रजिस्टर नंबर कोलेसर को भेजता है।
कोलेसर निर्धारित करता है कि वार्प द्वारा अनुरोधित 128 बाइट के लिए 4 लगातार सेक्टर अनुरोधों की आवश्यकता है। L1 कैश में प्रवेश करते समय, अनुरोध वर्चुअली एड्रेस्ड, 4-वे सेट-एसोसिएटिव कैश को भेजा जाता है। लुकअप चार स्लॉट के सेट में टैग की तुलना करता है। चूंकि यह पहली बार लोड है, यह मिस होता है और आगे नीचे जाता है। L1 हिट लगभग 15.4 नैनोसेकंड में लौटता है — 40 चक्र। वर्चुअल मेमोरी को SM छोड़ने से पहले अनुवाद की आवश्यकता होती है। SM की TLB में सोलह हाल के अनुवाद होते हैं। मिस की लागत लगभग 4.4 नैनोसेकंड होती है — ग्यारह चक्र।
अनुवाद के बाद, प्रत्येक 128-बाइट लाइन के लिए एक अनुरोध क्रॉसबार के पार अपने भौतिक पते के फ़ंक्शन द्वारा चुने गए 36 में से एक 2 MiB L2 स्लाइस की ओर बढ़ता है। सभी स्लाइस समानांतर में सेवा कर सकते हैं, जिससे कुल बैंडविड्थ एक सिंगल स्लाइस का 36 गुना हो जाता है।
स्रोत: Hacker News · HeadlinesBriefing द्वारा सारांशित