④ 스케줄링

블록이 SM(공장)에 배정된 뒤, 그 안에서 일꾼 조(warp)를 "어떤 순서로 어떻게 굴리는가" — 그리고 여러 커널이 한 GPU를 두고 어떻게 경쟁하는가. 이게 성능(빠르냐)과 응답성(급한 일을 제때 하냐)의 핵심입니다.

4.1 Warp 스케줄링 — "느린 메모리를 다른 일로 가린다"

①에서 살짝 본 핵심을 제대로 풀어볼게요. SM 안의 warp 스케줄러(작업 반장, 보통 4명)는 매 순간(클럭 사이클마다) "지금 일할 준비가 된 warp"를 골라 명령을 내립니다.

포인트는 여기예요. 어떤 warp가 창고(global memory)에서 데이터를 기다리며 멈추면(수백 사이클이나 걸림!), 반장은 그 조를 기다리지 않고 즉시 다른 준비된 warp에게 일을 시킵니다. 이 조 바꾸기에 비용이 거의 0이에요(CPU의 문맥 전환과 달리).

warp A: [메모리 대기...........]        ← A는 창고 기다리는 중
warp B:      [계산][계산]               ← 그동안 B를 시킴
warp C:            [계산][계산]          ← C도 시킴
warp D:                  [계산]          ← D도 → 계산기(ALU)가 안 놂!

그럼 반장(warp 스케줄러)은 매 사이클 정확히 무엇을 할까요? 4단계가 무한 반복됩니다:

4.2 Occupancy 튜닝 — 실제로 어떻게 올리나

①에서 본 occupancy를 코드에서 끌어올리는 실전 방법입니다. occupancy를 막는 건 보통 "일꾼이 자원을 너무 많이 쓰는" 경우예요:

막히는 원인	올리는 법
스레드가 레지스터(메모장)를 너무 많이 씀	변수를 줄이거나, __launch_bounds__로 컴파일러에 레지스터 상한을 힌트
블록이 shared memory(작업대)를 너무 많이 씀	타일 크기·버퍼를 줄임
블록 크기가 부적절	128~256 스레드 등으로 조정 (NVIDIA의 occupancy calculator 활용)

일일이 계산하기 귀찮다면, CUDA가 직접 계산해주는 API도 있습니다:

int minGrid, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGrid, &blockSize, myKernel);
// → "이 커널은 블록 크기 blockSize일 때 occupancy가 최대"를 런타임이 알려줌
myKernel<<<(n + blockSize - 1) / blockSize, blockSize>>>(...);

커널의 레지스터·shared memory 사용량을 자동으로 고려해주므로, grid-stride loop와 함께 쓰면 "어떤 GPU에서든 알아서 적절한 구성"이 됩니다. 실제 사용량 확인은 nvcc --ptxas-options=-v로.

4.3 스트림 우선순위 — 급한 일 먼저

여러 스트림(줄)이 동시에 일할 때, "이 작업은 급하니 먼저 해줘"라고 표시하고 싶다면 우선순위 스트림을 만듭니다. 자원이 비었을 때, 우선순위가 높은 스트림의 블록이 먼저 SM에 올라가요.

int lo, hi;
cudaDeviceGetStreamPriorityRange(&lo, &hi);  // hi가 더 높은 우선순위(값은 더 작음)
cudaStream_t urgent;
cudaStreamCreateWithPriority(&urgent, cudaStreamNonBlocking, hi);
// 급한 커널을 urgent 스트림에 넣으면 우선 배정됨

주의: 이건 "다음에 올릴 블록"을 고를 때 영향을 줄 뿐, 이미 돌고 있는 걸 멈추진 못합니다(다음 절).

4.4 선점(preemption)이 약하다 — GPU의 중요한 한계

선점(preemption)이란 "돌고 있는 일을 잠깐 멈추고 더 급한 일을 끼워넣는" 것입니다. CPU는 이걸 아주 잘해요(운영체제가 수시로 스레드를 갈아끼움). 그런데 GPU는 이게 약합니다.

이유는 ①에서 본 구조 때문이에요. 블록은 한 번 SM에 올라가면 자원을 차지한 채 끝까지 실행(run-to-completion)됩니다. 그래서 거대한 커널의 블록들이 모든 SM을 꽉 채우고 있으면, 뒤늦게 온 급한 커널은 자리가 빌 때까지 그냥 기다려야 해요. (이걸 head-of-line blocking — 앞 차가 막아서 뒤가 못 가는 상황이라고 합니다.)

	CPU	GPU
선점	강함 — 급한 일이 즉시 끼어듦	약함 — 블록은 보통 끝까지 실행
우선순위 효과	실행 중인 것도 밀어냄	주로 "다음에 올릴 것"에만 영향

(최신 GPU(Pascal 이후)엔 명령 단위 부분 선점이 있긴 한데, 거칠고 비쌉니다.) 그래서 응답성이 필요하면 이렇게 대처합니다:

• 커널을 작게 쪼개기 — 블록이 빨리 끝나면 SM 자리가 자주 비어서, 급한 일이 곧 올라갈 수 있어요.
• 스트림 우선순위 — 자리 빌 때 급한 블록부터.
• oversubscribe 금지 — GPU가 감당할 양보다 일을 더 쌓지 않기.

4.5 MIG · MPS — GPU를 나눠 쓰기

그럼 여러 프로그램(프로세스)이 한 GPU에 동시에 일을 시키면 어떻게 될까요? 의외로 기본 동작은 "동시에"가 아닙니다 — GPU는 프로세스마다 별도의 컨텍스트(작업 환경)를 두고, 이들을 시간 분할(time-slicing)로 번갈아 실행합니다. A 프로세스의 커널이 도는 동안 B는 통째로 대기하고, 컨텍스트를 바꾸는 비용도 들어요. 그래서 각 프로세스가 GPU를 절반도 못 채우는 작은 커널을 던지면, GPU가 절반씩 놀면서도 서로 기다리는 비효율이 생깁니다.

이 문제를 푸는 두 가지 방법이 MPS와 MIG입니다:

	MPS (Multi-Process Service)	MIG (Multi-Instance GPU)
방식	여러 프로그램의 커널을 합쳐 한 GPU에서 동시 실행	GPU를 물리적으로 여러 조각으로 분할(SM·메모리 격리)
격리(간섭 차단)	약함 — 서로 성능 영향 줄 수 있음	강함 — 하드웨어로 완전 분리
쓰는 곳	작은 커널 여러 개의 GPU 활용률↑	"이 조각은 이 작업 전용" 같은 보장(멀티테넌시·QoS)
지원	대부분의 GPU	A100/H100 등 데이터센터 GPU

쉽게 말하면 — MPS는 "한 주방을 같이 쓰되 주문을 합쳐 처리", MIG는 "주방을 벽으로 나눠 각자 전용 주방"입니다.