GPGPU 아키텍처와 병렬 연산 프로그래밍

그래픽 파이프라인을 넘어선 GPU

GPU는 본래 그래픽 렌더링을 위해 설계됐지만, 다수의 단순 연산 코어를 동시에 활용할 수 있다는 특성 때문에 범용 연산 가속기로 쓰이기 시작했다. 이를 GPGPU(General-Purpose computing on Graphics Processing Units)라고 한다. 과학 계산, 인공지능, 영상 처리처럼 데이터 병렬성이 큰 작업에서 특히 의미가 크다.

초기의 GPU는 고정 기능 렌더링 파이프라인을 가진 그래픽 전용 하드웨어였다. 처리 가능한 연산도 특정 그래픽 작업으로 제한됐고 개발자가 파이프라인을 직접 프로그래밍할 수 없었다.

셰이더 프로그래밍이 도입되면서 상황이 바뀌었다. DirectX와 OpenGL을 통해 제한적인 범용 계산을 수행할 수 있었고, 그래픽 API를 우회해 연산 작업을 GPU에 배치하는 방식도 사용됐다. 현재는 CUDA와 OpenCL처럼 범용 연산을 직접 다루는 API, 통합 셰이더 아키텍처, 고수준 프로그래밍 모델을 바탕으로 딥러닝과 HPC까지 적용 범위가 넓어졌다.

GPU는 CPU와 설계 목표가 다르다. CPU가 복잡한 제어 로직과 큰 캐시를 바탕으로 낮은 레이턴시에 집중한다면, GPU는 단순한 제어 로직과 많은 코어를 통해 처리량을 높인다. GPU는 수천 개 코어를 동시에 실행하고, 레이턴시보다 대역폭을 중시한다. GDDR 또는 HBM 메모리를 사용하며 CPU 대비 10배 이상 메모리 대역폭을 제공하지만 메모리 레이턴시는 높다.

⤢✕

데이터 병렬성을 실행하는 SIMT 모델

SIMD는 Single Instruction, Multiple Data의 약자로, 하나의 명령으로 여러 데이터를 함께 처리하는 방식이다. 벡터 처리의 기반이 되며 데이터 병렬성을 활용한다.

GPU에서는 이와 가까우면서도 다른 SIMT(Single Instruction, Multiple Threads) 모델이 사용된다. CPU의 SIMD는 AVX나 SSE처럼 명시적인 벡터 명령어와 고정된 벡터 길이를 다룬다. 반면 GPU 프로그래머는 스칼라 명령어를 작성하고, 하드웨어가 여러 스레드를 묶어 실행한다. 스레드별 분기는 가능하지만, 같은 실행 묶음 안에서 경로가 갈리면 비용이 생긴다.

NVIDIA는 32개 스레드 묶음을 Warp라고 부른다. Warp는 동일 명령을 동시에 실행하는 SIMT의 기본 단위다. AMD의 Wavefront는 같은 역할을 하며 64개 스레드로 구성된다. Wavefront 크기는 아키텍처에 따라 달라질 수 있다.

⤢✕

같은 Warp의 스레드가 서로 다른 조건 경로로 들어가면 분기 발산(Branch Divergence)이 발생한다. 이때 각 경로는 직렬화돼 순차적으로 실행되므로 성능이 떨어진다.

// Warp 내 스레드들의 조건 분기
if (threadIdx.x % 2 == 0) {
    // 짝수 스레드
    doEvenWork();
} else {
    // 홀수 스레드
    doOddWork();
}
// 두 경로가 순차적으로 실행됨

분기 발산을 줄이려면 조건문을 산술 연산으로 바꾸거나, Warp 안의 스레드가 가능한 한 동일한 작업을 수행하도록 구성해야 한다. 데이터 배치를 조정해 서로 다른 분기가 한 Warp에 섞이지 않게 하는 방법도 있다.

셰이더 단계가 공유하는 연산 자원

그래픽 파이프라인은 정점 데이터가 Vertex Shader, Tessellation, Geometry Shader, Rasterization, Fragment Shader, ROP를 거쳐 Frame Buffer에 기록되는 흐름으로 구성된다.

⤢✕

Vertex Shader는 정점을 처리하며 변환, 조명, 정점 위치 계산을 담당한다. 이 단계의 병렬성은 정점 개수에서 나온다. Tessellation Shader는 Hull Shader와 Domain Shader를 이용해 표면을 세분화하고 LOD(Level of Detail)를 조정해 기하학적 디테일을 높인다.

Geometry Shader는 프리미티브 단위로 동작하며 정점을 생성하거나 삭제할 수 있고 전체 도형 정보에도 접근한다. 다만 성능 이슈로 사용이 제한적이다. Fragment Shader는 픽셀 색상, 텍스처 샘플링, 조명 계산을 수행하며 픽셀 개수만큼 병렬성을 확보한다.

통합 셰이더 아키텍처에서는 각 셰이더 단계가 동일한 프로세서에서 실행된다. 스케줄러가 정점 작업, 픽셀 작업, 연산 작업을 Shader Processor에 배정하므로 자원 활용 효율과 동적 부하 분산이 가능하다. 이 구조가 GPGPU를 가능하게 한 기반 기술이다.

⤢✕

CUDA Core와 Stream Processor의 계층

NVIDIA CUDA Core는 정수·부동소수점 연산을 위한 Scalar ALU와 FPU, 레지스터 파일, 단순한 제어 로직으로 구성된다. CUDA Core는 SM(Streaming Multiprocessor)에 포함되며, SM 하나에는 64~128개의 CUDA Core, 공유 메모리, Warp 스케줄러, 특수 함수 유닛(SFU)이 들어간다.

여러 SM은 GPC(Graphics Processing Cluster)를 구성한다. GPC에는 Rasterizer와 Texture Units도 포함된다.

⤢✕

AMD의 Stream Processor는 VLIW(Very Long Instruction Word) 또는 GCN(Graphics Core Next) 구조를 사용하며, Compute Unit(CU)에서 Wavefront를 실행한다. Wavefront는 64개 스레드로 구성된다.

GCN의 CU에는 4개의 SIMD 유닛이 있고 각 SIMD는 16-wide다. CU는 64KB LDS(Local Data Share)와 Scalar Unit도 가진다.

특성	NVIDIA CUDA Core	AMD Stream Processor
워크그룹 크기	32 (Warp)	64 (Wavefront)
실행 모델	SIMT	SIMT
ISA	PTX → SASS	GCN ISA
메모리 모델	CUDA 메모리 계층	OpenCL 호환

성능은 메모리 계층에서 갈린다

Global Memory는 수 GB~수십 GB의 가장 큰 용량을 제공하며 모든 스레드가 접근할 수 있다. 다만 레이턴시가 수백 사이클에 이르므로 Coalescing이 중요하다.

Shared Memory는 SM 내부 스레드가 공유하는 메모리다. 레이턴시는 수 사이클로 낮지만 용량은 48KB~128KB로 제한되고, 개발자가 명시적으로 관리해야 한다. Registers는 스레드 전용이며 가장 빠르게 접근할 수 있지만 개수가 제한돼 있어 Register Pressure를 주의해야 한다. Constant/Texture Memory는 읽기 전용이고 캐싱 최적화를 위한 특수 목적 메모리다.

⤢✕

CUDA와 OpenCL이 작업을 배치하는 방식

CUDA에서는 Grid가 전체 실행 단위, Block이 스레드 그룹, Thread가 개별 실행 단위다. 커널은 여러 Block과 Thread로 구성돼 GPU에서 실행된다.

// CUDA 커널 예시
__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 호스트 코드
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c, n);

호스트와 디바이스 사이의 메모리 이동과 할당·해제도 명시적으로 수행한다.

// 디바이스 메모리 할당
cudaMalloc(&d_data, size);

// 호스트 → 디바이스 복사
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);

// 디바이스 → 호스트 복사
cudaMemcpy(h_result, d_result, size, cudaMemcpyDeviceToHost);

// 메모리 해제
cudaFree(d_data);

OpenCL은 벤더 중립적인 표준으로 CPU, GPU, FPGA 등 다양한 디바이스를 지원한다. 플랫폼, 디바이스, 컨텍스트, 커맨드 큐를 기준으로 실행 환경을 구성하며, CUDA와 마찬가지로 명시적 메모리 관리를 사용한다.

커널 성능을 다룰 때는 SM당 활성 Warp 수를 높이는 Occupancy, 레지스터와 공유 메모리 사용량, Block 크기를 함께 봐야 한다. 메모리 접근에서는 Coalesced 접근, Shared Memory 활용, 메모리 뱅크 충돌 회피가 핵심이다. 연산 측면에서는 분기를 줄이고 특수 함수 활용(**sinf, **expf)과 Instruction-Level Parallelism을 고려한다.

병렬성이 큰 연산의 적용 대상

GPGPU는 기상 시뮬레이션, 분자 동역학, 유체 역학(CFD), 양자 화학 같은 HPC 작업에 쓰인다. 인공지능과 머신러닝에서는 딥러닝 학습·추론, 행렬 연산 가속, TensorCore 활용, 대규모 모델 학습에 활용된다.

실시간 영상 처리, 인코딩·디코딩, 컴퓨터 비전, 의료 영상 분석도 GPU 병렬 처리와 맞닿아 있다. 금융 공학에서는 몬테카를로 시뮬레이션, 리스크 분석, 옵션 가격 결정, 고빈도 거래가 적용 영역이다.

GPGPU를 효율적으로 사용하려면 대규모 병렬 처리 자체뿐 아니라 SIMT 실행 단위, 분기 발산, 메모리 계층, Occupancy를 함께 고려해야 한다. 그래픽 렌더링에서 출발한 GPU의 구조는 이제 AI, 과학 계산, 빅데이터 처리의 가속 기반으로 쓰인다.