CUDA 메모리 관리
개요 (Overview)
CUDA 메모리 관리는 호스트 DRAM과 디바이스 메모리 사이의 데이터 이동, 할당 수명, 주소 가시성을 함께 다루는 계층입니다. GPU 연산 성능이 높아질수록 커널 자체보다 메모리 배치와 전송 방식이 병목이 되기 쉬워서, cudaMalloc과 cudaMemcpy 같은 기본 API를 정확히 이해하는 일이 중요합니다.
CUDA는 명시적 복사 방식, Unified Memory 방식, stream-ordered allocator 방식까지 여러 모델을 제공합니다. 각 모델은 코드 복잡도, 성능 예측 가능성, 페이지 이동 오버헤드, 재사용 효율이 다르므로 워크로드 성격에 맞춰 골라 써야 합니다.
CUDA 메모리 흐름
핵심 개념
cudaMalloc / cudaFree
cudaMalloc은 디바이스 메모리, 즉 GPU가 직접 접근할 수 있는 메모리를 할당합니다. 이 메모리는 커널에서 쓰는 텐서, 버퍼, 중간 결과의 기본 저장소가 되며, 수명 관리는 cudaFree가 담당합니다.
이 방식은 가장 단순하고 예측 가능하지만, 호스트와 디바이스 사이의 데이터 이동은 자동으로 일어나지 않습니다. 따라서 CPU 쪽 데이터가 필요하면 별도의 복사 API를 직접 호출해야 합니다.
cudaMemcpy / cudaMemcpyAsync
cudaMemcpy는 호스트와 디바이스 사이, 또는 디바이스 간 데이터 복사를 수행합니다. 방향은 cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice, cudaMemcpyDefault처럼 구분됩니다.
비동기 복사인 cudaMemcpyAsync는 스트림과 함께 사용하면 연산과 전송을 겹칠 수 있습니다. 다만 실제 겹침을 얻으려면 일반적으로 페이지 고정(host pinned) 메모리가 필요합니다.
cudaMallocHost / cudaHostAlloc
호스트에서 페이지 고정 메모리를 할당하면 DMA 전송 경로를 더 안정적으로 사용할 수 있습니다. cudaMallocHost와 cudaHostAlloc은 이런 pinned host memory를 만들고, GPU가 직접 접근 가능한 매핑 메모리의 기반이 됩니다.
Unified Addressing을 지원하는 환경에서는 일부 호스트 메모리가 동일한 포인터 값으로 접근되며, cudaHostGetDevicePointer 없이도 같은 주소를 쓸 수 있는 경우가 있습니다. 반대로 등록 메모리나 write-combined 메모리는 예외가 될 수 있습니다.
cudaHostRegister 와 mapped host memory
이미 애플리케이션이 확보한 호스트 버퍼를 그대로 CUDA 전송 경로에 올리고 싶을 때는 cudaHostRegister로 기존 메모리 구간을 page-locked 상태로 등록할 수 있습니다. 별도 복사 버퍼를 다시 만들지 않아도 되기 때문에, 네트워크 수신 버퍼나 외부 라이브러리가 소유한 메모리를 CUDA와 연결할 때 자주 쓰입니다.
이 버퍼를 mapped host memory로 열면 GPU가 cudaHostGetDevicePointer를 통해 직접 접근할 수 있어 zero-copy 경로를 구성할 수 있습니다. 다만 zero-copy는 PCIe/NVLink 뒤의 호스트 메모리를 직접 읽는 방식이라서, 대역폭과 지연 측면에서는 여전히 디바이스 메모리 상주 버퍼보다 불리한 경우가 많습니다.
cudaMallocManaged 와 Unified Memory
cudaMallocManaged는 Unified Memory 시스템이 자동으로 관리하는 메모리를 할당합니다. 프로그램은 하나의 포인터를 사용하지만, 실제 페이지는 CPU와 GPU 사이를 이동할 수 있습니다.
이 모델은 cudaMemPrefetchAsync로 미리 이동시키고, cudaMemAdvise로 접근 패턴을 힌트로 줄 수 있습니다. cudaPointerGetAttributes는 포인터가 host/device/managed 중 무엇인지 확인하는 데 사용됩니다.
cudaMallocAsync / cudaFreeAsync 와 memory pool
cudaMallocAsync와 cudaFreeAsync는 stream-ordered semantics를 따르는 할당자입니다. 할당과 해제가 스트림 순서 안에서만 유효하므로, 일반적인 cudaMalloc/cudaFree보다 재사용이 빠르고 동기화 부담이 적습니다.
이 API는 내부 memory pool에서 메모리를 가져오며, cudaMemPoolCreate, cudaMemPoolSetAttribute, cudaMemPoolTrimTo 같은 함수로 동작을 조정할 수 있습니다. 대형 반복 워크로드에서 할당 오버헤드를 줄이는 데 유리합니다.
비교/분석
| API 가족 | 용도 | 장점 | 주의점 |
|---|---|---|---|
cudaMalloc / cudaFree |
디바이스 메모리 직접 할당 | 단순하고 예측 가능함 | 호스트 데이터 이동을 직접 관리해야 함 |
cudaMemcpy / cudaMemcpyAsync |
호스트-디바이스 복사 | 전송 방향이 명확함, 스트림과 결합 가능 | 비동기 이득은 pinned host memory에 의존함 |
cudaMallocHost / cudaHostAlloc |
페이지 고정 호스트 메모리 | 전송과 중첩이 쉬움 | 호스트 메모리 사용이 무거워질 수 있음 |
cudaHostRegister / mapped host memory |
기존 호스트 버퍼 등록, zero-copy | 기존 버퍼를 재활용하고 직접 매핑 가능 | 직접 접근 지연이 커서 대용량 반복 접근에는 불리함 |
cudaMallocManaged |
Unified Memory | 코드 단순화, 페이지 단위 자동 이동 | 페이지 폴트와 migration 비용이 생김 |
cudaMallocAsync / cudaFreeAsync |
stream-ordered 할당 | 재사용 효율이 높고 할당 오버헤드가 낮음 | 스트림 순서 보장이 필요함 |
동작 원리
기본 모델에서는 CPU가 호스트 메모리를 채우고, cudaMalloc으로 확보한 디바이스 메모리로 cudaMemcpy를 통해 데이터를 옮깁니다. 이후 커널은 디바이스 메모리만 사용하므로, 데이터 배치와 복사 시점을 애플리케이션이 직접 결정합니다.
Pinned memory를 쓰면 전송 엔진이 더 직접적으로 동작할 수 있어서, 복사와 커널 실행을 스트림 단위로 겹치기 쉬워집니다. 이때 cudaMemcpyAsync는 계산과 전송의 파이프라이닝을 만드는 핵심 도구입니다.
cudaHostRegister로 등록한 mapped host memory는 복사 단계를 줄이고 GPU가 호스트 버퍼를 직접 읽게 만들 수 있습니다. 하지만 이 경로는 접근 지연과 링크 대역폭 제약을 그대로 받으므로, 제어 정보나 소량의 스트리밍 데이터에는 적합하지만 대형 텐서 본문에는 보통 device memory staging이 더 효율적입니다.
Unified Memory는 포인터 주소 공간을 단일화하고, 접근 시점에 필요한 페이지를 옮깁니다. 이 구조는 편하지만, 접근 패턴이 불규칙하면 page fault와 migration이 성능을 흔들 수 있습니다. 그래서 cudaMemPrefetchAsync와 cudaMemAdvise가 중요한 보조 수단이 됩니다.
Stream-ordered allocator는 할당 수명을 스트림 실행 순서에 묶어서 관리합니다. 같은 스트림 안에서 할당한 뒤 바로 사용하고, 마지막 사용 뒤에 해제하면 메모리 풀 재사용이 빨라집니다. 큰 반복 루프, 추론 서버, 그래프 캡처와 같이 할당 패턴이 반복되는 코드에 특히 잘 맞습니다.
장단점
장점
- 명시적 복사는 제어가 단순하고 성능 예측이 쉽습니다.
- Unified Memory는 코드 이식성과 작성 편의성이 좋습니다.
- Memory pool과 async alloc은 반복 할당 비용을 줄입니다.
- pinned memory와 비동기 복사는 전송-연산 중첩을 만들기 좋습니다.
단점
- 명시적 복사는 코드가 길어지고 관리 지점이 많아집니다.
- Unified Memory는 page fault와 migration 비용이 숨겨진 병목이 될 수 있습니다.
- pinned memory는 호스트 쪽 메모리 압박을 키울 수 있습니다.
- stream-ordered allocator는 순서 보장을 잘못 쓰면 디버깅이 어려워집니다.
관련 기술 / 참고 문헌
| 자료 | 링크 | 연결점 |
|---|---|---|
| CUDA Runtime API - Memory Management | https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html | cudaMalloc, cudaMemcpy, cudaHostAlloc, cudaMallocManaged |
| CUDA Runtime API - Stream Ordered Memory Allocator | https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY__POOLS.html | cudaMallocAsync, cudaFreeAsync, memory pool |
| CUDA Runtime API - Unified Addressing | https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__UNIFIED.html | cudaPointerGetAttributes, cudaMemcpyDefault |
| CUDA C++ Programming Guide - Unified Memory / Zero-Copy | https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html | mapped pinned memory, zero-copy, prefetch 운용 배경 |
| CUDA Programming Guide | https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html | Unified Memory, asynchronous execution, memory model |
| GPU 메모리 아키텍처 기초 | gpu_0005_gpu_memory_architecture_basics.html | HBM/GDDR 배경과 대역폭 관점 |
| vAttention Analysis | ../llm/llm_0020_vattention_analysis.html | CUDA VMM 기반 메모리 관리 예시 |
| KV Cache Offloading Analysis | ../llm/llm_0040_kv_cache_offloading_analysis.html | 메모리 티어링과 오프로딩 흐름 |
| Virtual Memory | ../system/system_0010_virtual_memory.html | 주소 공간, 매핑, 페이지 이동 배경 |
핵심 정리
CUDA 메모리 관리는 단순한 malloc 대체가 아니라, 호스트와 디바이스 사이의 데이터 이동 전략입니다. cudaMemcpy 계열은 가장 직접적이고, cudaMallocManaged는 가장 편리하며, cudaMallocAsync 계열은 반복 할당에 강합니다.
실전에서는 하나의 방식만 고집하기보다, 커널 데이터는 device memory에 두고 입력/출력은 pinned host memory로 옮기며, 큰 공유 버퍼는 Unified Memory나 memory pool로 정리하는 식의 혼합이 자주 쓰입니다. 워크로드의 접근 패턴과 동기화 비용이 어떤지 먼저 보는 것이 가장 중요합니다.
구현 예시
// 기본 디바이스 버퍼와 비동기 전송 예시
float *d_x = nullptr;
cudaMalloc(&d_x, bytes);
cudaMemcpyAsync(d_x, h_x, bytes, cudaMemcpyHostToDevice, stream);
// Unified Memory 예시
float *u_x = nullptr;
cudaMallocManaged(&u_x, bytes);
cudaMemPrefetchAsync(u_x, bytes, device, stream);
// Stream-ordered allocation 예시
float *p_x = nullptr;
cudaMallocAsync(&p_x, bytes, stream);
cudaFreeAsync(p_x, stream);