병렬 프로세서: 클라이언트에서 클라우드까지

크게 휘두른다. 가진 모든 것으로. 크게 맞히거나, 크게 빗나가거나. 어차피 살 거면 크게 살래. — Babe Ruth

6.1 도입 — 멀티프로세서, 그러니까 같이 하자

1장부터 5장까지 우리는 “한 명의 슈퍼스타 프로세서”를 어떻게 더 잘 굴릴 것인가에 매달렸다. 파이프라이닝, 분기 예측, 슈퍼스칼라, 더 큰 캐시, 더 빠른 메모리. 다 한 사람을 빠르게 만드는 기술이다. 그런데 어느 순간부터 “이 사람을 더 빠르게 만들기”가 물리적으로 한계에 부딪히기 시작했다. 클럭을 더 올리면 발열이 감당 안 되고, 명령어 수준 병렬성(ILP)도 짜낼 만큼 짜냈다. 그래서 업계가 내린 결론은 단순하다. “혼자 빠르게 못 하면, 여럿이 같이 하자.”

이게 바로 멀티프로세서(multiprocessor)의 시대다. 칩 안에 코어를 두 개, 네 개, 여덟 개, 수백 개씩 박아 넣고 동시에 일을 시킨다. 목표는 단순하다 — 프로세서 수가 늘어나면 성능도 같이 늘어나야 한다. 당연한 얘기 같지만, 뒤에서 보겠지만 이게 정말 정말 어려운 약속이다. 사람 8명이 모이면 일이 8배 빨라진다고 누가 그러던가? 회식 자리에서 “언제 출발하지?”만 30분 토론하는 게 인간이다.

작업 수준 병렬성 vs 병렬 처리 프로그램

여러 프로세서를 가지고 “병렬로 일을 시킨다”는 말에는 두 가지 다른 뜻이 숨어 있다. 이걸 헷갈리면 평생 이상한 곳에서 화가 난다.

• 작업 수준 병렬성(task-level parallelism): 서로 상관없는 작업들을 그냥 다른 코어에 던져 놓는다. 카톡 켜고, 음악 들으면서, 브라우저로 유튜브 보고, 백그라운드에서 백업 돌리는 그런 거. 각자 알아서 자기 일 하면 된다. 이건 OS가 알아서 분배해주니까 프로그래머가 별로 신경 쓸 게 없다.
• 병렬 처리 프로그램(parallel processing program): 한 가지 일을 여러 프로세서에 쪼개서 맡기는 것. 행렬 곱셈을 8개 코어가 동시에 처리한다거나, 영상 렌더링을 GPU 수천 개 스레드로 나눠 돌린다거나. 이쪽이 진짜 “어려운” 병렬이다. 협업이 필요하고, 동기화가 필요하고, 데이터 의존성을 고려해야 한다.

이 장의 주된 관심사는 두 번째다. “8명이서 단톡방으로 한 편의 기사를 같이 쓰는 일” — 이게 병렬 처리 프로그램의 본질이고, 또 본질적인 골칫거리다.

멀티코어와 SMP

멀티코어 마이크로프로세서(multicore microprocessor)는 말 그대로 한 칩 안에 여러 개의 프로세서(코어)가 들어 있는 구조다. 요즘 노트북에 든 칩은 코어 8개, 16개가 기본이다. 데스크톱 Threadripper 같은 괴물은 64코어, 128스레드까지 간다.

이런 멀티코어 시스템의 가장 흔한 형태가 SMP(Shared Memory Processor 또는 Symmetric Multiprocessor)다. 모든 코어가 하나의 물리적 주소 공간을 공유한다는 뜻이다. 이게 무슨 의미인지는 6.5절에서 진하게 다루지만, 일단 “모든 코어가 같은 화이트보드를 보고 일한다”고 생각하면 된다. 누구든 같은 주소에 쓰면 다른 코어도 그걸 읽을 수 있다. 편한 동시에 위험하다.

직렬/병렬 × 순차/동시 — 의외로 헷갈리는 사분면

같은 “여러 일을 한다”는 말도 두 축으로 나눠 보면 뜻이 달라진다.

“동시(concurrent)”는 “같이 일어나는 것처럼 보이는” 것이고, “병렬(parallel)”은 “물리적으로 동시에 일어나는” 것이다. 단일 코어 OS는 동시지만 병렬은 아니다. 멀티코어가 되어야 진짜 병렬이 된다. 자취방 식탁 하나에서 혼자서 책을 읽다 카톡을 답하다 빨래를 개는 건 동시이지만 병렬은 아니다. 식구가 셋이면 동시에 셋이 다른 일을 진짜로 할 수 있다 — 이게 병렬이다.

6.2 병렬 프로그래밍이 어려운 이유

“코어 8개니까 8배 빨라지겠지” — 만약 인생이 그렇게 단순했다면 이 장은 두 페이지로 끝났을 거다. 현실은 잔인하다. 그리고 그 잔인함을 가장 깔끔하게 표현한 게 1장에서도 만난 암달의 법칙(Amdahl's Law)이다.

암달의 법칙, 다시 만나다

암달의 법칙은 이렇게 말한다.

전체 속도 향상 = 1 / ( (1 - F) + F / S )

여기서 F는 병렬화 가능한 부분의 비율, S는 그 부분이 얼마나 빨라지는지다. 이걸 좀 더 풀어 쓰면, “아무리 병렬화 부분을 무한히 빠르게 만들어도, 직렬로 남는 (1-F) 부분이 전체 시간을 잡아먹는다”는 뜻이다. 직렬 부분은 절대로 줄어들지 않는 단단한 시멘트 바닥이다.

속도 향상 도전 — 100개 코어로 90배 빠르게?

구체적으로 가보자. 100개의 프로세서로 90배 속도 향상을 얻고 싶다고 하자. 즉 S = 100이고 원하는 전체 속도 향상은 90이다. 그러면 직렬 비율은 얼마여야 할까?

90 = 1 / ( (1 - F) + F / 100 )

풀어 보면:

(1 - F) + F / 100 = 1 / 90 ≈ 0.0111
1 - F + 0.01F     = 0.0111
1 - 0.99F         = 0.0111
0.99F             = 0.9889
F                 ≈ 0.9989

즉 코드의 99.89%가 완벽하게 병렬화되어 있어야 한다. 직렬로 남을 수 있는 부분은 겨우 0.11%다. 0.1% 남짓. 이게 얼마나 잔인한 숫자인지 감이 잘 안 온다면, 1만 줄짜리 프로그램에서 단 11줄만 직렬로 남아도 100개 코어가 90배를 못 낸다는 뜻이다.

그리고 1만 개 코어로 9000배 빠르게? 그러면 직렬 부분은 0.0011%여야 한다. 100만 줄 코드에서 11줄. 이쯤 되면 “병렬 프로그래밍은 사실상 직렬 부분을 어떻게 줄이느냐의 게임”이라는 말이 그냥 슬로건이 아니다.

강한 스케일링 vs 약한 스케일링

“문제의 크기”에 따라 속도 향상 이야기는 또 갈린다.

• 강한 스케일링(strong scaling): 문제 크기를 고정해 놓고 코어 수를 늘리며 얼마나 빨라지는지 보는 것. 1024×1024 행렬을 1코어로는 100초, 8코어로는 15초 같은 식. 가장 솔직하고 가혹한 비교다. 직렬 부분이 그대로 보이기 때문이다.
• 약한 스케일링(weak scaling): 코어 수와 함께 문제 크기도 같이 키우는 것. 1코어가 1024×1024를 100초에 처리한다면, 8코어는 1024×1024×8을 100초에 처리할 수 있는지 본다. “더 큰 문제를 같은 시간에” 식의 평가다.

슈퍼컴퓨터 업계에서는 약한 스케일링이 자연스러운 평가법이다. 어차피 슈퍼컴은 작은 문제 풀려고 사는 게 아니니까. 반면 데이터센터 응답 시간 같은 경우엔 강한 스케일링이 의미가 있다. 사용자는 “요청”이라는 고정된 크기의 일을 빠르게 처리해 주길 바라기 때문이다.

로드 밸런싱 — 한 명만 느려도 게임 끝

병렬 프로그램의 또 다른 무서운 함정은 로드 밸런싱(load balancing)이다. 일이 모든 코어에 골고루 나뉘지 않으면, 가장 늦은 코어가 끝날 때까지 다른 코어들은 대기해야 한다.

예를 들어 코어 100개에 일을 나누는데 99개 코어는 1초에 자기 몫을 끝내고, 운 나쁜 1개 코어가 2배의 일을 받아 2초가 걸린다고 하자. 전체 실행 시간은? 2초다. 가장 느린 놈을 기다려야 하니까. 99개 코어가 1초 동안 손가락 빨고 있다는 뜻이다. 이때의 속도 향상은 50배 정도지, 100배가 절대 아니다.

즉 “한 사람만 느려도 전체가 느려진다.” 이래서 일을 잘 쪼개서 골고루 나누는 일이 단순한 분배가 아니라 예술의 영역이 된다. 야구의 더블플레이를 떠올려 보자. 2루수와 유격수와 1루수가 각각 자기 동작을 한 박자도 안 어긋나게 해야 두 명을 잡는다. 누구 한 명이라도 1초 늦으면 그 더블플레이는 안 된다.

6.3 SISD, MIMD, SIMD, SPMD 그리고 벡터

병렬 컴퓨터를 분류하는 가장 오래된 표가 플린 분류(Flynn's Taxonomy)다. 마이클 플린이 1966년에 만들었으니, 거의 환갑이 다 된 분류법이지만 여전히 잘 통한다. 두 축은 단순하다 — “명령어 스트림이 하나냐 여럿이냐”와 “데이터 스트림이 하나냐 여럿이냐”.

이 분류는 너무 거칠어서 실무에서는 더 세분화된 용어들을 쓴다. 그 중에서 가장 자주 보는 게 SPMD(Single Program Multiple Data)다. 같은 프로그램의 복사본을 여러 노드에서 돌리되, 각자 다른 데이터를 처리하는 방식. MPI나 OpenMP 같은 거의 모든 실용 프레임워크가 SPMD다. “같은 코드를 돌리지만 자기 ID에 따라 다른 부분을 담당”하는 식이다.

SIMD: x86의 멀티미디어 확장 — MMX → SSE → AVX

SIMD는 한 명령어로 여러 데이터를 동시에 처리한다. 가장 익숙한 형태가 x86 CPU의 멀티미디어 확장 명령어들이다. 1990년대에 인텔이 멀티미디어(MP3 디코딩, MPEG 비디오 등)를 가속하려고 만든 게 MMX(MultiMedia eXtension)다. 64비트 레지스터 안에 8비트 짜리 8개를 동시에 더하거나 뺀다. 사진의 픽셀 8개를 한 명령어로 한 번에 밝기 조정할 수 있다는 얘기다.

그 뒤로는 비대해지는 길만 갔다.

• SSE(Streaming SIMD Extensions): 128비트 레지스터. 32비트 부동소수점 4개를 동시에.
• SSE2/3/4: 정수 연산도, 더 다양한 연산도.
• AVX(Advanced Vector eXtensions): 256비트 레지스터. 32비트 부동소수점 8개 동시에.
• AVX-512: 512비트 레지스터. 32비트 부동소수점 16개 동시에. 칩이 거의 화염방사기가 됨.

이런 식으로 레지스터 폭만 두 배씩 키우는 것도 SIMD의 한 갈래다. 이걸 “서브워드 병렬성(subword parallelism)” 또는 “패킹된 SIMD(packed SIMD)”라고 부른다. 큰 레지스터를 작은 조각으로 패킹해서 동시에 굴리는 거니까. 단순하고 효과적이다 — 단점은 매번 폭이 바뀔 때마다 새 명령어 셋을 추가해야 한다는 것. 그래서 x86 ISA가 점점 부풀어 올랐다.

벡터 아키텍처 — Cray 시절의 우아함

SIMD의 또 다른 갈래, 어쩌면 더 우아한 갈래는 벡터 아키텍처(vector architecture)다. 1970년대 Cray-1이 시작점이었다. 핵심 아이디어는 단순하다 — “벡터(배열)을 하나의 기본 데이터 타입으로 삼는 ISA를 만들자.” 그래서 등장한 게:

• 벡터 레지스터(vector register): 한 레지스터가 64개, 128개의 원소를 담는다.
• 벡터 명령어(vector instruction): vadd v1, v2, v3 한 줄이 64개 덧셈을 의미.
• 벡터 길이 레지스터(VL): “이번 연산은 앞에서부터 몇 개만 처리해 줘”를 동적으로 지정.
• 벡터 레인(vector lane): 실제 하드웨어에서 병렬로 처리하는 연산기들. 8개 레인이면 매 사이클 8개씩 흘려보냄.

DAXPY — 스칼라 vs 벡터 비교

벡터의 위력을 보여주는 고전 예제가 DAXPY다. Y = a × X + Y, X와 Y가 길이 N짜리 부동소수점 벡터인 연산이다. (DAXPY는 “Double A times X Plus Y”의 약어다.)

스칼라 RISC-V에서는 대충 이런 식이다.

// 스칼라 RISC-V — N=64라고 하자. 매번 한 원소씩 처리
        fld     fa0, 0(x10)        // a 로드 (한 번만)
        addi    t0, x0, 64        // 카운터 = 64
loop:   fld     ft0, 0(x11)        // X[i]
        fld     ft1, 0(x12)        // Y[i]
        fmul.d  ft2, fa0, ft0     // a * X[i]
        fadd.d  ft3, ft2, ft1     // + Y[i]
        fsd     ft3, 0(x12)        // 다시 Y[i]에 저장
        addi    x11, x11, 8
        addi    x12, x12, 8
        addi    t0, t0, -1
        bne     t0, x0, loop

한 번 도는 데 명령어가 9개쯤. 64번 돌면 약 580개 명령어가 실행된다. 분기 예측이며 캐시 미스며 별별 걱정거리가 다 있다.

이걸 벡터 RISC-V(가상의)로 쓰면:

// 벡터 RISC-V — N=64를 한 번에
        fld     fa0, 0(x10)        // a 로드
        vsetvli t0, x0, e64      // 벡터 길이 설정 (최대치)
        vle.v   v0, 0(x11)         // X 전체 로드
        vle.v   v1, 0(x12)         // Y 전체 로드
        vfmul.vf v2, v0, fa0     // a * X (전체)
        vfadd.vv v3, v2, v1      // + Y
        vse.v   v3, 0(x12)         // 결과 저장

명령어 수 7개. 약 80배 명령어 감소다. 분기도 없고, 카운터도 없고, 의존성 추적도 하드웨어가 알아서 한다. 한 명령어가 “64개”를 의미하니, 명령어 페치 부담도 거의 사라진다. 이게 벡터의 우아함이다.

벡터 vs 멀티미디어 SIMD

벡터 아키텍처와 x86식 멀티미디어 SIMD는 둘 다 “한 명령어로 여러 데이터”를 처리하지만, 결정적인 차이가 있다.

멀티미디어 SIMD는 “레지스터 안 비트를 잘게 쪼개서 쓴다”는 단순한 발상이지만, 폭이 바뀔 때마다 ISA가 부풀어 오른다는 단점이 있다. 벡터 아키텍처는 “벡터 길이”라는 추상으로 그 문제를 우아하게 피한다. 그래서 ARM의 SVE나 RISC-V의 V 확장이 21세기에 다시 벡터 스타일을 채택한 것이다. 옛것이 새것이 되는 컴퓨터 구조의 흔한 풍경.

벡터 레인의 비밀

“벡터 명령어 한 줄이 64개를 한 번에”라고 하지만, 실제 하드웨어가 정말 64개의 가산기를 가지고 있는 건 아니다. 보통은 4개나 8개의 레인을 두고, 각 레인이 매 사이클 한 원소씩 처리해서 16~8 사이클에 64개를 끝낸다. 레인을 늘리면 처리량이 비례해서 늘어난다 — 같은 ISA, 같은 코드, 단지 레인 수만 다른 구현체로. 이게 벡터의 또 다른 강점이다. 같은 프로그램을 실리콘 예산만 늘려서 더 빠르게 만들 수 있다.

6.4 하드웨어 멀티스레딩

프로세서를 늘리는 게 아니라, 한 프로세서가 여러 스레드를 빠르게 번갈아 처리하게 만드는 길도 있다. 이걸 하드웨어 멀티스레딩(hardware multithreading)이라고 한다. 핵심 아이디어는 “한 스레드가 캐시 미스나 의존성으로 멈춰 있을 때, 다른 스레드를 바로 돌리자”다. 비싼 함수 유닛을 놀게 두지 말자는 거다.

스레드 vs 프로세스

잠깐 용어 정리. 프로세스(process)는 자기만의 주소 공간을 가진 프로그램 인스턴스다. 크롬과 워드는 다른 프로세스. 스레드(thread)는 한 프로세스 안에서 여러 갈래로 갈라진 실행 흐름이다. 같은 메모리를 공유한다. 카톡 한 개 프로세스 안에서 “채팅 받는 스레드”와 “알림 띄우는 스레드”가 따로 도는 식.

하드웨어 멀티스레딩은 이 “스레드”들을 한 코어가 여러 개 돌리게 한다. 각 스레드는 자기 PC와 자기 레지스터 셋을 갖는다. 캐시와 ALU는 공유한다. 그래서 “여러 명이 한 책상을 같이 쓰는데, 각자 자기 노트와 자기 펜은 따로 쓴다” 정도의 그림이다.

세 가지 멀티스레딩 스타일

스레드 전환 타이밍에 따라 세 가지로 나뉜다.

1. Fine-grained Multithreading (세립도): 매 사이클마다 다른 스레드. 사이클별로 라운드 로빈. 모든 스레드가 골고루 진도가 나가지만, 단일 스레드 성능은 떨어진다. 한 스레드가 멈추든 안 멈추든 어차피 다음 사이클엔 다른 스레드 차례니까. 한 스레드가 빨리 끝나야 할 때는 손해다.
2. Coarse-grained Multithreading (조립도): 캐시 미스 같은 큰 정지가 발생할 때만 스레드 전환. 평소엔 한 스레드가 코어를 독점하니 단일 스레드 성능은 좋다. 단점은 “전환 비용”이 있어서, 짧은 정지엔 효과가 별로 없다. 그리고 파이프라인을 비웠다 채우는 시간이 든다.
3. Simultaneous Multithreading (SMT, 동시 멀티스레딩): 같은 사이클에 여러 스레드의 명령어를 동시에 발행한다. 슈퍼스칼라가 4개 명령어를 발행할 수 있으면, 이 중 2개는 스레드 A에서, 2개는 스레드 B에서. 함수 유닛을 가장 알차게 쓰는 방식. 인텔의 “하이퍼스레딩(Hyper-Threading)”이 바로 이것이다.

이슈 슬롯으로 그려 보면

슈퍼스칼라 코어가 사이클마다 4개의 이슈 슬롯을 갖는다고 하자. 4개의 스레드(빨/파/노/초)가 같은 시간 동안 어떻게 슬롯을 채우는지 보자. 비어 있는 슬롯은 “.”으로 표시.

// 단일 스레드 슈퍼스칼라 — 빨강 한 명만
사이클1:  빨 빨 .  .
사이클2:  빨 .  .  .       // 캐시 미스로 정지
사이클3:  .  .  .  .
사이클4:  빨 빨 빨 .

// Fine-grained — 사이클마다 갈아탐
사이클1:  빨 빨 .  .
사이클2:  파 파 파 .
사이클3:  노 노 .  .
사이클4:  초 초 초 초

// Coarse-grained — 정지 때만 갈아탐
사이클1:  빨 빨 .  .
사이클2:  빨 .  .  .       // 정지 → 전환 비용
사이클3:  파 파 파 .
사이클4:  파 파 .  .

// SMT — 한 사이클 안에서 섞음
사이클1:  빨 빨 파 파
사이클2:  빨 노 파 초
사이클3:  노 노 초 파
사이클4:  빨 빨 노 초

SMT 그림에서 빈 슬롯이 거의 없다는 게 핵심이다. 함수 유닛이 알차게 일한다. 인텔이 발표한 데이터에 따르면, i7에서 하이퍼스레딩(2-way SMT)으로 평균 1.31× 성능 향상, 1.07× 에너지 효율을 봤다고 한다. 트랜지스터 5%만 더 써서 30% 성능을 더 짜내는 셈이니, 이쯤이면 거의 공짜 점심이다.

6.5 멀티코어와 공유 메모리 멀티프로세서

드디어 본론. 한 칩 안에 진짜 여러 코어를 박아 놓고, 모두가 같은 메모리를 공유하는 SMP 시스템이다. 데스크톱부터 서버까지 가장 흔하게 만나는 형태다.

단일 물리 주소 공간 — 좋은 점, 나쁜 점

공유 메모리의 가장 큰 장점은 “주소 0x1000은 모두에게 같은 0x1000”이라는 점이다. 코어 A가 그 주소에 값을 쓰면 코어 B가 읽을 수 있다. 데이터를 명시적으로 “전송”할 필요가 없다. 같은 메모리니까. 이건 프로그래밍하기 정말 편하다.

그런데 동시에 끔찍하게 위험하기도 하다. A와 B가 동시에 같은 변수를 읽고 더하고 쓰면? 갱신이 사라질 수 있다. 카페에서 두 명이 한 노트에 동시에 1을 더하다 펜이 부딪히는 그림이다. 그래서 등장하는 게 동기화(synchronization)와 캐시 일관성(cache coherence)이다 — 둘 다 5장과 이번 장의 단골 주제.

UMA vs NUMA

공유 메모리도 한 가지가 아니다. 모든 코어가 메모리에 접근하는 시간이 똑같으면 UMA(Uniform Memory Access), 코어마다 다르면 NUMA(Non-Uniform Memory Access)다.

• UMA: “모든 코어가 메모리까지 같은 거리.” 작은 SMP, 단일 소켓 데스크톱이 대체로 이 모양. 단순하고 예측 가능. 단, 코어 수가 많아지면 메모리 버스가 병목이 된다.
• NUMA: 코어들이 여러 노드로 묶여 있고, 각 노드가 자기 “로컬 메모리”를 가진다. 자기 로컬은 빠르고, 다른 노드 메모리는 느리다. 큰 서버나 멀티소켓 시스템에서 일반적. 프로그래머가 “데이터를 쓸 코어 근처에 놓는 일”을 신경 써야 빨라진다.

락으로 동기화

공유 메모리에서 두 코어가 같은 변수를 안전하게 갱신하려면 락(lock)이 필요하다. 5장에서 살펴본 LR/SC 같은 원자적 연산이 그 토대다. 한 코어가 락을 잡고 임계 구역(critical section)에 들어가 변수 갱신을 마친 뒤 락을 푼다. 다른 코어는 그동안 기다린다.

화장실 한 칸을 공유하는 거랑 똑같다. 들어가면 문 잠그고, 끝나면 풀고. 락이 길수록 다른 사람들이 바깥에서 하품한다 → 직렬 부분이 늘어남 → 암달의 법칙이 우리를 비웃음. 그래서 “락은 짧게, 자주 가져가지 않기”가 병렬 프로그래밍 황금률 중 하나다.

병렬 합 — 트리 환원의 미학

구체적인 예로, 길이 N짜리 배열 A의 합을 P개의 코어로 구한다고 하자. 단순한 방법은:

1. 각 코어가 자기 몫의 부분합을 따로 계산 (병렬, N/P 시간)
2. 모든 부분합을 하나로 모음 (이게 직렬이면 P 시간)

2단계가 직렬이면 큰 P에서 손해다. 그래서 환원 트리(reduction tree)를 쓴다. P=8일 때 이런 식.

단계 0: [s0] [s1] [s2] [s3] [s4] [s5] [s6] [s7]   // 각자 부분합
단계 1: [s0+s1] [s2+s3] [s4+s5] [s6+s7]            // 4쌍 동시에
단계 2: [s0..3]      [s4..7]                        // 2쌍 동시에
단계 3: [s0..7]                                     // 최종

합치는 단계가 log₂(P)이다. P=1024라도 단계는 10. 직렬 1024단계가 아니라. 이게 트리 환원의 힘이다. OpenMP의 reduction 절도 내부적으론 이런 트리를 쓴다.

OpenMP — 프라그마로 병렬화

SMP에서 가장 인기 있는 병렬 프로그래밍 모델이 OpenMP다. C/C++/Fortran 코드에 프라그마(pragma)를 툭툭 붙이면 컴파일러가 알아서 스레드를 만들어 굴려준다. 가장 흔한 패턴.

// 1. 병렬 영역
#pragma omp parallel
{
    // 여기 코드 블록은 모든 스레드가 실행
    printf("안녕 from %d", omp_get_thread_num());
}

// 2. 병렬 for — 가장 자주 쓰는 패턴
#pragma omp parallel for
for (int i = 0; i < N; i++) {
    A[i] = B[i] * C[i];   // 반복마다 독립 → 자동 분배
}

// 3. reduction — 합/곱/최대 같은 환원
double sum = 0;
#pragma omp parallel for reduction(+:sum)
for (int i = 0; i < N; i++) {
    sum += A[i];          // 트리 환원으로 안전하게 합쳐짐
}

프라그마는 컴파일러에게 보내는 메모다. “이 루프 반복은 서로 독립이니 나눠 돌려도 돼” 같은 정보를 제공한다. 잘못 붙이면(예: 반복 간 의존이 있는 루프에 parallel for를 붙이면) 결과가 틀릴 수 있다. OpenMP는 안전을 보장해주지 않는다. “네가 옳다고 보증한 거니 나는 시키는 대로 돌릴게” 모드.

6.6 GPU 입문 — 게임 덕분에 진화한 괴물

GPU(Graphics Processing Unit)는 본래 “화면 그리기”의 보조 일꾼이었다. 그런데 1990년대 후반~2000년대 초반, 플레이스테이션과 X박스, 그리고 PC 게임의 폭발적인 성장이 GPU에게 어마어마한 진화 압력을 가했다. “3D 게임을 60프레임으로 돌려야 한다”는 단순한 요구가, “수천 개의 부동소수점 픽셀 셰이더 코어”라는 결과를 가져온 거다. 그 후로 누군가가 “이거 그래픽 말고 일반 계산에도 쓸 수 있지 않나?” 라고 했고, GPGPU(General Purpose GPU computing)와 NVIDIA CUDA의 시대가 열렸다. 지금은 딥러닝 학습의 거의 전부가 GPU에서 돌아간다.

CPU vs GPU — 같은 트랜지스터, 다른 철학

핵심은 “레이턴시 vs 처리량”의 다른 선택이다. CPU는 한 사람의 응답 시간을 최대한 짧게 만든다(분기 예측, OoO, 큰 캐시). GPU는 한 명 한 명은 좀 느려도, 동시에 수천 명을 돌려서 전체 처리량을 끝없이 끌어올린다. “ATM 한 대 vs 식권 자판기 1000대”의 차이.

CUDA — GPU 위의 C 같은 언어

NVIDIA가 만든 CUDA(Compute Unified Device Architecture)는 GPU 위에서 C 비슷한 코드를 돌리게 해주는 환경이다. 핵심 개념은 “커널(kernel)”이라는 함수를 정의하고, 그걸 “스레드 그리드” 위에서 실행시키는 것.

// CUDA 커널 — 모든 스레드가 이 함수를 실행
__global__ void daxpy(int n, double a,
                      double* x, double* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// 호스트(CPU)에서 호출
daxpy<<<blocks, threads_per_block>>>(n, 2.5, x, y);

한 스레드가 한 원소만 처리하지만, 수천~수십만 개 스레드가 동시에 켜져서 합쳐서 일을 끝낸다. 이게 GPU 프로그래밍의 멘탈 모델이다 — “한 명이 한 픽셀을 본다, 단 한 번에 백만 명이 같이.”

NVIDIA Fermi 용어 사전

GPU 용어가 회사마다 살짝 다른데, NVIDIA의 Fermi 세대를 기준으로 정리하면:

Multithreaded SIMD Processor

한 묶음의 SIMD 레인 + 자체 스케줄러를 갖춘 단위. 패터슨·헤네시는 이 용어를 쓰지만 NVIDIA 마케팅 용어로는 “Streaming Multiprocessor(SM)”이다. Fermi GPU 한 개에 SM이 16개쯤 들어 있다.

SIMD Thread

NVIDIA 용어로 “Warp.” 32개 스레드가 한 묶음으로 같은 명령어를 실행한다. 전형적인 SIMD 동작이지만, 각 스레드는 자기 데이터를 가진다.

SIMD Lane

NVIDIA 용어로 “CUDA Core.” 워프 내 한 스레드가 차지하는 실행 슬롯. 한 SM에 SIMD 레인이 32개~64개쯤.

Thread Block Scheduler

스레드 블록을 SM에 배정하는 디스패처. GPU 단위의 “감독관.”

Local Memory (Shared Memory)

한 SM 안에서 같은 블록의 스레드들이 공유하는 빠른 메모리. SRAM 기반. 수십 KB.

GPU Memory

GPU 카드 보드에 붙어 있는 큰 외부 DRAM(GDDR, HBM). 수~수십 GB. 모든 SM이 접근 가능. 대역폭은 어마어마하지만 레이턴시는 길다.

하드웨어 멀티스레딩이 GPU의 비밀 무기다. 한 워프가 메모리를 기다리고 있으면, SM은 즉시 다른 워프를 실행한다. CPU처럼 “캐시 미스 = 큰 손해”가 아니다. 워프가 충분히 많으면 메모리 레이턴시가 그냥 사라지는 것처럼 보인다. 그래서 GPU는 일부러 캐시를 작게 두고, 대신 더 많은 스레드와 더 큰 메모리 대역폭에 실리콘을 쓴다.

6.7 클러스터, 창고급 컴퓨터, 메시지 패싱

공유 메모리 SMP가 “한 칩/한 보드 안에서 같이 일하기”라면, 클러스터(cluster)는 “여러 독립 컴퓨터를 네트워크로 묶어서 같이 일하기”다. 같은 메모리를 공유하지 않는다. 각 노드가 자기만의 주소 공간을 가진다. 그래서 데이터를 주고받으려면 명시적으로 “보내고 받아야” 한다.

메시지 패싱 — send와 receive

메시지 패싱(message passing)은 두 가지 기본 동작으로 이루어진다.

• send(dest, data): 다른 노드로 데이터를 보냄.
• receive(src, buf): 다른 노드에서 데이터를 받음.

가장 널리 쓰는 라이브러리가 MPI(Message Passing Interface)다. 슈퍼컴퓨팅 세계의 공용 언어다.

장점과 단점

공유 메모리는 “편한 회의실” 같다 — 화이트보드 하나에 다같이 적으면 되니까. 메시지 패싱은 “편지로 소통”하는 회사다 — 누구한테 무엇을 보낼지 다 적어야 하지만, 회사가 무한히 커도 굴러간다. 슈퍼컴퓨팅과 클라우드 서비스가 메시지 패싱을 선택한 건, 결국 “스케일이 무한해야 하니까”다.

창고급 컴퓨터(WSC) — 구글, 아마존이 사는 동네

창고급 컴퓨터(Warehouse-Scale Computer)는 말 그대로 창고만 한 데이터센터에 수만~수십만 대의 서버를 채워 넣은 거대 클러스터다. 구글 검색, 아마존 쇼핑, 페이스북 피드, 넷플릭스 스트리밍 — 우리가 매일 쓰는 거의 모든 인터넷 서비스가 WSC 위에서 돌아간다.

WSC를 설계할 때 가장 중요한 지표는 “퍼포먼스/와트”와 “퍼포먼스/달러”다. 노드 한 대의 절대 성능보다, 전체 창고 단위에서 전기 요금과 장비 가격 대비 얼마나 일하느냐가 핵심이다. 그래서 WSC는 “비싼 슈퍼컴퓨터”가 아니라 “싼 노드 수만 대”의 모양을 띤다. 한 대가 자주 죽지만, 시스템 전체는 항상 살아 있어야 한다. 소프트웨어가 “고장은 일상이다”를 전제로 깔린다. (구글 SRE 책이 두꺼운 이유.)

6.8 멀티프로세서 네트워크 토폴로지

프로세서가 여럿이면 그들 사이에 어떻게 “선”을 깔지가 중요해진다. 이게 네트워크 토폴로지(topology)다. 하나의 회사가 어떻게 부서 간에 의사소통할지를 정하는 조직도와 비슷하다.

주요 토폴로지

• 버스(Bus): 한 줄의 공유 매체. 모두가 듣고, 한 번에 한 명만 말함. 단순하지만 노드가 많아지면 줄 서기 지옥. 작은 SMP에서 흔히 보임.
• 링(Ring): 동그랗게 연결. 옆 노드와만 직접 통신. 멀리 가려면 한 칸씩 돌아야 함. 인텔의 코어 간 인터커넥트가 한때 링이었다.
• 2D 메시(Mesh): 격자. 가로/세로 이웃과 연결. P 노드면 평균 거리는 √P 정도. 매니코어 칩(Intel Xeon Phi 등)에서 자주 본다.
• 토러스(Torus): 메시의 가장자리를 반대 끝과 연결해 도넛 모양. 평균 거리가 줄어든다. 슈퍼컴퓨팅에서 인기.
• 하이퍼큐브(Hypercube): n차원 큐브. P=2ⁿ개의 노드, 각 노드는 n개의 이웃을 가짐. 평균 거리가 log P. 깔끔하지만 차원이 늘면 와이어가 많아진다.
• 완전 연결(Fully connected): 모든 노드가 모든 노드와 직접 연결. 평균 거리 1. 이상적이지만 연결 수가 P²이라 작은 시스템에서만 가능.

이등분 대역폭(Bisection Bandwidth)

토폴로지를 비교하는 가장 좋은 한 가지 지표가 이등분 대역폭(bisection bandwidth)이다. 네트워크를 정확히 두 절반으로 잘랐을 때, 그 절단선을 가로지르는 링크들의 총 대역폭을 잰다. “최악의 경우 절반-대-절반 통신 속도”라고 보면 된다. 이게 클수록 알 수 없는 패턴의 트래픽도 잘 처리한다.

좋은 토폴로지는 “비용은 P log P 정도로 자라면서, 이등분 대역폭은 충분히 큰” 균형을 찾는 일이다. 하이퍼큐브가 이 점에서 인기 있는 이유다. 큰 데이터센터에서는 “팻 트리(fat tree)”나 “Clos 네트워크” 같은 변형이 표준이다.

6.10 벤치마크와 루프라인 모델

병렬 프로세서를 평가할 때 “이건 빠른가?”를 어떻게 측정할까? 슈퍼컴 업계의 LINPACK, 머신러닝의 MLPerf, 그래픽의 3DMark 등 분야별 벤치마크가 있지만, 그것들 너머에 한 가지 강력한 분석 도구가 있다. 루프라인 모델(Roofline Model)이다.

두 천장: 연산과 대역폭

어떤 코드든 둘 중 하나가 천장을 친다.

• 연산 천장(Peak FLOPS): 프로세서가 초당 낼 수 있는 부동소수점 연산의 최대치.
• 대역폭 천장(Peak Memory Bandwidth): 메모리에서 초당 끌어올 수 있는 바이트의 최대치.

어느 쪽이 천장이 되느냐는 코드의 산술 강도(arithmetic intensity, AI)가 결정한다. AI는 “바이트 한 개당 몇 번의 부동소수점 연산을 하는가”다. 단위는 FLOPS/byte. 행렬 곱셈은 AI가 크고 (한 바이트로 여러 곱셈을 만들어 냄), 단순 벡터 더하기는 AI가 작다(메모리에서 가져온 바이트당 한 번씩만 씀).

루프라인 그림은 가로축에 산술 강도, 세로축에 달성 가능한 FLOPS를 둔다.

달성 가능 FLOPS = min( Peak FLOPS,  AI × Peak Bandwidth )

  ↑ FLOPS
  |          ____ Peak FLOPS (수평 천장)
  |        /
  |      /
  |    /  ← 기울기 = Peak Bandwidth
  |  /
  |/
  +----------------------→ AI(산술 강도)
   대역폭 한계 영역  | 연산 한계 영역

AI가 작으면 직선(대역폭) 영역에 묶이고, AI가 충분히 크면 수평선(연산) 천장에 닿는다. 그 사이의 “꺾이는 점”에서 두 천장이 만난다. 이 점이 “이 머신에서 연산 한계에 도달하기 위해 필요한 최소 AI”다.

루프라인의 강력함은, “지금 내 코드가 어느 천장에 부딪히고 있나?”를 한 그림으로 보여준다는 점이다. 만약 내 점이 대역폭 직선 위에 놓여 있다면 “FPU 더 빠른 칩 사봐야 소용없음. 메모리 대역폭이 문제”라는 진단이 바로 나온다. 반대 경우면 “캐시 친화 최적화는 더 해도 의미 없음. 부동소수점 마이크로아키텍처가 한계.”

6.11 실제 비교: i7 960 vs Tesla

교과서가 들고 있는 벤치마크 중 하나는 인텔 Core i7 960(4코어 멀티스레딩 SMP)과 NVIDIA Tesla GPU의 여러 커널 비교다. 결론을 한 줄로 정리하면 — “많은 데이터에 같은 연산을 산술 강도가 어느 정도 있는 상태로 하면 GPU가 훨씬 강하다. 분기가 많고 데이터가 적은 워크로드는 CPU가 더 잘한다.”

벤치 결과는 대체로 GPU가 1.4×~14× 정도 빠른 경향. 하지만 이건 워크로드와 최적화 정도에 매우 민감하다. 같은 코드를 단순 포팅만 하면 GPU 이점이 작고, 데이터 레이아웃·메모리 합치기·블록 사이즈 튜닝까지 다 해야 제대로 된 격차가 나온다. 즉 “GPU가 빠르다”는 명제 뒤에는 “충분히 잘 짠 GPU 코드가 빠르다”는 단서가 항상 붙는다.

더 흥미로운 관찰은, 같은 워크로드라도 “데이터 크기”가 작아지면 GPU 이점이 사라진다는 거다. 커널 시작 오버헤드, 호스트-디바이스 메모리 전송 비용 같은 고정 비용이 작은 작업에서 부각되기 때문이다. 그래서 GPU는 “큰 일을 통째로 보내야” 진가를 발휘한다. 우체국에서 편지 한 통은 비행기 띄울 일이 아닌 셈.

6.13 오류와 함정

병렬 컴퓨팅 세계에 자주 출몰하는 함정들을 모았다.

• 피크 성능을 실제 성능으로 착각하기. “이 GPU는 30 TFLOPS다”라는 카탈로그 숫자는 모든 유닛이 이상적인 명령어를 매 사이클 받을 때나 나오는 환상의 숫자다. 실제 워크로드에서는 그 절반도 어렵다. 루프라인은 이 환상을 부순다.
• 암달의 법칙 무시. “코어 16개니까 16배”는 거의 절대 안 일어난다. 직렬 부분이 1%만 있어도 16배는 못 간다. 측정하지 않으면 항상 실망한다.
• 스트롱 스케일링과 위크 스케일링을 섞어 보고하기. “1024개 노드에서 1024×의 속도 향상” 같은 슬로건이 있다면 — 거의 항상 위크 스케일링이다. 솔직한 강한 스케일링 그래프를 요청하자.
• 락을 너무 길게 잡기. 임계 구역이 큰 락 안에 들어가면 그 부분 전체가 직렬이 된다. 락은 짧고 자주 바뀌어야 한다. 그리고 가능하면 락을 안 쓰는 자료구조(원자 연산 기반)를 고려하자.
• 거짓 공유(false sharing). 두 코어가 다른 변수를 갱신하지만, 그 두 변수가 같은 캐시 라인에 있으면 캐시 일관성 프로토콜이 라인을 핑퐁시키며 성능을 죽인다. 정렬과 패딩으로 변수가 다른 라인에 있게 하자.
• 메모리 일관성 모델을 가정으로 가지기. “쓰기는 다른 코어에 즉시 보일 거야”는 일반적으로 거짓이다. 약한 메모리 모델에서는 fence/barrier 명령어 없이는 순서가 보장되지 않는다.
• GPU에 잘 안 맞는 워크로드 강제 이식. 분기 발산이 큰 코드, 의존이 긴 시퀀셜 코드, 데이터가 작은 코드는 GPU에서 오히려 손해다. “GPU = 항상 빠름”은 미신.
• 네트워크 토폴로지 무시. 클러스터 위에서 노드 간 통신이 자주 일어나는 패턴이라면 이등분 대역폭이 진짜 병목이다. 노드 한 대 더 빠른 게 답이 아니다.

6.14 결론 — 병렬은 피할 수 없는 길

클럭 속도의 잔치는 끝났다. 무어의 법칙도 고전적 의미에서는 휘청이고 있다. 미래의 성능 향상은 점점 더 “병렬을 더 잘하기”에서 나올 것이다 — 이미 그렇다. 노트북 안의 8코어 16스레드, GPU의 수천 개 SIMD 레인, 클라우드 안의 수십만 노드까지, 우리는 이미 어마어마하게 병렬화된 세상에서 산다.

그런데 병렬 프로그래밍은 여전히 어렵다. 암달의 법칙은 잔인하고, 로드 밸런싱은 까다롭고, 캐시 일관성은 미묘하고, 거짓 공유는 음흉하고, 분기 발산은 GPU에서 골치다. 새로 들어오는 학생/엔지니어가 “병렬은 어려워” 라며 한숨 쉬는 건, 지금까지의 모든 위대한 병렬 컴퓨팅 연구자들도 똑같이 했던 한숨이다. 다행이라면, 그 한숨이 고스란히 노하우로 쌓여 있다는 것.

다섯 가지 큰 흐름

1. SIMD/벡터의 부활. ARM SVE, RISC-V V 확장이 “벡터 길이 가변”의 우아함을 다시 들고 온다. AI 추론도 이 흐름에 잘 맞는다.
2. 이종 컴퓨팅(heterogeneous computing). CPU + GPU + NPU + DSP가 한 시스템에 같이 있는 게 표준이 됐다. 어떤 일을 어디서 돌릴지 정하는 게 아키텍처와 컴파일러의 새 과제.
3. 도메인 특화 가속기(DSA). 텐서 연산만 잘하는 칩, 그래프만 잘하는 칩 — 일반 목적 칩의 효율 한계를 도메인 지식으로 깬다. TPU가 대표 예.
4. 창고급 컴퓨터의 본격화. 클라우드는 이제 “남의 컴퓨터”가 아니라 “주된 컴퓨팅 환경”이다. WSC를 잘 쓰는 능력 자체가 산업 경쟁력.
5. 병렬 친화 프로그래밍 모델의 진화. OpenMP, MPI, CUDA에 이어 SYCL, OpenACC, Triton, JAX/XLA, MLIR — 새 추상이 계속 나온다. 그리고 이런 추상이 “직렬 vs 병렬” 선택을 자동화하려고 한다.

이 책의 마지막 장이다. 1장에서 “컴퓨터는 빠르고, 일은 어려워”로 시작했던 우리는 이제 6장의 끝에서 “빠른 컴퓨터를 모아 더 어려운 일을 같이 한다”로 도착했다. 다음 컴퓨터 구조의 이야기는 여러분이 직접 쓰는 시대다. 코어 1024개를 잘 굴리는 사람이 되시길.