제2부 CUDA 프로그래밍

2.5 메모리 모델 II — 전송 최적화

2.5.1 이 장에서 배우는 것 / 선수 지식

CUDA 커널이 빠르게 실행된다고 프로그램이 빠른 것은 아니다. 많은 경우 실제 병목은 호스트(CPU)와 디바이스(GPU) 사이 데이터 이동(PCIe)이다. 이 장은 그 전송을 최적화한다 — 전송 시간 산수, pinned 메모리, 통합 메모리 튜닝, zero-copy. 복사와 연산을 겹치는 비동기 오버랩은 개념까지만, 실제 패턴은 2.10.

배우는 것:

• 호스트-디바이스 전송이 왜 PCIe에서 병목인지, PCIe 대역폭으로 전송 시간을 산수로 추정한다.
• 핀드(pinned) 메모리가 왜 빠르고 왜 비동기 복사의 전제인지.
• 통합 메모리를 cudaMemPrefetchAsync·cudaMemAdvise로 튜닝한다.
• 제로카피(zero-copy) 가 언제 이득이고 언제 독인지.

선수 지식: 1.9~1.10(HBM 대역폭·메모리 바운드), 2.1~2.3(커널·실행 구성), 2.4(메모리 공간). 비동기 오버랩은 2.10.

정직성: PCIe 대역폭·통합 메모리 동작은 GPU 세대·PCIe·드라이버·OS 의존 — 시점·플랫폼 확인 필요.

2.5.2 왜 데이터 이동이 문제인가

GPU 내부 대역폭(HBM·온칩)은 TB/s 단위인데, GPU 밖에서 안으로 데이터를 넣는 경로(PCIe)는 한~두 자릿수 느리다.

• 데이터를 한 번 적재해 두고 GPU에서 여러 번 재사용하면 PCIe 비용이 분할상환된다.
• 매 연산마다 호스트로 내렸다가 다시 적재하면 PCIe가 병목 — 초보 CUDA 코드가 CPU보다 느린 가장 흔한 이유다.

루프라인(1.10) 관점: 연산량 대비 PCIe 전송량이 낮으면, 커널을 아무리 최적화해도 PCIe가 천장이다.

2.5.2.1 전송 시간을 산수로

PCIe 세대별 이론 단방향 대역폭(x16, 시점 의존·확인 필요):

예 1 — $4096\times4096$ FP32 = $4096\cdot4096\cdot4 \approx67.1 \,\mathrm{MB}$. Gen4(실측 ~25 GB/s 가정): $67.1 \mathrm{MB} / 25 \,\mathrm{GB/s} \approx2.7 \,\mathrm{ms}$.

예 2 — 가치 판단. 같은 행렬로 GEMM($2N^3\approx1.37\mathrm{e}11 \,\mathrm{FLOP}$)을 50 TFLOPS에 실행하면 $\approx2.7 \,\mathrm{ms}$. 입력 두 장 전송(~5.4 ms) 대비 연산도 비슷한 자릿수 — 데이터를 한 번 적재해 여러 연산을 묶지 않으면 전송이 연산만큼 비용이 크다.

예 3 — 비효율 패턴. 벡터 덧셈(N=10⁷, FP32). 전송 $3\cdot40 \mathrm{MB} = 120 \,\mathrm{MB} \to\sim4.8 \,\mathrm{ms}$. 연산 $10^7 \mathrm{FLOP} / 5\mathrm{e}13 \approx0.2 \,\mu\mathrm{s}$. 전송이 연산보다 2만 배 비용이 크다. 이런 연산은 GPU에 단독으로 이전할 가치가 없다 — 더 큰 파이프라인의 일부로 데이터가 이미 GPU에 있을 때만 의미.

2.5.3 호스트-디바이스 전송 — pageable vs pinned

2.5.3.1 pageable의 숨은 복사

malloc/new로 잡은 메모리는 pageable — OS가 언제든 디스크로 스왑할 수 있다. GPU DMA 엔진은 물리 주소가 고정된 메모리만 안전하게 읽으므로, 드라이버는 staging 복사를 삽입한다:

추가 호스트 복사 한 번이 개입한다. 그래서 pageable 전송은 (1) 느리고, (2) staging이 끝나야 DMA가 시작되므로 실질적 비동기가 불가능 — cudaMemcpyAsync를 pageable 버퍼에 써도 동기적으로 동작한다.

2.5.3.2 pinned(page-locked) 메모리

cudaMallocHost(또는 cudaHostAlloc)로 잡은 메모리는 OS가 스왑하지 못하게 고정된다. DMA 엔진이 직접 읽으니 staging 복사가 사라진다:

효과: 더 높은 PCIe 대역폭(staging 제거 → 이론치에 근접), 실질적 비동기 가능(복사와 커널·다른 복사를 겹침, 2.10). 단 물리 메모리를 고정하므로 과하게 잡으면 시스템 전체 저하, 할당이 malloc보다 느리니 재사용한다.

float* h_pinned;
cudaMallocHost(&h_pinned, bytes);   // page-locked
// ... 사용 ...
cudaFreeHost(h_pinned);

2.5.4 통합 메모리 튜닝 — prefetch와 advise

cudaMallocManaged 포인터는 호스트·디바이스가 같은 주소로 접근하고, 런타임이 페이지를 자동 마이그레이션한다(2.2). 장점은 생산성, 단점은 페이지 폴트 마이그레이션이 보이지 않는 비용 — 잘못 쓰면 PCIe를 오가며 스래싱한다. 기본 설정으로 쓰면 편하지만 빠르게 쓰려면 힌트를 줘야 한다.

• cudaMemPrefetchAsync(ptr, bytes, dstDevice, stream) — 미리 옮긴다. 커널이 폴트로 한 장씩 끌어오는 대신 큰 단위를 한 번에 마이그레이션 → 폴트 제거. dstDevice에 GPU id 또는 cudaCpuDeviceId.
• cudaMemAdvise(ptr, bytes, advice, device) — 접근 패턴 힌트(데이터 자체를 옮기지는 않음).
  • SetReadMostly — 여러 프로세서가 읽기만 → 마이그레이션 대신 복제.
  • SetPreferredLocation — 선호 위치 지정 → 다른 쪽이 접근해도 가급적 거기 둠.
  • SetAccessedBy — 자주 접근하는 프로세서에 매핑을 만들어 폴트 회피.

통합 메모리 동작·성능은 GPU 아키텍처·OS·드라이버에 강하게 의존(페이지 폴트 마이그레이션 지원 여부, Grace-Hopper 같은 NVLink-C2C 일관성 시스템에서의 동작). 시점·플랫폼 의존, 확인 필요.

2.5.5 제로카피(zero-copy / mapped pinned)

cudaHostAlloc(&p, bytes, cudaHostAllocMapped)로 잡은 pinned 메모리를 디바이스 포인터로 매핑하면, 커널이 호스트 메모리를 PCIe 너머로 직접 읽고 쓴다(복사 없음).

• 이득: 정확히 한 번 접근하고 마는 데이터, 또는 데이터의 극히 일부만 GPU가 보는 경우.
• 독: 커널이 반복 접근하면 매 접근이 PCIe를 타므로 글로벌(HBM)보다 자릿수 느리다. "복사를 없앴다"가 곧 "빠르다"가 아니다.
• 적합처: 통합 메모리가 없거나 부적합한 환경, 통합 칩(Jetson, 9.1)처럼 CPU·GPU가 물리 메모리를 공유해 PCIe가 없는 경우.

2.5.6 워크스루 — pageable vs pinned 대역폭

// bw.cu — 빌드: nvcc -O2 -o bw bw.cu
#include <cstdio>
#include <cuda_runtime.h>
#define CK(x) do{cudaError_t e=(x); if(e){printf("err %s:%d: %s\n",__FILE__,__LINE__,cudaGetErrorString(e));return 1;}}while(0)
int main(){
  const size_t N = 256*1024*1024; const int iters = 20;
  void *d; CK(cudaMalloc(&d, N));
  void *h_page = malloc(N);                 // pageable
  void *h_pin;  CK(cudaMallocHost(&h_pin, N));  // pinned
  cudaEvent_t a,b; CK(cudaEventCreate(&a)); CK(cudaEventCreate(&b));
  auto bench = [&](void* src, const char* tag){
    CK(cudaMemcpy(d,src,N,cudaMemcpyHostToDevice));         // warmup
    CK(cudaEventRecord(a));
    for(int i=0;i<iters;i++) CK(cudaMemcpy(d,src,N,cudaMemcpyHostToDevice));
    CK(cudaEventRecord(b)); CK(cudaEventSynchronize(b));
    float ms=0; CK(cudaEventElapsedTime(&ms,a,b));
    printf("%-8s  %.2f ms/iter   %.1f GB/s\n", tag, ms/iters, (double)N*iters/(ms/1e3)/1e9);
    return 0;   // CK는 오류 시 1을 반환 → 람다 반환형(int)과 일치시킴
  };
  bench(h_page,"pageable"); bench(h_pin,"pinned");
  cudaFree(d); free(h_page); cudaFreeHost(h_pin);
  return 0;
}

예상(개략 — 절대 수치는 GPU·PCIe 세대·메인보드 의존):

pageable  ~14.x ms/iter   ~18 GB/s
pinned    ~10.x ms/iter   ~25 GB/s

읽는 법: pinned가 pageable보다 빠른 비율(여기선 ~1.4×)이 핵심이고, 절대값은 PCIe 세대 이론 대역폭에 얼마나 근접하는지로 본다(2.5.2.1과 대조). pinned가 pageable과 거의 같다면 PCIe가 이미 다른 이유로 막혔거나 통합 칩(staging 경로 없음)을 의심.

통합 메모리 + prefetch도 같은 방식으로 측정한다(Nsight Systems, 2.13): prefetch 없는 커널은 타임라인에 다수의 page fault/migration이 보이고, prefetch 있으면 마이그레이션이 커널 전 하나로 빠진다. 단 PCIe를 건너는 총량 자체는 prefetch로 줄지 않는다 — 줄이려면 데이터를 GPU에 두고 재사용(2.5.2).

2.5.7 트레이드오프 — 무엇을 어떻게 옮길까

견해: 기본은 ②. 성능 경로는 pinned + async로 가고 2.10의 스트림 오버랩으로 복사를 연산 뒤에 숨긴다. 개발 속도·복잡 자료구조가 우선이면 ③으로 시작해 프로파일러가 폴트를 보여주면 prefetch/advise로 튜닝. ①은 프로토타입까지, ④는 "왜 복사를 없애도 되는가"를 명확히 답할 수 있을 때만.

비엔비디아 정직: 통합·핀드 메모리의 대응물은 다른 가속기에도 있다(host-pinned/DMA, SoC 공유 메모리, 통합 가상 주소 공간). API는 다르지만 "PCIe(또는 그에 준하는 인터커넥트)를 건너는 비용을 분할상환하고, 고정 메모리로 DMA를 가능케 한다" 원리는 공통. NVLink-C2C로 CPU·GPU를 일관성 있게 묶는 시스템(4.8)에서는 이 장의 PCIe 병목 전제 자체가 약해진다 — 플랫폼을 먼저 확인.

2.5.8 중간 점검 (체크포인트)

• [ ] 버퍼 크기·PCIe 세대로 H2D 전송 시간 하한을 산수하고, 같은 데이터의 연산 시간과 비교해 "이전할 가치"를 판단한다.
• [ ] malloc→cudaMallocHost가 왜 빨라지는지(staging 제거)와 왜 비동기의 전제인지 설명한다.
• [ ] 통합 메모리의 페이지 폴트 스래싱 원인을 짚고 prefetch/advise로 고친다.
• [ ] 제로카피가 1회·부분 접근에서만 이득임을 안다.

증상: pinned가 pageable과 차이 없음 → PCIe가 다른 곳에서 막혔거나 통합 칩. prefetch를 넣어도 느림 → 매 반복 CPU가 데이터를 다시 만져 마이그레이션 반복(스래싱), 소유권을 한쪽으로 고정했는지 확인.

2.5.9 흔한 함정

• 과도한 H2D/D2H 전송. 루프 안 매 반복 복사 → PCIe가 전부 점유한다. 데이터를 한 번 적재하고 GPU에서 끝까지 처리하도록 재구성. "CPU보다 느린 GPU 코드"의 1순위.
• cudaMemcpyAsync를 pageable 버퍼에. async라 써도 pageable이면 동기로 떨어진다. 실질적 비동기는 pinned + non-default 스트림(2.10).
• 통합 메모리 오남용. "복사가 사라졌다"를 "비용이 사라졌다"로 착각. 폴트 마이그레이션은 보이지 않을 뿐 PCIe를 탄다. prefetch/advise로 소유권을 한쪽으로 고정.
• pinned 과다 할당. page-lock은 물리 메모리를 묶는다. 거대하게 잡으면 시스템 전체가 느려진다. 버퍼 재사용.
• 측정 함정. 비동기 호출 뒤 동기화 없이 시간을 재면 전송이 안 끝난 채 측정. 대역폭 측정은 반드시 동기화 후.

2.5.10 연습문제

연습 1 (전송 시간 산수). $8192\times8192$ FP16 두 장을 Gen5(실측 ~50 GB/s) PCIe로. (a) 총 바이트·전송 시간 하한. (b) 같은 두 행렬로 GEMM($2N^3$)을 200 TFLOPS에 실행할 때 연산 시간과 비교, "이전할 가치"를.

연습 2 (pageable→pinned). malloc만 cudaMallocHost로 바꾸면 대역폭이 오르는 이유를 staging 관점에서. 왜 pinned가 "비동기의 전제"인가?

연습 3 (통합 메모리 스래싱). 통합 메모리 반복 커널이 매 iteration마다 CPU에서 결과를 읽어 다음 입력으로 쓴다. 왜 느린가? 어떻게 완화?

연습 4 (제로카피 판단). 4 GB 호스트 배열 중 커널이 무작위 100개만 1회 읽는다. 전체 복사 vs 제로카피? 그 100개를 1000번 반복 읽으면?

2.5.11 요약 · 다음 장

• 실제 병목은 흔히 커널이 아니라 PCIe를 건너는 데이터 이동이다. 전송 시간을 산수하고 연산 시간과 비교해 "이전할 가치"를 먼저 판단하라.
• pinned 메모리는 staging을 없애 대역폭을 높이고 비동기 복사의 전제다. 통합 메모리는 코드를 단순화하되 폴트 마이그레이션을 prefetch·advise로 다스려야 한다. 제로카피는 1회·부분 접근이나 통합 칩에서만 이득.
• 절대 수치는 GPU 세대·PCIe·드라이버·OS 의존 — 시점·플랫폼 확인 필요.

다음 장: 2.6 공유 메모리 I — 뱅크·충돌·동기화. 블록 범위 온칩 스크래치패드를 직접 채우고 비우는 법, __syncthreads()의 정확한 위치, 그리고 뱅크 충돌을 타일링 맥락에서 다룬다.