제1부 GPU 아키텍처

1.8 온칩 메모리 II — L1/L2·캐시·메모리 병합

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

1.7에서 레지스터·공유 메모리·뱅크를 봤다. 이 장은 글로벌 메모리 접근의 병합(coalescing), L2 캐시와 영속성 힌트, 그리고 병합/비병합 접근의 성능 차이를 직접 측정하는 워크스루를 다룬다. 칩 밖 HBM 대역폭 자체는 1.9, 메모리바운드 판정은 1.10으로 잇는다.

학습 목표:

• 글로벌 메모리 접근의 병합 원리를 이해하고 병합/비병합의 유효 대역폭 차이를 추정한다.
• 정렬(alignment) 과 AoS vs SoA 레이아웃이 병합에 주는 영향을 안다.
• L1 캐시가 글로벌 로드를 언제 가로채는지, __ldg·읽기전용 캐시 경로를 안다(1.8.10).
• L2 캐시의 위치와 영속성(persistence) 힌트를 안다.
• 병합/비병합 복사 커널을 빌드해 유효 대역폭을 측정한다.

선수 지식: 1.2(워프가 메모리 접근 단위), 1.5(LD/ST 유닛), 1.7(온칩 계층). 대역폭 자체의 산수는 1.9.

정직성: 지연·대역폭·L2 용량은 세대 의존 개략값(확인 필요). 코드의 예상 출력도 환경 의존.

1.8.2 글로벌 메모리 병합(coalescing)

공유 메모리의 단위가 뱅크라면, 글로벌 메모리의 단위는 메모리 트랜잭션(보통 32바이트 섹터, 캐시라인 128바이트)이다. 한 워프의 32 스레드가 글로벌 메모리를 접근할 때, 하드웨어는 32개 주소를 모아 필요한 캐시라인/섹터만 가져온다.

• 32 스레드가 연속된 128바이트(=32×4B) 한 라인에 떨어지면 → 1 트랜잭션(완전 병합), 가져온 바이트가 전부 쓰임 → 효율 100%.
• 주소가 흩어지면(스트라이드·랜덤) → 여러 라인을 건드려 트랜잭션이 늘고, 각 라인에서 쓰는 바이트보다 가져온 바이트가 많아 대역폭 낭비.

정렬도 중요. 시작 주소가 캐시라인 경계에 정렬되지 않으면 연속 접근이라도 두 라인에 걸쳐 트랜잭션이 하나 더 생긴다. cudaMalloc은 충분히 정렬된 포인터를 주고, 구조체엔 __align__/float4 같은 벡터 타입으로 정렬·벡터화를 돕는다.

AoS vs SoA. 병합이 깨지는 흔한 실무 원인은 구조체 배열(Array of Structs) 이다.

// AoS — 비병합: 워프가 .x 만 읽으면 stride = sizeof(Particle)
struct Particle { float x, y, z, w; };
Particle p[N];   ...  p[i].x ...   // .x 가 16B 간격으로 흩어짐 → 라인 낭비

// SoA — 병합 친화: 같은 필드가 연속
struct Particles { float *x, *y, *z, *w; };  ...  ps.x[i] ...  // 4B 연속 → 1 라인에 32개 ✓

견해: GPU 데이터 레이아웃의 1번 원칙은 "워프가 같은 명령에서 만지는 주소를 연속으로" — 그래서 일반적으로 SoA를 기본으로 둔다. CPU 캐시 친화 코드(AoS)를 그대로 옮기면 병합이 깨진다. (메모리 모델·전송은 2.4~2.5.)

병합된 트랜잭션이 실제로 어디를 거쳐 오는지 — L1 캐시의 역할 — 는 1.8.10(번호는 이 장 끝에 붙지만 내용은 여기 다음 자리)에서 잇는다. 여기서는 먼저 그다음 층인 L2를 본다.

1.8.3 L2 캐시와 영속성(persistence) 힌트

1.8.3.1 L2의 위치

L2는 모든 SM이 공유하는 디바이스 범위 캐시로, 글로벌 메모리(HBM/GDDR) 앞단에 놓인다. 한 SM이 채운 라인을 다른 SM이 재사용할 수 있고, 디바이스 원자연산·여러 커널 간 데이터 재사용의 거점이다. 용량은 세대에 따라 수 MB~수십 MB로 커지는 추세다(확인 필요).

1.8.3.2 L2 영속성 — 자주 쓰는 데이터를 L2에 붙들기

Ampere 이후, L2의 일부를 영속(persisting) 으로 떼어 특정 글로벌 주소 범위(추론 가중치, 자주 읽는 룩업 테이블)가 L2에 더 오래 머물도록 힌트를 줄 수 있다. 스트리밍성 데이터가 L2를 휩쓸어 hot 데이터를 축출하는 것을 막는다.

// 1) L2 중 영속용 크기 설정
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, persistSize);
// 2) 특정 데이터 영역에 access policy window 설정
cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr  = hotData;       // 자주 읽는 영역
attr.accessPolicyWindow.num_bytes = hotBytes;
attr.accessPolicyWindow.hitRatio  = 0.6f;
attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);
// 3) 다 쓰면 윈도우 해제 + L2 영속 리셋 (안 하면 다음 커널에 잔류)
attr.accessPolicyWindow.num_bytes = 0;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);
cudaCtxResetPersistingL2Cache();

견해·함정: 영속 힌트는 재사용되는 작은 hot 영역에서만 이득. 무차별 적용하면 L2 일반 캐시 몫이 줄어 역효과. 리셋을 빠뜨리면 후속 커널이 영속 영역에 갇혀 느려진다. 효과는 Nsight Compute의 L2 hit rate로 검증(2.14). API·필드는 세대·CUDA 버전 의존.

1.8.4 워크스루 — 병합/비병합 복사 커널 비교

복붙해서 빌드·실행 가능한 예제. 같은 양의 데이터를 병합 vs stride 접근으로 복사해 시간을 잰다.

// coalesce.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void copyCoalesced(const float* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  // 연속 → 병합
    if (i < n) out[i] = in[i];
}
__global__ void copyStrided(const float* in, float* out, int n, int S) {
    int t = blockIdx.x * blockDim.x + threadIdx.x;
    int i = (t * S) % n;                            // stride 접근(병합 깨짐)
    if (t < n) out[t] = in[i];
}
int main() {
    const int n = 1 << 24;                 // 16M floats = 64 MB
    size_t bytes = (size_t)n * sizeof(float);
    float *d_in, *d_out;
    cudaMalloc(&d_in, bytes); cudaMalloc(&d_out, bytes);
    dim3 block(256), grid((n + 255) / 256);
    cudaEvent_t s, e; cudaEventCreate(&s); cudaEventCreate(&e);
    float ms; double gb = 2.0 * bytes / 1e9;  // read + write

    copyCoalesced<<<grid, block>>>(d_in, d_out, n);  // 워밍업
    cudaDeviceSynchronize();

    cudaEventRecord(s); copyCoalesced<<<grid, block>>>(d_in, d_out, n);
    cudaEventRecord(e); cudaEventSynchronize(e); cudaEventElapsedTime(&ms, s, e);
    printf("coalesced : %7.3f ms,  %6.1f GB/s\n", ms, gb / (ms / 1e3));

    cudaEventRecord(s); copyStrided<<<grid, block>>>(d_in, d_out, n, 32);
    cudaEventRecord(e); cudaEventSynchronize(e); cudaEventElapsedTime(&ms, s, e);
    printf("strided/32: %7.3f ms,  %6.1f GB/s (유효 대역폭)\n", ms, gb / (ms / 1e3));

    cudaFree(d_in); cudaFree(d_out); return 0;
}

nvcc -arch=sm_80 -O3 coalesce.cu -o coalesce   # -arch 는 GPU 세대에 맞게
./coalesce
# 예상 출력(개략, 환경 의존):
#   coalesced :   0.5xx ms,   ~수천 GB/s   ← HBM 대역폭에 근접
#   strided/32:   수 ms,      훨씬 낮음     ← 라인당 4B만 써서 유효 대역폭 급락

핵심: stride 32는 워프가 매 접근마다 다른 128B 라인을 건드려 가져온 바이트 대부분을 버린다. 유효 대역폭이 병합 대비 크게(흔히 한 자릿수 배) 낮아진다. 정확한 배수는 GPU·캐시 거동에 따라 다르니 Nsight Compute의 sector/request 효율 지표로 확인하라(2.14).

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

• [ ] 글로벌 접근식을 보고 병합인지 판별하고, 비병합이면 SoA·벡터화로 고치는 방향을 제시한다.
• [ ] 병합 효율을 캐시라인·트랜잭션 수로 추정한다(stride 1=100%, stride 2=50%, stride 32≈3%).
• [ ] L2의 위치(전 SM 공유, HBM 앞단)와 영속 힌트의 용도·리셋 필요성을 안다.
• [ ] 병합/비병합 복사 커널의 유효 대역폭 차이를 측정·해석한다.
• [ ] L1이 SM별 캐시임을(L2와 달리 전 SM 공유 아님), 캐싱 여부가 세대 의존임을, __ldg가 언제 도움되는지 안다(1.8.10).

증상: 글로벌 바운드 커널의 유효 대역폭이 HBM 한계의 일부에 불과 → 비병합/정렬 불량(낮은 sector efficiency) 의심(2.14). 추론 커널에서 가중치 재읽기가 글로벌로 샌다 → L2 영속 윈도우·재사용 패턴 점검.

1.8.6 트레이드오프 — carveout 의사결정

온칩 SRAM은 한정 자원이고 공유 메모리/L1이 경쟁한다(1.7). 결정은 보통 셋 중 하나:

1. 공유 메모리 많이 (carveout↑) — 타일링·스텐실·리덕션처럼 블록 내 재사용이 큰 커널. 단 블록당 공유 메모리가 크면 상주 블록↓ → 점유율↓(1.6). 누구에게: 재사용이 분명한 BLAS류·컨볼루션 타일.
2. L1 많이 (carveout↓) — 재사용 패턴이 불규칙해 하드웨어 캐시에 맡기는 편이 나은 커널(랜덤·간접 인덱싱). 누구에게: 그래프·희소·불규칙 워크로드.
3. 둘 다 적게, 레지스터·점유율 우선 — 산술 강도가 높아 ALU가 병목인 커널. 누구에게: compute-bound 커널.

비엔비디아 비교(정직): 사용자 관리 스크래치패드는 엔비디아 전유물이 아니다 — AMD LDS, 인텔 SLM이 같은 역할. 뱅크 충돌·병합·carveout의 사고방식은 이식 가능하지만 뱅크 수·라인 크기·API는 벤더·세대마다 다르니 수치를 그대로 옮기지 말 것.

1.8.7 흔한 함정

• AoS를 그대로 GPU로. CPU 친화 레이아웃이 워프 병합을 깬다. SoA 또는 벡터 타입 검토.
• 정렬 무시. 비정렬 시작 주소는 연속 접근에도 추가 트랜잭션을 만든다. cudaMalloc·float4로 정렬·벡터화.
• L2 영속 리셋 누락. 윈도우/영속을 안 풀면 후속 커널이 영향받는다.
• 유효 대역폭을 이론 대역폭으로 착각. 비병합이면 가져온 바이트의 일부만 쓴다 — Nsight의 sector efficiency로 실효를 확인(2.14).
• 수치 단정. 지연·대역폭·L2 용량은 세대 의존 개략값. 마이크로벤치 없이 절대값 단정 금지.

1.8.8 연습문제

연습 1 (병합 효율). float 배열을 워프가 stride 2(=8바이트 간격)로 읽는다. 한 워프가 건드리는 128B 라인 수와 효율(유효/가져온 바이트)을 추정하라.

연습 2 (AoS→SoA). struct {float x,y,z;} p[N];에서 워프가 p[i].y만 읽는다. 왜 비병합이고 어떻게 고치나?

연습 3 (L2 영속 판단). 추론 서버가 같은 임베딩 룩업 테이블(4 MB)을 매 요청 커널에서 읽는다. L2 영속을 적용할 가치가 있나? 주의점 두 가지는?

연습 4 (직접 해보기). 1.8.4 커널을 stride 1·2·8·32로 실행해 유효 대역폭을 표로 만들고, 어느 지점에서 급락하는지 관찰하라. Nsight Compute의 sector efficiency와 대조하라.

1.8.9 요약 · 다음 장

• 글로벌 접근은 워프 단위로 병합되어야 캐시라인을 알차게 쓴다 — 연속·정렬·SoA가 핵심 요인. 비병합은 유효 대역폭을 한 자릿수 배 낮춘다.
• L1은 SM별 캐시로 글로벌 로드를 가로챌 수 있다(세대·컴파일 의존, 1.8.10). L2는 전 SM 공유 디바이스 캐시이며, 영속 힌트로 작은 hot 영역을 붙들 수 있다(리셋 필수).
• 병합/비병합 복사 커널로 유효 대역폭 차이를 측정하고 Nsight로 검증한다.

이어서(번호는 이 장 끝에 붙지만 내용은 1.8.2 다음 자리에 속한다) 1.8.10 L1 캐시를 본 뒤 1.9로 이어진다.

1.8.10 L1 캐시 — 글로벌 로드의 하드웨어 캐싱

1.8.2는 워프의 글로벌 접근이 몇 개의 트랜잭션으로 묶이는지(병합)를 다뤘다. 그 트랜잭션이 실제로 어디를 거쳐 오는지가 이 절이다.

위치와 공유 SRAM: L1 데이터 캐시는 SM마다 있다 — L2(전 SM 공유)와 달리 다른 SM과 공유되지 않는다. 1.7의 온칩 SRAM을 공유 메모리와 나눠 쓰는 그 자원이 L1이다(carveout, 1.8.6). L2/HBM 앞에서 워프의 글로벌 로드를 가로채 — 이 SM의 어느 스레드가 최근 읽은 라인이 아직 L1에 있으면 L2·HBM까지 안 가고 바로 서빙된다.

히트/미스는 세대·컴파일 의존이 크다. 글로벌 로드를 L1에 기본으로 캐싱할지는 GPU 세대·컴파일 옵션에 따라 달라져 왔다 — 예컨대 과거 일부 세대는 글로벌 로드를 L1에 기본으로 캐싱하지 않고(로컬 메모리·레지스터 스필 전용으로 L1을 쓰고, 글로벌 캐싱은 별도 옵션으로 켜야 했음) 이후 세대에서 기본 캐싱으로 돌아온 이력이 있다 — 정확한 세대 경계·현재 기본값은 확인 필요(시점 의존). 최신 세대라고 가정하지 말고, 히트율은 Nsight Compute로 실측한다(2.14).

읽기 전용 데이터 — __ldg·const __restrict__: 커널 안에서 그 포인터로 절대 쓰지 않는 것이 보장된 글로벌 데이터(가중치·룩업 테이블 등)는 __ldg(ptr) 내장 함수나, 포인터를 const __restrict__로 선언해 컴파일러가 자동으로 읽기 전용 캐시 경로를 쓰게 유도할 수 있다. __restrict__(별칭 없음 보장)가 있어야 컴파일러가 이 최적화를 확신하고 적용한다.

캐시 연산자(고급, PTX 수준 개념): ld.global 명령엔 .ca(cache at all — 기본, L1+L2)·.cg(cache at global — L1 우회, L2만)·.cs(streaming — 재사용 안 할 데이터, evict-first) 같은 캐시 힌트가 있다(PTX ISA, 보통 인라인 PTX·컴파일러 경로로만 접근 — CUDA C++에서 직접 쓸 일은 드묾). 개념만 알아둔다: 한 번 쓰고 버릴 스트리밍 데이터를 기본(.ca)으로 읽으면 재사용될 다른 데이터를 L1에서 축출할 수 있다.

정직: L1 캐싱 정책(기본 on/off·세대별 차이·__ldg 적용 조건)은 세대·컴파일러 버전 의존이 크다 — 여기선 원리만, 실제 히트율·효과는 대상 GPU에서 Nsight Compute로 검증한다.

다음 장: 1.9 HBM과 대역폭 산수. 칩 밖 글로벌 메모리(GDDR vs HBM)의 물리 구조, "대역폭 = 폭 × 레이트", 그리고 LLM 추론에서 용량(GB) 이 배치·KV 캐시를 어떻게 제약하는지를 본다.