0Pricing
CUDA Academy · درس

العنونة التسلسلية

خطوات خالية من التعارض في الذاكرة المشتركة.

العنونة التسلسلية درس مجاني في 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 يتضمن محرر أكواد مدمج، لذا تكتب وتشغل أكواداً حقيقية مباشرة في متصفحك وتحصل على تعليقات فورية من الذكاء الاصطناعي — بدون إعداد محلي.

جميع الدروس في هذه الدورة

  1. فكرة شجرة الاختزال
  2. القضاء على تفرّع Warp
  3. العنونة التسلسلية
  4. الاختزال النهائي متعدد الكتل
← العودة إلى CUDA Academy