들어가며: 왜 병렬인가

1.1 이질 병렬 컴퓨팅: 천재와 군중

2005년쯤까지 프로그래머들은 행복한 망상에 빠져 살았다. "컴퓨터가 알아서 빨라진다"는 망상이다. 18개월마다 트랜지스터 수가 두 배가 되고(무어의 법칙), 클럭 주파수도 같이 올라가니, 같은 코드를 그냥 두기만 해도 매년 더 빨리 도는 마법이 있었다. 그러다 어느 순간 클럭이 3GHz 언저리에서 멈췄다. 전력 밀도와 발열의 벽이다. 더 이상 한 코어를 더 빠르게 돌릴 수 없게 되자 칩 설계자들은 코어 수를 늘리는 쪽으로 방향을 틀었다. 듀얼코어, 쿼드코어, 그다음엔 수십 개의 코어. 그리고 그 끝에 GPU가 있다. 수천 개의 작은 코어를 한 칩에 욱여넣은, 어떻게 보면 정신 나간 설계.

CPU와 GPU의 설계 철학을 한 문장으로 요약하면 이렇다. CPU는 한 명의 천재를, GPU는 잘 훈련된 군중을 지향한다. CPU 코어 한 개에는 분기 예측기, 비순차 실행기(out-of-order), 거대한 L1·L2 캐시, 깊은 파이프라인 같은 호화로운 부속이 붙어 있다. 단일 스레드를 가능한 한 빨리 끝내기 위해서다. 이걸 보통 "지연 지향(latency-oriented)" 설계라 부른다. 한 사람이 한 가지 어려운 일을 짧은 시간에 끝내는 데 최적화돼 있다.

GPU는 정반대다. 한 코어는 단순하다. 분기 예측 같은 화려한 장치도 없다. 대신 그런 코어를 수천 개 깔아 둔다. 각각은 느리지만 같은 시간에 수천 개의 일을 동시에 처리한다. 이걸 "처리량 지향(throughput-oriented)" 설계라 한다. 군중이 모여서 단순 작업을 같은 박자로 끝내는 모습이다. 한 명에게 미적분을 풀어 달라고 하면 GPU는 멍하니 있겠지만, 천 명에게 같은 모양의 곱셈을 시키면 CPU가 따라잡지 못한다.

"이질(heterogeneous)" 병렬 컴퓨팅이라는 말은 여기서 나온다. CPU와 GPU는 서로 다른 일을 잘하는 서로 다른 종류의 프로세서다. 운영체제, 입출력, 분기 가득한 제어 로직은 CPU가 한다. 행렬 곱, 합성곱, 픽셀 처리 같은 데이터 병렬 부하는 GPU가 받는다. 두 종족이 한 시스템에 공존하면서 각자 잘하는 걸 하는 구조 ─ 그것이 이 책 전체를 관통하는 모델이다.

1.2 왜 더 빨라야 하나

"지금도 충분히 빠른데?" 하는 사람이 있을 것이다. 그러나 우리가 풀려는 문제는 하드웨어보다 더 빠르게 자라고 있다. 몇 가지 풍경을 보자.

딥러닝. 2012년 AlexNet이 약 6천만 개의 파라미터로 ImageNet을 깼을 때 사람들은 놀랐다. 2020년의 GPT-3는 1750억 개. 2024년 이후의 프런티어 모델은 조 단위로 들어선다. 학습 한 번에 H100 수천 장을 몇 주 돌리는 일이 일상이 됐다. 이런 모델은 GPU 없이는 단순히 학습 시간이 길어지는 정도가 아니라, 인간 수명 안에 끝나지 않는다. 추론 단계에서도 마찬가지다. 사용자 한 명에게 응답 한 줄을 주려면 수십억 번의 부동소수점 곱이 필요하다.

과학 시뮬레이션. 기상 예보는 지구를 수 km 격자로 쪼개고 각 셀에서 운동 방정식과 열역학 방정식을 푼다. 격자를 두 배로 촘촘히 하면 셀 수는 8배(3차원이니까), 시간 간격까지 줄여야 하니 총 계산량은 16배가 된다. 같은 컴퓨터로 두 배 정확한 예보를 하려면 16배 빠른 기계가 필요하다는 뜻이다.

의료 영상. CT 한 번 촬영에 수백 장의 단면 이미지가 나온다. MRI 재구성, PET 영상의 노이즈 제거, 3D 볼륨 렌더링은 모두 픽셀(혹은 복셀) 단위로 같은 연산을 반복하는 전형적인 데이터 병렬 작업이다. 의사가 진단 화면을 기다리는 시간은 길어야 1~2초가 되어야 한다. 환자가 검사 침대 위에 누워 있는 동안 결과가 나와야 하기 때문이다.

자율주행. 차에 달린 카메라 6~12개, 라이다, 레이더가 매 프레임 수백 메가바이트의 데이터를 쏟아낸다. 30 FPS로 동작한다면 33ms 안에 객체 감지, 차선 추정, 경로 계획을 끝내야 한다. 사람의 안전이 ms 단위 응답에 걸려 있다. 이게 안 빠르면 차가 사람을 친다.

요약하면 이렇다. "빠른 컴퓨터가 있으니 큰 문제를 푸는 것"이 아니라 "풀어야 할 문제가 커서 빠른 컴퓨터를 만드는 것"이다. 그리고 그 "빠른 컴퓨터"의 답이 지금은 GPU 가속이다.

1.3 실제 응용 가속하기: Amdahl과 메모리 천장

GPU 가속이 만능 약은 아니다. 신비로운 100배 가속 신화에는 항상 조건이 붙는다. 그 조건을 가장 정직하게 말해 주는 것이 Amdahl의 법칙이다.

프로그램의 실행 시간 중에서 비율 p만큼이 병렬화 가능하고, 나머지 1−p는 순차적이라고 하자. 병렬 부분을 무한히 빠르게 만든다 해도 이론적인 최대 속도 향상은

Speedup ≤ 1 / (1 − p)

로 묶인다. p가 0.9면 최대 10배. p가 0.99여야 100배. p가 0.5밖에 안 된다면 아무리 GPU를 갈아 넣어도 2배다. 처음 GPU 포팅을 할 때 사람들이 흔히 겪는 좌절이 여기에서 온다. "커널은 50배 빨라졌는데 전체는 3배밖에 안 빨라요." 99% 같지만 99%가 아니었던 것이다. 직렬 부분 1%가 전체의 절반을 잡아먹고 있다.

두 번째 천장은 메모리 대역폭이다. GPU의 연산 성능은 TFLOPS 단위로 적혀 있지만, 그 숫자를 받쳐 줄 데이터를 충분히 빨리 공급할 수 있어야 의미가 있다. 가령 어떤 GPU의 단정밀도 연산이 20 TFLOPS, DRAM 대역폭이 1 TB/s라고 하자. 부동소수점 한 개가 4바이트니까 1 TB/s는 초당 250 G개의 부동소수점을 읽어들일 수 있다는 뜻이다. 이걸 20T번의 연산에 배분하면 한 부동소수점당 약 80번의 연산이 가능해야 대역폭이 천장이 되지 않는다. 이를 산술 강도(arithmetic intensity, FLOP/byte)라 부른다. 산술 강도가 낮으면(예: 단순 벡터 덧셈은 1 FLOP을 위해 12바이트를 읽어 0.083) 연산기가 데이터를 기다리며 논다. 이런 코드는 메모리 바운드라 부르고, 가속비는 연산 능력이 아니라 대역폭에 의해 결정된다.

그래서 "100× 가속이 가능한가?"라는 질문의 정직한 답은 조건부로 가능하다이다. 병렬 비율 p가 0.99 이상이고, 산술 강도가 충분히 높고, 메모리 접근 패턴이 좋으면 종종 도달한다. 그렇지 않으면 5배, 10배에서 천장에 부딪힌다. 이 책의 후반부 절반은 사실상 "어떻게 그 조건을 만들 것인가"를 다룬다.

1.4 병렬 프로그래밍의 어려움

병렬 코드를 짜는 일은 직렬 코드를 두 벌 짜는 것이 아니다. 사고방식 자체가 다르다. 몇 가지 큰 어려움을 미리 알려두면 마음의 준비가 된다.

알고리즘 사고의 전환. 직렬 알고리즘 책에 나오는 멋진 자료구조와 트릭 다수가 GPU에서는 그대로 안 먹는다. 분기 가득한 트리 순회, 연결 리스트 추적, 재귀 깊이가 수만인 알고리즘은 GPU의 강점인 SIMD 실행과 궁합이 안 맞는다. 대신 같은 일을 데이터 병렬로 다시 표현하는 능력이 필요하다. "정렬을 부분 합으로 푼다", "그래프 BFS를 행렬·벡터 곱으로 푼다" 같은 발상의 전환들이 자주 등장한다.

데이터 종속성. a[i] = a[i-1] + b[i] 같은 점화식은 i끼리 서로 기다려야 하니 단순한 병렬화가 안 된다. 이런 문제는 prefix sum 같은 영리한 변환을 거쳐야 병렬화된다. 9장에서 이걸 본격적으로 다룬다.

확장성과 성능 휴대성. 내 GTX 1080에선 잘 돌던 코드가 RTX 4090에서 느려질 수 있다. 워프 크기, SM 수, 공유 메모리 크기, L2 캐시 크기가 다 달라서다. "어디 가서나 좋은 성능"을 내려면 하드웨어 추상화에 대한 감각이 필요하다.

디버깅 난이도. 1만 개의 스레드 중 정확히 한 놈이 잘못된 인덱스로 메모리를 만지면, 증상은 무작위로 다른 곳에서 터진다. 직렬 코드의 printf 디버깅과는 차원이 다르다. cuda-gdb, compute-sanitizer, Nsight Compute 같은 도구를 일찍 익히는 게 정신 건강에 좋다.

1.5 다른 병렬 인터페이스들 (왜 CUDA로 시작하는가)

GPU 병렬 컴퓨팅에 입문하는 길은 여러 갈래다. 간단히 비교해 보면 다음과 같다.

• OpenMP: CPU의 멀티스레딩을 위한 컴파일러 디렉티브. #pragma omp parallel for 한 줄로 루프를 병렬화. 멀티코어 CPU 시대의 사실상 표준이지만 GPU 지원은 4.5 이후 들어온 후발주자.
• MPI: 노드 간 분산 메모리 모델. 슈퍼컴퓨터의 공용어. 한 노드 안에서는 여전히 CUDA나 OpenMP가 필요.
• OpenCL: NVIDIA, AMD, 인텔, 모바일 SoC까지 한 코드로 돌리겠다는 야심으로 출발. 모델이 CUDA와 거의 같지만 API가 더 장황하고, 실제 생태계의 모멘텀은 줄었다.
• SYCL: C++ 단일 소스로 이질 가속을 표현하려는 신세대 표준. 인텔의 oneAPI가 이걸 밀고 있다.
• HIP: AMD의 CUDA 호환 레이어. 코드 모양이 거의 같아서 hipify로 변환이 가능.

그러면 왜 CUDA인가? 첫째, 가장 성숙한 도구 사슬을 갖고 있다. nvcc, cuBLAS, cuDNN, Thrust, NCCL, Nsight ─ 이름만 들어도 풍성한 생태계다. 둘째, 산업 현장의 GPU 컴퓨팅 코드 중 압도적 다수가 CUDA로 쓰여 있다. 학습한 개념의 80~90%는 다른 인터페이스로 바로 옮겨 쓸 수 있다. 같은 모델이라 그렇다(블록, 스레드, 워프, 공유 메모리 …). 그러니 도구 하나로 깊게 파고, 나중에 필요하면 옆걸음으로 옮기는 전략이 학습 곡선상 합리적이다.

1.6 이 책이 데려갈 곳

이 책의 목표는 단순하다. 독자가 GPU에서 동작하는 의미 있는 커널을 직접 설계·구현·디버깅·최적화할 수 있게 만드는 것. "Hello, World" 수준에서 멈추지 않는다. 행렬 곱, 합성곱, prefix sum, 정렬, 히스토그램, 희소 행렬, 그래프 BFS, 머신러닝 추론까지 ─ GPU 컴퓨팅의 핵심 패턴들을 직접 손으로 구현해 본다. 그 과정에서 점유율, 메모리 합치기, 공유 메모리 뱅크 충돌, 워프 셔플, 텐서 코어 같은 단어들이 추상적인 술어가 아니라 키보드 위에서 찍어 본 코드와 결합된 감각이 된다.

1.7 책 구성 로드맵

책은 4부로 나뉜다.

그림 1.1 — 책의 4부 구성

Part I (1~6장)이 이 책의 토대다. 여기서 모델, 메모리 계층, 성능에 대한 직관이 잡힌다. 이 부분을 건너뛰면 후반부의 최적화 패턴들이 마술처럼 보일 것이다. Part II는 거의 모든 GPU 코드의 빌딩 블록인 패턴들이다. 한 번 익혀두면 새 알고리즘을 만났을 때 "아, 이건 reduction + scan이네" 식으로 분해해서 보게 된다. Part III와 IV는 실제 응용과 최신 기능이다. 이 책을 끝낼 즈음에 독자는 논문에서 GPU 구현 섹션을 읽으며 "이 사람이 왜 이걸 이렇게 짰는지"가 보이게 될 것이다.