당신의 시스템에는 CPU가 있습니다. 아마 GPU도 있을 것입니다. 사양표(spec sheet) 상으로는 둘 다 "코어(cores)"를 가지고 있고, 둘 다 특정 클록 속도(clock speed)로 작동하기 때문에 비슷해 보일 수 있지만, 이들은 완전히 다른 문제를 해결하기 위해 설계되었습니다. 그 _이유_를 이해하는 것이 CUDA에 관한 다른 모든 것을 이해하는 열쇠입니다.

왜 이런 하드웨어가 존재할까요?

CPU는 지연 시간 최적화(latency-optimized)되어 있습니다. CPU의 역할은 하나의 복잡한 작업을 가능한 한 빨리 끝내는 것이므로, 트랜지스터 예산의 큰 부분은 순수 산술 연산과는 상관없는 것들에 할당됩니다. 즉, if 문이 평가되기도 전에 어느 방향으로 갈지 예측하는 분기 예측기(branch predictors), 스톨(stalls)을 피하기 위해 명령어를 재정렬하는 비순차적 실행 엔진(out-of-order execution engines), 그리고 자주 사용되는 데이터를 코어 가까이에 유지하는 대규모 멀티 레벨 캐시(L1, L2, 때로는 L3) 등이 그것입니다. CPU는 예측 불가능하고, 분기되며, 순차적인 로직을 처리하도록 구축되었습니다.

A GPU는 CPU와 완전히 반대입니다. GPU는 워크로드(workload)가 예측 가능하고, 규칙적이며, 수학적으로 밀도가 높다고 가정합니다. 화면의 모든 픽셀이나 대규모 행렬의 모든 요소와 같이, 엄청난 양의 데이터에 동일한 명령어를 적용하는 방식입니다. 이러한 가정 때문에 GPU는 CPU가 의존하는 분기 예측 및 비순차적 실행 기계 장치의 대부분을 제거하고, 대신 트랜지스터 예산을 순수 산술 연산 장치에 투자합니다. 즉, 소수의 크고 복잡한 코어 대신 수천 개의 작고 단순한 코어를 갖게 됩니다.

CPU는 하나의 작업을 빠르게 끝내도록 최적화되어 있습니다. GPU는 수백만 개의 유사한 작업을 한 번에 끝내도록 최적화되어 있습니다.

행렬 곱셈 (Matrix multiplications), 컨볼루션 (Convolutions), 어텐션 메커니즘 (Attention mechanisms)과 같은 머신러닝 (Machine learning) 워크로드는 산술 수준에서 보면 거대한 텐서 (Tensors) 전반에 걸쳐 반복되는 몇 가지 동일한 연산들입니다. 이것이 바로 GPU가 설계된 목적이자, GPU가 신경망 (Neural Networks)을 학습시키고 실행하기 위한 기본 하드웨어가 된 이유인 규칙적이고 병렬적인 워크로드의 전형적인 형태입니다.

CUDA 이전: 그래픽을 렌더링하는 척하며 계산하기

GPU가 원래 임의의 프로그램을 실행하도록 설계되지 않았다는 사실은 잊기 쉽습니다. 2007년 이전까지 GPU의 존재 이유는 오로지 그래픽 파이프라인 (Graphics pipeline)뿐이었습니다. 삼각형으로 구성된 장면을 가져와서, 각 정점 (Vertex)에 **버텍스 셰이더 (Vertex shader)**를 실행하여 화면의 어디에 위치할지 파악한 다음, 결과로 나온 각 픽셀에 **프래그먼트 (픽셀) 셰이더 (Fragment (pixel) shader)**를 실행하여 최종 색상을 결정하고, 그 색상을 프레임버퍼 (Framebuffer)에 쓰는 방식이었습니다.

연구자들은 프래그먼트 셰이더가 사실 화면 전체에 걸쳐 픽셀당 한 번씩 병렬로 실행되는 작은 프로그램일 뿐이라는 중요한 사실을 알아차렸습니다. 만약 GPU가 당신의 _데이터 (Data)_를 마치 _사진 (Picture)_인 것처럼 취급하도록 속일 수 있다면, 일반적인 수치 계산에도 동일한 거대한 병렬성을 적용할 수 있었습니다. 이러한 우회 방법은 GPGPU —> _General Purpose computing on Graphics Processing Units_로 알려지게 되었으며, 이는 실제 수학 문제를 그래픽 파이프라인이 수용할 수 있는 형태로 변형하는 과정을 필요로 했습니다:

• 입력 데이터(예: 곱셈을 원하는 두 행렬)는 반드시 **텍스처 (textures)**로 인코딩되어야 했습니다. 왜냐하면 GPU가 읽는 방법을 알고 있는 유일한 대규모의 구조화된 입력이 텍스처였기 때문입니다.
• 실제로 연산을 트리거하려면, 뷰포트 전체를 덮는 두 개의 삼각형인 **풀스크린 쿼드 (full-screen quad)**를 렌더링해야 했습니다. 이는 순수하게 GPU가 출력 요소당 한 번씩 프래그먼트 셰이더 (fragment shader)를 호출하도록 강제하기 위한 구실에 불과했습니다. 렌더링되는 "이미지"는 시각적인 의미가 없었으며, 하드웨어가 수학 연산을 수행하도록 만들기 위한 변장이었습니다.
• 조명과 색상을 위해 설계된 셰이딩 언어로 작성된 프래그먼트 셰이더 (fragment shader)는 산술 연산을 위해 용도가 변경되어야 했습니다. 즉, 일반적인 메모리 읽기 대신 텍스처 룩업 (texture lookups)을 통해 입력을 읽고, 마치 픽셀 색상인 것처럼 수치 결과를 출력해야 했습니다.

픽셀 색상 연산으로 위장하여 a[i] + b[i]를 "계산"하는 극도로 단순화된 GLSL 프래그먼트 셰이더 (fragment shader)는 다음과 같은 모습이었습니다:

// 구식 GPGPU: 프래그먼트 셰이더로 위장한 덧셈
uniform sampler2D textureA; // 텍스처로 인코딩된 행렬 A
uniform sampler2D textureB; // 텍스처로 인코딩된 행렬 B
...

결과는 배열로 반환되지 않고 프레임버퍼 (framebuffer)에 담긴 이미지로 반환되었으며, 이를 다시 읽어 들여 숫자로 재해석해야 했습니다. 제어 흐름 (control flow)은 극도로 제한적이었고, 초기 셰이더 모델에서는 루프 (loops)와 조건문 (conditionals)이 취약하거나 지원되지 않았으며, 임의의 메모리 위치에 쓸 수 있는 방법 없이 오직 셰이더 호출이 발생한 특정 픽셀에만 쓸 수 있었습니다.

NVIDIA는 이러한 장벽을 제거하기 위해 2007년에 **Compute Unified Device Architecture (CUDA)**를 출시했습니다. CUDA는 실제 메모리 읽기 및 쓰기, 실제 제어 흐름을 갖춘 C 스타일의 프로그래밍 모델을 통해 GPU의 병렬 실행 유닛을 노출함으로써, 이 모든 과정이 그림인 척할 필요가 없도록 만들었습니다. 이것이 바로 GPU가 범용 도구 (general purpose tool)가 된 지점입니다.

GPU가 실행을 조직하는 방법

CUDA는 **SIMT (Single Instruction, Multiple Threads)**라고 불리는 모델을 사용합니다. 하나의 명령(instruction)이 발행되면, 수많은 스레드(threads)가 각자 자신만의 데이터를 가지고 이를 동시에 실행합니다. 이는 물리적 하드웨어에 직접 매핑되는 계층 구조로 조직됩니다:

• **스레드 (Thread)**는 실행의 가장 작은 단위로, 고유한 프로그램 카운터(program counter)와 개별 레지스터(private registers)를 가집니다.
• **워프 (Warp)**는 정확히 32개의 스레드로 구성된 그룹이며, 하드웨어가 실제로 스케줄링하는 단위입니다. 워프 내의 32개 스레드 모두는 동일한 클록 사이클(clock cycle)에 동일한 명령을 부여받습니다.
• **블록 (Block)**은 하나의 **스트리밍 멀티프로세서 (Streaming Multiprocessor, SM)**에 스케줄링되는 최대 1024개의 스레드입니다. 블록 내의 스레드들은 __syncthreads()를 통해 동기화할 수 있으며, 빠른 온칩 메모리 (on-chip memory)를 공유할 수 있습니다.
• **그리드 (Grid)**는 커널 (kernel)을 실행하는 데 필요한 모든 블록의 집합으로, 사용 가능한 모든 SM에 분산됩니다.

다음은 두 벡터를 더하는 아주 기본적인 CUDA 커널의 실제 모습입니다. 하나의 스레드가 하나의 요소를 처리합니다:

__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
...

위에서 본 GLSL의 모습과 비교해 보십시오. 텍스처(textures)나 프레임버퍼(framebuffers) 없이, 실제 배열 인덱싱(array indexing)과 실제 덧셈이 이루어집니다.

점유율 (Occupancy) 및 지연 시간 숨기기 (Latency hiding)

SM은 일반적으로 한 사이클에 실행할 수 있는 양보다 훨씬 더 많은 워프를 동시에 상주(resident)시킵니다. 워프가 (대개 느린 전역 메모리 (global memory) 읽기를 기다리며) 중단(stall)되면, SM의 스케줄러는 실행 준비가 된 다른 워프로 전환하고, 중단된 워프는 데이터가 도착하면 다시 재개됩니다. 이를 **지연 시간 숨기기 (latency hiding)**라고 하며, 이것이 GPU가 바쁘게 작동하기 위해 CPU 스타일의 캐시 (caches)나 투기적 실행 (speculative execution)을 필요로 하지 않는 주요 이유입니다. SM의 최대 워프 용량 중 실제로 상주하고 있는 비율을 **점유율 (occupancy)**이라고 부릅니다.

워프 발산 (Warp divergence)

워프 (warp) 내의 32개 스레드는 모두 단일 명령어 인출 (instruction fetch)을 공유하기 때문에, 반드시 코드를 함께 단계별로 실행해야 합니다.

if (data[i] > threshold) {
    // path A
} else {
...

만약 동일한 워프 내의 스레드들이 이 조건을 서로 다르게 평가하면, 워프는 **발산 (diverges)**합니다. 하드웨어는 경로 B를 타는 스레드들을 마스킹 (mask off) 처리하고, 활성화된 스레드들에 대해 경로 A를 실행한 뒤, 마스크를 전환하여 경로 B를 실행하며, 두 경로가 모두 끝나면 다시 재결합 (reconverges)합니다. 병렬로 실행될 수 있었던 작업이 이제 순차적으로 실행됩니다. 단순한 2-way 분기 (branch)만으로도 해당 워프의 유효 처리량 (throughput)을 대략 절반으로 줄일 수 있습니다.

핵심 요약 (Takeaway): 균일하고 분기가 없는 (branch-free) 워크로드는 GPU에서 효율적으로 실행됩니다. 워프의 동작을 분리하는 조건부 및 데이터 의존적 로직이 병렬 성능을 저하시키는 지점입니다.

진짜 병목 현상: 연산이 아닌 메모리

현대적인 GPU는 일반적으로 초당 메모리에서 가져올 수 있는 바이트 수보다 초당 수행할 수 있는 산술 연산 횟수가 훨씬 더 많습니다. 대부분의 실제 워크로드에서 메모리 대역폭 (memory bandwidth)이 제한 요인이 됩니다. 이는 **루프라인 모델 (roofline model)**로 공식화됩니다. 커널의 달성 가능한 성능은 피크 연산 처리량 (peak compute throughput) 또는 피크 메모리 대역폭 (peak memory bandwidth) 중, 커널의 연산 대비 바이트 비율 (math-to-bytes ratio)이 먼저 도달하는 제약 조건에 의해 제한됩니다.

효율적인 CUDA 코드는 글로벌 메모리 (global memory)로의 왕복 (round trips)을 최소화합니다. 행렬 곱셈 (matrix multiplication)에서 매우 빈번하게 사용되는 일반적인 패턴은 **타일링 (tiling)**입니다. 글로벌 메모리로부터 입력 데이터의 한 덩어리(chunk)를 공유 메모리 (shared memory)로 한 번 로드하고, 블록 내의 모든 스레드가 모든 산술 연산에 해당 타일 (tile)을 재사용하게 하며, 마지막에 최종 결과만을 글로벌 메모리에 다시 씁니다.

__global__ void matMulTiled(float *A, float *B, float *C, int n) {
    __shared__ float tileA[16][16];
    __shared__ float tileB[16][16];
...

각 값은 타일(tile)당 한 번씩 전역 메모리(global memory)에서 가져와지며, 매 곱셈마다 다시 가져오는 대신 공유 메모리(shared memory)에서 16번 재사용됩니다. 이는 이와 같이 연산 집약적인 커널(compute-heavy kernel)에서 사용할 수 있는 가장 큰 단일 최적화 방법인 경우가 많습니다.

메모리 병합 (Memory coalescing)

워프(warp)가 전역 메모리에서 데이터를 요청할 때, 하드웨어는 가능한 한 적은 수의 크고 연속적인 트랜잭션(transaction)으로 해당 요청을 처리하려고 시도합니다. 스레드 0이 인덱스 0을 요청하고, 스레드 1이 인덱스 1을 요청하는 방식 등으로 작동하면 효율적입니다. 이를 **병합된 액세스 (coalesced access)**라고 하며, 하나의 넓은 트랜잭션이 워프 전체의 요청을 충족할 수 있습니다.

// 병합된 액세스: 연속된 스레드가 연속된 메모리를 읽음
c[i] = a[i] + b[i];              // i = threadIdx.x + blockIdx.x * blockDim.x

...

만약 데이터가 연결 리스트(linked list), 파편화된 해시맵(fragmented hashmap), 또는 단순히 큰 스트라이드(stride)를 가진 것처럼 흩어져 있다면, 하드웨어는 이를 한 번의 트랜잭션으로 가져올 수 없으며 대신 스레드당(또는 작은 그룹당) 별도의 트랜잭션을 발행하게 됩니다. 이러한 **병합되지 않은 액세스 (uncoalesced access)**는 연산 자체는 변하지 않았더라도, 읽히는 데이터의 레이아웃(layout)만 바뀜으로써 유효 메모리 대역폭(effective memory bandwidth)을 수십 배까지 감소시킬 수 있습니다.

핵심 요약: 메모리 내 데이터의 레이아웃은 알고리즘 자체만큼이나 중요할 수 있습니다.

CUDA 코어에서 텐서 코어(Tensor Cores)까지

위의 모든 과정은 스레드당 클록 사이클당 하나의 스칼라 연산을 수행할 수 있는 단순한 ALU인 표준 **CUDA 코어 (CUDA core)**를 포함합니다. 이는 이미 범용 병렬 작업에 효과적입니다. 하지만 딥러닝 워크로드(workload)는 하나의 특정 연산, 즉 행렬 곱셈(matrix multiplication)이 지배적입니다. 완전 연결 계층(fully connected layer)은 행렬 곱셈이며, 컨볼루션(convolution) 또한 행렬 곱셈으로 재구성될 수 있고, 트랜스포머 어텐션(transformer attention) 내부의 핵심 연산 역시 행렬 곱셈입니다.

이를 특별히 가속하기 위해, NVIDIA는 2017년 Volta 아키텍처와 함께 **텐서 코어 (Tensor Core)**를 도입했습니다.

표준 CUDA 코어가 사이클당 하나의 스칼라 연산 (scalar operation)을 계산하는 반면, 텐서 코어 (Tensor Core)는 단일 클록 사이클 내에 작은 행렬 타일 (matrix tiles, 원래는 4×4)에 대해 완전한 혼합 정밀도 곱셈-누산 (mixed-precision multiply-accumulate) — D = A × B + C — 을 계산합니다. 이후의 아키텍처들은 이를 대폭 확장했습니다: 더 커진 유효 타일 크기, 그리고 더 많은 수치 정밀도(FP16, BF16, TF32, INT8) 및 추가적인 사이클을 절약하기 위해 알려진 0 값을 건너뛰는 구조적 희소성 (structured sparsity) 지원 등이 포함됩니다. 여기서 "혼합 정밀도 (Mixed precision)"란 속도를 위해 곱셈은 더 낮은 정밀도 형식으로 수행되지만, 정확도를 유지하기 위해 누산 (accumulation)은 더 높은 정밀도 형식으로 수행됨을 의미합니다.

이것이 바로 GPU가 AI에 유용하게 쓰이기도 하는 주요 그래픽 칩에서, 그래픽도 여전히 렌더링할 수 있는 AI 가속기 (AI accelerators)로 전환된 핵심적인 하드웨어적 이유입니다.

종합하기