제1부 GPU 아키텍처
1.7 온칩 메모리 I — 레지스터·공유 메모리·뱅크
GPU 성능은 연산기(ALU·텐서코어)가 아니라 데이터를 얼마나 빨리, 얼마나 가까이에서 공급하느냐로 결정되는 경우가 많다. 이 장은 SM 안의 온칩 메모리 중 레지스터·공유 메모리/L1과 그 뱅크 구조를 다룬다. L2·병합·대역폭은 1.8, 칩 밖 HBM은 1.9로 잇는다.
1.7.1 이 장에서 배우는 것 / 선수 지식
학습 목표:
- 온칩 메모리 계층(레지스터 → 공유/L1 → L2)을 지연·대역폭·용량·범위(스레드/블록/디바이스) 의 4축으로 구분한다.
- 레지스터 압박과 스필(spill)이 점유율·성능에 주는 영향을 안다.
- 공유 메모리/L1의 통합(unified) 구성(carveout)과 동적 공유 메모리 opt-in을 안다.
- 공유 메모리의 뱅크(bank) 구조를 이해하고, 어떤 접근이 뱅크 충돌(bank conflict) 을 일으키는지 그림으로 판별한다.
선수 지식: 1.2(SIMT·워프), 1.4~1.6(SM 내부·레지스터 파일·점유율). "워프 = 32 스레드가 한 명령을 함께 실행"을 전제로 쓴다.
정직성: 지연(사이클)·대역폭·용량 수치는 세대마다 다르고 마이크로벤치마크 의존이다. 모두 "시점·세대 의존, 확인 필요" 개략값. 정확한 수는 화이트페이퍼·측정으로.
1.7.2 왜 메모리 계층인가 — 큰 그림
GPU는 수천 스레드로 지연을 은닉(1.2)하지만, 은닉에는 한계가 있다. 대기 중인 워프가 충분히 많아야 하고 그만큼 레지스터·공유 메모리를 소비한다. 즉 온칩 메모리는 단순 "빠른 저장소"가 아니라 점유율을 좌우하는 자원이다(1.6).
계층의 핵심: 가까울수록 빠르고 작고 사적(私的)이다. 멀수록 느리고 크고 공유된다.
| 범위 | 메모리 | 지연·용량 |
|---|---|---|
| 스레드 | 레지스터 | 가장 빠름(~0지연), SM당 256KB / 스레드 ≤255개 |
| 블록 | 공유 메모리 / L1(통합) | 빠름(~20–30 사이클), SM당 수십~수백 KB |
| 디바이스 | L2 캐시(전 SM 공유) | 중간(~200 사이클), 수 MB~수십 MB |
| 디바이스 | 글로벌(HBM/GDDR, 칩밖) | 느림(~수백 사이클), 수십 GB (→ 1.9) |
지연 ≠ 대역폭에 주의: 지연은 한 요청이 돌아오는 시간, 대역폭은 파이프라인을 가득 채웠을 때의 처리량이다. GPU는 지연을 병렬성으로 은닉한 뒤 대역폭으로 승부하는 기계다.
1.7.3 레지스터 — 가장 빠르고 가장 사적인 저장소
레지스터는 SM 내부 레지스터 파일(1.4)에서 할당된다. 컴파일러(ptxas)가 지역 변수를 레지스터에 배치한다. 각 스레드는 자기만의 레지스터 집합을 본다 — 범위는 스레드, 수명은 커널 실행.
- 용량: SM당 흔히 64K개 32비트 레지스터(=256 KB, 세대 의존). 스레드 1개 상한은 255 레지스터.
- 지연: 연산 파이프라인에 직접 물려 거의 0지연.
레지스터 압박과 스필(spill). 레지스터는 SM 단위로 모든 상주 스레드가 나눠 쓴다(1.4). 스레드당 레지스터를 많이 쓰면 동시 상주 워프가 줄어 점유율이 떨어진다(1.6). 레지스터가 모자라면 컴파일러가 변수를 로컬 메모리로 방출한다(스필). 이름과 달리 로컬 메모리는 온칩이 아니라 글로벌 메모리 영역에 있고(L1·L2로 캐시될 뿐) 느리다.
nvcc -arch=sm_90 -Xptxas -v -c kernel.cu
# 출력 예(개략): ptxas info: Used 96 registers, 0 bytes spill stores, 0 bytes spill loads
# spill 이 0이 아니면 레지스터 압박. __launch_bounds__ 나 -maxrregcount 로 조절.견해: 점유율을 위해 레지스터를 맹목적으로 줄이면 스필로 오히려 느려진다(1.6). "스필 0을 유지하면서 점유율 확보"가 1차 목표다.
1.7.4 공유 메모리와 L1 — 통합된 온칩 스크래치패드
공유 메모리는 블록 안 모든 스레드가 함께 보는 온칩 스크래치패드다. 범위는 블록, 수명은 블록 실행. __shared__로 선언하며 프로그래머가 명시적으로 적재·동기화한다(자동 캐시가 아니다 — "사용자 관리 캐시"). 타일링(2.7)의 핵심 자원이다.
Volta 이후 공유 메모리와 L1 데이터 캐시는 같은 온칩 SRAM을 분할해 쓴다(통합 L1/공유 메모리).
carveout — 분할 비율 조절. 커널이 공유 메모리를 많이 쓰면 L1 몫이 줄고 반대도 성립한다. 런타임 힌트로 분할 선호를 줄 수 있다.
// 공유 메모리 선호 (L1 최소)
cudaFuncSetAttribute(myKernel, cudaFuncAttributePreferredSharedMemoryCarveout, 100); // 0~100(%)
// 블록당 공유 메모리가 기본 48KB 정적 상한을 넘으려면 동적 공유 메모리 + opt-in:
cudaFuncSetAttribute(myKernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 100*1024); // 예 100KB
myKernel<<<grid, block, 100*1024>>>(...); // 3번째 런치 인자 = 동적 공유 메모리 바이트함정: 블록당 정적 공유 메모리는 보통 48 KB로 제한, 그 이상은 동적 공유 메모리 + opt-in 필요. 한도·기본값은 세대마다 다르다(확인 필요).
1.7.5 공유 메모리 뱅크와 뱅크 충돌
1.7.5.1 뱅크 구조
공유 메모리는 32개 뱅크로 인터리브된다(워프 폭과 같음). 각 뱅크는 한 사이클에 32비트(4바이트) 워드 하나를 서비스한다. 연속한 4바이트 워드가 뱅크 0,1,…,31,0,1,… 순으로 라운드로빈 배치된다.
| 워드 인덱스 | 0 | 1 | 2 | 3 | … | 31 | 32 | 33 | … |
|---|---|---|---|---|---|---|---|---|---|
| 뱅크 번호 | 0 | 1 | 2 | 3 | … | 31 | 0 | 1 | … |
워드 i 는 bank (i mod 32).
규칙 (한 워프 32 스레드가 같은 명령에서 접근할 때):
- 32 스레드가 서로 다른 32개 뱅크에 흩어지면 → 1트랜잭션(충돌 없음).
- 여러 스레드가 같은 뱅크의 다른 워드를 노리면 → 뱅크 충돌: N-way 충돌 = N배 직렬화.
- 여러 스레드가 정확히 같은 워드를 읽으면 → 브로드캐스트(충돌 아님).
1.7.5.2 충돌 차수 예시와 패딩 해법
| 케이스 | 접근식 | 뱅크 매핑 | 결과 |
|---|---|---|---|
| 충돌 없음 (stride 1) | sh[threadIdx.x] | 스레드 t0..t31 → 워드 0..31 → 뱅크 0..31 (1:1) | 충돌 0 ✓ |
| 2-way 충돌 (stride 2) | sh[2*threadIdx.x] | 워드 0,2,4,...,62 → 뱅크 0,2,...,0,2 (t0와 t16 둘 다 bank0) | 2배 직렬화 ✗ |
| 최악 32-way (열 접근) | float sh[32][32]; sh[threadIdx.x][k] | 주소 t32+k → 뱅크 (t32+k)%32 = k (전원 같은 뱅크 k) | 32배 직렬화 ✗✗✗ |
| 패딩으로 해결 | float sh[32][33]; sh[threadIdx.x][k] | 주소 t33+k → 뱅크 (t33+k)%32 = (t+k)%32 (33 mod 32 = 1); t=0..31 이 뱅크 k,k+1,...,k+31 전부 다른 뱅크 | 충돌 0 ✓ |
비용: 32×33×4 = 4224 B (vs 4096 B, +128 B). 충돌 32배를 없애는 대가로 SRAM 3%를 더 쓴다 — 거의 항상 이득이다. 2.7의 타일 행렬곱에서 이 패딩을 실제 적용한다.
주의: 충돌 계산은 4바이트 워드 기준.
double(8B)·float4(16B)는 한 접근이 여러 뱅크를 가로질러 규칙이 달라진다(세대별 32/64비트 모드 차이도, 확인 필요).
1.7.6 중간 점검 (체크포인트)
- [ ] 임의의
__shared__접근식을 보고 뱅크 충돌 차수(1-way/…/32-way)를 직접 계산한다. - [ ]
[32][32]를[32][33]으로 패딩하면 왜 열 접근 충돌이 사라지는지 모듈러 산수로 설명한다. - [ ] 레지스터 스필이 왜 느린지(로컬=글로벌 영역),
-Xptxas -v의 어느 줄을 보는지 안다. - [ ] 공유 메모리/L1이 같은 SRAM을 carveout으로 나눔을, 48KB 벽과 동적 공유 메모리 opt-in을 안다.
증상: 공유 메모리 커널이 이론보다 느리다 → 뱅크 충돌 의심, 패딩·레이아웃 검토(Nsight shared conflicts, 2.14).
1.7.7 흔한 함정
- "공유 메모리는 항상 빠르다"는 착각. 뱅크 충돌이 있으면 32배까지 직렬화돼 글로벌보다 나을 것이 없을 수 있다.
- 열 접근 + 32 정렬 타일.
[K][K]에서 K가 32의 배수면 열 접근이 같은 뱅크로 몰린다. 패딩 한 칸으로 해결. - 로컬 메모리 ≠ 빠름. 이름과 달리 글로벌 영역. 레지스터 스필로 새어들면 느려진다.
- 48KB 벽. 블록당 정적 공유 메모리 상한을 모르고 큰 타일을 잡으면 런치 실패. 동적 공유 메모리 + opt-in 필요.
double·float4뱅크 계산. 충돌 규칙은 4바이트 워드 기준. 큰 타입은 분석이 달라진다.
1.7.8 연습문제
연습 1 (뱅크 충돌 차수). __shared__ float s[1024];에서 워프가 s[threadIdx.x * 3]을 읽는다. 충돌 차수는? s[threadIdx.x * k]의 일반 규칙은?
정답
stride 3은 32와 서로소(gcd(3,32)=1)라 뱅크 $(3\cdot t)\%32$가 0~31을 한 번씩 덮음 → 충돌 없음. 일반화: 충돌 차수 = gcd(k, 32). 홀수 k면 충돌 0, k=2면 2-way, k=32(배수)면 32-way. "stride를 홀수로"가 회피 휴리스틱.
연습 2 (패딩 산수). float t[32][33];에서 행 접근 t[k][tx]와 열 접근 t[tx][k]의 뱅크 분포를 계산하고 둘 다 충돌이 없는지 확인하라.
정답
행 t[k][tx]: 뱅크 (33k+tx)%32, tx가 0~31이면 뱅크 0~31 전부 → 충돌 0. 열 t[tx][k]: 뱅크 (33tx+k)%32 = (tx+k)%32, tx 0~31이면 k,k+1,…,k+31 전부 다름 → 충돌 0. 핵심은 33 mod 32 = 1.
연습 3 (carveout 의사결정). 한 커널이 블록당 공유 메모리 100 KB를 원한다. (a) 정적 __shared__로 되나? (b) 가능케 하는 API는? (c) 점유율 영향은?
정답
(a) 정적 상한 ~48KB라 불가. (b) cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 100*1024) opt-in 후 동적 공유 메모리(런치 3번째 인자). (c) 블록당 공유 메모리가 크면 SM 슬롯을 많이 먹어 동시 상주 블록↓ → 점유율↓(1.6). 재사용 이득과 점유율 손실을 저울질.
1.7.9 요약 · 다음 장
- 온칩 계층은 레지스터(스레드·~0지연) → 공유/L1(블록·수십 사이클) 로, 가까울수록 빠르고 작고 사적이며, 모두 점유율을 좌우하는 한정 자원이다.
- 레지스터 스필은 로컬(=글로벌) 메모리로 새어 느려진다 —
-Xptxas -v로 확인. - 공유 메모리/L1은 같은 SRAM을 carveout으로 분할하고, 48KB 초과는 동적 공유 메모리 opt-in이 필요하다.
- 공유 메모리는 32 뱅크로 나뉘며 같은 뱅크에 몰리면 최대 32배 직렬화 — 패딩·레이아웃으로 회피한다.
다음 장: 1.8 온칩 메모리 II — L1/L2·캐시·메모리 병합. 글로벌 접근의 병합(coalescing) 원리, L2 영속성 힌트, 그리고 병합/비병합 복사 커널을 직접 빌드해 유효 대역폭 차이를 측정한다.