- Authors
- Name
- Youngju Kim
- @fjvbn20031
- 들어가며
- 실행 모델: 논리적 계층 구조
- 하드웨어 구조: SM(Streaming Multiprocessor)
- 메모리 계층
- 워프 스케줄링과 점유율(Occupancy)
- 텐서코어(Tensor Core)
- 스트림과 비동기 실행
- 커널 예제: 타일드 행렬 곱셈
- 흔한 함정과 체크리스트
- 마치며
- 참고 자료
들어가며
GPU 프로그래밍을 처음 접하면 가장 헷갈리는 부분이 "내가 작성한 한 줄의 커널 코드가 실제 하드웨어에서 어떻게 수천 개로 펼쳐지는가"입니다. CPU는 코어 몇 개로 순차적인 작업을 빠르게 처리하도록 설계된 반면, GPU는 수천 개의 단순한 연산 유닛으로 같은 연산을 동시에 처리하도록 설계되어 있습니다. 이 차이를 이해하지 못하면 코드는 돌아가지만 성능은 CPU보다 못한 결과를 얻기 쉽습니다.
이 글에서는 CUDA의 논리적 실행 모델(그리드·블록·워프·스레드)부터 실제 하드웨어(SM, 메모리 계층, 텐서코어)까지를 다이어그램 중심으로 정리합니다. 추상적인 개념을 그림으로 옮겨 두면, 나중에 성능 튜닝을 할 때 "왜 이 코드가 느린가"를 머릿속에서 시뮬레이션할 수 있게 됩니다.
이 글에서 다루는 내용은 다음과 같습니다.
- 실행 모델: 그리드, 블록, 워프, 스레드의 계층 구조
- 하드웨어: SM(Streaming Multiprocessor) 내부 구조
- 메모리 계층: 레지스터, 공유 메모리, L1/L2 캐시, 글로벌 HBM
- 코얼레싱(coalesced) 메모리 접근 패턴
- 워프 스케줄링과 점유율(occupancy)
- 텐서코어(Tensor Core)와 혼합 정밀도 연산
- 스트림(stream)과 비동기 실행
- 간단한 커널 예제와 흔한 함정
아키텍처별(예: Hopper, Blackwell) 구체적인 수치는 세대와 칩에 따라 달라질 수 있으므로, 정확한 값은 공식 문서와 각 칩의 사양서를 확인하시기 바랍니다. 이 글에서는 개념적으로 일관된 부분을 중심으로 설명합니다.
실행 모델: 논리적 계층 구조
CUDA의 실행 모델은 철저히 계층적입니다. 커널(kernel)을 실행하면 수많은 스레드가 생성되는데, 이 스레드들은 무질서하게 흩어지는 것이 아니라 다음과 같은 계층으로 조직됩니다.
┌──────────────────────────────┐
│ GRID │
│ (커널 1회 실행 = 1 그리드) │
│ │
│ ┌────────┐ ┌────────┐ │
│ │ BLOCK │ │ BLOCK │ │
│ │ (0,0) │ │ (1,0) │ │
│ └────────┘ └────────┘ │
│ ┌────────┐ ┌────────┐ │
│ │ BLOCK │ │ BLOCK │ │
│ │ (0,1) │ │ (1,1) │ │
│ └────────┘ └────────┘ │
└──────────────────────────────┘
│
▼ (블록 하나를 확대)
┌──────────────────────────────┐
│ BLOCK │
│ 최대 1024 스레드 (예시) │
│ │
│ Thread Thread Thread ... │
│ (0) (1) (2) │
│ │
│ ── 32개씩 묶여 WARP 형성 ── │
│ [warp 0] [warp 1] ... │
└──────────────────────────────┘
│
▼ (워프 하나를 확대)
┌──────────────────────────────┐
│ WARP │
│ 32개 스레드가 같은 명령을 │
│ 동시에(SIMT) 실행 │
│ │
│ T0 T1 T2 ... T30 T31 │
│ └──── lockstep 실행 ────┘ │
└──────────────────────────────┘
그리드(Grid)
커널을 한 번 실행할 때 생성되는 전체 스레드 집합이 하나의 그리드입니다. 그리드는 1차원, 2차원, 3차원으로 구성할 수 있어 행렬이나 볼륨 데이터에 자연스럽게 매핑됩니다.
블록(Block)
그리드는 여러 개의 블록으로 나뉩니다. 한 블록 안의 스레드들은 다음 두 가지 특권을 가집니다.
- 공유 메모리(shared memory) 를 통해 빠르게 데이터를 주고받을 수 있습니다.
- 동기화(
__syncthreads()) 를 통해 서로의 진행 상황을 맞출 수 있습니다.
서로 다른 블록끼리는 기본적으로 직접 동기화할 수 없습니다. 이 제약이 바로 GPU의 확장성을 보장합니다. 블록은 서로 독립적이므로, 하드웨어는 가용한 SM에 블록을 자유롭게 배치할 수 있습니다.
워프(Warp)
하드웨어 입장에서 실제 스케줄링 단위는 스레드가 아니라 워프입니다. 워프는 32개의 스레드 묶음이며, 같은 명령어를 동시에 실행합니다. 이 방식을 SIMT(Single Instruction, Multiple Threads)라고 부릅니다.
SIMT 실행 방식 (1 명령어 → 32 스레드)
PC ──▶ [ ADD r1, r2, r3 ]
│
┌────────┼─────────────────────────────────┐
▼ ▼ ▼ ... ▼ ▼
T0 T1 T2 T30 T31
(각 스레드는 자신의 레지스터로 같은 명령을 수행)
여기서 핵심 함정이 등장합니다. 같은 워프 안의 스레드들이 서로 다른 분기를 타면(예: if-else), 하드웨어는 두 경로를 순차적으로 실행합니다. 이를 워프 다이버전스(warp divergence) 라고 하며 성능 저하의 주요 원인입니다.
워프 다이버전스 (분기로 인한 직렬화)
if (threadIdx.x < 16) { A() } else { B() }
사이클 1~k : T0..T15 → A() 실행, T16..T31 → 대기(마스크 off)
사이클 k+1~ : T0..T15 → 대기, T16..T31 → B() 실행
결과: A 경로 시간 + B 경로 시간 (병렬이 아니라 직렬)
인덱스 계산
각 스레드는 자신의 전역 위치를 다음과 같이 계산합니다. 이 공식은 거의 모든 CUDA 커널의 첫 줄에 등장합니다.
int idx = blockIdx.x * blockDim.x + threadIdx.x;
blockDim.x = 4 인 경우의 전역 인덱스
blockIdx.x: 0 1 2
┌────────┐ ┌────────┐ ┌────────┐
threadIdx.x: │0 1 2 3 │ │0 1 2 3 │ │0 1 2 3 │
global idx : │0 1 2 3 │ │4 5 6 7 │ │8 9 10 11│
└────────┘ └────────┘ └────────┘
하드웨어 구조: SM(Streaming Multiprocessor)
논리적 모델을 실제로 실행하는 하드웨어 단위가 SM(Streaming Multiprocessor) 입니다. GPU 한 장은 수십 개의 SM으로 구성되며, 블록은 SM에 배정되어 실행됩니다. 하나의 SM은 보통 여러 개의 블록을 동시에 수용할 수 있습니다.
┌─────────────────────────────────────────────────────────────┐
│ SM (Streaming Multiprocessor) │
│ │
│ ┌───────────────┐ ┌───────────────┐ │
│ │ Warp Scheduler│ │ Warp Scheduler│ (보통 4개 파티션) │
│ │ + Dispatch │ │ + Dispatch │ │
│ └───────┬───────┘ └───────┬───────┘ │
│ ▼ ▼ │
│ ┌──────────────┐ ┌──────────────┐ │
│ │ CUDA Cores │ │ CUDA Cores │ (FP32 / INT 연산) │
│ │ ████████████│ │ ████████████│ │
│ └──────────────┘ └──────────────┘ │
│ ┌──────────────┐ ┌──────────────┐ │
│ │ Tensor Cores │ │ Tensor Cores │ (행렬 곱셈 가속) │
│ └──────────────┘ └──────────────┘ │
│ ┌──────────────┐ ┌──────────────┐ │
│ │ LD/ST │ │ SFU │ (메모리 / 초월함수) │
│ └──────────────┘ └──────────────┘ │
│ │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ Register File (수만 개의 32비트 레지스터) │ │
│ └───────────────────────────────────────────────────────┘ │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ Shared Memory / L1 Cache (블록 단위로 분할) │ │
│ └───────────────────────────────────────────────────────┘ │
└─────────────────────────────────────────────────────────────┘
SM 내부의 주요 구성 요소는 다음과 같습니다.
- 워프 스케줄러(Warp Scheduler): 실행 준비가 된 워프를 골라 실행 유닛으로 보냅니다. 보통 한 SM은 여러 개의 스케줄러를 가진 파티션으로 나뉩니다.
- CUDA 코어: FP32/INT 같은 기본 산술 연산을 수행하는 ALU입니다.
- 텐서코어(Tensor Core): 작은 행렬 곱셈-누적(MMA)을 한 번에 처리하는 전용 유닛입니다.
- LD/ST 유닛: 메모리 로드/스토어를 담당합니다.
- SFU(Special Function Unit):
sin,sqrt,exp같은 초월 함수를 빠르게 계산합니다. - 레지스터 파일: 스레드별 지역 변수가 저장되는, 가장 빠른 저장소입니다.
- 공유 메모리/L1: SM 내에서 블록이 사용하는 빠른 온칩 메모리입니다.
여기서 중요한 통찰은, 레지스터 파일과 공유 메모리가 유한한 자원이라는 점입니다. 한 SM이 동시에 수용할 수 있는 블록과 워프의 수는 이 자원을 얼마나 쓰느냐로 결정됩니다. 이것이 바로 다음에 설명할 점유율의 본질입니다.
메모리 계층
GPU 성능의 8할은 메모리에서 결정된다고 해도 과언이 아닙니다. 연산 유닛은 매우 빠르지만, 데이터가 제때 공급되지 않으면 놀게 됩니다. CUDA의 메모리는 속도와 용량이 반비례하는 계층 구조를 이룹니다.
메모리 계층 (위로 갈수록 빠르고 작음)
▲ 빠름 / 작음
│ ┌──────────────────────────┐
│ │ Registers (스레드 전용) │ 수 사이클, 수십~수백 KB/SM
│ ├──────────────────────────┤
│ │ Shared Memory / L1 (블록) │ ~수십 사이클, 수십~수백 KB/SM
│ ├──────────────────────────┤
│ │ L2 Cache (전체 SM 공유) │ ~수백 사이클, 수~수십 MB
│ ├──────────────────────────┤
│ │ Global Memory (HBM, DRAM)│ ~수백 사이클, 수십 GB
│ └──────────────────────────┘
│ 느림 / 큼
▼
각 메모리 공간의 특성을 표로 정리하면 다음과 같습니다. (구체 수치는 아키텍처마다 다르므로 상대적 경향으로 이해하시기 바랍니다.)
| 메모리 공간 | 스코프 | 상대 지연 | 상대 대역폭 | 용량 | 비고 |
|---|---|---|---|---|---|
| 레지스터 | 스레드 | 가장 낮음 | 가장 높음 | 매우 작음 | 컴파일러가 할당 |
| 공유 메모리 | 블록 | 낮음 | 높음 | 작음 | 명시적 관리, 뱅크 구조 |
| L1 캐시 | SM | 낮음 | 높음 | 작음 | 공유 메모리와 자원 공유 |
| L2 캐시 | 전체 GPU | 중간 | 중간 | 중간 | 자동 캐싱 |
| 글로벌(HBM) | 전체 GPU | 가장 높음 | 절대값은 크나 상대적 느림 | 큼 | 호스트와 데이터 교환 |
| 상수 메모리 | 전체 GPU | 캐시 적중 시 낮음 | 브로드캐스트 유리 | 작음 | 읽기 전용 |
| 로컬 메모리 | 스레드 | 높음 | 낮음 | 글로벌에 위치 | 레지스터 스필 시 사용 |
핵심 전략은 명확합니다. 글로벌 메모리 접근을 최소화하고, 한 번 읽은 데이터를 공유 메모리나 레지스터에 담아 재사용하는 것입니다. 타일링(tiling) 기법이 바로 이 원리를 구현한 것입니다.
코얼레싱(Coalesced) 메모리 접근
글로벌 메모리는 한 번에 일정 크기의 메모리 세그먼트(예: 32, 64, 128바이트)를 트랜잭션으로 가져옵니다. 한 워프의 32개 스레드가 연속된 주소에 접근하면, 하드웨어는 이를 적은 수의 트랜잭션으로 합쳐 처리합니다. 이것이 코얼레싱입니다.
코얼레스드 접근 (좋음): 워프가 연속 주소에 접근
Thread: T0 T1 T2 T3 ... T31
Addr : 0 4 8 12 ... 124
└──────────────────────┘
→ 1번의 메모리 트랜잭션으로 처리
언코얼레스드 접근 (나쁨): 스레드가 흩어진 주소에 접근
Thread: T0 T1 T2 T3 ...
Addr : 0 512 1024 1536 ...
└─┐ └─┐ └─┐ └─┐
▼ ▼ ▼ ▼
→ 스레드마다 별도 트랜잭션 → 대역폭 낭비
코얼레싱을 깨는 대표적 원인은 다음과 같습니다.
- 잘못된 데이터 레이아웃(예: 행 우선/열 우선 불일치)
- 큰 스트라이드(stride)를 둔 접근
- 구조체 배열(AoS) 대신 배열 구조체(SoA)를 쓰지 않은 경우
AoS vs SoA
AoS (Array of Structs): [x0 y0 z0][x1 y1 z1][x2 y2 z2]...
→ x만 읽으려 해도 y,z가 끼어들어 비연속 접근
SoA (Struct of Arrays): [x0 x1 x2 ...][y0 y1 y2 ...][z0 z1 z2 ...]
→ x 배열이 연속 → 코얼레싱 유리
워프 스케줄링과 점유율(Occupancy)
GPU가 빠른 진짜 비결은 코어 개수가 아니라 지연 은닉(latency hiding) 에 있습니다. 한 워프가 글로벌 메모리를 기다리며 멈추면, 스케줄러는 즉시 실행 준비가 된 다른 워프로 전환합니다. 충분히 많은 워프가 대기 중이라면, 메모리 지연이 거의 보이지 않게 가려집니다.
워프 스케줄링 타임라인 (지연 은닉)
시간 ───────────────────────────────────────▶
Warp0: [실행]──[메모리 대기 ........]──[실행]
Warp1: [실행]──[메모리 대기 ......]──[실행]
Warp2: [실행]──[메모리 대기 ....]──
Warp3: [실행]──[메모리 대기]
→ 어느 한 워프가 멈춰도, 다른 워프가 실행 유닛을 채움
→ 실행 유닛이 놀지 않게 됨 (높은 사용률)
점유율(Occupancy)
점유율은 한 SM에서 동시에 활성화된 워프 수를 그 SM이 지원하는 최대 워프 수로 나눈 비율입니다. 점유율이 높을수록 지연을 가릴 워프가 많아집니다. 다만 점유율이 100%여야 최고 성능이 나오는 것은 아닙니다. 적당한 점유율로도 충분히 지연을 가릴 수 있고, 오히려 스레드당 레지스터를 넉넉히 써서 더 빨라지는 경우도 있습니다.
점유율을 제한하는 요인을 표로 정리하면 다음과 같습니다.
| 제한 요인 | 설명 | 줄이는 방법 |
|---|---|---|
| 스레드당 레지스터 수 | 레지스터가 많으면 동시 워프 수가 줄어듦 | 변수 단순화, 컴파일러 한계 지정 |
| 블록당 공유 메모리 | 공유 메모리가 많으면 동시 블록 수가 줄어듦 | 타일 크기 조정 |
| 블록당 스레드 수 | 너무 작으면 워프 부족, 너무 크면 자원 부족 | 128, 256 등 적절히 선택 |
| SM당 최대 블록/워프 수 | 하드웨어 상한 | 아키텍처 사양 확인 |
점유율 직관
레지스터/공유메모리를 적게 쓰는 커널:
SM 자원 ████████████████ → 많은 워프 동시 수용 → 높은 점유율
레지스터/공유메모리를 많이 쓰는 커널:
SM 자원 ████ → 적은 워프만 수용 → 낮은 점유율
실무에서는 NVIDIA의 Occupancy Calculator나 cudaOccupancyMaxPotentialBlockSize API로 적절한 블록 크기를 찾고, Nsight Compute로 실제 병목을 확인하는 것이 정석입니다.
텐서코어(Tensor Core)
딥러닝의 핵심 연산은 행렬 곱셈(GEMM)입니다. 일반 CUDA 코어가 스칼라 단위로 곱셈-덧셈을 처리하는 것과 달리, 텐서코어는 작은 행렬 블록의 곱셈-누적(MMA, Matrix Multiply-Accumulate)을 한 명령으로 처리합니다. 덕분에 처리량이 수 배에서 수십 배까지 올라갑니다.
텐서코어의 MMA 연산 (개념도)
A (M×K) B (K×N) C (M×N)
┌──────────┐ ┌──────────┐ ┌──────────┐
│ 타일 │ × │ 타일 │ + │ 누적기 │ ──▶ D = A·B + C
└──────────┘ └──────────┘ └──────────┘
(저정밀) (저정밀) (고정밀 누적)
한 워프가 협력하여 하나의 타일 MMA를 수행
입력은 FP16/BF16/INT8 등, 누적은 FP32로 정밀도 유지
텐서코어의 특징은 혼합 정밀도(mixed precision) 입니다. 입력은 FP16, BF16, TF32, INT8, 그리고 최신 세대에서는 FP8 같은 낮은 정밀도로 받되, 누적은 FP32 같은 높은 정밀도로 수행해 정확도 손실을 줄입니다.
| 입력 정밀도 | 특징 | 주요 용도 |
|---|---|---|
| FP16 | 반정밀도, 넓은 지원 | 학습/추론 |
| BF16 | FP32와 같은 지수 범위 | 학습 안정성 |
| TF32 | FP32 입력을 내부적으로 처리 | 코드 변경 최소화 |
| INT8 | 정수, 양자화 추론 | 저지연 추론 |
| FP8 | 최신 세대 지원 | 대규모 학습/추론 |
대부분의 경우 텐서코어를 직접 다루기보다는 cuBLAS, cuDNN, CUTLASS 같은 고수준 라이브러리를 통해 활용합니다. 직접 작성하려면 mma 계열의 워프 수준 API를 사용합니다.
Hopper, Blackwell 등 최신 세대는 텐서코어 처리량과 지원 정밀도(예: FP8), 비동기 데이터 이동 기능을 세대마다 확장해 왔습니다. 구체적인 지원 정밀도와 성능 수치는 세대와 칩에 따라 다르므로 공식 사양을 확인하시기 바랍니다.
스트림과 비동기 실행
기본적으로 CUDA 작업은 하나의 기본 스트림(default stream)에서 순차적으로 실행됩니다. 하지만 스트림(stream) 을 여러 개 사용하면, 서로 독립적인 작업을 겹쳐서(overlap) 실행할 수 있습니다. 대표적으로 데이터 전송과 커널 연산을 동시에 진행해 전체 시간을 줄입니다.
단일 스트림 (직렬): 전송과 연산이 차례로
H2D[==] Kernel[======] D2H[==]
─────────────────────────────────▶ 시간
다중 스트림 (오버랩): 청크 단위로 겹쳐서
S1: H2D[==] Kernel[====] D2H[==]
S2: H2D[==] Kernel[====] D2H[==]
S3: H2D[==] Kernel[====] D2H[==]
──────────────────────────────────────▶ 시간
→ 전송과 연산이 겹쳐 총 시간 단축
오버랩을 제대로 활용하려면 다음 조건이 필요합니다.
- 호스트 메모리가 페이지 고정(pinned) 되어 있어야 비동기 복사가 가능합니다.
cudaMemcpyAsync와 커널 실행에 명시적으로 스트림을 지정해야 합니다.- 이벤트(
cudaEvent)로 스트림 간 의존성을 관리할 수 있습니다.
cudaStream_t stream;
cudaStreamCreate(&stream);
cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream);
myKernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
cudaMemcpyAsync(h_out, d_out, bytes, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);
cudaStreamDestroy(stream);
최근 세대에서는 CUDA Graphs로 반복되는 작업 그래프를 미리 캡처해 실행 오버헤드를 줄이거나, 비동기 메모리 복사 기능을 통해 전송과 연산을 더 촘촘히 겹치는 기법이 많이 쓰입니다.
커널 예제: 타일드 행렬 곱셈
지금까지의 개념을 한데 모은 예제로, 공유 메모리를 활용한 타일드 행렬 곱셈을 살펴보겠습니다. 핵심 아이디어는 글로벌 메모리에서 작은 타일을 한 번 읽어 공유 메모리에 올린 뒤, 그 타일을 여러 번 재사용해 글로벌 접근을 줄이는 것입니다.
먼저 가장 단순한 벡터 덧셈부터 보겠습니다.
__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}
// 호스트 측 실행 예시
// int threads = 256;
// int blocks = (n + threads - 1) / threads;
// vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);
이제 공유 메모리 타일링을 적용한 행렬 곱셈입니다. C = A * B를 계산하며, 각 블록은 16x16 타일을 담당합니다.
#define TILE 16
__global__ void matMulTiled(const float* A, const float* B, float* C, int N) {
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];
int row = blockIdx.y * TILE + threadIdx.y;
int col = blockIdx.x * TILE + threadIdx.x;
float acc = 0.0f;
// K 차원을 타일 단위로 순회
for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
// 한 타일을 공유 메모리에 로드 (코얼레스드 접근)
int aCol = t * TILE + threadIdx.x;
int bRow = t * TILE + threadIdx.y;
As[threadIdx.y][threadIdx.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
Bs[threadIdx.y][threadIdx.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
__syncthreads(); // 타일 로드 완료 보장
// 공유 메모리에서 부분합 계산 (글로벌 접근 없음)
for (int k = 0; k < TILE; ++k) {
acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
__syncthreads(); // 다음 타일 로드 전 동기화
}
if (row < N && col < N) {
C[row * N + col] = acc;
}
}
이 커널이 빠른 이유를 그림으로 정리하면 다음과 같습니다.
타일링으로 글로벌 접근 재사용
순진한 버전: C[i][j] 계산마다 A행 N개 + B열 N개를 매번 글로벌에서 읽음
→ 같은 데이터를 수없이 중복 로드
타일드 버전: TILE×TILE 블록을 공유 메모리에 1회 로드 →
블록 내 모든 스레드가 그 타일을 재사용
→ 글로벌 트래픽이 약 TILE배 감소
┌─────────┐ load ┌──────────────┐ reuse ┌──────────┐
│ Global │ ───────▶ │ Shared (타일) │ ──────▶ │ 256 스레드 │
│ (HBM) │ 1회 │ 16×16 │ 여러번 │ 연산 │
└─────────┘ └──────────────┘ └──────────┘
__syncthreads()가 두 번 등장하는 이유에 주목해 주세요. 첫 번째는 모든 스레드가 타일 로드를 끝낼 때까지 기다려 데이터 경쟁을 막고, 두 번째는 다음 타일을 덮어쓰기 전에 현재 계산이 끝났음을 보장합니다. 이 동기화를 빠뜨리면 결과가 비결정적으로 망가집니다.
흔한 함정과 체크리스트
마지막으로 실무에서 자주 마주치는 함정을 정리합니다.
1. 언코얼레스드 메모리 접근
가장 흔하고 치명적입니다. 데이터 레이아웃을 SoA로 바꾸거나, 인덱싱을 연속 접근이 되도록 조정하세요. 프로파일러의 메모리 효율 지표를 먼저 확인하는 습관이 중요합니다.
2. 워프 다이버전스
워프 내부에서 데이터에 따라 분기가 갈리면 직렬화됩니다. 가능하면 분기 조건을 워프 경계와 정렬하거나, 분기 대신 산술 연산(predication)으로 바꾸는 것을 고려하세요.
3. 공유 메모리 뱅크 충돌
공유 메모리는 여러 뱅크로 나뉘어 있는데, 한 워프의 스레드들이 같은 뱅크에 동시에 접근하면 충돌이 일어나 직렬화됩니다. 패딩을 한 칸 추가해(예: [TILE][TILE+1]) 충돌을 피하는 기법이 자주 쓰입니다.
4. 과도한 레지스터 사용
레지스터를 너무 많이 쓰면 점유율이 떨어지고, 한계를 넘으면 느린 로컬 메모리로 스필됩니다. 커널을 잘게 나누거나 변수를 줄여 보세요.
5. 동기화 누락 또는 오용
__syncthreads()를 분기 안에서 일부 스레드만 호출하면 데드락이 발생합니다. 모든 스레드가 같은 동기화 지점을 통과하도록 작성하세요.
6. 에러 체크 생략
CUDA 호출의 반환값과 cudaGetLastError()를 확인하지 않으면 조용히 잘못된 결과를 얻습니다. 디버그 빌드에서는 매 호출마다 에러를 검사하는 것이 안전합니다.
| 증상 | 의심 원인 | 확인 도구 |
|---|---|---|
| 느린데 코어는 노는 듯 | 메모리 바운드, 언코얼레싱 | Nsight Compute |
| 점유율이 낮음 | 레지스터/공유 메모리 과다 | Occupancy Calculator |
| 분기 많은 커널이 느림 | 워프 다이버전스 | Nsight Compute |
| 결과가 가끔 틀림 | 동기화 누락, 경쟁 상태 | compute-sanitizer |
| 전송이 연산과 안 겹침 | pinned 메모리 미사용 | Nsight Systems |
마치며
CUDA를 잘 쓰는 핵심은 결국 두 가지로 요약됩니다. 첫째, 논리적 실행 모델(그리드·블록·워프)이 실제 하드웨어(SM·메모리 계층)에 어떻게 매핑되는지를 머릿속에 그릴 수 있어야 합니다. 둘째, 데이터를 어떻게 메모리 계층에 잘 배치하고 재사용하느냐가 성능의 대부분을 결정합니다.
이 글에서 다룬 다이어그램들을 떠올리며 코드를 작성하면, "왜 이 커널이 느린가"를 추측이 아니라 구조로 설명할 수 있게 됩니다. 그다음은 Nsight 같은 프로파일러로 가설을 검증하고 반복적으로 다듬는 과정입니다.
마지막으로, 아키텍처는 세대마다 빠르게 진화합니다. 텐서코어의 정밀도 지원, 비동기 메모리 이동, SM당 자원 같은 구체적 수치는 칩과 버전에 따라 달라지므로, 실제 튜닝 시에는 항상 해당 아키텍처의 공식 문서를 확인하시기 바랍니다.