③ CUDA 동작 방식 — 실행 모델

②에서 kernel<<<...>>>() 한 줄을 썼죠. 그런데 그 한 줄을 실행하면 실제로 무슨 일이 벌어질까요? CPU가 GPU에게 일을 시키고, 그 일이 일꾼에게 도착하기까지의 전 과정을 따라가 봅니다. 이걸 알면 성능 문제가 왜 생기는지, 어떻게 푸는지가 보입니다.

3.1 커널 launch의 실체 — "주문서를 큐에 넣는다"

먼저 비유로요. CUDA 커널 호출은 마치 식당에서 주방에 주문서를 넣는 것과 같습니다. 손님(CPU)은 주문서를 카운터에 올려놓고 음식이 나오기 전에 자기 할 일을 하러 갑니다. 주방(GPU)은 주문서를 받아 요리하죠. kernel<<<...>>>() 한 줄이 실행되면 내부에선 이런 단계가 일어납니다:

① CPU(손님):   "이 커널을 이 구성으로 실행해줘" 주문서 작성
② 드라이버:    주문서를 GPU의 하드웨어 큐(대기열)에 넣음(enqueue)
               ⚠️ 이 '넣는 행위' 자체가 ~µs(마이크로초)의 오버헤드
③ CPU:         음식 기다리지 않고 다음 줄로 진행 (← 비동기!)
④ GPU(주방):   큐에서 커널을 꺼내 grid(일감 전체)를 펼침
               → 'GigaThread 엔진'이 블록들을 각 SM(공장)에 분배
⑤ 각 SM:       배정받은 블록의 warp(32명 조)를 스케줄해 실제 실행

그림 — 커널 호출은 "큐에 주문 제출". CPU는 안 기다리고 다음 줄로 가고, GPU가 알아서 블록을 SM에 펼쳐 실행한다

⚠️ 가장 중요: 커널 실행은 "비동기"다

CPU는 커널이 끝나기를 기다리지 않습니다 — 주문서만 넣고 바로 다음 줄로 가요. 그래서 커널 바로 다음 줄에서 결과를 읽으려 하면 아직 요리가 안 끝났을 수 있습니다(버그!). 결과가 필요하면 "다 될 때까지 기다려"라고 명시해야 해요:

cudaMemcpy(...) — 결과를 복사해오는데, 이게 자동으로 완료를 기다려줍니다(그래서 ②장 예제는 안전했어요).
cudaDeviceSynchronize() — "GPU의 모든 일이 끝날 때까지 여기서 대기".

👤 CPU 개발자라면 — "함수 호출인데 왜 안 기다리지?"

CPU에서 foo()를 부르면, foo가 끝나야 다음 줄이 실행되죠(동기). 그래서 kernel<<<>>>() 다음 줄에서 결과를 읽으려다 "값이 이상해요" 하는 일이 흔합니다. CUDA 커널 호출은 함수 호출이 아니라 "비동기 작업 제출"이에요 — std::async나 작업 큐에 일을 던지는 것에 가깝습니다. CPU는 제출만 하고 바로 다음 줄로 가버려요. 그래서 결과가 필요한 시점엔 반드시 "끝날 때까지 기다려"를 명시해야 합니다(cudaMemcpy는 자동으로 기다려줌, 아니면 cudaDeviceSynchronize()). 또 하나 — 커널 안에서 난 에러는 호출한 그 줄이 아니라 나중에 동기화하는 줄에서 보고됩니다. 그래서 에러 체크를 cudaDeviceSynchronize() 직후에 하는 게 정석이에요.

💡 ②의 launch 오버헤드를 기억하세요

위 단계 ②의 ~µs 오버헤드는 큰 커널 하나엔 무시할 만하지만, 작은 커널을 수십 개 반복하면 이게 쌓여서 병목이 됩니다. 이 문제를 통째로 해결하는 게 ⑤ CUDA Graphs예요. 지금은 "launch엔 공짜가 아닌 비용이 있다"만 기억!

3.2 비동기 · 스트림 — 여러 줄을 동시에

커널이 비동기라는 성질을 무기로 쓸 수 있습니다. 바로 스트림(stream)이에요. 스트림은 쉽게 말하면 "주문서를 넣는 줄(대기열)"입니다.

같은 스트림(같은 줄)의 작업들은 순서대로 처리됩니다.
다른 스트림(다른 줄)끼리는 동시에 처리될 수 있어요 — 주방에 창구가 여러 개인 셈.

이걸 이용하면 "데이터 복사"와 "계산"을 겹쳐서(overlap) 시간을 아낄 수 있습니다. 한쪽 줄에서 복사하는 동안 다른 줄에서 계산하는 거죠:

cudaStream_t s1, s2;
cudaStreamCreate(&s1);  cudaStreamCreate(&s2);

// 데이터를 둘로 쪼개, 청크0은 s1 줄에서 청크1은 s2 줄에서 처리
cudaMemcpyAsync(dA0, hA0, sz, cudaMemcpyHostToDevice, s1);
kernel<<<b, t, 0, s1>>>(dA0);
cudaMemcpyAsync(dA1, hA1, sz, cudaMemcpyHostToDevice, s2);
kernel<<<b, t, 0, s2>>>(dA1);

그림 — 스트림 2개로 나누면 "복사"와 "계산"이 다른 장치에서 동시에 진행되어 전체 시간이 준다 (pinned memory 필수)

👤 CPU 개발자라면 — "스트림 = 스레드"가 아닙니다

스트림을 CPU의 워커 스레드처럼 생각하기 쉬운데, 다릅니다. 스트림은 일을 하는 주체가 아니라 "순서가 보장되는 작업 대기열(큐)"일 뿐이에요. 실제 일은 GPU의 SM들이 하고, 스트림은 그저 "이 큐에 넣은 것들은 순서대로, 다른 큐와는 병렬로 처리해도 됨"이라는 의존성 표시입니다. 그래서 "스트림 8개 = 8배 빠름"이 아니라, 스트림은 "복사하는 동안 계산도 같이 돌려서 빈 시간을 없애는" 용도예요(자원이 남을 때만 효과). CPU의 작업 큐 + 비동기 실행 모델에 가장 가깝습니다.

📌 조건

비동기 복사(cudaMemcpyAsync)가 진짜로 겹치려면, 호스트 메모리가 pinned memory여야 합니다. 일반 메모리면 효과가 제한돼요.

⚠️ "기본 스트림(default stream)"의 함정

스트림을 안 만들고 그냥 kernel<<<b, t>>>()로 실행하면 기본 스트림(스트림 0)에 들어갑니다. 그런데 이 기본 스트림은 보통 줄이 아니라 "다른 모든 스트림과 동기화되는" 특별한 줄이에요(legacy 동작). 기본 스트림에 작업이 들어가는 순간, 앞서 만든 s1·s2의 병렬 진행이 전부 멈췄다 가게 됩니다 — "스트림을 썼는데 overlap이 안 돼요"의 단골 원인이죠.

대처는 둘 중 하나 — ① overlap을 원하는 작업은 전부 명시적 스트림에 넣기(기본 스트림 안 쓰기), ② 컴파일 시 nvcc --default-stream per-thread로 기본 스트림을 "보통 스트림"으로 바꾸기.

3.3 이벤트 · 동기화 — 시간 재기와 줄 맞추기

여러 스트림(줄)이 돌아갈 때, "언제 끝났나" 측정하거나 "이 줄은 저 줄을 기다려라" 조율하려면 이벤트(event)를 씁니다. 이벤트는 작업 흐름 속의 "표시 깃발" 같은 거예요.

시간 측정 — 깃발 두 개를 꽂고 그 사이 GPU 시간을 잼. CPU 타이머보다 정확합니다.
스트림 간 의존성 — cudaStreamWaitEvent로 "B 줄은 A 줄의 이 깃발까지 기다려".
전체/부분 동기화 — cudaDeviceSynchronize()(모든 일), cudaStreamSynchronize(s)(한 줄만), cudaEventQuery(e)(기다리지 않고 "끝났어?"만 물어보기).
CPU 함수 끼워넣기 — cudaLaunchHostFunc(s, fn, arg)로 스트림 흐름 속에 "이 시점에 CPU 함수 fn을 불러줘"를 넣을 수 있어요(예: GPU 결과가 준비되는 순간 로그 남기기). 단 그 함수 안에서 CUDA 호출은 금지.

// 커널 실행 시간 재기
cudaEvent_t start, stop;
cudaEventCreate(&start);  cudaEventCreate(&stop);
cudaEventRecord(start);        // 시작 깃발
kernel<<<b, t>>>();
cudaEventRecord(stop);         // 끝 깃발
cudaEventSynchronize(stop);    // 끝 깃발까지 대기
float ms;  cudaEventElapsedTime(&ms, start, stop);  // 걸린 시간(ms)

그림 — 이벤트(🚩깃발)의 두 용도: 깃발 두 개 사이의 GPU 시간을 재거나(위), 다른 스트림이 그 깃발까지 기다리게 해 순서를 보장한다(아래)

3.4 블록 → SM 배정 — 일감을 공장에 나눠주기

앞 단계 ④에서 grid가 펼쳐지면, GPU 안의 GigaThread 엔진(=일감 배분 담당)이 블록들을 SM(공장)에 나눠줍니다. 규칙은 단순하지만 중요해요:

블록 하나는 한 SM에만 배정됩니다 — 절대 두 공장에 쪼개 보내지 않아요(그래서 같은 블록끼리만 공용 작업대를 쓸 수 있는 거죠).
SM 하나는 블록을 여러 개 동시에 받습니다 — 자원(레지스터·shared·warp 슬롯)이 허락하는 만큼.
블록이 끝나서 자원이 비면, 대기 중이던 다음 블록이 그 SM에 올라갑니다 (블록이 많아도 순차적으로 소화).

Grid의 블록들: [B0][B1][B2][B3][B4][B5]...   (블록 수 > SM 수일 수 있음)
                 │   │   │   │
                SM0 SM1 SM0 SM1   ← 자리 나는 대로 배정, 끝나면 다음 블록 투입

🧠 여기서 나오는 중요한 규칙

블록이 어느 SM에, 어떤 순서로 배정될지는 보장이 없습니다. 그래서 블록은 서로 독립적이어야 하고, "블록 0이 끝나야 블록 1이 일한다" 같은 가정을 하면 안 돼요. 블록끼리 협력이 꼭 필요하면 → 커널을 두 개로 나누거나(중간에 동기화), cooperative groups의 grid 동기화 같은 특수 기능을 써야 합니다. 그리고 이 배정된 일을 SM이 실제로 어떻게 굴리는지가 바로 다음 장, ④ 스케줄링입니다.

✅ ③ 장 정리

① 커널 호출 = 비동기로 주문서를 큐에 넣기(~µs 비용). CPU는 안 기다리니 결과 읽기 전엔 동기화 필요.
② 스트림(줄)으로 복사·계산을 겹쳐 시간 절약, 이벤트(깃발)로 측정·조율.
③ grid가 펼쳐지면 GigaThread가 블록을 SM에 분배, 블록은 독립적이어야 함.
다음은 SM이 그 블록 안 warp를 "어떤 순서로, 어떻게" 굴리는가입니다.