العنونة التسلسلية
خطوات خالية من التعارض في الذاكرة المشتركة.
العنونة التسلسلية درس مجاني في CUDA Academy على CoddyKit. هذا هو الدرس 3 من أصل 4. يمكنك قراءة الدرس كاملاً أدناه مجاناً — ثم تمرن عليه مباشرة في المتصفح باستخدام محرر أكواد مدمج ومدرس ذكاء اصطناعي متاح 24/7. هذا الدرس جزء من مسار التعلم في CUDA Academy، وتقدمك يتزامن عبر الويب وتطبيق CoddyKit. تتضمن دورة CUDA Academy 4 دروس في المجموع.
بعض أجزاء هذا الدرس لم تُترجم بعد وتظهر باللغة الإنجليزية.
Shared Memory Has Banks
Shared memory is split into 32 banks, one per warp lane. When 32 threads hit 32 different banks, all reads happen in a single fast cycle.
Bank Conflicts Slow You Down
If two threads in a warp touch the same bank, that is a bank conflict. The hardware serializes those accesses, costing extra cycles.
Interleaved Addressing
The previous reduction used interleaved addressing: stride starts small and doubles, so partners are close together in shared memory.
int index = 2 * s * tid;
data[index] += data[index + s];Why Interleaving Conflicts
With small, doubling strides, several lanes in a warp map onto the same bank. Those accesses can no longer happen in one cycle.
Flip the Stride Order
Sequential addressing starts the stride large and halves it each step, the reverse of interleaving. This single change removes the conflicts.
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s)
data[tid] += data[tid + s];
__syncthreads();
}Big Stride, Clean Banks
A large stride spreads partner addresses far apart, so each lane lands on its own bank. The warp reads conflict-free in one cycle.
The tid < s Guard
Only the lower half of threads work each step, written as tid < s. That keeps active threads contiguous, so warps stay non-divergent too.
Two Wins at Once
Sequential addressing fixes bank conflicts and avoids warp divergence in the same kernel. One layout change, two performance problems solved.
Still Sync Each Step
You still need a __syncthreads after each step. Threads must see the previous level's writes before they read for the next level.
Result Lands at Index 0
As the stride halves toward zero, all partial sums fold into data[0]. Thread 0 then writes that block's result back to global memory.
A Classic Optimization
This pattern comes straight from NVIDIA's famous reduction guide. Sequential addressing is a textbook step toward a conflict-free kernel.
Quick Check
Think about why a large, halving stride beats a small, doubling one.
Recap
You swapped interleaved for sequential addressing: stride starts large and halves, killing bank conflicts and divergence at once. Up next: multi-block sums. ✨
الأسئلة الشائعة
هل درس «العنونة التسلسلية» مجاني؟
نعم — نص درس «العنونة التسلسلية» كامل متاح مجاناً هنا على الويب. لتمرينه بشكل تفاعلي (محرر أكواد مدمج ومدرس ذكاء اصطناعي متاح 24/7) وفتح باقي دورة CUDA Academy، انتقل إلى CoddyKit PRO. تتضمن دورة CUDA Academy 4 دروس في المجموع.
ماذا ستتعلم في «العنونة التسلسلية»؟
خطوات خالية من التعارض في الذاكرة المشتركة. تتمرن على CUDA Academy مع أكواد عملية تشغلها مباشرة في المتصفح، ومدرس ذكاء اصطناعي متاح 24/7 يجيب على أسئلتك أثناء عملك.
هل أحتاج إلى خبرة سابقة لأبدأ CUDA Academy؟
لا تُشترط خبرة سابقة. CUDA Academy على CoddyKit منظم للمبتدئين حتى المتقدمين، لذا يمكنك البدء من هنا أو من البداية والتقدم بسرعتك الخاصة. هذا هو الدرس 3 من أصل 4.
كم من الوقت يستغرق درس «العنونة التسلسلية»؟
معظم دروس CoddyKit تستغرق حوالي 5–10 دقائق. كل منها موجز وتفاعلي، لذا تحرز تقدماً مستمراً وتستأنف من حيث توقفت عبر الويب والتطبيق.
هل يمكنني كتابة وتشغيل أكواد في درس CUDA Academy هذا؟
نعم. كل درس في CUDA Academy يتضمن محرر أكواد مدمج، لذا تكتب وتشغل أكواداً حقيقية مباشرة في متصفحك وتحصل على تعليقات فورية من الذكاء الاصطناعي — بدون إعداد محلي.
جميع الدروس في هذه الدورة
- فكرة شجرة الاختزال
- القضاء على تفرّع Warp
- العنونة التسلسلية
- الاختزال النهائي متعدد الكتل