GPU Memory Model
register, shared memory, L1/L2 cache, global memory, HBM, pinned memory, unified memory 읽기
GPU 성능은 연산 유닛의 개수만으로 정해지지 않는다. 데이터가 GPU memory에 있는지, SM 가까이에서 재사용되는지, CPU와 GPU 사이를 얼마나 자주 오가는지가 처리량을 크게 바꾼다.
GPU는 많은 thread를 동시에 실행한다. 그만큼 많은 데이터가 동시에 필요하다. 연산 유닛이 아무리 많아도 memory가 데이터를 제때 공급하지 못하면 SM은 기다린다.
{
"diagram": "html-diagram",
"variant": "explainer",
"title": "GPU memory map",
"width": 960,
"height": 500,
"mobileWidth": 920,
"regions": [
{
"id": "host-region",
"label": "1. Host memory\\nCPU RAM",
"x": 36,
"y": 51,
"width": 240,
"height": 288
},
{
"id": "device-region",
"label": "2. Device memory\\nGPU side",
"x": 416,
"y": 51,
"width": 508,
"height": 288
}
],
"nodes": [
{
"id": "pageable",
"kind": "host",
"label": "Pageable\\nmemory",
"caption": "일반 CPU allocation",
"x": 156,
"y": 167,
"width": 168,
"height": 68
},
{
"id": "pinned",
"kind": "host",
"label": "Pinned\\nmemory",
"caption": "page-locked host buffer",
"x": 156,
"y": 283,
"width": 168,
"height": 68
},
{
"id": "transfer",
"kind": "nic",
"label": "H2D copy",
"caption": "cudaMemcpy / to(\"cuda\")",
"x": 336,
"y": 225,
"width": 144,
"height": 76
},
{
"id": "global-in",
"kind": "memory",
"label": "Global memory\\nHBM / VRAM",
"caption": "input, weight, activation",
"x": 506,
"y": 167,
"width": 150,
"height": 72
},
{
"id": "cache",
"kind": "note",
"label": "L2 / L1\\ncache",
"caption": "hardware cache path",
"x": 678,
"y": 167,
"width": 130,
"height": 64
},
{
"id": "sm",
"kind": "gpu",
"label": "SM working\\nset",
"caption": "register, shared memory",
"x": 844,
"y": 167,
"width": 130,
"height": 74
},
{
"id": "global-out",
"kind": "memory",
"label": "Global memory\\noutput",
"caption": "D2H는 필요할 때만",
"x": 554,
"y": 283,
"width": 150,
"height": 66
},
{
"id": "spill",
"kind": "switch",
"label": "Local spill",
"caption": "device DRAM backing",
"x": 816,
"y": 283,
"width": 150,
"height": 66
},
{
"id": "managed",
"kind": "note",
"label": "Unified Memory",
"caption": "하나의 pointer처럼 보이지만 page migration은 runtime이 처리",
"x": 480,
"y": 405,
"width": 320,
"height": 70
}
],
"links": [
{
"id": "pageable-transfer",
"path": "M 240 167 C 270 167 270 195 264 200",
"tone": "muted",
"flow": {
"speed": "slow",
"emphasis": "muted"
}
},
{
"id": "pinned-transfer",
"path": "M 240 283 C 272 283 272 258 264 252",
"tone": "primary",
"width": 2.6,
"flow": {
"speed": "normal",
"emphasis": "strong"
}
},
{
"id": "transfer-global",
"path": "M 408 220 C 430 210 430 178 431 167",
"tone": "primary",
"width": 2.6,
"flow": {
"speed": "fast",
"count": 2,
"emphasis": "strong"
}
},
{
"id": "global-cache",
"points": [[581, 167], [613, 167]],
"tone": "primary",
"width": 2.6,
"flow": {
"speed": "fast",
"emphasis": "strong"
}
},
{
"id": "cache-sm",
"points": [[743, 167], [779, 167]],
"tone": "primary",
"width": 2.6,
"flow": {
"speed": "fast",
"emphasis": "strong"
}
},
{
"id": "sm-output",
"path": "M 844 204 C 812 244 706 274 630 281",
"tone": "primary",
"flow": {
"speed": "normal",
"emphasis": "soft"
}
},
{
"id": "sm-spill",
"path": "M 852 204 C 844 224 829 240 816 250",
"tone": "warning",
"dashed": true,
"flow": false
},
{
"id": "managed-host",
"path": "M 360 388 C 260 360 190 335 156 317",
"tone": "muted",
"dashed": true,
"opacity": 0.48,
"flow": false
},
{
"id": "managed-device",
"path": "M 602 388 C 590 330 538 250 506 205",
"tone": "muted",
"dashed": true,
"opacity": 0.48,
"flow": false
}
],
"callouts": [
{
"id": "host-note",
"tone": "muted",
"anchor": "top",
"x": 156,
"y": 352,
"width": 198,
"text": "둘 다 GPU 안의 memory 계층이 아니다."
},
{
"id": "managed-note",
"tone": "primary",
"anchor": "bottom",
"x": 770,
"y": 432,
"width": 236,
"text": "Managed model이지 별도의 빠른 hardware memory가 아니다."
}
]
}Memory Map
GPU memory를 계층처럼 그릴 때 먼저 분리해야 할 것이 있다. Register, shared memory, cache, global memory는 CUDA kernel이 device 쪽에서 만나는 memory space와 hardware resource에 가깝다. 반면 pageable memory와 pinned memory는 CPU RAM을 OS가 어떻게 다룰 수 있는지를 나타내는 host memory 상태다. Unified Memory는 별도의 hardware 층이라기보다 CUDA가 CPU와 GPU 사이의 memory 접근을 관리해주는 추상화다.
이 글의 그림은 가장 흔한 실행 경로를 기준으로 그렸다.
- CPU RAM의 pageable memory나 pinned memory에서 시작한다.
cudaMemcpy나tensor.to("cuda")같은 host-device transfer로 GPU global memory에 데이터가 놓인다.- Kernel이 global memory를 읽으면 cache path를 거쳐 SM working set으로 들어간다.
- Thread가 바로 쓰는 값은 register에, block 안에서 재사용하는 tile은 shared memory에 놓인다.
- 결과는 다시 global memory에 저장되고, CPU가 필요할 때만 host로 복사된다.
이 구분을 해두면 “어디에 저장되는가”와 “언제 이동하는가”를 섞어서 보지 않게 된다. Register와 shared memory는 SM 가까이에 있는 자원이다. Global memory는 CUDA programming model의 device memory space이고, 실제로는 GPU의 HBM이나 VRAM 같은 off-chip DRAM이 받친다. Pinned memory는 GPU 안에 있는 더 빠른 memory가 아니라, GPU와 데이터를 주고받기 좋게 page-locked로 잡아 둔 host memory다.
Host and Device
CUDA에서 memory를 볼 때 첫 구분은 host와 device다.
Host는 CPU 쪽이다. 일반적인 process memory, 즉 CPU RAM이 여기에 있다.
Device는 GPU 쪽이다. GPU가 kernel을 실행하면서 직접 읽고 쓰는 memory가 여기에 있다. 흔히 VRAM이라고 부르고, 데이터센터 GPU에서는 HBM을 쓰는 경우가 많다.
CUDA C++ 코드에서는 host memory와 device memory 사이를 명시적으로 오가는 흐름이 자주 나온다.
float* x_gpu;
cudaMalloc(&x_gpu, size);
cudaMemcpy(x_gpu, x_cpu, size, cudaMemcpyHostToDevice);
kernel<<<blocks, threads>>>(x_gpu);
cudaMemcpy(x_cpu, x_gpu, size, cudaMemcpyDeviceToHost);cudaMalloc은 device memory를 잡는다. cudaMemcpyHostToDevice는 CPU RAM의 데이터를 GPU memory로 복사한다. kernel은 GPU memory에 있는 데이터를 읽고 쓴다.
PyTorch에서는 이 흐름이 더 짧게 보인다.
import torch
x = torch.randn(1024, 1024)
x_gpu = x.to("cuda")
y_gpu = x_gpu * 2처음 만든 x는 CPU tensor다. x.to("cuda")를 호출하면 GPU memory에 같은 값을 담은 CUDA tensor가 생긴다. x_gpu * 2는 그 CUDA tensor를 입력으로 받아 GPU 쪽에서 실행된다.
Global Memory
CUDA 문서에서 global memory는 GPU의 큰 device memory 공간을 가리킨다. 모델 weight, input tensor, activation, output 같은 큰 데이터가 여기에 놓인다.
Global memory는 용량이 크고 모든 thread block에서 접근할 수 있다. 대신 SM 안의 register나 shared memory보다 멀다. 그래서 global memory 접근이 많고 재사용이 적으면 GPU가 계산보다 memory 접근을 기다리는 시간이 길어진다.
HBM은 global memory를 구성하는 실제 memory 기술로 볼 수 있다. High Bandwidth Memory라는 이름처럼 넓은 대역폭을 제공하지만, SM 바로 옆의 register나 shared memory와 같은 수준으로 가까운 공간은 아니다. 그래서 global memory는 CUDA가 보여주는 주소 공간의 이름이고, HBM이나 VRAM은 그 공간을 실제로 받치는 물리 memory라고 보는 편이 정확하다.
GPU code를 읽을 때 global memory와 HBM을 완전히 같은 층의 말로 보면 헷갈린다.
global memory: CUDA programming model에서 보는 주소 공간
HBM / VRAM: 그 주소 공간을 받치는 물리 memoryRegisters
Register는 CUDA thread가 쓰는 가장 가까운 저장 공간이다.
__global__ void add(float* x, float* y, float* out) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
float a = x[i];
float b = y[i];
out[i] = a + b;
}i, a, b 같은 값은 보통 compiler가 register에 배치한다. register는 빠르지만 thread마다 따로 잡히는 자원이다. 한 thread가 register를 많이 쓰면 SM에 동시에 올릴 수 있는 thread나 block 수가 줄어들 수 있다.
CUDA 최적화에서 occupancy 이야기가 나오는 이유도 여기에 있다. register 사용량, shared memory 사용량, thread block 크기가 함께 SM에 올라갈 수 있는 작업량을 결정한다.
Local Memory
이름 때문에 local memory를 register처럼 가까운 공간으로 오해하기 쉽다. CUDA에서 local memory는 thread마다 private하게 보이는 공간이지만, 물리적으로는 device memory 쪽에 놓일 수 있다.
Compiler가 register에 담기 어려운 값을 local memory로 내릴 수 있다. 예를 들어 큰 local array나 register가 부족한 상황이 여기에 걸린다.
__global__ void work(float* out) {
float tmp[128];
int j = threadIdx.x & 127;
tmp[j] = j * 2.0f;
out[threadIdx.x] = tmp[j];
}이런 코드가 항상 local memory로 간다는 뜻은 아니다. compiler와 architecture, access pattern에 따라 달라진다. 다만 “thread local variable이니까 항상 register처럼 싸다”라고 보면 안 된다.
Shared Memory
Shared memory는 같은 thread block 안의 thread들이 함께 쓰는 on-chip memory다.
Global memory에서 읽은 데이터를 shared memory에 올려두고 여러 thread가 재사용하면 global memory 접근을 줄일 수 있다. matrix multiplication에서 tile을 shared memory에 담는 패턴이 대표적이다.
__shared__ float tile[256];
tile[threadIdx.x] = x[i];
__syncthreads();
int next = (threadIdx.x + 1) % blockDim.x;
out[i] = tile[next] * 2;이 코드에서는 block 안의 thread들이 tile을 공유한다. 어떤 thread가 쓴 값을 다른 thread가 읽기 전에 __syncthreads()로 block 안의 실행 지점을 맞춘다.
Shared memory는 빠르지만 block 범위를 넘지 않는다. block A의 shared memory를 block B가 직접 읽을 수 없다.
Constant and Texture
CUDA에는 constant memory와 texture memory도 있다.
Constant memory는 kernel 실행 중 값이 바뀌지 않는 데이터를 읽는 데 쓰인다. 같은 warp의 thread들이 같은 constant 값을 읽는 경우 cache 효율이 좋다. 작은 coefficient, configuration value처럼 모든 thread가 공통으로 참조하는 값에 맞는다.
Texture memory는 원래 graphics workload에서 온 읽기 전용 memory path다. CUDA에서는 특정 access pattern에서 cache locality를 활용하는 용도로 볼 수 있다. 일반적인 딥러닝 코드를 읽을 때 매번 직접 마주치는 공간은 아니지만, CUDA memory space에는 이런 특수한 읽기 경로도 있다.
L1 and L2 Cache
GPU에도 cache가 있다. CPU cache와 이름은 같지만, 여기서는 GPU 안의 cache를 말한다.
L1 cache는 SM 쪽에 가깝다. L2 cache는 GPU 전체에서 공유되는 cache로 보면 된다. CUDA Best Practices Guide는 L2 cache가 global memory 접근보다 낮은 latency와 높은 bandwidth를 제공한다고 설명한다.
프로그래머가 모든 cache 동작을 직접 제어하는 것은 아니다. 하지만 access pattern은 cache 효율에 큰 영향을 준다.
같은 데이터를 여러 번 읽거나, 가까운 주소를 연속적으로 읽거나, warp 안의 thread들이 이어진 주소를 읽으면 memory system이 더 잘 처리할 수 있다.
Coalescing
memory coalescing은 warp 안의 thread들이 memory를 읽고 쓸 때 중요하다.
Warp 하나는 보통 32개 thread다. 이 thread들이 global memory의 이어진 주소를 읽으면 hardware가 더 적은 memory transaction으로 묶어 처리할 수 있다.
good:
thread 0 -> x[0]
thread 1 -> x[1]
thread 2 -> x[2]
...
bad:
thread 0 -> x[0]
thread 1 -> x[1024]
thread 2 -> x[2048]
...두 경우 모두 thread 수는 같다. 차이는 memory access pattern이다. GPU는 많은 thread를 동시에 실행하므로, 각 thread가 어디를 읽는지가 성능에 직접 드러난다.
PyTorch나 cuBLAS 같은 library를 쓸 때도 tensor layout, stride, contiguous 여부가 성능에 영향을 주는 이유가 여기에 있다.
Pinned Memory
CPU RAM에서 GPU memory로 데이터를 복사할 때 host memory가 항상 같은 성능을 내지는 않는다.
일반적인 CPU allocation은 pageable memory다. Operating system이 page를 옮기거나 swap할 수 있다. GPU DMA가 안정적으로 접근하려면 중간 staging이 필요할 수 있다.
Pinned memory, 또는 page-locked memory는 OS가 해당 page를 옮기지 못하게 고정한 host memory다. CUDA는 pinned memory를 host-device transfer에 더 효율적으로 쓸 수 있다.
float* x_host;
cudaMallocHost(&x_host, size);
float* x_gpu;
cudaMalloc(&x_gpu, size);
cudaMemcpy(x_gpu, x_host, size, cudaMemcpyHostToDevice);Pinned memory는 transfer 성능에 유리하지만 무한정 쓰면 안 된다. Host memory를 page-locked로 잡으면 OS가 memory를 관리하는 자유도가 줄어든다. Data loader나 input pipeline에서 필요한 범위로 제한해서 쓰는 편이 좋다.
중요한 점은 pinned memory가 GPU 안의 새로운 계층이 아니라는 것이다. 위치는 여전히 CPU RAM이다. 다만 page가 고정되어 있으므로 GPU와의 DMA transfer나 asynchronous copy에 더 적합한 host buffer가 된다.
Unified Memory
Unified Memory는 host와 device가 함께 접근할 수 있는 managed memory를 제공한다. CUDA에서는 cudaMallocManaged로 할당한다.
int* data;
cudaMallocManaged(&data, n * sizeof(int));
kernel<<<blocks, threads>>>(data);
cudaDeviceSynchronize();
printf("%d\n", data[0]);이 코드는 명시적인 cudaMemcpy 없이 CPU와 GPU가 같은 pointer를 쓰는 것처럼 보인다. 사용하기는 편하지만 데이터 이동이 사라지는 것은 아니다. 필요하면 page migration이나 synchronization이 발생한다.
Unified Memory는 memory 관리 코드를 단순하게 만들 수 있다. 하지만 성능을 예측해야 하는 상황에서는 page가 어느 쪽에 있고, 언제 이동하는지 봐야 한다.
따라서 Unified Memory도 register, shared memory, global memory와 같은 물리 계층으로 두면 헷갈린다. 더 정확하게는 CUDA가 제공하는 managed memory 모델이다. 개발자는 하나의 pointer처럼 다루지만, runtime과 driver는 page migration, coherence, prefetch 같은 일을 통해 host와 device 접근을 맞춘다.
Frameworks
PyTorch 같은 framework를 쓰면 register, shared memory, coalescing을 매번 직접 제어하지 않는다. 그래도 GPU memory model을 알면 성능 문제를 읽는 데 도움이 된다.
예를 들어 다음 현상들은 모두 memory model과 연결된다.
nvidia-smi에서 memory usage는 높은데 GPU utilization은 낮다.to("cuda")가 training loop 안에 반복해서 나온다.- tensor가 non-contiguous라 연산 전에 copy가 생긴다.
- batch size를 키우면 throughput이 오르다가 어느 순간 memory 부족이 난다.
- DataLoader가 느려 GPU가 다음 batch를 기다린다.
GPU가 느리게 보일 때 항상 연산 유닛이 부족한 것은 아니다. 데이터가 CPU RAM에 있거나, device memory에 있어도 access pattern이 나쁘거나, global memory에서 재사용 없이 계속 읽고 있을 수 있다.
정리
Register와 shared memory는 SM 가까이에 있는 on-chip 자원이다. L1과 L2 cache도 GPU 안의 memory hierarchy에 들어간다. Global memory는 CUDA device memory space이고, 실제 저장 매체는 GPU의 HBM이나 VRAM이다.
Pageable memory와 pinned memory는 GPU 쪽 계층이 아니다. 둘 다 CPU RAM에 있는 host memory이고, 차이는 OS가 page를 옮길 수 있는지에 있다. Pinned memory는 page-locked host buffer라서 host-device transfer에 유리하지만, 많이 잡으면 OS memory 관리에 부담을 준다.
Unified Memory는 별도의 빠른 hardware memory가 아니다. CPU와 GPU가 함께 접근할 수 있는 managed memory 모델이고, 편의성을 얻는 대신 page migration과 synchronization 비용을 의식해야 한다.
성능 문제를 볼 때는 이 분류 위에서 접근 pattern을 따로 본다. Warp 안의 thread들이 이어진 주소를 읽는지, shared memory에 올린 데이터를 재사용하는지, CPU와 GPU 사이의 transfer가 training loop 안에서 반복되는지 같은 질문이 여기에 들어간다.
References
- NVIDIA — CUDA C++ Best Practices Guide: Device Memory Spaces https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#device-memory-spaces
- NVIDIA — CUDA C++ Programming Guide: Understanding Memory https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/understanding-memory.html
- NVIDIA — CUDA C++ Programming Guide: Unified Memory https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-unified-memory-programming-hd
- NVIDIA — CUDA Runtime API: Memory Management https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html
- NVIDIA — CUDA C++ Best Practices Guide: Data Transfer Between Host and Device https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#data-transfer-between-host-and-device