CUDA - 5장. CUDA 최적화 기법

 

⚡ 5장. CUDA 최적화 기법

🎯 학습 목표

이 장에서는 CUDA 프로그램의 성능을 극대화하기 위한 핵심 전략을 학습합니다.
공유 메모리, 스레드 최적화, 메모리 정렬, 연산 최소화, 스트리밍, 텍스처 메모리 등 실전 중심의 최적화 기법을 익히고 적용합니다.

---

🧭 최적화 전략 미리보기

항목	설명	난이도
공유 메모리	블록 내 연산 캐싱	🟢 초급
메모리 정렬	연속된 주소 접근	🟢 초급
워프 최적화	분기 제거, 32개 단위 실행	🟡 중급
연산 최소화	중복 제거, 사전 계산	🟢 초급
스트리밍	복수 작업 동시 처리	🟡 중급
텍스처 메모리	이미지/2D 연산 특화 캐시	🔴 고급

---

5.1 메모리 최적화 🟢

✅ 공유 메모리 활용

공유 메모리는 블록 내 스레드 간 빠른 데이터 공유에 사용됩니다.

__shared__ float tile[256];

int idx = threadIdx.x;
tile[idx] = input[idx];
__syncthreads();

output[idx] = tile[idx] * 2.0f;

• 전역 메모리보다 수십 배 빠름
• 반복 접근, 블록 내 협업에 적합

---

✅ 메모리 정렬(Coalescing)

int idx = threadIdx.x + blockIdx.x * blockDim.x;
output[idx] = input[idx]; // ✅ 연속 주소 접근

• 32개 스레드(워프)가 연속 주소에 접근할 경우, 메모리 접근 병합(coalescing) 발생 → 성능 향상
• float4, int4 등의 벡터 타입도 정렬에 유리

---

5.2 스레드 병렬성 및 워프 최적화 🟡

🧠 워프(Warp) 개념

• 32개의 스레드 → 1 워프 단위로 동시에 실행
• 워프 내 분기(조건문 등)가 있으면 분기 실행(divergence) 발생 → 성능 저하

// ❌ 비효율적
if (threadIdx.x % 2 == 0) a[idx] += 1;
else a[idx] *= 2;

// ✅ 조건문 제거
int even = (threadIdx.x % 2 == 0);
a[idx] = even ? a[idx] + 1 : a[idx] * 2;

✅ 스레드 최적화 팁

항목	전략
블록 크기	32의 배수(128, 256 등)
워프 통일성	분기 없이 동일한 명령 실행 유도
연속 인덱싱	인접 스레드가 인접 데이터 처리

---

5.3 연산량 최소화 및 메모리 대역폭 최적화 🟢~🟡

➗ 중복 계산 제거

// ❌
for (...) sum += sqrt(a[i]) * sqrt(b[i]);

// ✅
float sa = sqrt(a[i]), sb = sqrt(b[i]);
sum += sa * sb;

---

📶 메모리 대역폭 최적화 전략

전략	설명
공유 메모리 캐싱	반복 접근 시 전역 → 공유 전환
불필요 복사 최소화	필요한 범위만 복사
__restrict__	읽기 전용 변수로 힌트 제공
정렬된 구조체	__align__, float4 활용

__global__ void compute(const float* __restrict__ in, float* out) {
    int i = threadIdx.x;
    out[i] = in[i] * 2.0f;
}

---

5.4 CUDA 스트리밍 (Streams) 활용 🟡

🧩 개념

• CUDA 스트림은 작업 단위(데이터 전송, 커널 실행 등)를 비동기 처리할 수 있도록 해줍니다.
• 여러 연산을 겹쳐서 실행(overlap) 가능

▶️ 예시 코드

cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

cudaMemcpyAsync(d1, h1, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d2, h2, size, cudaMemcpyHostToDevice, stream2);

kernel<<<blocks, threads, 0, stream1>>>(d1);
kernel<<<blocks, threads, 0, stream2>>>(d2);

cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

✅ 데이터 전송과 커널 실행을 병렬로 실행 가능

---

5.5 텍스처 메모리 최적화 🔴

🖼️ 특징

항목	설명
구조	전용 캐시 + 보간 기능
용도	2D/3D 이미지, 보간 LUT
속도	전역 메모리보다 빠름 (특히 2D 엑세스)

📌 사용 예시

texture<float, 2, cudaReadModeElementType> tex;

__global__ void kernel() {
    float val = tex2D(tex, x, y);
}

⚠️ 고급 기능이지만 영상/시뮬레이션 등에서 큰 성능 향상 가능

---

📈 최적화 전후 성능 비교 (예시)

적용 전략	Before	After	향상률
공유 메모리	3.2ms	0.9ms	↑ 3.6배
워프 최적화	2.5ms	1.4ms	↑ 1.8배
스트리밍	6.8ms	3.3ms	↑ 2배 이상

테스트 환경: RTX 3060, CUDA 12.3, Ubuntu 22.04 기준

---

📌 상황별 전략 요약

상황	전략 추천
연산량 많은 반복	중복 계산 제거, shared memory
이미지 처리	텍스처 메모리
데이터 전송 지연	스트림 + cudaMemcpyAsync()
병목이 메모리 쪽	메모리 정렬, read-only cache, Coalescing
조건문 복잡	워프 정렬, 조건 분기 제거

---

🧪 연습 문제

1. 공유 메모리를 사용해 vectorAdd 성능 개선 후 속도 비교
2. if-else 문을 제거한 워프 정렬 코드 작성
3. cudaMemcpyAsync() + 커널 실행으로 스트리밍 실습
4. Nsight Compute를 사용해 병목 지점 프로파일링

---

🔎 성능 분석 도구

도구	기능
Nsight Systems	전체 흐름 타임라인 보기
Nsight Compute	커널 단위 성능 상세 분석
cuda-memcheck	메모리 접근 오류 탐지
nv-nsight-cu-cli	CLI 기반 프로파일링

🎯 병목 지점을 확인하고, 해당 영역에 최적화 기법 적용 필요

---

✅ 최종 요약

항목	설명 요약
공유 메모리	빠르고 블록 내부 협업에 유리
Coalescing	연속 주소 접근 → 빠른 전역 메모리
워프 최적화	분기 방지, 블록 크기 조정
연산 최적화	불필요 계산 제거, 사전 처리
스트리밍	복수 연산 병렬 실행
텍스처 메모리	2D/보간 특화 캐시
프로파일링 도구	Nsight로 병목 구간 시각적 분석