Mixed Precision 훈련 -- 숫자의 정밀도를 줄여서 속도를 올리자

학습 목표

• •CUDA 메모리 계층 구조를 이해한다
• •Global, Shared, Constant, Local 메모리의 특징을 배운다
• •각 메모리 유형의 적절한 사용 시나리오를 파악한다
• •Shared Memory를 활용한 최적화 기법을 익힌다

CUDA 메모리 계층

GPU는 다양한 종류의 메모리를 가지며, 각각 속도와 용도가 다릅니다.

메모리 계층 개요

메모리 비교

Global Memory

Global Memory는 GPU의 메인 메모리입니다.

특징

장점:
- 가장 큰 용량 (8GB ~ 80GB)
- 모든 스레드에서 접근 가능
- Host와 데이터 교환 가능

단점:
- 가장 느린 메모리 (400-600 cycles)
- Off-chip에 위치

할당 및 사용

cuda
복사
// Global Memory 할당
float* d_data;
cudaMalloc(&d_data, size);

// Host에서 Global Memory로 복사
cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice);

// 커널에서 Global Memory 접근
__global__ void kernel(float* data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float value = data[idx];  // Global Memory 읽기
    data[idx] = value * 2;    // Global Memory 쓰기
}

Coalesced Access (합쳤된 접근)

좋은 접근 패턴 (Coalesced):
Thread 0 → data[0]
Thread 1 → data[1]
Thread 2 → data[2]
→ 한 번의 메모리 트랜잭션으로 처리

나쁜 접근 패턴 (Strided):
Thread 0 → data[0]
Thread 1 → data[32]
Thread 2 → data[64]
→ 여러 번의 메모리 트랜잭션 필요 (32배 느려질 수 있음)

Shared Memory

Shared Memory는 블록 내 스레드들이 공유하는 고속 메모리입니다.

특징

선언 방법

cuda
복사
// 정적 할당 (컴파일 시 크기 결정)
__shared__ float sharedData[256];

// 동적 할당 (런타임에 크기 결정)
extern __shared__ float dynamicShared[];

// 커널 호출 시 동적 크기 지정
kernel<<<grid, block, sharedMemSize>>>(args);

활용 예시: 타일 행렬 곱셈

cuda
복사
__global__ void matMulTiled(float* A, float* B, float* C, int N) {
    // Shared Memory에 타일 로드
    __shared__ float tileA[TILE_SIZE][TILE_SIZE];
    __shared__ float tileB[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < N / TILE_SIZE; t++) {
        // Global → Shared로 데이터 로드
        tileA[threadIdx.y][threadIdx.x] = A[row * N + t * TILE_SIZE + threadIdx.x];
        tileB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * N + col];

        __syncthreads();  // 모든 스레드가 로드 완료할 때까지 대기

        // Shared Memory에서 연산 (빠름!)
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        }

        __syncthreads();
    }

    C[row * N + col] = sum;
}

__syncthreads()

cuda
복사
// 블록 내 모든 스레드 동기화
__syncthreads();

// 언제 필요한가?
// 1. Shared Memory에 쓴 후, 다른 스레드가 읽기 전
// 2. 모든 스레드가 특정 작업을 완료해야 할 때

// 주의: 조건문 안에서 __syncthreads() 사용 시
// 모든 스레드가 같은 조건을 만족해야 함 (deadlock 방지)

Constant Memory

Constant Memory는 읽기 전용 캐시 메모리입니다.

특징

- 크기: 64KB 고정
- 모든 스레드가 같은 값을 읽을 때 최적
- 캐시되어 빠른 접근 가능
- Host에서 설정, Device에서 읽기만

선언 및 사용

cuda
복사
// 전역 범위에 선언
__constant__ float constData[256];

// Host에서 값 설정
float hostData[256];
cudaMemcpyToSymbol(constData, hostData, sizeof(hostData));

// 커널에서 읽기
__global__ void kernel() {
    float value = constData[threadIdx.x];  // 캐시된 읽기
}

적합한 용도

✓ 컨볼루션 필터 가중치
✓ 룩업 테이블
✓ 물리 상수
✓ 변환 행렬

✗ 스레드마다 다른 데이터 접근
✗ 동적으로 변하는 데이터

Local Memory (Register)

Local Memory와 Register는 각 스레드의 개인 저장소입니다.

Register

cuda
복사
__global__ void kernel() {
    // 레지스터에 저장 (가장 빠름)
    float localVar = 0.0f;    // 레지스터
    int index = threadIdx.x;  // 레지스터

    // 연산도 레지스터에서 수행
    localVar = localVar + 1.0f;
}

Local Memory (Register Spill)

cuda
복사
__global__ void kernel() {
    // 레지스터가 부족하면 Local Memory로 스필
    float largeArray[100];  // 큰 배열 → Local Memory

    // Local Memory는 실제로 Global Memory에 위치
    // → 느리므로 피해야 함
}

레지스터 사용량 확인

bash
복사
# 컴파일 시 레지스터 사용량 출력
nvcc --ptxas-options=-v kernel.cu

# 출력 예시:
# ptxas info: Used 32 registers, 48 bytes smem, 0 bytes cmem

메모리 최적화 전략

전략 1: Global Memory 접근 최소화

cuda
복사
// 나쁜 예: Global Memory 반복 접근
__global__ void bad(float* data) {
    for (int i = 0; i < 100; i++) {
        data[threadIdx.x] += 1.0f;  // 매번 Global 접근
    }
}

// 좋은 예: 레지스터 활용
__global__ void good(float* data) {
    float temp = data[threadIdx.x];  // 한 번만 읽기
    for (int i = 0; i < 100; i++) {
        temp += 1.0f;  // 레지스터에서 연산
    }
    data[threadIdx.x] = temp;  // 한 번만 쓰기
}

전략 2: Shared Memory로 데이터 재사용

cuda
복사
// 여러 스레드가 같은 데이터를 사용할 때
__shared__ float sharedData[BLOCK_SIZE];

// 1. 협력하여 데이터 로드
sharedData[threadIdx.x] = globalData[blockIdx.x * BLOCK_SIZE + threadIdx.x];
__syncthreads();

// 2. Shared Memory에서 반복 접근
for (int i = 0; i < BLOCK_SIZE; i++) {
    result += sharedData[i];  // 빠른 접근
}

전략 3: Coalesced Access 보장

cuda
복사
// 연속 스레드 → 연속 메모리
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float value = data[idx];  // Coalesced!

// 구조체 배열 (SoA)이 배열의 구조체 (AoS)보다 효율적
// AoS: struct Particle { float x, y, z; } particles[N];
// SoA: float x[N], y[N], z[N];  // 권장

정리

최적화 원칙

1. Global Memory 접근을 최소화
2. 데이터 재사용이 있으면 Shared Memory 활용
3. Coalesced Access 패턴 유지
4. 레지스터 스필을 피함 (배열 크기 제한)
5. __syncthreads()로 동기화 보장

핵심 포인트: 메모리 계층을 이해하고 적절한 메모리를 선택하는 것이 CUDA 성능 최적화의 핵심입니다.