지난 글에서는 2D 데이터 처리와 비동기 스트림(Asynchronous Streams)을 통해 Host와 Device 간 연산이 겹칠 수 있음을 확인했습니다. 이제는 GPU 메모리 계층을 살짝 들여다보고, 성능 최적화를 위해 왜 이러한 구조가 존재하는지 알아볼 차례입니다. 오늘 다룰 주제는 Shared Memory와 Constant Memory라는 특별한 메모리 공간들입니다. 이들은 단순한 글로벌 메모리(Global Memory) 접근보다 훨씬 더 빠르거나 특정 상황에서 효율적인 접근을 가능하게 해줍니다.
메모리 계층 구조 이해하기
GPU 메모리는 단순하지 않습니다. 성능을 극대화하기 위해 다양한 메모리 종류가 존재합니다.
- Global Memory(전역 메모리): 우리가 지금까지 cudaMalloc로 할당한 메모리입니다. 용량이 크지만, 접근 속도가 상대적으로 느립니다.
- Shared Memory(공유 메모리): 같은 블록 내의 스레드들이 공유하는 빠른 접근 가능한 메모리. 블록 내부 협업에 유용.
- Constant Memory(상수 메모리): 읽기 전용으로 빠른 캐싱이 가능한 메모리. 모든 스레드가 동일한 상수 데이터를 자주 참조할 때 유리.
- Register, Local Memory 등: 각 스레드별로 할당되는 레지스터나 지역 메모리도 있지만, 여기서는 깊게 다루지 않습니다.
이 중 오늘은 Shared Memory와 Constant Memory에 초점을 맞춥니다.
Shared Memory 맛보기
Shared Memory는 같은 블록 내 모든 스레드가 공유할 수 있는 고속 메모리 공간입니다. 글로벌 메모리보다 접근 속도가 훨씬 빠르므로, 스레드들이 공동으로 사용하는 중간 데이터나 자주 참조하는 중간 계산 결과를 보관하는데 유용합니다.
사용 예시
예를 들어 큰 배열에서 부분합(Partial Sum)을 구하는 커널을 생각해봅시다. 각 블록이 배열의 일부를 처리한 뒤 최종 결과를 Host로 가져온다면, 이 때 shared memory를 활용해 블록 내부의 스레드들이 협력하여 부분합을 계산할 수 있습니다.
__global__ void partial_sum_kernel(const float* input, float* block_sums, int N) {
extern __shared__ float sdata[];
// 런타임에 지정되는 shared memory 크기
int tid = threadIdx.x;
int i = blockIdx.x * (blockDim.x * 2) + tid;
// Global 메모리에서 데이터를 가져와 Shared Memory에 저장
sdata[tid] = (i < N ? input[i] : 0)
+ (i + blockDim.x < N ? input[i + blockDim.x] : 0);
__syncthreads();
// 공유 메모리에서 병렬 감소연산(Parallel Reduction) 수행
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
if (tid < stride) {
sdata[tid] += sdata[tid + stride];
}
__syncthreads();
}
// 블록별 결과를 output에 저장
if (tid == 0) {
block_sums[blockIdx.x] = sdata[0];
}
}
위 예제에서는 extern __shared__ float sdata[];를 통해 shared memory를 선언했습니다. 커널 런칭 시 <<<blocks, threads, shared_mem_size_in_bytes>>> 형태로 추가 인자를 줘서 공유 메모리 크기를 지정할 수 있습니다. 이 메모리에 글로벌 메모리에서 읽어온 데이터를 로드한 뒤, 블록 내의 모든 스레드가 빠르게 접근할 수 있어 병렬 연산을 더 효율적으로 할 수 있습니다.
장점
- Global Memory 접근 빈도를 줄이고, 캐시처럼 사용할 수 있습니다.
- 블록 내에서 공동 작업이 빈번한 연산(예: Reduction, Stencil 연산, Matrix 곱 부분 블록 캐싱)에 큰 성능 향상을 기대할 수 있습니다.
Constant Memory 맛보기
Constant Memory는 읽기 전용 메모리로, 모든 스레드가 동일한 상수 데이터를 자주 참조할 때 유용합니다. Constant Memory에 있는 데이터는 캐시를 통해 빠르게 접근 가능하고, 메모리 대역폭을 효율적으로 사용합니다.
사용 예시
예를 들어 필터 커널(Blur, Sharpen 등) 값이 고정된 경우 이 값을 Constant Memory에 저장해두면, 모든 스레드가 반복해서 이 필터 값을 Global Memory보다 빠르게 참조할 수 있습니다.
__constant__ float d_filter[9]; // 3x3 필터 예시
__global__ void apply_filter(const float* input, float* output, int M, int N) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < N && y < M) {
float sum = 0.0f;
for (int fy = -1; fy <= 1; ++fy) {
for (int fx = -1; fx <= 1; ++fx) {
int ix = min(max(x+fx,0),N-1);
int iy = min(max(y+fy,0),M-1);
float val = input[iy * N + ix];
float fval = d_filter[(fy+1)*3 + (fx+1)];
sum += val * fval;
}
}
output[y * N + x] = sum;
}
}
__constant__ 키워드로 선언한 변수는 Host 코드에서 cudaMemcpyToSymbol을 통해 값을 초기화할 수 있습니다.
float h_filter[9] = { ... }; // 필터 값 준비
cudaMemcpyToSymbol(d_filter, h_filter, 9 * sizeof(float));
이제 커널에서 d_filter 값은 빠른 접근이 가능하고, 모든 스레드가 동일한 필터값을 쉽게 참조할 수 있습니다.
최적화 맛보기: 왜 이런 메모리가 필요할까?
Global Memory 접근은 상대적으로 느리고, 병렬로 동작하는 수많은 스레드들이 동시에 Global Memory를 읽으면 메모리 대역폭 병목이 발생할 수 있습니다. Shared Memory나 Constant Memory는 이런 병목을 줄이고 연산을 가속하는 데 도움이 됩니다.
Before:
- 모든 데이터 접근을 Global Memory에 의존 → 더 느리고 비효율적
After: - Shared Memory를 활용한 데이터 재사용, Constant Memory를 통한 상수 데이터 고속 접근 → 향상된 성능
물론 이런 최적화를 적용하기 위해서는 코드 복잡도가 올라갈 수 있고, 하드웨어 특성과 데이터 접근 패턴을 잘 이해해야 합니다. 하지만 이런 계층 구조를 이해하는 것만으로도 앞으로 성능 튜닝을 위한 아이디어를 많이 얻을 수 있습니다.
다음 글 예고
다음 글(#7)에서는 이제까지 배운 내용들을 좀 더 체계적으로 빌드 시스템(CMake)과 결합하고, Modern C++ 특징(콘셉트, 람다, constexpr 등)을 CUDA 코드에 녹여내는 방법을 살펴볼 예정입니다. 지금까지는 기초 CUDA 문법과 메모리 계층, 비동기 처리 등을 다뤘다면, 이제 "현대적 C++ 스킬"을 적극적으로 CUDA 코드에도 적용해볼 시간입니다.
유용한 링크 & 리소스
- NVIDIA CUDA Programming Guide: Shared/Constant Memory 개념 정리
- NVIDIA CUDA Samples: Shared Memory 최적화 예제, Constant Memory 활용 예제
- Optimization Tricks on CUDA NVIDIA 개발자 블로그: 다양한 최적화 기법 소개
- Parallel Reduction 예제: Shared Memory 활용한 고전적인 예제