[GPU 엔지니어링 #2] GPU 구조와 SIMT

이슈 배경 및 포스팅 목표

최근 ChatGPT와 같은 거대 언어 모델(LLM)과 딥러닝이 폭발적으로 성장하면서, GPU는 단순한 그래픽 처리 장치를 넘어 'AI 시대의 심장'이 되었습니다. 하지만 많은 분들이 비싼 GPU 서버를 구축하고도 "왜 기대만큼 연산 속도가 나오지 않을까?" 혹은 "왜 메모리(VRAM)는 마음대로 증설할 수 없을까?"라는 엔지니어링 문제에 부딪히곤 합니다.

이러한 문제를 근본적으로 해결하려면 단순히 코드를 잘 짜는 것을 넘어, 코드가 실행되는 하드웨어 밑바닥의 작동 원리(아키텍처)를 이해해야 합니다.

오늘 글에서는 CPU와 GPU의 설계 철학 차이에서 출발해, GPU의 핵심 실행 모델인 SIMT, 성능 저하의 주범인 분기 발산(Warp Divergence), 그리고 VRAM의 물리적 한계까지 흩어져 있던 개념들을 하나의 흐름으로 완벽하게 정리해 보겠습니다.

 

GPU의 구조: SM, CUDA 코어, Warp

 

• SM (Streaming Multiprocessor) - "독립적인 작업장": GPU 칩 내부에는 수십 개에서 최상위 모델의 경우 140여 개 이상의 SM이 탑재되어 있습니다. SM은 단순한 코어의 묶음이 아니라, 그 자체로 명령어를 해석하고 메모리를 공유하며 수백 개의 연산 유닛을 품고 있는 완전한 독립 작업장입니다.
• CUDA 코어 (ALU) - "단순 반복 작업자": SM 내부에 존재하는 실질적인 연산 유닛(ALU)입니다. CPU의 ALU가 복잡한 논리 연산과 상황 대처에 능한 '만능 박사'라면, GPU의 CUDA 코어는 행렬 곱셈이나 덧셈 같은 단순 사칙연산에 극도로 특화된 '단순 반복 숙련공'입니다. 이들은 크기가 매우 작아 하나의 SM 안에 수백 개씩 들어갈 수 있습니다.
• 워프 (Warp) - "32명의 작업 조": 아무리 작업자가 많아도 개별적으로 지시를 내리면 통제가 불가능합니다. 그래서 NVIDIA는 하드웨어적으로 32개의 스레드(작업자)를 하나의 조(Warp)로 묶어버렸습니다. 워프에 속한 32개의 스레드는 단 하나의 명령어를 받아 **완벽하게 똑같은 연산(Lock-step)**을 동시에 수행합니다. 다루는 데이터의 위치만 다를 뿐, 움직임은 완전히 일치해야 합니다.

 

[핵심 원리: Lock-step 실행] 가장 중요한 작동 원리는 하나의 워프 스케줄러가 "단 1개의 명령어"를 가져오면, 워프에 속한 32개의 스레드가 각자의 CUDA 코어에 매핑되어 "정확히 똑같은 연산"을 동시에 수행한다는 것입니다. 다루는 데이터의 위치만 다를 뿐 행동은 완벽히 일치합니다.

 

워프와 워프 스케줄러의 심층 이해 (+ SM 파티셔닝)

이 워프들에게 일감을 던져주는 주체는 누구일까요? 바로 SM 내부에 존재하는 **워프 스케줄러(Warp Scheduler)**입니다.

• 워프(Warp) vs 워프 스케줄러(Warp Scheduler)의 차이:
  • 워프: 실행되는 스레드들의 묶음(논리적/물리적 대상)입니다. (예: 공장의 작업 조)
  • 워프 스케줄러: SM 내부에서 대기 중인 수십 개의 워프들 중, "지금 당장 메모리에서 데이터를 가져와 연산할 준비가 끝난 워프"를 찾아내어 명령어를 하달(Dispatch)하는 하드웨어 제어 장치입니다. (예: 공장의 작업 반장)

스케줄러의 가장 중요한 역할은 '지연 시간 은닉(Latency Hiding)'입니다. A 워프가 VRAM에서 데이터를 가져오느라 시간이 걸려 멍하니 서 있다면, 스케줄러는 즉시 계산 준비가 끝난 B 워프를 CUDA 코어에 투입시켜 하드웨어가 1초도 쉬지 않게 만듭니다.

 

[심화: SM 파티셔닝 (Sub-core Partitioning) 체제의 도입]

GPU 세대가 발전하면서 SM 하나에 들어가는 CUDA 코어의 수가 128개 이상으로 폭발적으로 늘어났습니다. 그러자 스케줄러(반장) 1명이 128명의 작업자에게 제때 명령어를 나눠주지 못하는 병목 현상이 발생했습니다. 이를 해결하기 위해 최신 아키텍처에서는 거대한 SM 내부를 4개의 구역(Sub-core)으로 쪼개는 파티셔닝(Partitioning) 설계를 도입했습니다. 각 구역마다 1명의 전담 워프 스케줄러와 전용 레지스터를 독립적으로 배치하여, 4명의 반장이 동시에 명령을 쏟아내도록 만들어 동시 처리량을 극대화한 것입니다.

 

SIMD에서 SIMT로

이러한 워프 기반의 실행 구조를 바탕으로 NVIDIA는 SIMT(Single Instruction, Multiple Threads)라는 혁신적인 실행 모델을 정립했습니다. 이를 이해하려면 과거의 SIMD와 비교해 보아야 합니다.

• SIMD (Single Instruction, Multiple Data): 하나의 명령어로 여러 데이터를 처리하는 하드웨어 방식입니다. 하지만 이 방식을 쓰려면 개발자가 코드 단에서 직접 "데이터 4개를 1세트로 묶어라"라고 명시하는 복잡한 벡터(Vector) 프로그래밍을 해야만 했습니다. 매우 귀찮고 오류가 나기 쉬운 작업입니다.
• SIMT (Single Instruction, Multiple Threads): 반면 SIMT 모델에서는 개발자가 그냥 "스레드 1개가 데이터 1개를 처리한다"는 스칼라(Scalar) 코드를 아주 단순하게 작성하기만 하면 됩니다. 코드를 컴파일하고 실행하면, GPU 안의 '워프 스케줄러'가 알아서 스레드 32개를 모아 하드웨어 레벨에서 SIMD 방식으로 묶어서 실행해 줍니다. 개발자의 편의성과 하드웨어의 병렬 처리 성능을 완벽하게 조화시킨 패러다임입니다.

 

분기 발산 (Warp Divergence): if문이 초래하는 치명적 병목

SIMT 구조는 32개의 스레드가 무조건 똑같은 명령어를 실행해야 한다는 강력한 전제 조건이 있습니다.

만약, 이 스레드들이 if-else 같은 조건문을 만나 실행 경로가 엇갈리게 되면 큰 병목현상이 발생합니다.

 

GPU는 CPU처럼 분기를 예측하거나 두 갈래 길을 동시에 처리할 수 있는 복잡한 회로가 없습니다. 따라서 워프 내에서 경로가 엇갈리면, 하드웨어는 울며 겨자 먹기로 각 경로를 순차적으로 실행(직렬화)해야만 합니다.

// [문제 상황] 데이터 값에 따라 스레드의 동작이 달라지는 분기문 발생
__global__ void process_data(int *data) {
    int tid = threadIdx.x; // 워프 내 각 스레드의 고유 ID (0~31)

    // 워프 내 32개의 스레드가 이 조건문을 만납니다.
    if (data[tid] > 0) {
        // [경로 A] 16개의 스레드가 이 연산을 수행할 때,
        // 나머지 16개는 유용한 작업을 하지 못하고 마스킹(대기) 상태가 됩니다.
        data[tid] = data[tid] * 2; 
    } else {
        // [경로 B] A 경로가 끝나면, 이번엔 반대로 A 스레드들이 멈추고
        // 나머지 16개 스레드가 순차적으로 이 연산을 수행합니다.
        data[tid] = data[tid] * -1;
    }
}

 

이처럼 실행 경로가 갈라지는 현상을 워프 발산(Warp Divergence)이라고 부릅니다. 이 현상이 발생하면 1번 사이클에 끝날 일이 2번 사이클로 늘어나고, 매 시간마다 연산 유닛(CUDA 코어)의 50%가 아무것도 하지 않고 자원을 낭비하게 됩니다. 조건문이 중첩될수록 효율은 기하급수적으로 떨어지며, 이는 GPU 프로그래밍에서 성능을 갉아먹는 1순위 주범입니다.

 

VRAM의 물리적 한계: 왜 마음대로 꽂을 수 없을까?

수천 개의 코어가 워프 단위로 엄청난 연산을 쏟아내려면, 그에 걸맞은 속도로 데이터를 공급해 주는 메모리가 필수적입니다. 여기서 일반적인 시스템 RAM과 GPU 전용인 VRAM의 아키텍처 차이가 명확해집니다.

• 시스템 RAM (지연 시간 최적화): CPU가 다양한 상황에 대처하기 위해 필요한 소규모의 데이터를 최대한 빨리(응답 속도 위주) 가져오는 데 특화되어 있습니다. 도로로 치면 좁지만 제한 속도가 무제한인 고속도로입니다. (메모리 버스 폭: 보통 64-bit)
• VRAM (대역폭 극대화): GPU의 수많은 코어에 엄청난 양의 데이터를 한 번에 때려 부어야 하므로, 응답 속도보다는 데이터가 지나가는 통로 자체를 무식할 정도로 넓힌 메모리입니다. 도로로 치면 차선이 수천 개인 초대형 고속도로입니다. (메모리 버스 폭: 256-bit ~ 5120-bit 이상)

Q. 왜 컴퓨터 RAM처럼 슬롯에 꽂아서 추가하지 못할까요?

초당 수백 기가바이트(GB/s)에서 테라바이트(TB/s)의 막대한 데이터를 전송하려면 수백, 수천 가닥의 전기 배선(Pin)이 필요합니다.

데이터 전송 클럭이 이렇게 극단적으로 높을 때, 칩과 메모리 사이의 거리가 단 몇 밀리미터(mm)라도 멀어지면 찰나의 시간 차이로 인해 전기 신호가 훼손되고 노이즈가 발생합니다.

따라서 VRAM은 물리적인 슬롯(DIMM) 형태를 거치지 않고, 반드시 GPU 칩셋 바로 옆 기판에 바짝 붙여 납땜(Soldered)해야만 신호 무결성을 유지할 수 있습니다. 특히 최신 데이터센터용 AI 칩(H100 등)은 이 거리마저도 아깝다고 판단하여, 아예 HBM(High Bandwidth Memory)을 칩 다이(Die) 옆에 2.5D 패키징으로 하나의 덩어리로 만들어버립니다. 이 때문에 사용자는 최초에 구매한 VRAM 용량을 영원히 안고 가야 하는 물리적 한계를 가지게 됩니다.

 

따라서, VRAM은 증설이 힘드니, Out of Memory 에러가 발생하지 않도록 모델의 Batch Size를 세밀하게 튜닝해야 합니다.