استخدام __shfl_down_sync للاختزال
نفّذ اختزالات Warp دون حواجز.
استخدام __shfl_down_sync للاختزال درس مجاني في CUDA Academy على CoddyKit. هذا هو الدرس 2 من أصل 4. يمكنك قراءة الدرس كاملاً أدناه مجاناً — ثم تمرن عليه مباشرة في المتصفح باستخدام محرر أكواد مدمج ومدرس ذكاء اصطناعي متاح 24/7. هذا الدرس جزء من مسار التعلم في CUDA Academy، وتقدمك يتزامن عبر الويب وتطبيق CoddyKit. تتضمن دورة CUDA Academy 4 دروس في المجموع.
تنقل Shuffle السجلات
تتيح shuffle لـ lane واحد قراءة قيمة سجل lane آخر مباشرةً. فتنتقل البيانات بين الخيوط دون ذاكرة مشتركة ودون حاجز بينهما.
تعرّفوا على shfl_down_sync
الأداة الأساسية للاختزالات هي shfl_down_sync. إذ يجلب كل lane قيمة من lane يقع أعلى منه بعدد ثابت من المواضع في warp.
float v = __shfl_down_sync(mask, val, offset);قراءة الوسائط
يأخذ الاستدعاء قناعًا نشطًا، والقيمة المراد مشاركتها، وإزاحة. ويتلقى lane i القيمة الموجودة في lane i زائد الإزاحة.
أضيفوا ما تتلقونه
لا يحتاج جمع warp إلا إلى إضافة القيمة التي جرى تبديلها. فيأخذ lane i رقم جاره ويدمجه في المجموع الجاري الخاص به.
val += __shfl_down_sync(mask, val, offset);تنصيف الإزاحة
مثل الاختزال الشجري، تبدأ الإزاحة عند 16 ثم تنخفض إلى النصف في كل خطوة: 16، 8، 4، 2، 1. وبعد خمس خطوات، يُجمع warp بأكمله.
for (int o = 16; o > 0; o >>= 1)
val += __shfl_down_sync(mask, val, o);سبب كفاية خمس خطوات
يضم warp عددًا قدره 32 من الـlanes، و2 مرفوعة إلى القوة الخامسة تساوي 32. لذلك تكفي خمس خطوات تنصيف لدمج القيم الـ32 كلها في قيمة واحدة.
تصل الإجابة إلى lane 0
بعد انتهاء الحلقة، يوجد مجموع warp الكامل في lane 0. أما الـlanes الأخرى فتحوي أجزاء غير مفيدة، لذا ينبغي لـ lane 0 وحده كتابة النتيجة.
لا حاجة إلى الحواجز
بما أن التبادل يحدث عبر السجلات بالتزامن، يمكنكم تخطي syncthreads بالكامل. فلقد تولت لاحقة sync تنسيق الـlanes المقنّعة.
أسرع من الذاكرة المشتركة
يتفوق اختزال warp باستخدام shuffle على الإصدار المعتمد على الذاكرة المشتركة: تعليمات أقل، ولا تعارضات بنوك، ولا حالات توقف بسبب حاجز. إنه المسار السريع للـ32 عنصرًا الأخيرة.
مرّروا القناع الصحيح
إذا كانت بعض الـlanes قد خرجت بالفعل، فسيكون استخدام القناع الكامل خطأً. التقطوا المجموعة النشطة باستخدام activemask كي تلامس كل عملية shuffle الـlanes الموجودة فقط.
اختزالات من مستويين
غالبًا ما تُجري الاختزالات الكبيرة اختزالًا لكل warp باستخدام عمليات shuffle، ثم تدمج نتائج warps في الذاكرة المشتركة. وتتعامل عمليات shuffle بكفاءة كبيرة مع المستوى الداخلي البسيط.
تحقق سريع
فكّروا في أيّ مسار يحتفظ بالإجابة عند انتهاء الحلقة.
مراجعة
جمعتم عناصر warp باستخدام shfl_down_sync: خفّضتم الإزاحة إلى النصف خمس مرات، فأصبح المسار 0 يحتفظ بالمجموع، من دون الحاجة إلى حواجز. التالي: الاقتراع والتصويت. ✨
الأسئلة الشائعة
هل درس «استخدام __shfl_down_sync للاختزال» مجاني؟
نعم — نص درس «استخدام __shfl_down_sync للاختزال» كامل متاح مجاناً هنا على الويب. لتمرينه بشكل تفاعلي (محرر أكواد مدمج ومدرس ذكاء اصطناعي متاح 24/7) وفتح باقي دورة CUDA Academy، انتقل إلى CoddyKit PRO. تتضمن دورة CUDA Academy 4 دروس في المجموع.
ماذا ستتعلم في «استخدام __shfl_down_sync للاختزال»؟
نفّذ اختزالات Warp دون حواجز. تتمرن على CUDA Academy مع أكواد عملية تشغلها مباشرة في المتصفح، ومدرس ذكاء اصطناعي متاح 24/7 يجيب على أسئلتك أثناء عملك.
هل أحتاج إلى خبرة سابقة لأبدأ CUDA Academy؟
لا تُشترط خبرة سابقة. CUDA Academy على CoddyKit منظم للمبتدئين حتى المتقدمين، لذا يمكنك البدء من هنا أو من البداية والتقدم بسرعتك الخاصة. هذا هو الدرس 2 من أصل 4.
كم من الوقت يستغرق درس «استخدام __shfl_down_sync للاختزال»؟
معظم دروس CoddyKit تستغرق حوالي 5–10 دقائق. كل منها موجز وتفاعلي، لذا تحرز تقدماً مستمراً وتستأنف من حيث توقفت عبر الويب والتطبيق.
هل يمكنني كتابة وتشغيل أكواد في درس CUDA Academy هذا؟
نعم. كل درس في CUDA Academy يتضمن محرر أكواد مدمج، لذا تكتب وتشغل أكواداً حقيقية مباشرة في متصفحك وتحصل على تعليقات فورية من الذكاء الاصطناعي — بدون إعداد محلي.
جميع الدروس في هذه الدورة
- Warps وLanes وMasks
- استخدام __shfl_down_sync للاختزال
- دوال Ballot وVote
- Cooperative Groups