CUDA 상세

개요

CUDA(Compute Unified Device Architecture)는 NVIDIA가 2007년에 도입한 병렬 컴퓨팅 플랫폼이자 프로그래밍 모델로, GPU의 범용 연산 자원을 C/C++ 기반 코드와 전용 런타임으로 활용할 수 있게 해준다. 개발자는 같은 프로그램 안에서 호스트 CPU 코드와 디바이스 GPU 커널을 함께 작성하며, AI 훈련/추론, 과학 계산, 데이터 분석, 미디어 처리처럼 데이터 병렬성이 큰 작업을 GPU로 오프로딩할 수 있다.

CUDA의 핵심은 GPU 하드웨어 구조를 소프트웨어 계층으로 노출한다는 점이다. NVIDIA GPU는 다수의 스트리밍 멀티프로세서(SM)로 구성되고, 각 SM은 워프 스케줄러, 레지스터 파일, 공유 메모리, 연산 유닛을 바탕으로 많은 스레드를 동시에 처리한다. CUDA는 이를 스레드, 스레드 블록, 그리드, 워프, 스트림 같은 추상화로 제공해 개발자가 하드웨어 세부 구조를 이해하면서도 이식 가능한 병렬 프로그램을 작성하게 돕는다.

핵심 개념

CUDA 프로그래밍 모델

CUDA 프로그래밍 모델은 계층적인 스레드 조직과 메모리 접근 방법을 정의한다.

스레드 계층 구조:
- 스레드(Thread): 가장 작은 실행 단위. 각 스레드는 독립적인 프로그램 카운터, 레지스터 상태, 스레드 인덱스를 가진다.
- 스레드 블록(Thread Block): 여러 스레드가 묶인 협업 단위. 같은 블록 안의 스레드는 __shared__ 메모리로 데이터를 교환하고 __syncthreads()로 동기화할 수 있다.
- 그리드(Grid): 하나의 커널 실행에서 생성되는 전체 블록 집합. 블록 간 실행 순서는 보장되지 않으므로 블록 간 직접 동기화 없이 동작하도록 설계해야 한다.

워프(Warp):
- 32개 스레드가 이루는 기본 스케줄링 단위
- 하드웨어는 같은 워프의 스레드들을 SIMT(Single Instruction, Multiple Thread) 방식으로 실행한다.
- 워프 내부에서 분기 경로가 갈라지면 같은 워프가 여러 경로를 순차적으로 처리하므로 처리 효율이 낮아질 수 있다.

커널(Kernel):
- GPU에서 실행되는 함수. __global__ 키워드로 선언
- 호스트(Host)에서 호출하면 디바이스(Device)에서 병렬 실행

점유율(Occupancy):
- 한 SM 안에 동시에 상주하는 활성 워프 수의 비율
- 레지스터 사용량, 블록당 공유 메모리 사용량, 블록 크기가 점유율을 함께 결정한다.
- 점유율이 높다고 항상 가장 빠른 것은 아니지만, 메모리 지연 시간을 숨길 수 있는 워프 수를 확보하는 데 중요하다.

스트리밍 멀티프로세서(SM)

NVIDIA GPU의 핵심 연산 유닛인 SM은 워프 실행, 데이터 이동, 온칩 메모리 재사용을 위한 여러 리소스를 포함한다:

SM 스케줄링 동작:
1. 워프 스케줄러가 의존성이 해소된 활성 워프를 선택한다.
2. 선택된 워프의 명령어를 적절한 실행 파이프에 발행한다.
3. 연산 결과는 레지스터에 기록되고, 필요 시 공유 메모리나 글로벌 메모리로 이동한다.
4. 메모리 대기나 분기 때문에 멈춘 워프 대신 다른 워프를 실행해 지연 시간을 숨긴다.

CUDA 메모리 계층

CUDA는 다양한 특성을 가진 메모리 계층을 제공한다:

메모리 접근 최적화 기법:
- Coalescing: 같은 워프의 스레드들이 인접 주소에 접근하면 하나의 메모리 트랜잭션으로 처리
- 공유 메모리 뱅크 컨플릭트: 32개 뱅크로 구성된 공유 메모리에서 인접 스레드가 다른 뱅크에 접근해야 효율적
- 비동기 복사와 TMA: cp.async, cuda::memcpy_async, TMA를 이용해 데이터 이동과 계산을 겹칠 수 있음

스케줄링 메커니즘

CUDA의 스케줄링은 여러 수준에서 동작한다:

블록 스케줄링:
- 그리드 내 블록을 SM에 배치하는 단계
- 스케줄러는 사용 가능한 SM의 수와 블록의 리소스 요구량을 고려하여 배치
- 블록은 독립적으로 실행되므로 순서 보장 불가

워프 스케줄링:
- SM 내에서 활성 워프 중 다음에 실행할 워프를 선택
- 구현 세부 정책은 세대와 마이크로아키텍처에 따라 달라지지만, 연구와 최적화 문헌에서는 주로 LRR(Loose Round Robin), GTO(Greedy-Then-Oldest), Two-Level 같은 정책으로 설명한다.
- 개발자는 정책 이름 자체보다 워프 수, 메모리 대기, 분기, 명령어 수준 병렬성이 실제 성능에 미치는 영향을 더 중요하게 봐야 한다.

지연 시간 은닉(Latency Hiding):
- 메모리 지연 시간을 숨기기 위해 여러 워프를 번갈아 실행
- 한 워프가 메모리 대기 중일 때 다른 워프를 실행하여 SM 활용도 극대화

비교/분석

GPU 아키텍처 세대별 CUDA 변경사항

메모리 접근 패턴 비교

스케줄링 정책 비교

동작 원리

커널 실행 흐름

1. 호스트에서 커널 호출:
   c kernel<<<gridDim, blockDim>>>(args);
   - gridDim: 그리드 크기 (블록 수)
   - blockDim: 블록 크기 (스레드 수)

2. 블록 배치:
   - GPU의 스케줄러가 각 블록을 사용 가능한 SM에 배치
   - 하나의 SM에는 여러 블록이 할당될 수 있음 (리소스 허용 시)

3. 스레드 생성 및 실행:

   • 각 SM은 할당된 블록의 스레드를 워프 단위로 생성
   • 32개 스레드씩 그룹화되어 SIMT 방식으로 실행

4. 동기화 및 완료:
   - 블록 내 스레드 동기화: __syncthreads()
   - 커널 완료 대기: cudaDeviceSynchronize()

메모리 접근 동작

글로벌 메모리 Coalescing:
- 워프의 32개 스레드가 연속된 128B 주소에 접근하면 단일 128B 메모리 트랜잭션으로 처리
- 32개 스레드가 각각 4B(FP32)씩 접근 시 128B = 32 × 4B

공유 메모리 뱅크 컨플릭트:
- 32개 뱅크로 구성, 각 뱅크는 4B 폭
- 같은 워프의 스레드가 같은 뱅크에 접근하면 접근이 직렬화될 수 있다.
- 패드 추가(shared float pad[N+1])로 컨플릭트 해소

점유율과 파이프라인 겹치기:
- 블록 크기와 공유 메모리 사용량이 크면 동시에 상주 가능한 워프 수가 줄어든다.
- 반대로 적절한 점유율과 cp.async 또는 TMA 기반 프리페치를 조합하면 메모리 로드와 계산을 겹쳐 실효 처리량을 높일 수 있다.

텐서 코어 활용

Tensor Core는 행렬 곱셈-누산(MMA) 연산을 하드웨어 수준에서 가속화한다:

WGMMA(Warp Group Matrix Multiply-Accumulate):
- Hopper에서 도입된 워프 그룹 단위 MMA 명령어
- 4개 워프(128 스레드)가 협력하여 큰 타일 행렬 연산을 수행
- TMA와 조합하면 글로벌 메모리에서 공유 메모리로 가져온 타일을 효율적으로 소비할 수 있다.

WMMA 프로그래밍 인터페이스:

// 16x16x16 행렬 곱셈 예시
wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

장단점

장점

• 풍부한 생태계: cuBLAS, cuDNN, NCCL, TensorRT 등 최대 규모의 GPU 가속 라이브러리
• 범용성: AI, HPC, 그래픽스, 데이터 분석 등 다양한 도메인 지원
• 세밀한 제어: 메모리 접근 패턴, 스케줄링, 동기화를 프로그래밍 수준에서 제어 가능
• 도구 체인: Nsight Systems, Nsight Compute, CUPTI 등 풍부한 프로파일링 도구
• 호환성: NVIDIA GPU 전체 라인업에서 동일한 CUDA 코드 실행 가능
• 진화적 발전: 세대별로 TMA, 분산 공유 메모리 등 새로운 기능 지속 추가

단점

• NVIDIA 종속: NVIDIA GPU에서만 실행 가능, AMD/Intel GPU는 지원 불가
• 메모리 관리 복잡성: 다양한 메모리 계층의 적절한 활용 필요
• 스케줄링 오버헤드: 블록/워프 스케줄링으로 인한 간접 비용
• 동기화 제약: 호스트-디바이스 간 동기화로 인한 잠재적 병목
• 전력 소비: 대규모 병렬 실행 시 높은 전력 소모

관련 기술

참고 표준 및 자료

• NVIDIA CUDA Programming Guide (최신 버전)
• NVIDIA CUDA C++ Best Practices Guide
• NVIDIA Hopper Tuning Guide
• NVIDIA GPU Architecture Whitepapers (Volta, Ampere, Hopper)
• ACM/IEEE GPU 스케줄링 및 메모리 최적화 연구 문헌

관련 문서

• CUDA 메모리 관리
• 텐서 코어 아키텍처
• NVIDIA GPU 상세
• GPU 메모리 아키텍처 기초
• 혼합 정밀도 하드웨어
• AI 가속기 개요
• ROCm 상세

핵심 정리

1. CUDA는 NVIDIA GPU의 범용 병렬 연산 능력을 프로그래밍할 수 있는 플랫폼으로, 스레드 블록/그리드/워프의 계층적인 프로그래밍 모델을 제공한다.
2. SM(Streaming Multiprocessor)은 CUDA의 핵심 연산 유닛으로, 워프 스케줄러, 레지스터 파일, 공유 메모리, 연산 파이프를 함께 제공한다.
3. CUDA 메모리 계층은 레지스터, 공유 메모리, 글로벌 메모리 등으로 구성되며, 각각 다른 접근 속도와 범위를 가진다.
4. 효율적인 CUDA 프로그래밍을 위해서는 메모리 Coalescing, 공유 메모리 뱅크 컨플릭트 회피, 워프 스케줄링 활용이 중요하다.
5. CUDA는 지속적으로 발전하여 TMA, 분산 공유 메모리, WGMMA 등 새로운 기능을 추가하며 AI/HPC 워크로드의 성능을 향상시키고 있다.