GPU 메모리 최적화 기법
개요 (Overview)
GPU 성능은 연산 유닛의 수만큼이나 메모리 접근 방식에 크게 좌우됩니다. 같은 커널이라도 글로벌 메모리 접근이 흩어져 있거나, 같은 데이터를 여러 번 다시 읽거나, 호스트와 디바이스 사이의 이동이 잦으면 처리량이 급격히 떨어집니다. 그래서 GPU 메모리 최적화는 단순히 더 큰 버퍼를 쓰는 일이 아니라, 이동할 바이트 수를 줄이고, 같은 데이터를 더 많이 재사용하고, 접근을 연속적으로 만드는 작업입니다.
실무에서는 병목을 먼저 찾고, 그 다음 접근 패턴을 정리한 뒤, Shared Memory 타일링으로 재사용을 늘리고, 필요하면 압축이나 정밀도 축소로 메모리 사용량 자체를 낮춥니다. 남는 지연 시간은 prefetch와 비동기 복사로 가리는 방식이 자주 쓰입니다. 이 문서는 이런 순서를 기준으로 GPU 메모리 최적화의 핵심을 정리합니다.
최적화 흐름
핵심 개념
병목부터 찾기
GPU 최적화의 첫 단계는 연산이 아니라 메모리 병목을 확인하는 일입니다. Nsight Compute나 Nsight Systems로 DRAM 대기, 낮은 대역폭 사용률, 불규칙한 warp stall을 확인하면 어떤 커널이 메모리 중심인지 빠르게 구분할 수 있습니다. 계산량이 많아 보여도 실제로는 메모리 접근 때문에 SM이 쉬는 경우가 많습니다.
coalesced access
warp 안의 thread들이 인접한 주소를 읽고 쓸수록 global memory transaction 수가 줄어듭니다. 반대로 stride가 크거나 구조체 배치가 좋지 않으면 같은 양의 데이터를 옮겨도 실제 효율은 크게 떨어집니다. 최신 GPU는 예전보다 misalignment에 관대하지만, strided access 자체는 여전히 피하는 편이 좋습니다.
Shared Memory 타일링
Shared Memory는 thread block이 함께 쓰는 on-chip 버퍼입니다. 자주 재사용되는 데이터 블록을 global memory에서 한 번 읽어 Shared Memory에 올려두면, 이후에는 block 내부 thread들이 더 빠르게 접근할 수 있습니다. 행렬 곱셈, transpose, reduction, attention 계열 커널에서 이 방식이 특히 강합니다. 다만 bank conflict가 생기면 성능이 흔들리므로 padding이나 인덱싱 조정이 필요합니다.
occupancy와 on-chip 자원 균형
Shared Memory와 레지스터를 많이 쓰면 데이터 재사용은 좋아지지만, 동시에 실행 가능한 block 수가 줄어 occupancy가 떨어질 수 있습니다. 그래서 타일 크기를 무작정 키우기보다, 한 번의 global memory load가 얼마나 많은 계산에 재사용되는지와 그 대가로 줄어드는 병렬성을 함께 봐야 합니다. 실무에서는 profiler에서 achieved occupancy, shared memory usage, register pressure를 같이 보면서 sweet spot을 찾습니다.
메모리 압축과 정밀도 축소
같은 내용을 더 작은 표현으로 담으면 옮겨야 할 바이트 수가 줄어듭니다. FP32 대신 FP16/BF16을 쓰거나, INT8/INT4처럼 더 낮은 정밀도를 쓰거나, 구조적으로 sparse한 데이터를 압축해 저장하는 방식이 여기에 들어갑니다. 이 방법은 bandwidth와 용량을 동시에 아낄 수 있지만, 복원 비용과 정확도 손실 여부를 함께 봐야 합니다.
prefetch와 비동기 복사
메모리 지연을 없앨 수 없다면 미리 당겨오면 됩니다. 호스트-디바이스 이동은 cudaMemcpyAsync로 겹칠 수 있고, 커널 내부의 global-to-shared 이동은 cp.async 계열과 pipeline으로 숨길 수 있습니다. 반복 루프에서 다음 타일을 미리 준비해 두면 계산과 이동이 겹치면서 실효 대역폭이 좋아집니다.
비교/분석
| 기법 | 주로 줄이는 비용 | 잘 맞는 상황 | 주의점 |
|---|---|---|---|
| coalesced access | global memory transaction | warp 단위로 연속 데이터 접근 | 데이터 배치가 어긋나면 효과가 급감함 |
| Shared Memory 타일링 | 같은 데이터의 재읽기 | 행렬/타일/블록 재사용이 큰 커널 | bank conflict와 shared memory 용량 제한 |
| occupancy 조정 | 지연 은닉 실패 | 타일 크기와 on-chip 버퍼가 큰 커널 | shared memory와 register 사용이 지나치면 동시 실행 수가 감소함 |
| 압축/정밀도 축소 | 저장·전송 바이트 수 | 정확도 여유가 있고 데이터가 크면 유리 | 복원 비용과 수치 오차를 함께 고려해야 함 |
| prefetch / async copy | 메모리 지연 시간 | 반복 구조와 독립 iteration이 있는 루프 | 레지스터, shared memory, 동기화 비용이 늘 수 있음 |
동작 원리
GPU 메모리 최적화는 보통 다음 순서로 진행됩니다. 먼저 profiler로 병목 커널을 찾고, global memory 접근이 warp 단위로 묶이는지 확인합니다. 그다음 반복해서 읽는 데이터를 block 단위 타일로 묶어 Shared Memory에 올리고, thread들은 그 타일 안에서 계산하도록 바꿉니다.
이 단계에서는 bank conflict가 없는지도 같이 확인해야 합니다. tile을 잘게 나눴더라도 여러 thread가 같은 bank에 몰리면 Shared Memory 지연이 다시 커질 수 있기 때문입니다. transpose처럼 열 방향 접근이 많은 커널은 padding을 한 칸 넣는 것만으로도 충돌을 크게 줄일 수 있습니다.
이후에도 bandwidth가 부족하면 데이터 표현 자체를 줄입니다. 예를 들어 FP16이나 BF16으로 저장하면 같은 용량에 더 많은 데이터를 담을 수 있고, sparse representation은 실제 사용되는 값만 옮겨서 전송량을 낮춥니다. 마지막으로 계산과 이동을 분리하지 말고 겹치도록 구성합니다. cudaMemcpyAsync는 호스트와 디바이스 사이 이동을, cp.async는 device 내부 이동을 숨기는 데 자주 쓰입니다.
다만 각 단계는 독립 최적화가 아니라 서로 영향을 줍니다. 타일을 키우면 재사용은 늘지만 shared memory 사용량이 커지고, async copy를 깊게 걸면 파이프라인은 좋아지지만 register pressure와 동기화 비용이 늘 수 있습니다. 따라서 최종 튜닝은 단일 기법의 최대화가 아니라, 대역폭 절감과 occupancy 유지 사이의 균형을 맞추는 과정에 가깝습니다.
핵심은 한 번 읽은 데이터를 여러 번 쓰는 구조로 바꾸는 것입니다. 이 구조가 되면 global memory는 단순 공급원 역할을 하고, 빠른 on-chip memory가 실제 작업 공간이 됩니다.
장단점
장점
- global memory 접근량을 직접 줄일 수 있습니다.
- Shared Memory와 타일링은 재사용이 큰 커널에서 효과가 큽니다.
- 압축과 정밀도 축소는 메모리 용량과 bandwidth를 동시에 절약합니다.
- 비동기 복사는 계산과 이동을 겹쳐 지연을 숨기기 좋습니다.
- profiler 기반으로 병목 유형별 대응책을 나눌 수 있어 튜닝 우선순위를 세우기 쉽습니다.
단점
- 접근 패턴을 바꾸는 작업은 코드 복잡도를 높입니다.
- Shared Memory는 bank conflict와 용량 제한을 신경 써야 합니다.
- 압축과 낮은 정밀도는 정확도 손실이나 복원 비용을 가져올 수 있습니다.
- prefetch와 비동기 복사는 동기화와 리소스 배분을 잘못 잡으면 오히려 느려질 수 있습니다.
- 타일 크기와 버퍼를 공격적으로 키우면 occupancy 하락으로 지연 은닉 능력이 떨어질 수 있습니다.
관련 기술 / 참고 문헌
| 자료 | 링크 | 연결점 |
|---|---|---|
| CUDA C++ Best Practices Guide - Memory Optimizations | https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#memory-optimizations | coalescing, shared memory, pinned memory, async transfer |
| CUDA Programming Guide | https://docs.nvidia.com/cuda/cuda-programming-guide/index.html | shared memory, pipelines, async copies, memory model |
| Using Shared Memory in CUDA C/C++ | https://developer.nvidia.com/blog/using-shared-memory-cuda-cc/ | block 단위 Shared Memory와 __syncthreads() |
| How to Access Global Memory Efficiently in CUDA C/C++ Kernels | https://developer.nvidia.com/blog/how-access-global-memory-efficiently-cuda-c-kernels/ | coalesced access, strided access, tiling |
| Boosting Application Performance with GPU Memory Prefetching | https://developer.nvidia.com/blog/boosting-application-performance-with-gpu-memory-prefetching/ | prefetch, loop unrolling, async copy |
| CUDA Occupancy Calculator | https://docs.nvidia.com/cuda/archive/11.4.1/cuda-occupancy-calculator/index.html | shared memory, register usage, block size에 따른 occupancy 점검 |
| GPU 메모리 아키텍처 기초 | gpu_0005_gpu_memory_architecture_basics.html | GDDR/HBM 대역폭과 메모리 계층 배경 |
| CUDA 메모리 관리 | gpu_0010_cuda_memory_management.html | cudaMemcpyAsync, cudaMallocManaged, pinned memory |
| FlashAttention Analysis | ../llm/llm_0023_flashattention_analysis.html | tiling과 메모리 재사용의 대표 사례 |
| KV Cache Quantization Analysis | ../llm/llm_0015_kv_cache_quantization_analysis.html | 정밀도 축소와 압축형 메모리 절약 사례 |
핵심 정리
GPU 메모리 최적화는 더 빠른 메모리를 찾는 일이 아니라, 더 적게 읽고 더 많이 재사용하는 구조를 만드는 일입니다. coalesced access는 기본기이고, Shared Memory 타일링은 재사용을 키우는 핵심 수단입니다.
압축과 정밀도 축소는 용량과 bandwidth를 동시에 줄일 수 있지만, 정확도와 복원 비용을 함께 봐야 합니다. 남는 지연은 prefetch와 async copy로 숨기고, 최종적으로는 profiler가 보여주는 병목에 맞춰 조합을 선택하는 것이 가장 효과적입니다.
특히 Shared Memory와 레지스터를 많이 쓰는 최적화는 재사용 이득만큼 occupancy 저하 위험도 함께 가져옵니다. 그래서 좋은 GPU 메모리 튜닝은 한 기법을 극단적으로 밀기보다, transaction 수 감소, bank conflict 회피, occupancy 유지의 균형점을 찾는 과정으로 이해하는 편이 정확합니다.