⑥ 성능 최적화

드디어 ①~⑤의 지식이 실전 무기가 되는 장입니다. 핵심 원칙 하나만 먼저 — 대부분의 GPU 커널은 "계산"이 아니라 "메모리 기다림"에서 느려집니다. 그래서 메모리 접근 패턴을 손보는 게 가장 효과가 큽니다. 효과 큰 순서대로 정리할게요.

6.1 병목 진단 — 내 커널은 memory-bound인가 compute-bound인가

기법으로 들어가기 전에, "이 커널의 한계는 어디서 오는가"를 판단하는 사고법부터. 커널의 속도는 결국 둘 중 하나에 막힙니다 — 계산기(ALU)의 속도(compute-bound)거나, 창고에서 데이터를 나르는 속도(memory-bound)거나. 어느 쪽인지에 따라 처방이 완전히 달라져요.

판단 도구가 연산 강도(arithmetic intensity)입니다 — "메모리에서 1바이트 가져올 때마다 계산을 몇 번 하는가"(FLOP/byte):

// 예: 벡터 덧셈 C[i] = A[i] + B[i]
// 원소 하나당 — 읽기 8바이트(A,B) + 쓰기 4바이트(C) = 12바이트 이동, 덧셈 1번
연산 강도 = 1 FLOP / 12 byte ≈ 0.08

// GPU의 균형점 (예: A100) — 계산 19,500 GFLOP/s ÷ 메모리 1,555 GB/s ≈ 12.5 FLOP/byte
// 0.08 ≪ 12.5  →  벡터 덧셈은 극단적 memory-bound!
//                 계산기는 거의 놀고, 성능은 100% 메모리 대역폭이 결정

이 "균형점"보다 연산 강도가 낮으면 memory-bound(대부분의 커널이 여기), 높으면 compute-bound입니다. 이 관계를 그래프로 그린 게 유명한 roofline 모델이에요 — 지붕(roof)처럼 생긴 성능 상한선에서, 왼쪽 비탈(메모리 한계)에 있느냐 천장(계산 한계)에 닿았느냐를 보는 거죠.

진단	의미	처방
memory-bound	계산기가 데이터를 기다리며 놂	메모리 이동량 줄이기 — 병합(6.2), 재사용(6.3)이 약. 계산 최적화는 효과 없음
compute-bound	계산기가 풀가동 중	계산량 줄이기 — 더 싼 수식, 낮은 정밀도(FP16 등), Tensor 코어 활용
latency-bound	둘 다 놀고 있음 (일감 부족)	병렬성 늘리기 — occupancy·ILP 확보, 더 큰 문제를 한 번에

흥미로운 점 — 연산 강도는 알고리즘 구현에 따라 바뀝니다. 행렬 곱을 그냥 짜면 같은 원소를 N번씩 창고에서 다시 읽어 memory-bound지만, 타일링(6.3)으로 한 번 가져온 데이터를 우려먹으면 "바이트당 연산 횟수"가 타일 크기만큼 올라가 compute-bound로 바뀌어요. 6.2~6.3의 기법들은 결국 "연산 강도를 끌어올리는 기술"입니다. 내 커널이 어느 쪽인지는 추측 말고 Nsight Compute(6.7)가 직접 알려줍니다.

6.2 메모리 병합 (Coalescing) — 효과 가장 큼

①에서 "32명이 창고의 연속된 칸을 한꺼번에 가져오면 효율적"이라고 했죠. 그게 coalescing(메모리 병합)입니다. 한 warp(32 스레드)가 연속된 주소를 읽으면, GPU가 그걸 하나의 큰 운반으로 묶어 처리해요. 반대로 주소가 흩어져 있으면 운반을 여러 번 쪼개야 해서 대역폭(메모리 통로)이 낭비됩니다.

// 👍 좋음: 스레드 i가 배열 i를 읽음 → 0,1,2,3...연속 → 한 번에 운반(병합)
int i = blockIdx.x*blockDim.x + threadIdx.x;
C[i] = A[i] + B[i];

// 👎 나쁨: 스레드가 stride(간격)로 띄엄띄엄 접근 → 운반이 쪼개짐
C[i * stride] = A[i * stride];

6.3 Shared Memory 타일링 — "한 번 가져와 우려먹기"

창고(global memory)는 느립니다. 그러니 여러 번 쓸 데이터는 한 번만 가져와서 공용 작업대(shared memory)에 올려두고 재사용하면 됩니다. 이걸 타일링(tiling)이라 하고, 대표 예가 행렬 곱이에요. (행렬 곱은 같은 값을 여러 번 읽거든요.)

__global__ void matmul(float* A, float* B, float* C, int N) {
    __shared__ float tA[16][16], tB[16][16];   // 공용 작업대(타일)
    int row = blockIdx.y*16+threadIdx.y, col = blockIdx.x*16+threadIdx.x;
    float sum = 0;
    for (int t = 0; t < N/16; t++) {
        tA[threadIdx.y][threadIdx.x] = A[row*N + t*16+threadIdx.x];  // ① 타일을 작업대로
        tB[threadIdx.y][threadIdx.x] = B[(t*16+threadIdx.y)*N + col];
        __syncthreads();                       // ② 다 채울 때까지 대기
        for (int k=0;k<16;k++)              // ③ 작업대에서 16번 우려먹기
            sum += tA[threadIdx.y][k]*tB[k][threadIdx.x];
        __syncthreads();                       // ④ 덮어쓰기 전 대기
    }
    C[row*N+col] = sum;
}

각 원소를 매번 창고에서 읽는 대신, 16×16 타일을 작업대에 한 번 올려 16번 재사용합니다. 창고 접근이 수십 배 줄어 성능이 크게 올라가요. (__syncthreads()가 왜 필요한지는 ②장을 떠올리세요 — "다 채운 뒤에 쓰자"를 보장.)

6.4 Bank Conflict — 작업대의 숨은 함정

공용 작업대(shared memory)에도 함정이 하나 있습니다. shared memory는 내부적으로 32개의 "뱅크(bank)"로 나뉘어 있어요 — 32개의 출납 창구라고 생각하세요. 한 warp의 여러 스레드가 같은 창구(뱅크)의 서로 다른 물건을 동시에 달라고 하면, 창구가 하나뿐이라 줄을 서야 합니다(직렬화 → 느려짐). 이게 bank conflict예요.

• 충돌 없음 — 스레드들이 서로 다른 창구에 접근하거나, 모두 같은 한 물건을 요청(이건 방송처럼 한 번에 처리=브로드캐스트).
• 흔한 회피법: 패딩 — 배열 폭에 1을 더해 창구 정렬을 어긋나게 만듭니다:

__shared__ float tile[16][16];      // 👎 열 접근 시 bank conflict 가능
__shared__ float tile[16][16+1];    // 👍 +1 패딩으로 충돌 회피

6.5 Warp 셔플과 Reduction — 작업대조차 건너뛰기

shared memory보다 더 빠른 협력 수단이 하나 더 있습니다. 같은 warp의 32 스레드는 어차피 한 몸처럼 움직이니, 서로의 레지스터 값을 직접 주고받을 수 있어요 — 이게 warp 셔플(shuffle)입니다. 작업대(shared memory)에 놓고 가져가는 단계조차 생략하고, 손에서 손으로 건네는 거죠(메모리 접근 0회, 동기화 불필요).

대표 활용처가 reduction(합·최댓값 등 하나로 모으기)입니다. 배열 전체의 합을 구한다고 해봅시다:

// warp 안 32개 값의 합 — 5번의 셔플로 끝 (반씩 접어 더하기)
float val = ...;                       // 각 스레드가 가진 값
for (int offset = 16; offset > 0; offset >>= 1)
    val += __shfl_down_sync(0xffffffff, val, offset);
// offset 16: 0번←16번, 1번←17번… → 8 → 4 → 2 → 1
// 끝나면 레인 0번이 32개의 합을 가짐

블록 전체·배열 전체의 합은 이걸 계층으로 쌓습니다:

// 전형적인 3단 reduction
① 각 warp가 셔플로 자기 32개를 합침            (레지스터끼리, 가장 빠름)
② warp 대표들이 shared memory에 모여 다시 합침  (블록당 합 완성)
③ 각 블록 대표가 atomicAdd로 전역 결과에 누적   (블록 수만큼만 — 경쟁 적음)

"수천만 개를 atomicAdd로 한 칸에 몰아넣기"(대기열 대참사)와 비교하면, 아토믹 호출이 원소 수 → 블록 수로 줄어듭니다. 이 패턴은 평균·분산·내적·softmax 등 어디에나 등장하니 형태를 기억해두세요. 직접 짜기 귀찮으면 CUB 라이브러리가 극한 최적화 버전을 제공합니다.

6.6 Occupancy · 분기 · 그 밖의 작은 기법들

• Occupancy 확보 — ④에서 본대로, SM에 warp를 충분히 채워 메모리 지연을 가립니다. 단 100%가 목표는 아니고 "지연을 가릴 만큼만".
• Warp divergence 줄이기 — 같은 warp 안에서 if/else로 갈라지면 느려진다고 했죠. 데이터를 미리 정렬하거나, 분기 없는(branchless) 코드로 바꿔 한 warp가 같은 길을 가게 합니다.
• 전송 숨기기 — 스트림 overlap + pinned memory로 복사와 계산을 겹쳐 복사 시간을 가립니다.

그리고 비용 대비 효과가 좋은 "한 줄짜리" 기법들:

기법	방법	왜 빨라지나
벡터화 적재	float 4개씩 float4로 읽기	메모리 명령 수가 1/4 — 같은 대역폭을 더 적은 명령으로
__restrict__	포인터 인자에 붙이기	"이 포인터들은 안 겹침"을 컴파일러에 보장 → 적극적 재배치·캐싱 허용
Constant memory	모든 스레드가 읽는 계수·설정값을 __constant__에	①장의 공지문 — 브로드캐스트로 한 번에 전달
-use_fast_math	컴파일 옵션	sin·exp·나눗셈을 빠른 근사 버전으로 (⚠️ 정밀도 손실 — 허용될 때만)

6.7 프로파일링 — 추측 말고 측정

최적화의 진짜 첫걸음은 "어디가 느린지 찾는 것"입니다. 눈으로 코드를 보고 추측하면 거의 틀려요. NVIDIA가 무료로 주는 측정 도구 두 가지:

도구	무엇을 보여주나
Nsight Systems (nsys)	전체 타임라인 — CPU/GPU/스트림/복사가 언제 도는지, overlap이 잘 되는지
Nsight Compute (ncu)	커널 하나 심층 분석 — occupancy, 메모리 처리량, 병합률, bank conflict 등

nsys profile ./app          # 전체 타임라인 수집 → 병목 구간 파악
ncu --set full ./app        # 커널별 상세 메트릭

6.8 최적화 체크리스트 (우선순위 순)

1. 프로파일링으로 병목 확인 — memory-bound인가 compute-bound인가 latency-bound인가. 처방이 달라짐.
2. CPU↔GPU 복사 줄이기 — 데이터는 한 번 올려서 GPU 안에서 끝내기(PCIe가 제일 느린 길).
3. 메모리 병합(coalescing) — 효과 가장 큼. 이웃 스레드가 이웃 메모리를.
4. shared memory 재사용(타일링) + bank conflict 회피, reduction엔 warp 셔플
5. occupancy 확보 — 지연 가릴 만큼 (cudaOccupancyMaxPotentialBlockSize 활용)
6. 분기 줄이기 + 전송 overlap, float4·__restrict__ 같은 작은 기법
7. 반복 워크로드는 CUDA Graphs로 launch 비용 제거
8. 라이브러리로 대체 가능한지 늘 확인 (cuBLAS·CUB·CUTLASS…)