GPU 아키텍처
GPU(Graphics Processing Unit)는 이름 그대로 화면의 수많은 픽셀을 한꺼번에 계산하려고 만들어진 프로세서입니다. 픽셀마다 하는 계산은 비슷하고 서로 독립이라서, 단순한 연산 유닛을 아주 많이 두고 같은 계산을 수많은 데이터에 동시에 적용하는 구조가 알맞았습니다. CPU가 몇 개의 크고 영리한 코어로 한 작업을 빨리 끝내는 데 맞춰져 있다면, GPU는 단순한 코어를 수천 개 두고 많은 작업을 한꺼번에 처리하는 처리량에 맞춰져 있습니다. 행렬 곱셈 같은 계산이 픽셀 계산과 닮아 있어, 오늘날 GPU는 그래픽뿐 아니라 과학 계산과 딥러닝에도 널리 쓰입니다.
NVIDIA GPU에서는 스레드 32개가 워프(warp)로 묶여 한 번에 같은 명령어를 실행합니다. 이 방식을 SIMT(Single Instruction, Multiple Threads)라고 부릅니다. CPU의 SIMD와 비슷하지만 스레드마다 분기할 수 있다는 점이 다릅니다. 대신 한 워프 안에서 분기가 갈리면 각 경로가 차례로 실행되어 효율이 떨어집니다. 또 GPU에도 캐시가 있지만 CPU처럼 큰 캐시에 면적을 쓰지 않고, 메모리를 기다리는 워프를 두고 다른 워프를 실행하는 식으로 지연을 숨깁니다. 이 토픽에서는 CPU와 GPU의 구조 차이, 워프와 분기, 지연 숨기기를 봅니다.
핵심 개념
지연 지향과 처리량 지향
CPU는 한 스레드의 지연을 줄이려고 큰 캐시, 분기 예측, 비순차 실행 같은 장치에 칩 면적을 씁니다. GPU는 그 면적을 연산 유닛에 더 많이 써서 같은 시간에 처리하는 일의 양(처리량)을 키웁니다.
핵심 개념
SIMT와 워프
NVIDIA GPU는 스레드 32개를 묶은 워프 단위로 명령어를 하나씩 실행합니다. 워프 안의 스레드는 같은 명령을 각자의 데이터로 실행하고, 분기가 갈리면 각 경로를 차례로 실행하며 해당 경로가 아닌 스레드는 꺼 둡니다.
핵심 개념
스레드 계층과 SM
CUDA에서는 스레드를 블록으로, 블록을 그리드로 묶어 실행합니다. 블록은 GPU의 스트리밍 멀티프로세서(SM) 하나에 배정되어 그 안의 워프들이 번갈아 실행됩니다.
핵심 개념
GPU의 메모리 구조
모든 스레드가 접근하는 전역 메모리, 블록 안의 스레드가 공유하는 공유 메모리, 스레드마다 쓰는 레지스터로 나뉩니다. 공유 메모리는 전역 메모리보다 훨씬 가까워서 데이터를 재사용할 때 쓰입니다.
핵심 개념
지연 숨기기
한 워프가 메모리를 기다리면 SM이 곧바로 다른 워프를 실행합니다. 실행할 수 있는 워프가 충분히 많으면 기다리는 시간이 다른 워프의 계산으로 채워져 연산 유닛이 쉬지 않습니다.
핵심 개념
호스트와 장치 사이의 전송
CPU(호스트)의 메모리와 GPU(장치)의 메모리는 따로이고, 데이터를 PCIe 같은 연결로 옮겨야 합니다. 계산이 이 전송 시간보다 충분히 크지 않으면 GPU로 옮겨도 이득이 작을 수 있습니다.
GPU의 배경
GPU는 3차원 그래픽의 계산을 빠르게 하려는 전용 하드웨어에서 출발했습니다. 정점과 픽셀을 하나씩 처리하는 일은 서로 독립이라서 병렬화하기 쉬웠고, 그래픽 처리 장치는 점점 프로그래밍할 수 있게 되었습니다. 그래픽이 아닌 계산에도 GPU를 쓰는 범용 GPU 컴퓨팅(GPGPU)이 CUDA 같은 프로그래밍 환경과 함께 퍼지면서, 과학 계산과 딥러닝으로 쓰임이 넓어졌습니다.
숫자로 확인하기
자주 하는 오해
면접에서 이렇게 물어요
면접 질문
CPU와 GPU의 구조 차이는 무엇인가요?
CPU는 큰 캐시와 복잡한 제어로 한 스레드를 빠르게 처리하는 지연 지향이고, GPU는 단순한 코어를 수천 개 두고 같은 연산을 많은 데이터에 동시에 적용하는 처리량 지향입니다.
면접 질문
워프와 워프 분기란 무엇인가요?
NVIDIA GPU에서 스레드 32개 묶음이 한 번에 같은 명령어를 실행하는 단위가 워프입니다. 한 워프 안에서 조건문으로 경로가 갈리면 각 경로가 차례로 실행되어 성능이 떨어지는 것이 워프 분기입니다.
면접 질문
GPU는 메모리 지연을 어떻게 다루나요?
GPU에도 캐시가 있지만 지연을 주로 많은 워프를 번갈아 실행해 숨깁니다. 한 워프가 메모리를 기다리는 동안 다른 워프가 계산하므로 연산 유닛이 쉬지 않습니다.
면접 질문
GPU에 알맞지 않은 작업은?
조건 분기가 많아 스레드마다 경로가 갈리거나 앞 결과에 의존하는 순차 작업, 데이터가 작아 CPU와 GPU 사이 전송 시간이 더 큰 작업입니다.
실제 시스템에서는
함께 보면 좋은 토픽
참고 자료
CUDA 벡터 덧셈 — 스레드 하나가 원소 하나
// 커널: 스레드마다 원소 하나를 더함 __global__ void add(const float *a, const float *b, float *c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; } // 호스트: 블록 256개 스레드, 원소 수에 맞춰 블록 수를 계산해 실행 int threads = 256; int blocks = (n + threads - 1) / threads; add<<<blocks, threads>>>(d_a, d_b, d_c, n);
커널은 모든 스레드가 같은 코드를 실행하고, blockIdx와 threadIdx로 자기가 맡을 원소 번호 i를 계산합니다. 원소마다 하는 일이 같고 서로 독립이라서 GPU의 데이터 병렬 구조에 잘 맞습니다. 원소 수 n이 블록 크기의 배수가 아닐 수 있어서 마지막 블록의 남는 스레드를 if (i < n)으로 막습니다. n이 32의 배수가 아니면 경계에 걸친 워프 하나에서만 분기가 생기고, 나머지 워프는 모두 같은 경로를 탑니다. <<<blocks, threads>>>는 CUDA의 실행 구성 문법으로, 한 블록의 스레드 수는 하드웨어 한계(보통 1024개) 이하여야 합니다.
워프 분기 — 같은 워프가 다른 경로를 탈 때
// 나쁜 예: 같은 워프 안에서 홀수/짝수 스레드가 서로 다른 경로 if (threadIdx.x % 2 == 0) { pathA(); } else { pathB(); } // 나은 예: 워프(32개) 단위로 경로가 갈림 → 워프 안에서는 모두 같은 경로 if ((threadIdx.x / 32) % 2 == 0) { pathA(); } else { pathB(); }
첫 조건은 같은 워프의 스레드 32개 중 절반은 pathA, 절반은 pathB로 갈립니다. 워프는 두 경로를 차례로 실행하고 해당하지 않는 스레드는 꺼 두므로 실행이 두 번으로 늘어납니다. 둘째 조건은 스레드 번호를 32로 나눈 몫으로 경로를 정해 한 워프의 스레드가 모두 같은 경로를 탑니다. 워프 안에서는 분기가 갈리지 않으므로 한 경로만 실행됩니다. 이처럼 데이터나 스레드 번호를 정리해 같은 워프가 같은 경로를 타게 만드는 것이 GPU 코드를 빠르게 하는 일반적인 방법입니다.
PyTorch — 텐서를 GPU로 옮겨 계산
import torch device = "cuda" if torch.cuda.is_available() else "cpu" a = torch.randn(4096, 4096).to(device) b = torch.randn(4096, 4096).to(device) c = a @ b # 행렬 곱셈이 GPU에서 실행됨 torch.cuda.synchronize() # GPU 작업이 끝날 때까지 기다림(시간 측정 시)
행렬 곱셈은 원소마다 독립적인 곱셈과 덧셈이 많아 GPU에 알맞은 계산입니다. to(device)는 CPU 메모리의 텐서를 GPU 메모리로 옮기는 호출로, 이 전송에는 시간이 듭니다. 작은 텐서에서는 전송과 실행 시작 비용이 계산보다 커서 CPU가 더 빠를 수 있으므로 크기가 충분히 큰지 확인해야 합니다. CUDA 연산은 호출이 먼저 반환되고 GPU에서 비동기로 실행되기 때문에, 시간을 잴 때는 torch.cuda.synchronize()로 작업 완료를 기다린 뒤 재야 정확합니다.
Q1
GPU가 그래픽·행렬 연산에서 CPU보다 유리한 이유는?
Q2
GPU가 상대적으로 불리한 작업은?
Q3
NVIDIA GPU에서 수십 개(32개)의 스레드가 묶여 같은 명령을 동시에 실행하는 단위의 이름은?