تابعت تدوينتنا السابقة نواة جمع المتجهات من 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 بايت، وسطر ذاكرة تخزين مؤقت واحد، وترجمة عنوان واحدة، وعبورًا واحدًا عبر جهاز التقاطع (crossbar)، وأحد الشرائح الستة والثلاثين من L2، وعندما يفشل في كل مكان، تنشيطًا واحدًا وأربع قراءات أعمدة في شريحة DRAM. يعيش الورب الخاص بنا على أحد الأقسام الفرعية الأربعة لـ SM إلى جانب أحد عشر وربًا مقيمًا آخر. تُصدر التعليمة إلى وحدة التحميل/التخزين (LSU)، التي ترسل رمز العملية وقناع المسارات النشطة والعناوين المحسوبة ورقم السجل إلى المُجمِّع (coalescer).
يحدد المُجمِّع أن هناك حاجة إلى 4 طلبات قطاعات متجاورة للـ 128 بايت التي طلبها الورب. عند الدخول إلى ذاكرة التخزين المؤقت L1، يُرسل الطلب إلى ذاكرة التخزين المؤقت ذات العنونة الافتراضية والمترابطة بالمجموعات رباعية الاتجاهات. تقارن عملية البحث العلامات في مجموعات من أربع خانات. وبما أن هذا تحميل لأول مرة، فإنه يفشل وينزل أكثر. تعود إصابة L1 في حوالي 15.4 نانوثانية — 40 دورة. تتطلب الذاكرة الافتراضية ترجمة قبل مغادرة SM. تحتفظ TLB الخاصة بـ SM بست عشرة ترجمة حديثة. تكلف الأخطاء حوالي 4.4 نانوثانية — إحدى عشرة دورة.
بعد الترجمة، يتقدم طلب واحد لكل سطر بحجم 128 بايت عبر جهاز التقاطع إلى إحدى شرائح L2 البالغة 2 MiB والبالغ عددها 36، يتم اختيارها بواسطة دالة تعتمد على عنوانها الفيزيائي. يمكن لجميع الشرائح الخدمة بالتوازي، مما يمنح عرض النطاق الترددي الإجمالي 36 ضعف شريحة واحدة.
المصدر: Hacker News · لخّصه HeadlinesBriefing