CUDA Academy · 강의

레지스터와 로컬 메모리

스레드별로 가장 빠른 저장 공간과 그 초과 사용을 알아봅니다.

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

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

The Fastest Memory You Have

Every thread gets its own private registers, the quickest storage on the chip. They live right next to the math units, so reads cost almost nothing.

Plain Variables Become Registers

When you write a normal local variable in a kernel, the compiler usually keeps it in a register. No special syntax is needed at all. 🙂

__global__ void k() {
    int x = 5;   // x lives in a register
    float y = x * 2.0f;
}

Private to One Thread

A register value belongs to exactly one thread. Thread 7 cannot see thread 8's register, which is why per-thread work stays cleanly separated.

Registers Are a Shared Pool

Each streaming multiprocessor has a fixed register file. All resident threads split it, so using fewer registers per thread lets more threads run at once.

When You Run Out

If a kernel needs more registers than are available, the extra values get pushed elsewhere. This event is called a register spill.

Where Spills Go: Local Memory

Spilled values land in local memory, which despite the name is not on-chip. It actually lives in slow off-chip DRAM.

Local Is Still Per-Thread

Local memory is private to each thread just like registers. The difference is purely speed, since local memory is far slower to reach.

Big Local Arrays Spill

Declaring a large per-thread array often forces it into local memory, because there simply are not enough registers to hold every element.

__global__ void k() {
    float buf[64]; // likely lives in local memory
}

Spills Hurt Performance

Every spill turns a free register access into a slow DRAM trip. Cutting register pressure is one of the easiest optimization wins you can make.

Seeing Your Register Count

You can ask the compiler to report registers per thread. Pass --ptxas-options=-v to nvcc and it prints usage for each kernel.

nvcc --ptxas-options=-v kernel.cu

Capping Registers on Purpose

The __launch_bounds__ qualifier hints the compiler to limit registers, trading a little per-thread speed for more threads running together.

__global__ void __launch_bounds__(256)
myKernel() { /* ... */ }

Quick Check

You declared a big per-thread array and performance dropped. What likely happened?

Recap: Fast and Private

You learned that registers are the fastest per-thread storage, that the pool is limited, and that overflow spills into slow local memory. Keep register pressure low. 🎯

자주 묻는 질문

“레지스터와 로컬 메모리” 강의는 무료인가요?

네 — “레지스터와 로컬 메모리” 전체 내용을 이 웹사이트에서 무료로 읽을 수 있습니다. 인터랙티브하게 실습하려면(내장 코드 에디터와 24/7 AI 튜터), CoddyKit PRO로 업그레이드하면 CUDA Academy 강의 전체를 잠금 해제할 수 있습니다. CUDA Academy 강의에는 총 4개의 강의가 포함되어 있습니다.

“레지스터와 로컬 메모리”에서 뭘 배우나요?

스레드별로 가장 빠른 저장 공간과 그 초과 사용을 알아봅니다. 브라우저에서 직접 실행하는 실습 코드로 CUDA Academy을(를) 배우며, 24/7 AI 튜터가 강의를 진행하면서 질문에 답변해줍니다.

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

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

“레지스터와 로컬 메모리” 강의는 얼마나 걸리나요?

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

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

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

이 강의의 모든 강의

← CUDA Academy(으)로 돌아가기