제2부 CUDA 프로그래밍

2.4 메모리 모델 I — 종류

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

CUDA에서 한 변수·버퍼는 여러 "메모리 공간(memory space)" 중 하나에 산다. 공간마다 누가 보는가(범위)·얼마나 사는가(수명)·어디에 캐시되는가·쓰기 가능한가가 다르다. 이 장은 그 공간들을 분류하고 "어떤 데이터를 어디에 둘지" 판단하는 토대를 세운다. 호스트-디바이스 전송 최적화는 2.5, 공유 메모리 활용(타일링)은 2.6~2.7로 잇는다.

배우는 것:

• 글로벌·상수·텍스처/표면·로컬·레지스터·공유 공간을 범위·수명·캐시 경로·용도로 구분한다.
• 어떤 읽기 전용 데이터를 상수(작고·워프 공통)에, 어떤 것을 텍스처(2D 지역성·보간)에 둘지 판단한다.
• "로컬 메모리"가 이름과 달리 느린 HBM(스필·동적 인덱스)임을 안다.

선수 지식: 1.7~1.8(온칩 메모리·병합·L2), 1.10(메모리 바운드), 2.1~2.3(커널·실행 구성). 공유 메모리의 실전 활용은 2.6~2.7.

정직성: 상수 메모리 한도·캐시 동작은 세대 의존(확인 필요).

2.4.2 CUDA 메모리 공간 분류

2.4.2.1 레지스터와 로컬 메모리

• 레지스터: 가장 빠름. 스레드 단위 범위, 스레드 수명. SM 레지스터 파일(1.4)에 있고 점유율(1.6)에 직접 영향. 스칼라 지역 변수는 기본 레지스터.
• 로컬 메모리: 이름과 달리 물리적으로 글로벌과 같은 HBM에 있고 L1/L2로 캐시된다. 스레드 단위 범위지만 느리다. 컴파일러가 (a) 레지스터가 부족해 스필하거나, (b) 동적 인덱스로 접근하는 지역 배열(컴파일 타임에 인덱스를 모르면 레지스터에 못 둠)을 만나면 로컬로 떨어진다. "왜 내 커널이 느리지?"의 흔한 범인. nvcc --ptxas-options=-v로 스필 바이트 확인.

2.4.2.2 글로벌 메모리

• 가장 일반적인 R/W 공간. 그리드 전체 범위(모든 스레드가 봄), 앱이 해제할 때까지 수명. 물리적으로 디바이스 HBM, L2(전 SM 공유) + L1/텍스처 캐시 경유.
• cudaMalloc으로 할당하고 cudaMemcpy로 채운다. 접근 성능의 핵심은 병합(coalescing)(1.8).

2.4.2.3 상수 메모리

• 디바이스에서 읽기 전용, 그리드 전체 범위. 총 64 KB(전용 상수 메모리, 아키텍처 의존·확인 필요), 그중 커널이 쓰는 부분이 상수 캐시(SM당 약 8 KB)에 올라간다.
• 강점은 브로드캐스트: 워프 내 모든 스레드가 같은 주소를 읽을 때 한 번 읽어 32스레드에 뿌린다 — 글로벌보다 효율적. 반대로 워프 내 스레드가 서로 다른 주소를 읽으면 직렬화되어 손해.
• 용도: 모든 스레드가 공유하는 작은 읽기 전용 파라미터(필터 계수·변환 행렬·하이퍼파라미터). 선언 __constant__ float coeff[256];, 채우기 cudaMemcpyToSymbol(coeff, host_ptr, bytes);.

__constant__ float K[5];     // 디바이스 상수
// 호스트: float k[5]={...}; cudaMemcpyToSymbol(K, k, sizeof(k));
__global__ void conv1d(const float* x, float* y, int n){
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if(i<2 || i>=n-2) return;
  // K[0..4]: 워프 내 모든 스레드가 같은 주소 → 상수 캐시 브로드캐스트
  y[i]= K[0]*x[i-2]+K[1]*x[i-1]+K[2]*x[i]+K[3]*x[i+1]+K[4]*x[i+2];
}

K[t] 인덱스 t가 워프 전체에서 동일하므로 브로드캐스트가 성립. 인덱스가 스레드마다 달랐다면 상수 메모리는 직렬화되어 손해 — 이때는 글로벌·공유가 낫다.

2.4.2.4 텍스처·표면 메모리

• 글로벌 메모리를 텍스처 캐시라는 별도 읽기 경로로 보는 방식. 2D/3D 공간 지역성에 최적화된 캐시, 하드웨어 보간(filtering)·정규화 좌표·경계 처리(clamp/wrap) 무료. 읽기 전용(텍스처), 읽기·쓰기(표면).
• 용도: 이미지·볼륨 처리·스텐실·공간적으로 뭉친 불규칙 접근. 단순 읽기 전용 캐싱이면 __ldg() 또는 const __restrict__ 포인터로 충분한 경우가 많아, 텍스처는 보간·좌표 기능이 필요할 때 주로 쓴다.

2.4.2.5 공유 메모리(짧게)

• 블록 단위 범위, 블록 수명. 온칩, 매우 빠름. 블록 내 스레드 협력·재사용의 핵심 도구. 뱅크·병합 메커니즘은 1.7, 타일링 활용은 2.6~2.7.

2.4.2.6 한눈에 비교

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

• [ ] 여섯 공간을 범위·수명·캐시·R/W로 구분한다.
• [ ] 어떤 읽기 전용 데이터를 상수(작고·워프 공통)에 둘지, 텍스처(2D 지역성·보간)에 둘지 판단한다.
• [ ] "로컬 메모리 ≠ 빠름"(스필·동적 인덱스 → HBM)을 알고 --ptxas-options=-v로 확인한다.
• [ ] 상수 메모리가 워프 공통 주소일 때만 이득(브로드캐스트), 분기 접근이면 직렬화임을 안다.

증상: 커널이 갑자기 느려졌는데 글로벌 접근은 그대로라면 → 레지스터 스필(→로컬→HBM) 의심, --ptxas-options=-v로 spill 확인(1.6). 상수로 바꿨더니 더 느려짐 → 워프 내 스레드가 서로 다른 주소를 읽음(브로드캐스트 깨짐).

2.4.4 흔한 함정

• 상수 메모리 오해. (a) 큰 데이터를 못 넣는다(총 64 KB, 확인 필요). (b) 워프 내 분기 접근(스레드마다 다른 주소)은 브로드캐스트가 깨져 직렬화 — 이때 상수는 글로벌보다 느리다. 상수는 "작고·읽기전용·워프 공통"일 때만 이득.
• 로컬 메모리 스필을 글로벌로 오인. --ptxas-options=-v로 spill stores/loads 확인(1.6).
• 텍스처를 단순 캐싱에 과용. 보간·좌표 기능이 필요 없으면 __ldg()/const __restrict__로 충분.
• 동적 인덱스 지역 배열. 컴파일 타임에 인덱스를 모르는 지역 배열은 로컬(HBM)로 떨어진다. 가능하면 인덱스를 상수화하거나 공유 메모리로.

2.4.5 연습문제

연습 1 (상수 vs 글로벌). 다음 중 상수 메모리에 두면 이득인 것과 손해인 것을 고르고 이유를: (a) 모든 스레드가 읽는 3×3 변환 행렬, (b) 스레드 i가 table[i]만 읽는 1M 룩업테이블, (c) 모든 스레드가 같은 16개 양자화 스텝.

연습 2 (로컬 메모리 진단). 커널에 float tmp[64];를 두고 tmp[someVar](런타임 인덱스)로 접근하니 느려졌다. 왜인가? 어떻게 확인하나?

연습 3 (공간 선택). 1D 컨볼루션에서 (a) 입력 신호, (b) 5-탭 필터 계수, (c) 블록 내 재사용되는 이웃 윈도우를 각각 어느 공간에 두나?

2.4.6 요약 · 다음 장

• CUDA 메모리 공간은 범위·수명·캐시 경로·R/W로 갈린다: 레지스터(스레드) → 공유(블록) → 글로벌·상수·텍스처(그리드).
• 상수는 워프 공통 읽기 전용, 텍스처는 공간 지역성·보간, 로컬은 이름과 달리 느린 HBM(스필·동적 인덱스).
• 어떤 데이터를 어디에 둘지는 범위·재사용·접근 패턴으로 정한다.

다음 장: 2.5 메모리 모델 II — 전송 최적화. 실제 병목인 PCIe를 건너는 호스트-디바이스 전송 — pageable vs pinned, 통합 메모리 prefetch/advise, zero-copy를 산수와 함께 본다.