CUDA Academy · 강의

순차 주소 지정

공유 메모리에서 충돌 없는 보폭을 사용합니다.

순차 주소 지정은(는) CoddyKit의 무료 CUDA Academy 강의입니다. 이것은 4개 중 3번째 강의입니다. 아래에서 전체 강의를 무료로 읽을 수 있으며, 내장 코드 에디터와 24/7 AI 튜터와 함께 브라우저에서 직접 실습할 수 있습니다. 이 강의는 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 AI 튜터), CoddyKit PRO로 업그레이드하면 CUDA Academy 강의 전체를 잠금 해제할 수 있습니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.

“순차 주소 지정”에서 뭘 배우나요?

공유 메모리에서 충돌 없는 보폭을 사용합니다. 브라우저에서 직접 실행하는 실습 코드로 CUDA Academy을(를) 배우며, 24/7 AI 튜터가 강의를 진행하면서 질문에 답변해줍니다.

CUDA Academy을(를) 시작하는 데 경험이 필요한가요?

사전 경험은 필요하지 않습니다. CoddyKit의 CUDA Academy은(는) 초급자부터 고급 학습자까지를 위해 구성되어 있으므로, 여기서 시작하거나 처음부터 시작할 수 있으며 자신의 속도대로 진행할 수 있습니다. 이것은 4개 중 3번째 강의입니다.

“순차 주소 지정” 강의는 얼마나 걸리나요?

대부분의 CoddyKit 강의는 약 5~10분이 소요됩니다. 각 강의는 간결하고 인터랙티브하여 꾸준한 진행이 가능하며, 웹과 앱에서 중단한 부분부터 바로 시작할 수 있습니다.

이 CUDA Academy 강의에서 코드를 작성하고 실행할 수 있나요?

네. 모든 CUDA Academy 강의에는 내장 코드 에디터가 포함되어 있으므로, 브라우저에서 바로 실제 코드를 작성하고 실행한 후 즉시 AI 피드백을 받을 수 있습니다 — 로컬 설정이 필요 없습니다.

이 강의의 모든 강의

← CUDA Academy(으)로 돌아가기