CUDA Academy · 강의

shared 배열 선언하기

블록마다 빠른 임시 메모리를 사용합니다.

__shared__ 배열 선언하기은(는) CoddyKit의 무료 CUDA Academy 강의입니다. 이것은 4개 중 1번째 강의입니다. 아래에서 전체 강의를 무료로 읽을 수 있으며, 내장 코드 에디터와 24/7 AI 튜터와 함께 브라우저에서 직접 실습할 수 있습니다. 이 강의는 CUDA Academy 학습 경로의 일부이며, 진행 상황이 웹과 CoddyKit 앱에 동기화됩니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.

이 강의의 일부는 아직 번역되지 않았으며 영어로 표시됩니다.

A Scratchpad Per Block

Every block gets a tiny, fast pool of on-chip memory called shared memory. Think of it as a scratchpad the whole block can write to and read from together.

Why It Is So Fast

Shared memory sits right on the streaming multiprocessor, so it is roughly 100x faster than global memory. Use it to avoid hammering slow off-chip DRAM. 🚀

The shared Keyword

You declare it with the __shared__ qualifier inside a kernel. This one array is then shared by every thread in the block.

__global__ void kern() {
    __shared__ float tile[256];
}

Fixed Size at Compile Time

When you give a size in brackets, that is static shared memory. The compiler must know the size, so it has to be a constant, not a runtime value.

__shared__ int counts[128];

One Copy, Not One Per Thread

This is the key idea: a __shared__ array is created once per block, not once per thread. All threads see the exact same array.

Threads Cooperate Through It

Because every thread sees the same data, shared memory lets threads cooperate. One thread can stash a value and a neighbor can pick it up.

Block Scope, Not Beyond

Its lifetime matches the block. The array is born when the block starts and gone when it ends, so it has block scope only. 🧱

Each Thread Owns a Slot

A common pattern is one slot per thread, indexed by threadIdx.x. Each thread loads its element into shared memory in parallel.

__shared__ float s[256];
s[threadIdx.x] = input[i];

A Tiny but Precious Resource

Shared memory is small, often just 48 to 100 KB per SM. Asking for too much per block reduces how many blocks can run at once.

The Classic Use: Staging Tiles

The most common job is staging a tile of global data so the block can reuse it many times without going back to slow DRAM.

Not Visible to Other Blocks

Remember the boundary: shared memory is private to its block. Two different blocks each get their own separate copy and cannot peek at each other.

Quick Check

Let us check how shared memory is scoped.

Recap

You learned that __shared__ gives each block a fast on-chip scratchpad, created once per block and ideal for staging reusable data. Next: keeping threads in step. 🎯

자주 묻는 질문