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 включает встроенный редактор кода, поэтому ты пишешь и запускаешь реальный код прямо в браузере и получаешь моментальную обратную связь от AI — локальная установка не требуется.

Все уроки этого курса

← Назад к CUDA Academy