CUDA 성능 최적화

학습 목표

• •Occupancy 개념을 이해하고 최적화한다
• •Memory Coalescing의 중요성을 배운다
• •Branch Divergence를 이해하고 최소화한다
• •Nsight 도구를 활용한 프로파일링 기법을 익힌다

Occupancy (점유율)

Occupancy는 SM의 자원이 얼마나 효율적으로 사용되는지를 나타냅니다.

Occupancy 정의

Occupancy = 활성 Warp 수 / 최대 가능 Warp 수

예: SM의 최대 Warp = 64, 실제 활성 = 32
Occupancy = 32/64 = 50%

Occupancy에 영향을 주는 요소

Occupancy 계산 예시

SM 자원 (예: Ampere 아키텍처):
- 최대 레지스터: 65,536개/SM
- 최대 Shared Memory: 164KB/SM
- 최대 스레드: 2,048개/SM (64 Warps)
- 최대 블록: 32개/SM

커널 사용량:
- 레지스터: 스레드당 64개
- Shared Memory: 블록당 48KB
- 블록 크기: 256 스레드

계산:
- 레지스터 제한: 65536 / 64 = 1024 스레드 가능
- Shared 제한: 164KB / 48KB = 3.4 → 3블록, 768 스레드
- 블록 크기 제한: 최대 32블록, 각 256 = 8192 (최대 초과)

실제 Occupancy = min(1024, 768) / 2048 = 37.5%

Occupancy 향상 기법

cuda
복사
// 방법 1: 레지스터 사용량 제한
__global__ void __launch_bounds__(256, 4) kernel() {
    // 256 스레드/블록, 최소 4블록/SM 보장
    // 컴파일러가 레지스터 사용 최적화
}

// 방법 2: 지역 변수 줄이기
__global__ void bad() {
    float data[100];  // 레지스터 스필 → 느림
}

__global__ void good() {
    float data[8];    // 레지스터에 유지 가능
}

높은 Occupancy가 항상 좋은가?

꼭 그렇지 않음!
- Occupancy 50%도 충분히 빠를 수 있음
- 메모리 대역폭에 의해 성능이 제한되는 경우
- ILP(Instruction Level Parallelism)이 높은 경우

최적화 순서:
1. 먼저 알고리즘 최적화
2. 메모리 접근 패턴 최적화
3. 그 다음 Occupancy 고려

Memory Coalescing (메모리 합치기)

Coalescing은 여러 스레드의 메모리 접근을 하나의 트랜잭션으로 합치는 것입니다.

Coalesced vs Non-coalesced

Coalesced Access 패턴

cuda
복사
// 좋은 패턴: 연속 접근
__global__ void coalesced(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[idx];  // Coalesced!
}

// 나쁜 패턴: Stride 접근
__global__ void strided(float* data, int stride) {
    int idx = threadIdx.x * stride;  // stride > 1이면 비효율
    float val = data[idx];  // Non-coalesced!
}

구조체 배열 vs 배열의 구조체

cuda
복사
// AoS (Array of Structures) - 비효율
struct Particle {
    float x, y, z;
    float vx, vy, vz;
};
Particle particles[N];
// Thread i가 particles[i].x 접근 → Non-coalesced

// SoA (Structure of Arrays) - 효율적
struct ParticleData {
    float x[N], y[N], z[N];
    float vx[N], vy[N], vz[N];
};
// Thread i가 x[i] 접근 → Coalesced!

2D 배열 접근

cuda
복사
// 행 우선 순서 (C/C++ 기본) - 행 방향 접근이 Coalesced
float data[HEIGHT][WIDTH];

// 좋은 패턴: 같은 행의 연속 열 접근
__global__ void good(float* data, int width) {
    int row = blockIdx.y;
    int col = threadIdx.x + blockIdx.x * blockDim.x;
    float val = data[row * width + col];  // Coalesced
}

// 나쁜 패턴: 같은 열의 연속 행 접근
__global__ void bad(float* data, int width) {
    int col = blockIdx.x;
    int row = threadIdx.x + blockIdx.y * blockDim.y;
    float val = data[row * width + col];  // Non-coalesced
}

Branch Divergence (분기 발산)

Branch Divergence는 Warp 내 스레드들이 서로 다른 분기를 실행할 때 발생합니다.

SIMT 실행 모델

Warp (32 threads) 특성:
- 모든 스레드가 같은 명령어를 실행
- 분기가 있으면 양쪽을 순차 실행
- 비활성 스레드는 대기

if (condition) {
    A;  // 일부 스레드만 실행
} else {
    B;  // 나머지 스레드만 실행
}
// 실제로는 A와 B 모두 실행 (시간 = A + B)

Branch Divergence 예시

cuda
복사
// Divergent (나쁨)
__global__ void divergent(float* data) {
    int idx = threadIdx.x;
    if (idx % 2 == 0) {
        data[idx] = expf(data[idx]);  // 짝수 스레드
    } else {
        data[idx] = logf(data[idx]);  // 홀수 스레드
    }
}
// Warp 내에서 절반씩 다른 분기 → 2배 느림

// Non-divergent (좋음)
__global__ void nonDivergent(float* data) {
    int idx = threadIdx.x;
    int warpId = idx / 32;
    if (warpId % 2 == 0) {
        // Warp 0, 2, 4, ... 전체가 같은 분기
        data[idx] = expf(data[idx]);
    } else {
        // Warp 1, 3, 5, ... 전체가 같은 분기
        data[idx] = logf(data[idx]);
    }
}
// 각 Warp가 같은 분기 → Divergence 없음

Divergence 최소화 기법

cuda
복사
// 1. Warp 단위로 조건 맞추기
if (threadIdx.x / 32 == targetWarp) { ... }

// 2. 데이터 정렬로 분기 제거
// 조건에 따라 데이터를 미리 정렬하여
// 같은 조건의 데이터가 같은 Warp에 오도록

// 3. 조건 연산 대신 산술 연산
// 나쁜 예
if (x > 0) y = a; else y = b;

// 좋은 예 (Divergence 없음)
float mask = (float)(x > 0);
y = mask * a + (1 - mask) * b;

Shared Memory Bank Conflict

Bank 구조

Shared Memory는 32개의 Bank로 구성
- 각 Bank는 4바이트 단위
- 연속 4바이트가 연속 Bank에 매핑

Bank 0  Bank 1  Bank 2  ...  Bank 31
[0-3]   [4-7]   [8-11]       [124-127]
[128-131] ...

Bank Conflict

cuda
복사
__shared__ float data[32];

// Conflict 없음: 각 스레드가 다른 Bank
float val = data[threadIdx.x];

// 2-way Conflict: 2개 스레드가 같은 Bank
float val = data[threadIdx.x * 2];  // 0,2,4... → Bank 0,2,4...
                                     // 실제론 0,8,16.. → 같은 Bank!

// 32-way Conflict (최악): 모든 스레드가 같은 Bank
float val = data[0];  // 모두 Bank 0

Bank Conflict 해결

cuda
복사
// 패딩으로 해결
__shared__ float data[32][33];  // 32 + 1 패딩
// 열 접근 시 각 행이 다른 Bank로 오프셋됨

NVIDIA 프로파일링 도구

Nsight Systems

bash
복사
# 시스템 전체 프로파일링
nsys profile ./my_cuda_app

# 결과를 GUI로 확인
nsys-ui report.qdrep

• •CPU/GPU 활동 타임라인 시각화
• •커널 실행 시간, 메모리 전송 분석
• •병목 지점 식별

Nsight Compute

bash
복사
# 커널 상세 분석
ncu --set full ./my_cuda_app

# 특정 커널만 분석
ncu --kernel-name myKernel ./my_cuda_app

• •Occupancy 분석
• •메모리 처리량 측정
• •Warp 효율성 분석
• •하드웨어 카운터 수집

주요 분석 메트릭

최적화 체크리스트

메모리 최적화

□ Global Memory 접근이 Coalesced인가?
□ AoS를 SoA로 변환했는가?
□ Shared Memory를 활용해 데이터 재사용?
□ Bank Conflict가 없는가?
□ 불필요한 Host-Device 전송이 없는가?

실행 최적화

□ 적절한 블록 크기를 사용하는가? (256 권장)
□ Occupancy가 충분한가?
□ Branch Divergence가 최소화되었는가?
□ 스트림을 활용해 연산/전송 오버랩?

알고리즘 최적화

□ 병렬화에 적합한 알고리즘인가?
□ 불필요한 동기화가 없는가?
□ 워크로드가 균등하게 분배되는가?

정리

최적화 우선순위

1. 알고리즘 선택 (가장 큰 영향)
2. 메모리 접근 패턴 (Coalescing)
3. Branch Divergence 최소화
4. Occupancy 튜닝
5. 미세 최적화 (Bank Conflict 등)

핵심 포인트: 측정 없이 최적화하지 말고, Nsight 도구로 병목을 파악한 후 체계적으로 최적화하세요.