GPU 메모리 코일레싱 (Memory Coalescing)
개요 (Overview)
GPU의 병렬 연산 성능은 메모리 서브시스템의 효율성에 크게 좌우됩니다. Memory Coalescing은 워프(warp) 내 32개 스레드가 각각 요청하는 글로벌 메모리 접근을 하드웨어가 자동으로 합쳐 최소한의 메모리 트랜잭션으로 처리하는 최적화 기법입니다. 비용이 높은 글로벌 메모리 접근에서 스트라이드 패턴이나 비정렬 접근이 발생하면 대역폭이 크게 저하되지만, 코일레싱이 제대로 이루어지면 하나의 트랜잭션(32바이트 단위)으로 여러 스레드의 요청을 한번에 서비스할 수 있습니다.
코일레싱의 효율은 Global Memory Load/Store Efficiency 지표로 측정하며, 이 값이 높을수록 실제 사용된 대역폭 대비 필요한 대역폭의 비율이 좋음을 의미합니다. NVIDIA CUDA Best Practices Guide에서는 코일레싱을 High Priority 최적화 항목으로 지정하고 있으며, 특히 GDDR 메모리에서 ECC가 활성화된 경우 분산 접근 시 ECC 전송 오버헤드가 증가하므로 더욱 중요합니다.
실무에서는 코일레싱을 단순히 "연속 주소 접근"으로만 이해하면 부족합니다. 같은 워프 안에서도 데이터 구조가 AoS(Array of Structures)인지 SoA(Structure of Arrays)인지, 블록 크기가 32의 배수인지, 2차원 타일을 어떤 축으로 읽는지에 따라 요청이 서로 다른 세그먼트로 분산될 수 있습니다. 따라서 커널 작성 시에는 메모리 레이아웃, 인덱싱 방식, Shared Memory 재배치를 함께 설계해야 실제 대역폭 이득을 안정적으로 얻을 수 있습니다.
핵심 개념
워프와 메모리 트랜잭션
CUDA에서 최소 실행 단위인 워프(warp)는 32개의 스레드로 구성됩니다. 각 스레드는 하나의 메모리 주소에 접근하지만, 하드웨어는 워프 단위로 접근 요청을 모아 처리합니다. Compute Capability 6.0 이상에서는 워프의 동시 접근이 32바이트 단위 트랜잭션으로 합쳐지며, 하나의 32바이트 세그먼트 내에서 스레드가 어떤 위치에 접근하든 동일한 수의 트랜잭션이 수행됩니다.
정렬 접근 (Aligned Access)
스트라이드가 1이고 32바이트 경계에 정렬된 접근이 가장 효율적입니다. 예를 들어 32개 스레드가 각각 인접한 float(4바이트) 요소에 접근하면, 총 128바이트가 필요하며 이는 4개의 32바이트 트랜잭션으로 처리됩니다. cudaMalloc()으로 할당된 메모리는 최소 256바이트 경계로 정렬이 보장되므로, 적절한 크기의 블록 크기를 사용하면 자연스럽게 정렬 접근이 가능합니다.
미정렬 순차 접근 (Misaligned Sequential Access)
스트라이드가 1이지만 32바이트 경계와 정렬되지 않은 접근입니다. 이 경우 128바이트 데이터가 5개의 32바이트 세그먼트에 걸쳐 분산되어 트랜잭션이 하나 더 발생합니다. V100 기준으로 offset=0일 때 약 790 GB/s인 대역폭이 미정렬 시 약 10~25% 저하되지만, 인접 워프들이 캐시 라인을 재사용하기 때문에 이론적 손실보다는 적게 나타납니다.
스트라이드 접근 (Strided Access)
워프 내 스레드들이 일정한 간격(stride)을 두고 메모리에 접근하는 패턴입니다. 다차원 배열이나 행렬 연산에서 흔히 발생하며, 스트라이드가 커질수록 효율이 급격히 떨어집니다. 스트라이드가 2이면 50% 효율(8개 세그먼트), 스트라이드가 32이면 32개 세그먼트가 필요하여 거의 피크 대역폭의 1/8 수준으로 성능이 저하됩니다.
Shared Memory를 활용한 코일레싱
글로벌 메모리에서 스트라이드 접근이 불가피한 경우, Shared Memory를 활용하여 코일레싱된 접근으로 변환할 수 있습니다. 행렬 곱셈 C = A × B의 경우, 행렬 A의 타일을 Shared Memory에 로드할 때 워프가 인접한 열에 접근하도록 설계하면 글로벌 메모리 읽기가 코일레싱됩니다. C = A × A^T 연산에서는 Shared Memory를 사용하지 않으면 스트라이드 읽기로 인해 12.8 GB/s에 불과하던 대역폭을 140.2 GB/s로 개선할 수 있습니다(약 11배 향상).
데이터 레이아웃과 워프 매핑
코일레싱은 주소 자체보다도 "워프의 lane ID가 어떤 주소열로 매핑되는가"에 더 직접적으로 좌우됩니다. 예를 들어 float4처럼 벡터화된 로드는 명령 수를 줄일 수 있지만, 구조체 필드 배치가 lane별로 흩어져 있으면 여전히 비효율적인 세그먼트 접근이 발생합니다. 반대로 SoA 레이아웃으로 각 스레드가 같은 필드의 연속 원소를 읽게 만들면, 동일한 계산이라도 트랜잭션 수를 줄여 코일레싱 효율을 높일 수 있습니다.
비교/분석
| 접근 패턴 | 트랜잭션 수 (워프당) | 대역폭 효율 | 비고 |
|---|---|---|---|
| 정렬 순차 (stride=1) | 4 | 100% | 이상적인 패턴 |
| 미정렬 순차 | 5 | ~80% | 캐시 재사용으로 완화 |
| Stride = 2 | 8 | 50% | Shared Memory 필요 |
| Stride = N (큰 값) | 32 | ~12.5% | Shared Memory 필수 |
| 랜덤 접근 | 32 | 최악 | 회피 필요 |
| 최적화 단계 | 대역폭 (V100 기준) | 비고 |
|---|---|---|
| 비최적화 (stride 읽기) | 12.8 GB/s | C = A × A^T |
| Shared Memory로 코일레싱 | 140.2 GB/s | 11배 향상 |
| Bank conflict 해소 | 199.4 GB/s | 15.6배 향상 |
동작 원리
하드웨어 코일레싱 메커니즘
워프가 글로벌 메모리에 접근하면, 하드웨어는 32개 스레드의 요청 주소를 분석하여 32바이트 세그먼트 단위로 그룹화합니다. Compute Capability 6.0 이상에서는 L1 캐시와 무관하게 32바이트 트랜잭션을 기준으로 동작합니다. 요청된 워드가 32바이트 세그먼트 내에 일부만 포함되더라도 전체 세그먼트가 페치되므로, 인접 워프들이 동일한 캐시 라인의 나머지 데이터를 재사용할 수 있습니다.
블록 크기와 정렬
스레드 블록 크기는 워프 크기(32)의 배수여야 합니다. 블록 크기가 32의 배수가 아니면 후속 블록의 첫 워프가 비정렬된 주소에서 접근을 시작할 수 있습니다. cudaMalloc()이 256바이트 경계로 정렬을 보장하므로, 첫 번째 블록의 정렬은 보장되지만 후속 블록은 블록 크기에 따라 달라집니다.
Shared Memory 활용 흐름
- 글로벌 메모리에서 데이터를 Shared Memory에 코일레싱된 접근으로 로드
__syncthreads()로 동기화- Shared Memory에서 계산에 필요한 데이터를 재사용
- 결과를 글로벌 메모리에 코일레싱된 접근으로 저장
이 흐름은 글로벌 메모리의 접근 빈도를 줄이고, 스트라이드 접근 패턴을 Shared Memory 내부에서 해결하여 전체 대역폭 효율을 크게 향상시킵니다.
커널 프로파일링 단계에서는 Global Memory Load Efficiency, Global Memory Store Efficiency, gld_transactions, gst_transactions 같은 지표를 함께 보면 좋습니다. 효율 수치가 낮은데도 실행 시간이 예상보다 작게 나오는 경우는 캐시 재사용이나 연산 병목이 숨어 있을 수 있으므로, 단일 지표만 보지 말고 워프당 트랜잭션 수와 stall 원인을 함께 해석해야 합니다.
장단점
장점
- 글로벌 메모리 대역폭을 최대한 활용하여 커널 성능 향상
- 하드웨어가 자동으로 처리하므로 프로그래머의 추가 노력이 적음
cudaMalloc()의 정렬 보장과 적절한 블록 크기 설정만으로 기본적인 코일레싱 달성 가능- Shared Memory와 결합하면 스트라이드 접근 패턴도 효율적으로 변환 가능
단점
- 데이터 레이아웃이 접근 패턴과 불일치하면 코일레싱이 어려움
- 2차원/3차원 배열의 열 방향 접근은 본질적으로 스트라이드 패턴
- Shared Memory 활용 시 Shared Memory 리소스 제한으로 occupancy 저하 가능
- Bank conflict가 발생하면 Shared Memory 접근 자체가 병목이 됨
관련 기술
| 자료 | 링크 | 연결점 |
|---|---|---|
| CUDA C++ Best Practices Guide — Coalesced Access to Global Memory | docs.nvidia.com/cuda/cuda-c-best-practices-guide | 코일레싱의 공식 기술 사양과 최적화 가이드 |
| CUDA C++ Programming Guide — Memory Hierarchy | docs.nvidia.com/cuda/cuda-c-programming-guide | 글로벌 메모리, Shared Memory 등 메모리 계층 전체 이해 |
| GPU 메모리 아키텍처 기초 | gpu_0005_gpu_memory_architecture_basics.html | GPU 메모리 계층 구조와 대역폭 비교 |
| CUDA 메모리 관리 | gpu_0010_cuda_memory_management.html | cudaMalloc, Unified Memory 등 CUDA 메모리 API |
| GPU 메모리 최적화 기법 | gpu_0015_gpu_memory_optimization.html | Occupancy, Shared Memory, bank conflict 등 최적화 기법 |
핵심 정리
GPU 메모리 코일레싱은 워프 내 스레드들의 글로벌 메모리 접근을 32바이트 단위 트랜잭션으로 자동 합치는 하드웨어 기반 최적화입니다. 정렬된 순차 접근(stride=1)이 최대 효율을 보이며, 스트라이드가 커질수록 대역폭이 급격히 저하됩니다. CUDA Best Practices Guide에서는 코일레싱을 High Priority 항목으로 지정하고 있습니다.
실제 성능 차이는 매우 큽니다. 행렬 곱셈 C = A × A^T에서 Shared Memory를 활용한 코일레싱 최적화는 대역폭을 11배 향상시키며(12.8 → 140.2 GB/s), bank conflict까지 해소하면 15.6배까지 개선됩니다. 적절한 블록 크기(32의 배수)와 cudaMalloc()의 정렬 보장만으로도 기본적인 코일레싱이 가능하지만, 복잡한 접근 패턴에서는 Shared Memory를 활용한 변환이 필수적입니다.
코일레싱 효율은 Global Memory Load/Store Efficiency 지표로 측정할 수 있으며, 커널 프로파일링 시 이 값을 확인하여 최적화 여부를 판단해야 합니다.