الدرس 15: تشخيص وإصلاح الوصول ذي الـ stride
في الدرس السابق رأينا ما هو الـ coalescing: عندما تلمس threads متتالية داخل warp عناوين متتالية، تخدمها جميعًا معاملة واحدة، وعندما تقفز بمقدار stride كبير يستهلك الـ warp معاملات كثيرة مهدورة. الآن نتعلم المهارة العملية — رصد وصول غير موحَّد داخل kernel، وإصلاحه. العلامة الكلاسيكية هي فهرس يضرب thr
تخيّل 4 عمّال يقطفون التفاح من صفوف الأشجار. في البداية يركض كل عامل إلى عمود مختلف في طرف بعيد من البستان — 4 ركضات طويلة. أما إذا أخذ كل عامل التفاحة التالية في صفٍّ واحد متلاصق، فسيعملون جميعًا جنبًا إلى جنب في مرور واحد. لم نغيّر البستان — بل فقط مَن-يقطف-ماذا.
- الوصول ذو الـ stride
- نمط تقفز فيه threads متتالية بمقدار stride كبير (على سبيل المثال data[threadIdx.x * N])، فتكون عناوينها متباعدة ويقع كلٌّ منها في معاملة ذاكرة مختلفة.
- موحَّد مقابل غير موحَّد
- موحَّد: عناوين threads المتجاورة متتالية، معاملة واحدة لكامل الـ warp. غير موحَّد: العناوين مبعثرة، ويستهلك الـ warp معاملات كثيرة، معظمها مهدور.
- إصلاح إعادة تعيين الفهرس
- إصلاح الوصول المبعثر: اجعل threads المتجاورة تلمس عناوين متجاورة — أعطِ الـ thread رقم t العنصر base + t بدلًا من base + t*N، كل ذلك في الذاكرة العامة.
- ترتيب row-major
- تخطيط تكون فيه عناصر الصف نفسه متجاورة في الذاكرة. المرور صفًّا صفًّا يجعل threads المتجاورة تلمس عناوين متجاورة (موحَّد)؛ أما المرور عمودًا عمودًا فيقفز بمقدار width (مبعثر).
