الدرس 25: scan متوازٍ — kernel المضاعفة
في الدرس السابق رأينا ما الذي يحسبه مجموع البادئات (prefix sum) وفكرة المضاعفة: في كل جولة يضيف كل عنصر جاره الذي يبعد d، وتتضاعف d (1، 2، 4...). هنا نبنيه ككود kernel حقيقي في CUDA. نعمل داخل block واحد باستخدام مصفوفة في الذاكرة المشتركة (__shared__) حيث يملك thread واحد موقعًا واحدًا. الصعوبة الح
يملك كل thread موقعًا واحدًا في المصفوفة المشتركة. في كل خطوة يقرأ أولًا القيمة الموجودة على بُعد d مواقع إلى يساره ويضعها جانبًا، وينتظر حتى يقرأ الجميع (__syncthreads)، وعندها فقط يضيفها إلى نفسه، ثم ينتظر مرة أخرى حتى يكتب الجميع قبل الخطوة التالية. الانتظارات هي ما يحافظ على الترتيب ويمنع الكتابة فوق القيم.
- Hillis-Steele
- kernel الـ scan بالمضاعفة في log2(n) خطوة: في الخطوة d، يضيف كل thread ذو tid>=d القيمة من tid-d. تتضاعف d في كل خطوة.
- tile في الذاكرة المشتركة
- مصفوفة __shared__ لكامل الـ block. يحمّل الـ scan المدخلات إليها، ويعمل عليها في مكانها عبر جميع الخطوات، وفي النهاية يكتب إلى out.
- حواجز القراءة/الكتابة
- __syncthreads() بعد القراءة (كي يقرأ الجميع قبل أي كتابة) وبعد الكتابة (كي تكون الكتابات مرئية للخطوة التالية).
- عدد الخطوات log2(n)
- الحلقة for(d=1; d<n; d*=2) تعمل log2(n) مرة، لأن d تتضاعف في كل خطوة حتى تصل إلى n.
