Featured image of post GPU의 초병렬 아키텍처와 CUDA의 물리: SIMT, 워프, 텐서 코어의 계산 원리
아키텍처 기술

GPU의 초병렬 아키텍처와 CUDA의 물리: SIMT, 워프, 텐서 코어의 계산 원리

높은 처리량을 극한까지 추구하는 GPU의 내부 설계. SM, 워프 스케줄링, 텐서 코어, 공유 메모리 최적화의 진수.

GPU의 초병렬 아키텍처와 CUDA의 물리: SIMT, 워프, 텐서 코어의 계산 원리

현대의 고도화된 계산 과학, 인공지능, 딥러닝, 그리고 고해상도 컴퓨터 그래픽스를 뒷받침하는 근간 기술, 그것이 바로 GPU(Graphics Processing Unit)입니다. 본 문서에서는 GPU의 아키텍처와 그 위에서 동작하는 병렬 계산 기반인 CUDA(Compute Unified Device Architecture)의 물리적, 하드웨어적 측면을 깊이 파헤칩니다. 단순한 프로그래밍 문법이 아니라, 하드웨어가 ‘왜 그렇게 설계되었는지’, ‘어떻게 극한의 계산 처리량을 낼 수 있는지’를 스트리밍 멀티프로세서(SM), SIMT 실행 모델, 워프 스케줄링, 텐서 코어, 그리고 메모리 계층의 관점에서 철저하게 해부합니다.

제1장: CPU와 GPU의 설계 사상의 분기점

1.1 짧은 지연 시간 추구 vs 높은 처리량 추구

범용 프로세서인 CPU(Central Processing Unit)와 병렬 계산에 특화된 GPU는 그 탄생 배경에서부터 설계 사상이 근본적으로 다릅니다. CPU는 ‘어떻게 하면 하나의 작업(스레드)을 빠르게 끝낼 것인가’라는 ‘낮은 레이턴시(지연 시간 최소화)‘를 지상 과제로 삼아 진화해 왔습니다. 반면 GPU는 ‘대량의 작업을 묶어 전체적으로 단위 시간당 얼마나 많은 처리를 완료할 수 있는가’라는 ‘높은 스루풋(처리량 극대화)‘을 추구합니다.

CPU는 운영체제 제어, 복잡한 분기 조건이 수반되는 애플리케이션 실행, 사용자의 무작위적인 인터럽트 처리 등 예측 불가능한 작업을 신속하게 처리해야 합니다. 이를 위해 고도의 분기 예측 회로, 비순차적 실행(Out-of-Order Execution, 명령어의 순서를 바꿔서 실행하는 메커니즘), 거대한 L1/L2/L3 캐시 메모리를 탑재하여 메모리 접근 지연을 숨기면서 단일 스레드의 성능을 극한까지 높이고 있습니다.

이에 반해 GPU는 원래 화면상의 수백만 픽셀에 대해 동일한 셰이딩 연산을 적용하는 것과 같이 고도로 병렬화 가능한 작업을 처리하기 위해 탄생했습니다. 복잡한 제어 회로나 거대한 캐시에 다이(Die) 면적을 할애하는 대신, 단순한 연산기(ALU: Arithmetic Logic Unit)를 한계까지 꽉 채우는 선택을 했습니다.

1.2 다이 면적에서 캐시, 제어 회로, ALU의 할당 비율

실리콘 다이(반도체 칩)의 제한된 면적(트랜지스터 예산)을 어떻게 배분할 것인가가 두 아키텍처의 차이를 결정짓습니다.

  • CPU의 다이 면적 배분: 다이의 절반 이상이 대용량 캐시 메모리(SRAM)와 고도의 제어 회로(분기 예측, 명령어 페치, 디코드, 스케줄링 등)로 채워져 있습니다. 실제 연산을 수행하는 ALU가 차지하는 비율은 상대적으로 작습니다.
  • GPU의 다이 면적 배분: 캐시 메모리나 제어 회로는 필요 최소한으로 억제되어 있으며, 다이의 대부분이 수천에서 수만 개에 달하는 ALU(CUDA 코어)로 채워져 있습니다.

GPU는 메모리 접근 지연(레이턴시)을 캐시로 숨기는 것이 아니라, ‘컨텍스트 스위칭’을 통해 숨깁니다. 특정 스레드 그룹이 메모리에서 데이터가 도착하기를 기다리는 동안, 즉각적으로 다른 스레드 그룹의 연산을 실행함으로써 연산기를 항상 가동 상태(높은 점유율: Occupancy)로 유지합니다. 이것이 GPU에서의 ‘높은 처리량 추구’의 물리적 구현입니다. 하드웨어 수준의 멀티스레딩(Hardware Multithreading)이 매우 가볍게 이루어지기 때문에 수천~수만 개의 동시 실행 스레드가 존재한다는 것을 전제로 합니다.

제2장: SIMT 실행 모델의 본질

2.1 SIMD와 SIMT의 차이

병렬 처리의 분류로 플린의 분류학(Flynn’s taxonomy)이 있지만, GPU의 실행 모델은 종종 SIMD(Single Instruction, Multiple Data)와 비교됩니다. CPU의 벡터 확장 명령어(AVX 등)는 순수한 SIMD이며, 하나의 명령어로 여러 데이터(예: 256비트 폭의 레지스터에 저장된 8개의 32비트 부동소수점 수)를 동시에 처리합니다. SIMD에서는 데이터 요소마다 다른 분기(if-else)를 수행하는 것이 매우 어렵습니다.

반면, NVIDIA가 제창한 CUDA의 실행 모델은 **SIMT(Single Instruction, Multiple Threads)**라고 불립니다. SIMT에서는 여러 독립적인 ‘스레드’가 그룹(후술할 ‘워프’)을 형성하고, 같은 명령어를 공유하여 실행합니다. 그러나 SIMD와 달리 SIMT의 각 스레드는 **독립적인 레지스터 상태와 명령어 주소 카운터(프로그래밍 모델 상)**를 가지고 있습니다. 이를 통해 프로그래머는 마치 각 스레드가 독립적으로 동작하는 것처럼 코드를 작성할 수 있습니다.

2.2 32스레드 단위의 ‘워프(Warp)’

GPU 하드웨어는 스레드를 개별적으로 스케줄링하는 것이 아니라, **32개의 스레드를 하나로 묶은 ‘워프(Warp)’**라는 단위로 관리하고 실행합니다. (AMD의 GPU에서는 Wavefront라고 불리며, 64스레드 단위 등이 채택되기도 합니다).

스트리밍 멀티프로세서(SM) 내의 명령어 페치·디코드 유닛은 워프 단위로 하나의 명령어를 페치하고, 워프 내의 32개 스레드 모두에게 같은 명령어를 발행(디스패치)합니다. 즉, 워프 내의 32스레드는 물리적으로 완전히 동시에, 같은 명령어를 각자가 가진 다른 데이터에 대해 실행합니다. 이것이 SIMT의 핵심입니다.

2.3 워프 다이버전스(Warp Divergence, 분기 불일치)의 물리적 페널티

각 스레드가 독립적인 프로그램 카운터를 가진 것처럼 행동할 수 있다고는 하지만, 물리적으로는 워프 내의 모든 스레드가 동일한 명령어를 실행해야 합니다. 그렇다면 코드 내에 if-else와 같은 조건 분기가 있고, 워프 내의 스레드들 사이에서 분기 조건의 참/거짓이 나뉘면 어떻게 될까요?

이 현상을 **워프 다이버전스(Warp Divergence: 분기 불일치)**라고 부릅니다.

워프 다이버전스가 발생하면 하드웨어는 다음 단계로 처리를 수행합니다.

  1. 먼저 if 조건이 참이 된 스레드(활성 스레드)에 대해서만 명령어를 실행합니다. 이때 조건이 거짓이 된 스레드는 ‘마스킹(비활성화)‘되어 연산 결과가 기록되지 않습니다.
  2. 다음으로 else 조건(또는 조건이 거짓일 경우의 경로)으로 전환하여, 이번에는 조금 전 마스킹되었던 스레드를 활성화하고 참이었던 스레드를 마스킹하여 명령어를 실행합니다.

즉, 분기 경로가 여러 개인 경우 하드웨어는 해당 경로들을 병렬이 아닌 직렬(순차적)로 실행할 수밖에 없게 됩니다. 극단적인 예로, 워프 내의 32스레드가 32가지의 각기 다른 분기 경로를 거치게 되면 실행 시간은 32배로 뛰어오릅니다. 워프 다이버전스는 GPU의 계산 처리량을 격감시키는 가장 큰 요인 중 하나이며, 알고리즘 설계에서 가장 피해야 할 안티 패턴입니다. 물리적으로는 ALU가 전력을 소비하고 있음에도 불구하고, 마스킹되어 유효한 계산 결과를 생성하지 못하는 ‘헛된 사이클’이 발생하고 있음을 의미합니다.

제3장: 스트리밍 멀티프로세서(SM)의 하드웨어 해부

GPU는 다수의 **스트리밍 멀티프로세서(SM: Streaming Multiprocessor)**의 집합체로 구성되어 있습니다. SM이야말로 GPU의 진정한 계산 엔진입니다. 최신 아키텍처(예: Hopper H100)에서는 하나의 GPU 다이에 100개 이상의 SM이 탑재되어 있습니다.

3.1 SM 내부의 파이프라인 구성

SM은 내부에 다시 여러 개의 서브 파티션(보통 4개)으로 분할되어 있으며, 각각이 독립적인 워프 스케줄러와 디스패치 유닛을 가지고 있습니다.

  • 워프 스케줄러(Warp Scheduler): 실행 가능한 상태(레지스터나 메모리가 준비된 상태)에 있는 워프를 선택합니다. GPU 스케줄러는 제로 오버헤드로 워프를 전환할 수 있으며, 이것이 메모리 접근 지연을 숨기는 열쇠가 됩니다.
  • 디스패치 유닛(Dispatch Unit): 스케줄링된 워프에 대해 명령어를 발행합니다.
  • CUDA 코어(INT32 / FP32 / FP64 ALU): 실제 정수 연산이나 부동소수점 연산을 수행하는 유닛입니다.
  • 로드/스토어 유닛(LD/ST Unit): 메모리에 대한 읽기/쓰기를 담당합니다.
  • 특수 함수 유닛(SFU: Special Function Unit): sin, cos, exp, 역수 등의 초월 함수를 빠르게 계산하는 전용 하드웨어입니다.

명령어 파이프라인은 매우 깊게 설계되어 있으며, 페치, 디코드, 스케줄링, 레지스터 읽기, 실행(다수 사이클), 라이트 백의 각 단계를 거칩니다. FP32의 FMA(Fused Multiply-Add) 연산 지연 시간은 보통 수 사이클에서 십여 사이클 정도 걸리지만, 매 사이클마다 다른 워프에서 명령어를 발행함으로써 파이프라인을 항상 가득 찬 상태로 유지합니다.

3.2 거대한 레지스터 파일과 레지스터 압박

SM에는 CPU와는 비교할 수 없을 정도로 거대한 레지스터 파일이 탑재되어 있습니다(예: 1SM당 64KB~256KB의 SRAM). 이는 SM에서 동시 실행되는 수천 개 스레드의 컨텍스트를 모두 유지하기 위해서입니다.

컨텍스트 스위칭이 제로 사이클에 완료되는 이유는 스레드의 레지스터 상태를 메모리로 대피(스필, spill)시킬 필요가 없기 때문입니다. 하지만 스레드당 사용하는 레지스터 수가 증가하면, SM 내에서 동시에 기동할 수 있는 워프의 수(점유율)가 떨어집니다. 이를 **레지스터 압박(Register Pressure)**이라고 부릅니다. 레지스터가 고갈되면 데이터는 느린 로컬 메모리(물리적으로는 글로벌 메모리의 일부)로 스필되며, 치명적인 성능 저하를 일으킵니다.

3.3 공유 메모리(Shared Memory)와 뱅크 충돌

SM에는 프로그래머가 명시적으로 제어할 수 있는 초고속 온칩 메모리인 **공유 메모리(Shared Memory)**가 존재합니다. L1 캐시와 동일한 물리 SRAM 영역을 공유하지만 명시적인 데이터 캐시로 기능하며, 블록 내 스레드 간 데이터 공유 및 동기화에 사용됩니다.

공유 메모리의 물리적 구조는 **메모리 뱅크(Memory Banks)**라고 불리는 여러 개의 독립된 모듈(보통 32개)로 나뉘어 있습니다. 연속되는 32비트 주소는 다른 뱅크에 인터리브(할당)됩니다.

워프 내의 32스레드가 서로 다른 뱅크에 동시에 접근할 경우, 접근은 완전히 병렬로(1사이클 만에) 처리됩니다. 이를 뱅크 충돌 없음(Bank Conflict-Free)이라고 합니다. 그러나 여러 스레드가 같은 뱅크의 서로 다른 주소에 동시에 접근하려고 하면 요청이 직렬화되어 페널티(지연)가 발생합니다. 이를 **뱅크 충돌(Bank Conflict)**이라고 부릅니다. 예를 들어, 2-way 뱅크 충돌이라면 접근 시간은 2배가 되고, 최악의 경우 32-way 충돌에서는 32배로 지연됩니다. 행렬 전치와 같은 알고리즘에서는 스트라이드(보폭) 접근으로 인해 심각한 뱅크 충돌이 발생하므로, 패딩(의미 없는 데이터를 삽입하여 메모리 주소를 비켜가게 하는 기법)을 사용하여 충돌을 회피하는 고도화된 최적화가 필수적입니다.

제4장: 텐서 코어(Tensor Core)의 적화연산 파이프라인

Volta 아키텍처에서 처음 도입되어 이후 GPU의 성능을 비약적으로 끌어올린 혁명적인 하드웨어가 바로 **텐서 코어(Tensor Core)**입니다. AI와 딥러닝의 폭발적인 발전은 텐서 코어 없이는 설명할 수 없습니다.

4.1 행렬 적화연산(MMA)의 하드웨어 구현

딥러닝 연산의 대부분은 신경망의 가중치 행렬과 입력 데이터의 행렬곱(GEMM: General Matrix Multiply)입니다. 수식으로는 $D = A \times B + C$ ($A, B$는 입력 행렬, $C$는 누산기(accumulator) 행렬)로 표현됩니다.

기존의 CUDA 코어에서는 이 행렬곱을 1개 요소씩 FMA(Fused Multiply-Add) 명령어를 사용하여 계산했습니다. 반면 텐서 코어는 작은 행렬(예: 4x4나 16x16)의 곱셈 및 덧셈 연산을 하드웨어 레벨에서 1사이클(또는 수 사이클) 만에 실행하는 전용 회로입니다.

물리적으로는 수십 개에서 수백 개의 곱셈기와 거대한 덧셈 트리(Adder Tree)를 와이어로 직결하여, 중간 결과를 레지스터에 다시 쓰지 않고 단번에 적화연산을 완료합니다. 이로 인해 일반 CUDA 코어에 비해 면적당 연산 처리량(TFLOPS)이 비교할 수 없을 만큼 높습니다.

4.2 혼합 정밀도(Mixed-Precision)의 진수

텐서 코어의 또 다른 핵심은 혼합 정밀도(Mixed-Precision) 연산 지원입니다. 딥러닝에서는 계산 과정에서 높은 정밀도(FP32/FP64)를 필요로 하지 않는 상황이 자주 있습니다. 텐서 코어는 입력 행렬 $A$와 $B$를 낮은 정밀도(FP16, BF16, 또는 더 낮은 FP8, INT8, INT4)로 읽어들여 내부 곱셈을 낮은 정밀도로 수행한 후, 덧셈(누산) 과정을 더 높은 정밀도(FP32나 INT32)로 수행하는 파이프라인을 갖추고 있습니다.

  • FP16 / BF16: 학습의 표준. BF16(Bfloat16)은 지수부가 FP32와 같은 8비트라서 다이내믹 레인지가 넓어 기울기 소실을 방지하기 쉽습니다.
  • FP8 / INT8 / INT4: 추론(Inference) 고속화의 비장의 카드. 데이터 전송량(메모리 대역폭)도 감소하므로 처리량이 극적으로 향상됩니다.

Hopper 아키텍처에서는 Transformer 모델의 계산을 획기적으로 가속하는 ‘FP8 Tensor Core’가 도입되어, FP32 대비 이론상 수십 배의 스루풋을 달성했습니다. 소프트웨어 측(CUDA)에서는 wmma(Warp-Level Matrix Multiply and Accumulate) API나 mma.sync PTX 명령어를 통해 텐서 코어를 직접 구동하고, 워프 내의 스레드가 협력하여 행렬의 단편을 레지스터로 로드·연산·스토어하는 매우 복잡한 컬렉티브(집단) 처리를 수행합니다.

제5장: CUDA 메모리 계층과 최적화 기법

GPU의 계산 능력이 아무리 높더라도, 데이터 공급이 병목이 되면 성능이 나오지 않습니다(메모리 월(Memory Wall) 문제). CUDA 프로그래밍 최적화의 9할은 ‘메모리 접근 최적화’라고 해도 과언이 아닙니다.

5.1 글로벌 메모리의 코얼레싱(Coalescing) 접근

GPU의 메인 메모리(HBM이나 GDDR)인 글로벌 메모리는 매우 넓은 대역폭(예를 들어 수 TB/s)을 가지지만, 지연 시간도 수백 사이클로 매우 큽니다.

글로벌 메모리에 대한 접근 효율을 극대화하는 절대 원칙이 바로 **코얼레싱(Coalescing: 병합)**입니다. GPU의 메모리 컨트롤러는 메모리에 대해 32바이트, 64바이트 또는 128바이트 단위의 트랜잭션으로 접근을 수행합니다. 워프 내의 32스레드가 메모리에 접근할 때, 그 메모리 주소들이 연속된 영역(정렬된 128바이트 경계 내)에 들어맞을 경우, 하드웨어는 이 요청들을 **한 번의 메모리 트랜잭션으로 병합(Coalesce)**하여 처리합니다.

반대로 스레드가 무작위 주소에 접근하거나 스트라이드(간격이 벌어진) 접근을 하게 되면 병합이 이루어지지 않아 여러 트랜잭션이 발생합니다. 이를 ‘비병합(Non-Coalesced) 접근’이라고 부르며, 유효 메모리 대역폭을 10분의 1 이하로 떨어뜨리는 치명적인 성능 버그가 됩니다.

5.2 CUDA C++ 코드 예제: 행렬 전치 최적화와 공유 메모리

아래는 비병합 접근을 회피하고 공유 메모리를 활용하여 성능을 획기적으로 개선하는 행렬 전치(Matrix Transpose)의 최적화된 커널 코드 예시입니다.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35

// 공유 메모리를 사용한 최적화된 행렬 전치 커널
// TILE_DIM = 32, BLOCK_ROWS = 8 로 설정
__global__ void transposeSharedOptimized(float *odata, const float *idata, int width, int height) {
    // 공유 메모리 선언. 뱅크 충돌을 피하기 위해 '+ 1' 패딩을 넣음
    __shared__ float tile[TILE_DIM][TILE_DIM + 1];

    // 입력 행렬상의 글로벌 인덱스(읽기용)
    int xIndex = blockIdx.x * TILE_DIM + threadIdx.x;
    int yIndex = blockIdx.y * TILE_DIM + threadIdx.y;

    // 출력 행렬상의 글로벌 인덱스(쓰기용)
    // 블록의 X와 Y를 맞바꾸어 쓰기 시의 코얼레싱을 확보
    int xIndex_out = blockIdx.y * TILE_DIM + threadIdx.x;
    int yIndex_out = blockIdx.x * TILE_DIM + threadIdx.y;

    // 1. 글로벌 메모리에서 공유 메모리로 읽기 (코얼레싱 접근)
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        if (xIndex < width && (yIndex + j) < height) {
            // 스레드는 연속된 주소를 읽어들임
            tile[threadIdx.y + j][threadIdx.x] = idata[(yIndex + j) * width + xIndex];
        }
    }

    // 블록 내 모든 스레드의 읽기 완료를 동기화
    __syncthreads();

    // 2. 공유 메모리에서 글로벌 메모리로 쓰기 (코얼레싱 접근)
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        if (xIndex_out < height && (yIndex_out + j) < width) {
            // 공유 메모리에서는 전치된 위치에서 읽어냄.
            // [TILE_DIM+1] 패딩으로 인해 열 방향 접근 시에도 뱅크 충돌이 발생하지 않음
            odata[(yIndex_out + j) * height + xIndex_out] = tile[threadIdx.x][threadIdx.y + j];
        }
    }
}

이 코드의 핵심은 3가지입니다.

  1. 읽기 시 코얼레싱: idata에서의 읽기는 threadIdx.x가 연속되는 X 방향으로 이루어지므로 완벽하게 코얼레싱됩니다.
  2. 쓰기 시 코얼레싱: odata에 쓸 때에도 블록의 좌표를 맞바꿈으로써 threadIdx.x 방향으로 연속되도록 설계되어 코얼레싱됩니다.
  3. 공유 메모리 패딩: tile[TILE_DIM][TILE_DIM + 1]처럼 1요소 분량을 엇갈리게(패딩) 만듦으로써, 쓸 때 열 방향(tile[threadIdx.x][threadIdx.y + j])으로 접근할 때의 뱅크 충돌을 완전히 배제했습니다.

5.3 캐시 계층과 특수 메모리

  • L1/L2 캐시 정책: 최근 GPU 아키텍처에서는 프로그래머가 PTX 명령어(.ca, .cg, .cs 등)를 사용하여 캐시 동작을 힌트로 제어할 수 있습니다. 예를 들어, 한 번만 접근하는 데이터는 L2 캐시를 우회하여(스트리밍 접근) 캐시 오염을 막을 수 있습니다.
  • 텍스처 메모리 / 상수 메모리: 이미지 처리에 특화된 텍스처 메모리는 2D 공간 지역성을 가지는 접근에 대해 전용 캐시를 활용합니다. 상수(Constant) 메모리는 모든 스레드가 동일한 상수를 읽어들이는 브로드캐스트 접근에 대해 극도로 높은 효율성을 자랑합니다.

제6장: 딥러닝 시대에서의 GPU의 미래

단일 GPU의 성능 향상뿐만 아니라, 시스템 전체로서의 스케일링이 현재 계산 과학의 프론티어입니다.

6.1 NVLink와 NVSwitch에 의한 초고속 상호 연결

거대한 LLM(대규모 언어 모델)은 단일 GPU의 메모리(예를 들어 80GB나 144GB)에는 다 담을 수 없습니다. 모델 병렬화(텐서 병렬화나 파이프라인 병렬화)를 수행하기 위해서는 GPU 간에 테라바이트급의 데이터를 초당 주고받아야 합니다. 기존의 PCIe(PCI Express) 버스로는 이 대역폭을 감당할 수 없기 때문에 NVIDIA는 NVLink라는 독자적인 고속 인터커넥트를 개발했습니다. 나아가 NVSwitch라는 스위치 칩을 매개함으로써 8기나 256기와 같은 GPU들이 완벽한 논블로킹(non-blocking) 크로스바 스위치로 결합되어, 마치 하나의 거대한 GPU인 것처럼 동작하는 클러스터를 구축하는 것이 가능해졌습니다.

6.2 Transformer Engine과 FP8 생태계

자연어 처리뿐만 아니라 이미지와 음성 인식에서도 사실상의 표준이 된 Transformer 아키텍처에 최적화하기 위해, Hopper 아키텍처에서는 Transformer Engine이라 불리는 전용 하드웨어와 소프트웨어 협조 메커니즘이 탑재되었습니다. 이는 텐서의 통계 정보를 동적으로 감시하고, FP8과 FP16의 계산 정밀도를 레이어별로 자동 전환(Dynamic Scaling)함으로써 정확도 저하를 방지하면서도 극한의 계산 속도와 메모리 대역폭 절약을 실현하는 구조입니다.

6.3 GPU 클러스터의 스케일링 법칙과 미래 전망

OpenAI의 ‘스케일링 법칙(Scaling Laws)‘이 보여주듯, 모델의 파라미터 수와 계산량을 늘리면 늘릴수록 AI 성능은 계속 향상되고 있습니다. 이에 따라 GPU는 단순한 프로세서에서 수만 기를 광섬유로 연결한 ‘데이터 센터 그 자체가 한 대의 거대한 GPU(슈퍼컴퓨터)‘로 진화하고 있습니다.

향후 아키텍처의 진화는 실리콘 포토닉스(광 인터커넥트) 도입, CPO(Co-Packaged Optics), 그리고 SRAM에서 HBM으로 이어지는 3D 적층 기술의 추가적인 고도화를 향할 것입니다. 그러나 ‘병렬 처리에 의한 처리량 극대화’라는, GPU 탄생 초기부터 변하지 않는 DNA는 앞으로도 계산 과학의 최전선을 계속 개척해 나갈 것입니다.

【추가 논고】 GPU의 스케줄링 및 점유율(Occupancy)에 대한 수리적 분석

title: “그래픽스 연산 프로세서의 초병렬 아키텍처와 CUDA의 물리: SIMT, 워프, 텐서 코어의 계산 원리” description: “높은 처리량을 극한까지 추구하는 그래픽스 연산 프로세서의 내부 설계. SM, 워프 스케줄링, 텐서 코어, 공유 메모리 최적화의 진수.” slug: “gpu-architecture-cuda-parallel-computing” date: “2026-10-03T05:00:00+09:00” categories: [“architecture”, “technology”] tags: [“gpu”, “cuda”, “parallel-computing”, “hardware”] image: “eyecatch.jpg”

그래픽스 연산 프로세서의 초병렬 아키텍처와 CUDA의 물리: SIMT, 워프, 텐서 코어의 계산 원리

현대의 고도화된 계산 과학, 인공지능, 딥러닝, 그리고 고해상도 컴퓨터 그래픽스를 뒷받침하는 근간 기술, 그것이 바로 그래픽스 연산 프로세서(Graphics Processing Unit)입니다. 본 문서에서는 그래픽스 연산 프로세서의 아키텍처와 그 위에서 동작하는 병렬 계산 기반인 CUDA(Compute Unified Device Architecture)의 물리적, 하드웨어적 측면을 깊이 파헤칩니다. 단순한 프로그래밍 문법이 아니라, 하드웨어가 ‘왜 그렇게 설계되었는지’, ‘어떻게 극한의 계산 처리량을 낼 수 있는지’를 스트리밍 멀티프로세서(SM), SIMT 실행 모델, 워프 스케줄링, 텐서 코어, 그리고 메모리 계층의 관점에서 철저하게 해부합니다.

추보 제1장의 보충: 범용 연산 프로세서와 그래픽스 연산 프로세서의 설계 사상 분기점

1.1 짧은 지연 시간 추구 vs 높은 처리량 추구

범용 프로세서인 범용 연산 프로세서(Central Processing Unit)와 병렬 계산에 특화된 그래픽스 연산 프로세서는 그 탄생 배경에서부터 설계 사상이 근본적으로 다릅니다. 범용 연산 프로세서는 ‘어떻게 하면 하나의 작업(스레드)을 빠르게 끝낼 것인가’라는 ‘낮은 레이턴시(지연 시간 최소화)‘를 지상 과제로 삼아 진화해 왔습니다. 반면 그래픽스 연산 프로세서는 ‘대량의 작업을 묶어 전체적으로 단위 시간당 얼마나 많은 처리를 완료할 수 있는가’라는 ‘높은 스루풋(처리량 극대화)‘을 추구합니다.

범용 연산 프로세서는 운영체제 제어, 복잡한 분기 조건이 수반되는 애플리케이션 실행, 사용자의 무작위적인 인터럽트 처리 등 예측 불가능한 작업을 신속하게 처리해야 합니다. 이를 위해 고도의 분기 예측 회로, 비순차적 실행(Out-of-Order Execution, 명령어의 순서를 바꿔서 실행하는 메커니즘), 거대한 L1/L2/L3 캐시 메모리를 탑재하여 메모리 접근 지연을 숨기면서 단일 스레드의 성능을 극한까지 높이고 있습니다.

이에 반해 그래픽스 연산 프로세서는 원래 화면상의 수백만 픽셀에 대해 동일한 셰이딩 연산을 적용하는 것과 같이 고도로 병렬화 가능한 작업을 처리하기 위해 탄생했습니다. 복잡한 제어 회로나 거대한 캐시에 다이(Die) 면적을 할애하는 대신, 단순한 연산기(ALU: Arithmetic Logic Unit)를 한계까지 꽉 채우는 선택을 했습니다.

1.2 다이 면적에서 캐시, 제어 회로, ALU의 할당 비율

실리콘 다이(반도체 칩)의 제한된 면적(트랜지스터 예산)을 어떻게 배분할 것인가가 두 아키텍처의 차이를 결정짓습니다.

  • 범용 연산 프로세서의 다이 면적 배분: 다이의 절반 이상이 대용량 캐시 메모리(SRAM)와 고도의 제어 회로(분기 예측, 명령어 페치, 디코드, 스케줄링 등)로 채워져 있습니다. 실제 연산을 수행하는 ALU가 차지하는 비율은 상대적으로 작습니다.
  • 그래픽스 연산 프로세서의 다이 면적 배분: 캐시 메모리나 제어 회로는 필요 최소한으로 억제되어 있으며, 다이의 대부분이 수천에서 수만 개에 달하는 ALU(CUDA 코어)로 채워져 있습니다.

그래픽스 연산 프로세서는 메모리 접근 지연(레이턴시)을 캐시로 숨기는 것이 아니라, ‘컨텍스트 스위칭’을 통해 숨깁니다. 특정 스레드 그룹이 메모리에서 데이터가 도착하기를 기다리는 동안, 즉각적으로 다른 스레드 그룹의 연산을 실행함으로써 연산기를 항상 가동 상태(높은 점유율: Occupancy)로 유지합니다. 이것이 그래픽스 연산 프로세서에서의 ‘높은 처리량 추구’의 물리적 구현입니다. 하드웨어 수준의 멀티스레딩(Hardware Multithreading)이 매우 가볍게 이루어지기 때문에 수천~수만 개의 동시 실행 스레드가 존재한다는 것을 전제로 합니다.

추보 제2장의 보충: SIMT 실행 모델의 본질

2.1 SIMD와 SIMT의 차이

병렬 처리의 분류로 플린의 분류학(Flynn’s taxonomy)이 있지만, 그래픽스 연산 프로세서의 실행 모델은 종종 SIMD(Single Instruction, Multiple Data)와 비교됩니다. 범용 연산 프로세서의 벡터 확장 명령어(AVX 등)는 순수한 SIMD이며, 하나의 명령어로 여러 데이터(예: 256비트 폭의 레지스터에 저장된 8개의 32비트 부동소수점 수)를 동시에 처리합니다. SIMD에서는 데이터 요소마다 다른 분기(if-else)를 수행하는 것이 매우 어렵습니다.

반면, NVIDIA가 제창한 CUDA의 실행 모델은 **SIMT(Single Instruction, Multiple Threads)**라고 불립니다. SIMT에서는 여러 독립적인 ‘스레드’가 그룹(후술할 ‘워프’)을 형성하고, 같은 명령어를 공유하여 실행합니다. 그러나 SIMD와 달리 SIMT의 각 스레드는 **독립적인 레지스터 상태와 명령어 주소 카운터(프로그래밍 모델 상)**를 가지고 있습니다. 이를 통해 프로그래머는 마치 각 스레드가 독립적으로 동작하는 것처럼 코드를 작성할 수 있습니다.

2.2 32스레드 단위의 ‘워프(Warp)’

그래픽스 연산 프로세서의 하드웨어는 스레드를 개별적으로 스케줄링하는 것이 아니라, **32개의 스레드를 하나로 묶은 ‘워프(Warp)’**라는 단위로 관리하고 실행합니다. (AMD의 그래픽스 연산 프로세서에서는 Wavefront라고 불리며, 64스레드 단위 등이 채택되기도 합니다).

스트리밍 멀티프로세서(SM) 내의 명령어 페치·디코드 유닛은 워프 단위로 하나의 명령어를 페치하고, 워프 내의 32개 스레드 모두에게 같은 명령어를 발행(디스패치)합니다. 즉, 워프 내의 32스레드는 물리적으로 완전히 동시에, 같은 명령어를 각자가 가진 다른 데이터에 대해 실행합니다. 이것이 SIMT의 핵심입니다.

2.3 워프 다이버전스(Warp Divergence, 분기 불일치)의 물리적 페널티

각 스레드가 독립적인 프로그램 카운터를 가진 것처럼 행동할 수 있다고는 하지만, 물리적으로는 워프 내의 모든 스레드가 동일한 명령어를 실행해야 합니다. 그렇다면 코드 내에 if-else와 같은 조건 분기가 있고, 워프 내의 스레드들 사이에서 분기 조건의 참/거짓이 나뉘면 어떻게 될까요?

이 현상을 **워프 다이버전스(Warp Divergence: 분기 불일치)**라고 부릅니다.

워프 다이버전스가 발생하면 하드웨어는 다음 단계로 처리를 수행합니다.

  1. 먼저 if 조건이 참이 된 스레드(활성 스레드)에 대해서만 명령어를 실행합니다. 이때 조건이 거짓이 된 스레드는 ‘마스킹(비활성화)‘되어 연산 결과가 기록되지 않습니다.
  2. 다음으로 else 조건(또는 조건이 거짓일 경우의 경로)으로 전환하여, 이번에는 조금 전 마스킹되었던 스레드를 활성화하고 참이었던 스레드를 마스킹하여 명령어를 실행합니다.

즉, 분기 경로가 여러 개인 경우 하드웨어는 해당 경로들을 병렬이 아닌 직렬(순차적)로 실행할 수밖에 없게 됩니다. 극단적인 예로, 워프 내의 32스레드가 32가지의 각기 다른 분기 경로를 거치게 되면 실행 시간은 32배로 뛰어오릅니다. 워프 다이버전스는 그래픽스 연산 프로세서의 계산 처리량을 격감시키는 가장 큰 요인 중 하나이며, 알고리즘 설계에서 가장 피해야 할 안티 패턴입니다. 물리적으로는 ALU가 전력을 소비하고 있음에도 불구하고, 마스킹되어 유효한 계산 결과를 생성하지 못하는 ‘헛된 사이클’이 발생하고 있음을 의미합니다.

추보 제3장의 보충: 스트리밍 멀티프로세서(SM)의 하드웨어 해부

그래픽스 연산 프로세서는 다수의 **스트리밍 멀티프로세서(SM: Streaming Multiprocessor)**의 집합체로 구성되어 있습니다. SM이야말로 그래픽스 연산 프로세서의 진정한 계산 엔진입니다. 최신 아키텍처(예: Hopper H100)에서는 하나의 그래픽스 연산 프로세서 다이에 100개 이상의 SM이 탑재되어 있습니다.

3.1 SM 내부의 파이프라인 구성

SM은 내부에 다시 여러 개의 서브 파티션(보통 4개)으로 분할되어 있으며, 각각이 독립적인 워프 스케줄러와 디스패치 유닛을 가지고 있습니다.

  • 워프 스케줄러(Warp Scheduler): 실행 가능한 상태(레지스터나 메모리가 준비된 상태)에 있는 워프를 선택합니다. 그래픽스 연산 프로세서의 스케줄러는 제로 오버헤드로 워프를 전환할 수 있으며, 이것이 메모리 접근 지연을 숨기는 열쇠가 됩니다.
  • 디스패치 유닛(Dispatch Unit): 스케줄링된 워프에 대해 명령어를 발행합니다.
  • CUDA 코어(INT32 / FP32 / FP64 ALU): 실제 정수 연산이나 부동소수점 연산을 수행하는 유닛입니다.
  • 로드/스토어 유닛(LD/ST Unit): 메모리에 대한 읽기/쓰기를 담당합니다.
  • 특수 함수 유닛(SFU: Special Function Unit): sin, cos, exp, 역수 등의 초월 함수를 빠르게 계산하는 전용 하드웨어입니다.

명령어 파이프라인은 매우 깊게 설계되어 있으며, 페치, 디코드, 스케줄링, 레지스터 읽기, 실행(다수 사이클), 라이트 백의 각 단계를 거칩니다. FP32의 FMA(Fused Multiply-Add) 연산 지연 시간은 보통 수 사이클에서 십여 사이클 정도 걸리지만, 매 사이클마다 다른 워프에서 명령어를 발행함으로써 파이프라인을 항상 가득 찬 상태로 유지합니다.

3.2 거대한 레지스터 파일과 레지스터 압박

SM에는 범용 연산 프로세서와는 비교할 수 없을 정도로 거대한 레지스터 파일이 탑재되어 있습니다(예: 1SM당 64KB~256KB의 SRAM). 이는 SM에서 동시 실행되는 수천 개 스레드의 컨텍스트를 모두 유지하기 위해서입니다.

컨텍스트 스위칭이 제로 사이클에 완료되는 이유는 스레드의 레지스터 상태를 메모리로 대피(스필, spill)시킬 필요가 없기 때문입니다. 하지만 스레드당 사용하는 레지스터 수가 증가하면, SM 내에서 동시에 기동할 수 있는 워프의 수(점유율)가 떨어집니다. 이를 **레지스터 압박(Register Pressure)**이라고 부릅니다. 레지스터가 고갈되면 데이터는 느린 로컬 메모리(물리적으로는 글로벌 메모리의 일부)로 스필되며, 치명적인 성능 저하를 일으킵니다.

3.3 공유 메모리(Shared Memory)와 뱅크 충돌

SM에는 프로그래머가 명시적으로 제어할 수 있는 초고속 온칩 메모리인 **공유 메모리(Shared Memory)**가 존재합니다. L1 캐시와 동일한 물리 SRAM 영역을 공유하지만 명시적인 데이터 캐시로 기능하며, 블록 내 스레드 간 데이터 공유 및 동기화에 사용됩니다.

공유 메모리의 물리적 구조는 **메모리 뱅크(Memory Banks)**라고 불리는 여러 개의 독립된 모듈(보통 32개)로 나뉘어 있습니다. 연속되는 32비트 주소는 다른 뱅크에 인터리브(할당)됩니다.

워프 내의 32스레드가 서로 다른 뱅크에 동시에 접근할 경우, 접근은 완전히 병렬로(1사이클 만에) 처리됩니다. 이를 뱅크 충돌 없음(Bank Conflict-Free)이라고 합니다. 그러나 여러 스레드가 같은 뱅크의 서로 다른 주소에 동시에 접근하려고 하면 요청이 직렬화되어 페널티(지연)가 발생합니다. 이를 **뱅크 충돌(Bank Conflict)**이라고 부릅니다. 예를 들어, 2-way 뱅크 충돌이라면 접근 시간은 2배가 되고, 최악의 경우 32-way 충돌에서는 32배로 지연됩니다. 행렬 전치와 같은 알고리즘에서는 스트라이드(보폭) 접근으로 인해 심각한 뱅크 충돌이 발생하므로, 패딩(의미 없는 데이터를 삽입하여 메모리 주소를 비켜가게 하는 기법)을 사용하여 충돌을 회피하는 고도화된 최적화가 필수적입니다.

추보 제4장의 보충: 텐서 코어(Tensor Core)의 적화연산 파이프라인

Volta 아키텍처에서 처음 도입되어 이후 그래픽스 연산 프로세서의 성능을 비약적으로 끌어올린 혁명적인 하드웨어가 바로 **텐서 코어(Tensor Core)**입니다. AI와 딥러닝의 폭발적인 발전은 텐서 코어 없이는 설명할 수 없습니다.

4.1 행렬 적화연산(MMA)의 하드웨어 구현

딥러닝 연산의 대부분은 신경망의 가중치 행렬과 입력 데이터의 행렬곱(GEMM: General Matrix Multiply)입니다. 수식으로는 $D = A \times B + C$ ($A, B$는 입력 행렬, $C$는 누산기(accumulator) 행렬)로 표현됩니다.

기존의 CUDA 코어에서는 이 행렬곱을 1개 요소씩 FMA(Fused Multiply-Add) 명령어를 사용하여 계산했습니다. 반면 텐서 코어는 작은 행렬(예: 4x4나 16x16)의 곱셈 및 덧셈 연산을 하드웨어 레벨에서 1사이클(또는 수 사이클) 만에 실행하는 전용 회로입니다.

물리적으로는 수십 개에서 수백 개의 곱셈기와 거대한 덧셈 트리(Adder Tree)를 와이어로 직결하여, 중간 결과를 레지스터에 다시 쓰지 않고 단번에 적화연산을 완료합니다. 이로 인해 일반 CUDA 코어에 비해 면적당 연산 처리량(TFLOPS)이 비교할 수 없을 만큼 높습니다.

4.2 혼합 정밀도(Mixed-Precision)의 진수

텐서 코어의 또 다른 핵심은 혼합 정밀도(Mixed-Precision) 연산 지원입니다. 딥러닝에서는 계산 과정에서 높은 정밀도(FP32/FP64)를 필요로 하지 않는 상황이 자주 있습니다. 텐서 코어는 입력 행렬 $A$와 $B$를 낮은 정밀도(FP16, BF16, 또는 더 낮은 FP8, INT8, INT4)로 읽어들여 내부 곱셈을 낮은 정밀도로 수행한 후, 덧셈(누산) 과정을 더 높은 정밀도(FP32나 INT32)로 수행하는 파이프라인을 갖추고 있습니다.

  • FP16 / BF16: 학습의 표준. BF16(Bfloat16)은 지수부가 FP32와 같은 8비트라서 다이내믹 레인지가 넓어 기울기 소실을 방지하기 쉽습니다.
  • FP8 / INT8 / INT4: 추론(Inference) 고속화의 비장의 카드. 데이터 전송량(메모리 대역폭)도 감소하므로 처리량이 극적으로 향상됩니다.

Hopper 아키텍처에서는 Transformer 모델의 계산을 획기적으로 가속하는 ‘FP8 Tensor Core’가 도입되어, FP32 대비 이론상 수십 배의 스루풋을 달성했습니다. 소프트웨어 측(CUDA)에서는 wmma(Warp-Level Matrix Multiply and Accumulate) API나 mma.sync PTX 명령어를 통해 텐서 코어를 직접 구동하고, 워프 내의 스레드가 협력하여 행렬의 단편을 레지스터로 로드·연산·스토어하는 매우 복잡한 컬렉티브(집단) 처리를 수행합니다.

추보 제5장의 보충: CUDA 메모 계층과 최적화 기법

그래픽스 연산 프로세서의 계산 능력이 아무리 높더라도, 데이터 공급이 병목이 되면 성능이 나오지 않습니다(메모리 월(Memory Wall) 문제). CUDA 프로그래밍 최적화의 9할은 ‘메모리 접근 최적화’라고 해도 과언이 아닙니다.

5.1 글로벌 메모리의 코얼레싱(Coalescing) 접근

그래픽스 연산 프로세서의 메인 메모리(HBM이나 GDDR)인 글로벌 메모리는 매우 넓은 대역폭(예를 들어 수 TB/s)을 가지지만, 지연 시간도 수백 사이클로 매우 큽니다.

글로벌 메모리에 대한 접근 효율을 극대화하는 절대 원칙이 바로 **코얼레싱(Coalescing: 병합)**입니다. 그래픽스 연산 프로세서의 메모리 컨트롤러는 메모리에 대해 32바이트, 64바이트 또는 128바이트 단위의 트랜잭션으로 접근을 수행합니다. 워프 내의 32스레드가 메모리에 접근할 때, 그 메모리 주소들이 연속된 영역(정렬된 128바이트 경계 내)에 들어맞을 경우, 하드웨어는 이 요청들을 **한 번의 메모리 트랜잭션으로 병합(Coalesce)**하여 처리합니다.

반대로 스레드가 무작위 주소에 접근하거나 스트라이드(간격이 벌어진) 접근을 하게 되면 병합이 이루어지지 않아 여러 트랜잭션이 발생합니다. 이를 ‘비병합(Non-Coalesced) 접근’이라고 부르며, 유효 메모리 대역폭을 10분의 1 이하로 떨어뜨리는 치명적인 성능 버그가 됩니다.

5.2 CUDA C++ 코드 예제: 행렬 전치 최적화와 공유 메모리

아래는 비병합 접근을 회피하고 공유 메모리를 활용하여 성능을 획기적으로 개선하는 행렬 전치(Matrix Transpose)의 최적화된 커널 코드 예시입니다.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35

// 공유 메모리를 사용한 최적화된 행렬 전치 커널
// TILE_DIM = 32, BLOCK_ROWS = 8 로 설정
__global__ void transposeSharedOptimized(float *odata, const float *idata, int width, int height) {
    // 공유 메모리 선언. 뱅크 충돌을 피하기 위해 '+ 1' 패딩을 넣음
    __shared__ float tile[TILE_DIM][TILE_DIM + 1];

    // 입력 행렬상의 글로벌 인덱스(읽기용)
    int xIndex = blockIdx.x * TILE_DIM + threadIdx.x;
    int yIndex = blockIdx.y * TILE_DIM + threadIdx.y;

    // 출력 행렬상의 글로벌 인덱스(쓰기용)
    // 블록의 X와 정렬을 맞바꾸어 쓰기 시의 코얼레싱을 확보
    int xIndex_out = blockIdx.y * TILE_DIM + threadIdx.x;
    int yIndex_out = blockIdx.x * TILE_DIM + threadIdx.y;

    // 1. 글로벌 메모리에서 공유 메모리로 읽기 (코얼레싱 접근)
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        if (xIndex < width && (yIndex + j) < height) {
            // 스레드는 연속된 주소를 읽어들임
            tile[threadIdx.y + j][threadIdx.x] = idata[(yIndex + j) * width + xIndex];
        }
    }

    // 블록 내 모든 스레드의 읽기 완료를 동기화
    __syncthreads();

    // 2. 공유 메모리에서 글로벌 메모리로 쓰기 (코얼레싱 접근)
    for (int j = 0; j < TILE_DIM; j += BLOCK_ROWS) {
        if (xIndex_out < height && (yIndex_out + j) < width) {
            // 공유 메모리에서는 전치된 위치에서 읽어냄.
            // [TILE_DIM+1] 패딩으로 인해 열 방향 접근 시에도 뱅크 충돌이 발생하지 않음
            odata[(yIndex_out + j) * height + xIndex_out] = tile[threadIdx.x][threadIdx.y + j];
        }
    }
}

이 코드의 핵심은 3가지입니다.

  1. 읽기 시 코얼레싱: idata에서의 읽기는 threadIdx.x가 연속되는 X 방향으로 이루어지므로 완벽하게 코얼레싱됩니다.
  2. 쓰기 시 코얼레싱: odata에 쓸 때에도 블록의 좌표를 맞바꿈으로써 threadIdx.x 방향으로 연속되도록 설계되어 코얼레싱됩니다.
  3. 공유 메모리 패딩: tile[TILE_DIM][TILE_DIM + 1]처럼 1요소 분량을 엇갈리게(패딩) 만듦으로써, 쓸 때 열 방향(tile[threadIdx.x][threadIdx.y + j])으로 접근할 때의 뱅크 충돌을 완전히 배제했습니다.

5.3 캐시 계층과 특수 메모리

  • L1/L2 캐시 정책: 최근 그래픽스 연산 프로세서 아키텍처에서는 프로그래머가 PTX 명령어(.ca, .cg, .cs 등)를 사용하여 캐시 동작을 힌트로 제어할 수 있습니다. 예를 들어, 한 번만 접근하는 데이터는 L2 캐시를 우회하여(스트리밍 접근) 캐시 오염을 막을 수 있습니다.
  • 텍스처 메모리 / 상수 메모리: 이미지 처리에 특화된 텍스처 메모리는 2D 공간 지역성을 가지는 접근에 대해 전용 캐시를 활용합니다. 상수(Constant) 메모리는 모든 스레드가 동일한 상수를 읽어들이는 브로드캐스트 접근에 대해 극도로 높은 효율성을 자랑합니다.

추보 제6장의 보충: 딥러닝 시대에서의 그래픽스 연산 프로세서의 미래

단일 그래픽스 연산 프로세서의 성능 향상뿐만 아니라, 시스템 전체로서의 스케일링이 현재 계산 과학의 프론티어입니다.

6.1 NVLink와 NVSwitch에 의한 초고속 상호 연결

거대한 LLM(대규모 언어 모델)은 단일 그래픽스 연산 프로세서의 메모리(예를 들어 80GB나 144GB)에는 다 담을 수 없습니다. 모델 병렬화(텐서 병렬화나 파이프라인 병렬화)를 수행하기 위해서는 그래픽스 연산 프로세서 간에 테라바이트급의 데이터를 초당 주고받아야 합니다. 기존의 PCIe(PCI Express) 버스로는 이 대역폭을 감당할 수 없기 때문에 NVIDIA는 NVLink라는 독자적인 고속 인터커넥트를 개발했습니다. 나아가 NVSwitch라는 스위치 칩을 매개함으로써 8기나 256기와 같은 그래픽스 연산 프로세서들이 완벽한 논블로킹(non-blocking) 크로스바 스위치로 결합되어, 마치 하나의 거대한 그래픽스 연산 프로세서인 것처럼 동작하는 클러스터를 구축하는 것이 가능해졌습니다.

6.2 Transformer Engine과 FP8 생태계

자연어 처리뿐만 아니라 이미지와 음성 인식에서도 사실상의 표준이 된 Transformer 아키텍처에 최적화하기 위해, Hopper 아키텍처에서는 Transformer Engine이라 불리는 전용 하드웨어와 소프트웨어 협조 메커니즘이 탑재되었습니다. 이는 텐서의 통계 정보를 동적으로 감시하고, FP8과 FP16의 계산 정밀도를 레이어별로 자동 전환(Dynamic Scaling)함으로써 정확도 저하를 방지하면서도 극한의 계산 속도와 메모 대역폭 절약을 실현하는 구조입니다.

6.3 그래픽스 연산 프로세서 클러스터의 스케일링 법칙과 미래 전망

OpenAI의 ‘스케일링 법칙(Scaling Laws)‘이 보여주듯, 모델의 파라미터 수와 계산량을 늘리면 늘릴수록 AI 성능은 계속 향상되고 있습니다. 이에 따라 그래픽스 연산 프로세서는 단순한 프로세서에서 수만 기를 광섬유로 연결한 ‘데이터 센터 그 자체가 한 대의 거대한 그래픽스 연산 프로세서(슈퍼컴퓨터)‘로 진화하고 있습니다.

향후 아키텍처의 진화는 실리콘 포토닉스(광 인터커넥트) 도입, CPO(Co-Packaged Optics), 그리고 SRAM에서 HBM으로 이어지는 3D 적층 기술의 추가적인 고도화를 향할 것입니다. 그러나 ‘병렬 처리에 의한 처리량 극대화’라는, 그래픽스 연산 프로세서 탄생 초기부터 변하지 않는 DNA는 앞으로도 계산 과학의 최전선을 계속 개척해 나갈 것입니다.

결어: 계산 과학의 극한으로

GPU 아키텍처는 인류가 지금까지 만들어낸 가장 복잡하고, 또한 가장 처리량에 특화된 계산 엔진입니다. CPU가 ‘한 대의 초고성능 F1 머신’이라면, GPU는 ‘수만 대의 덤프트럭이 통제된 움직임으로 동시에 물자를 나르는 거대한 물류 시스템’에 비유할 수 있습니다.

SIMT에 의한 워프 단위의 명령어 실행, 수천 개의 스레드를 제로 사이클로 전환하는 하드웨어 스케줄링, 한계까지 대역폭을 끌어내는 코얼레싱 접근, 그리고 딥러닝의 돌파구를 견인한 텐서 코어의 파이프라인. 이 모든 것은 ‘물리 법칙의 한계(광속, 열, 전력, 실리콘 미세화의 한계) 속에서, 어떻게 부동소수점 연산의 총량을 극대화할 것인가’라는 엔지니어들의 광기라고도 할 수 있는 집념의 결정체입니다.

앞으로의 소프트웨어 엔지니어, AI 연구자, HPC 연구자에게 있어 GPU 아키텍처를 이해하는 것은 단순한 교양이 아닙니다. 프레임워크(PyTorch나 TensorFlow) 이면에서 무슨 일이 일어나고 있는지를 직관적으로 파악하고, 하드웨어의 능력을 극한까지 끌어내기 위한 ‘필수 과목’인 것입니다. 메모리의 뱅크 충돌을 피하고, 워프 다이버전스를 배제하며, 텐서 코어의 파이프라인을 데이터로 계속 채우는 것. 그 최적화의 끝에, 과거에는 슈퍼컴퓨터로 수개월이 걸렸던 계산이 책상 위 몇 장의 GPU로 몇 시간 만에 완료되는 미래가 지금 바로 현실이 되고 있습니다.

우리는 지금 인류 역사상 가장 흥미진진한 컴퓨터 아키텍처의 황금시대를 살고 있습니다. CUDA의 물리와 GPU 초병렬 아키텍처의 진수를 이해하고, 차세대 혁신을 만들어내는 것은 이 글을 읽고 있는 여러분 자신일지도 모릅니다.

전문 용어 해설(Glossary)

  • SM (Streaming Multiprocessor): GPU의 주요 연산 블록. CPU의 코어에 해당하지만, 그 내부에 다수의 CUDA 코어, 워프 스케줄러, 공유 메모리 등을 내포한다.
  • SIMT (Single Instruction, Multiple Threads): 워프 내의 모든 스레드가 동일한 명령어를 공유하면서 독립적인 데이터에 대해 연산을 수행하는 GPU 특유의 실행 모델.
  • Warp (워프): 32개 스레드의 집합체. 하드웨어에 의한 스케줄링과 명령어 발행의 최소 단위.
  • Warp Divergence (워프 다이버전스): 워프 내의 스레드 간에 분기 조건이 나뉘어, 실행 경로가 직렬화되면서 처리량이 떨어지는 현상.
  • Tensor Core (텐서 코어): 행렬 적화연산(MMA)을 하드웨어 레벨에서 단번에 처리하는 전용 회로. 딥러닝 고속화에 특화.
  • Coalesced Access (코얼레싱 접근): 워프 내의 스레드가 연속된 메모리 주소에 접근했을 때, 하드웨어가 이를 1개의 트랜잭션으로 병합하여 넓은 대역폭을 실현하는 구조.
  • Shared Memory (공유 메모리): SM 내부에 탑재된 프로그래머 제어 가능한 초고속 L1 스크래치패드 메모리.
  • Bank Conflict (뱅크 충돌): 공유 메모리에서 여러 스레드가 동일 뱅크의 다른 주소에 동시에 접근하여, 접근이 직렬화되는 페널티.
  • Occupancy (오큐펀시 / 점유율): SM 상에서 동시에 활성화할 수 있는 워프 수의 이론적 최대치에 대한 실제 비율. 높을수록 메모리 접근 지연을 숨기기 쉽다.
  • Register Spilling (레지스터 스필링): 스레드가 사용하는 레지스터 수가 하드웨어 상한을 초과하여, 넘친 데이터가 느린 메모리(로컬 메모리)로 대피되는 현상.

참고문헌 및 권장 읽기 목록

  1. NVIDIA CUDA C++ Programming Guide: 모든 CUDA 프로그래머가 반드시 읽어야 할 공식 문서. 메모리 접근 패턴 및 최적화 모범 사례가 망라되어 있다.
  2. NVIDIA Ampere / Hopper Architecture Whitepaper: 텐서 코어의 파이프라인이나 비동기 메모리 전송, Transformer Engine의 하드웨어 구현 세부 사항이 기재된 공식 백서.
  3. Computer Architecture: A Quantitative Approach (John L. Hennessy, David A. Patterson): 컴퓨터 아키텍처의 고전적 명저. CPU와 GPU의 설계 사상 차이, 캐시 계층, 명령어 수준 병렬성에 대해 깊이 배울 수 있다.
  4. Programming Massively Parallel Processors: A Hands-on Approach (David B. Kirk, Wen-mei W. Hwu): CUDA 프로그래밍을 알고리즘 설계 관점에서 해설한 교과서. 공유 메모리의 타일링 기법이나 리덕션, 프리픽스 섬 등의 구현이 상세히 설명되어 있다.
  5. Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking: 학술 논문. NVIDIA가 공개하지 않은 캐시 레이턴시나 텐서 코어의 정확한 처리량을 마이크로 벤치마크로 밝혀낸 걸작.

본 문서에서 해설한 아키텍처 지식은 하드웨어의 진화와 함께 진부화되는 부분도 있을 수 있지만, ‘대역폭을 극대화하고 병렬성을 끌어내며 레이턴시를 숨긴다’는 근본적인 물리 원칙은 컴퓨터 과학의 보편적 진리로 계속 남을 것입니다.

comments powered by Disqus