제2부 CUDA 프로그래밍
2.6 공유 메모리 I — 뱅크·충돌·동기화
2.6.1 이 장에서 배우는 것 / 선수 지식
공유 메모리로 데이터를 재사용해 글로벌 메모리 트래픽을 줄이는 타일링(tiling) 의 개념과 메커니즘을 익힌다. 왜 공유 메모리가 필요한지(naive 행렬곱의 한계), 타일 한 단계의 적재·동기화·계산·동기화 흐름, 뱅크 충돌 회피 — 까지. 실제 타일 행렬곱 코드·측정은 2.7로 잇는다.
배우는 것:
- 왜 naive 행렬곱이 메모리 바운드인지 산수로 설명한다(연산-바이트 비율).
- 공유 메모리 타일이 글로벌 접근을 타일 크기에 비례해 줄이는 원리.
__syncthreads()를 정확한 위치(적재 후·계산 후)에 두는 이유와 누락 시 race.- 뱅크 충돌을 패딩으로 회피하는 법(하드웨어 메커니즘은 1.7).
선수 지식: 1.7(공유 메모리 뱅크·충돌·패딩의 하드웨어), 1.10(메모리 바운드·루프라인), 2.1~2.3(커널·실행 구성), 2.4(메모리 공간). 이 장은 1.7의 뱅크 메커니즘을 활용하는 단계다.
환경 셋업은 0.3·A.1. 코드는 CUDA Toolkit 12.x 계열
nvcc가정(버전 시점 의존).
2.6.2 왜 공유 메모리인가 — naive 행렬곱의 한계
온칩 계층(1.7)을 재사용 관점에서 본다. 핵심은 접근 비용의 자릿수 차이:
| 위치 | 범위 | 상대 지연(개략) | 비고 |
|---|---|---|---|
| 레지스터 | 스레드 | ~1 | 컴파일러 할당 |
| 공유 메모리 | 블록(SM당) | ~수십 사이클 | 프로그래머가 명시 제어 |
| L2 캐시 | GPU 전체 | ~수백 사이클 | 자동 |
| 글로벌(HBM) | GPU 전체 | ~수백 사이클+ | 대역폭 한정 |
공유 메모리는 프로그래머가 명시적으로 채우고 비우는, 블록 단위로 공유되는 빠른 스크래치패드다. 같은 데이터를 블록 안 여러 스레드가 여러 번 읽을 때, 글로벌에서 한 번만 가져와 공유 메모리에 적재해 두고 거기서 반복 읽으면 글로벌 트래픽이 준다. 이것이 타일링의 전부다.
naive 행렬곱은 왜 느린가. $C = A \times B$, 모두 $N\times N$. naive 커널은 출력 원소 하나를 스레드 하나가 맡고 내적 한 줄을 순회한다:
for k in 0..N: sum += A[row][k] * B[k][col]한 스레드가 글로벌에서 읽는 양: A의 한 행(N) + B의 한 열(N) = 2N. 출력은 $N^2$개이므로 전체 글로벌 읽기 = $2N^3$ 원소($2N^3\cdot4$ 바이트, FP32). 연산은 곱·합 $2N^3$ FLOP.
연산-바이트 비율 = $2N^3 / (2N^3\times4)$ = 0.25 FLOP/byte. 형편없이 낮다. 현대 데이터센터 GPU는 FP32에서 수십~수백 FLOP/byte는 공급해야 연산기를 채운다(1.10 루프라인). 즉 naive 행렬곱은 연산기가 유휴 상태인 메모리 바운드 커널이다.
문제의 본질은 재사용 부족이 아니라 재사용을 글로벌에서 한다는 데 있다. A[row][k] 한 원소는 C[row][*] 전체(N개 출력)를 계산할 때 N번 쓰인다. naive에선 그 N번이 전부 글로벌 접근이다. 그 재사용을 공유 메모리로 끌어내리는 것이 타일링이다.
2.6.3 타일링의 메커니즘
2.6.3.1 출력을 타일로 분할한다
출력 C를 $T\times T$ 타일로 나눈다. 블록 하나가 출력 타일 하나를 맡고, 블록 안 스레드($T\times T$개)가 각자 그 타일의 원소 하나를 계산한다. 한 출력 타일을 계산하려면 A의 가로 띠와 B의 세로 띠가 필요하고, 이 띠를 $T\times T$ 타일로 분할해 K 방향으로 한 타일씩 공유 메모리에 적재하며 부분합을 누적한다.
한 단계(phase)에서는 A 한 타일과 B 한 타일을 공유 메모리로 적재하고, 그 안에서 T번 곱-합으로 부분합을 누적한 뒤 다음 타일로 넘어간다.
K 방향으로 N/T 단계를 돌면 한 출력 타일이 완성된다.
2.6.3.2 한 단계의 흐름 — 적재·동기화·계산·동기화
각 단계에서 동기화 두 번이 반드시 들어간다.
- ② 적재 후 동기화: 내 스레드는
As[ty][tx]만 채웠지만 ③에서는As[ty][0..T-1]전부를 읽는다. 그 칸들은 다른 스레드가 채운다. 모두 적재를 끝낼 때까지 기다리지 않으면 안 채워진 칸을 읽는다(race). - ④ 계산 후 동기화: 다음 단계가
As/Bs를 덮어쓴다. 누가 아직 ③을 읽는데 덮어쓰면 race. 계산이 다 끝난 뒤 덮어쓴다.
두 동기화 중 하나라도 빠지면 결과가 간헐적으로 틀린다(어떤 실행은 우연히 맞는다). 가장 잡기 어려운 버그다 — 2.7에서 재현·진단한다.
2.6.4 뱅크 충돌과 패딩 (타일링 맥락)
공유 메모리는 32개 뱅크로 나뉜다(1.7). 한 워프(32 스레드)가 같은 뱅크의 서로 다른 주소를 동시에 접근하면 뱅크 충돌로 직렬화된다(모두 같은 주소면 브로드캐스트, 충돌 아님).
뱅크 충돌은 한 워프가 공유 타일을 열(column) 방향으로 가를 때 생긴다. __shared__ float T[TILE][TILE](TILE=16, float 4바이트 → 행 stride=16워드)에서 워프가 같은 열 c의 여러 행 T[t][c](스레드가 행 t를 가름)를 동시에 읽는다고 하자. 주소가 16워드씩 떨어져 뱅크 $(16\cdot t + c) \operatorname{mod} 32$가 단 2개(c, c+16 계열)만 번갈아 나온다 → 타일 16행이 뱅크당 8개씩 몰려 8-way 충돌로 직렬화된다.
해법: 행에 패딩 한 칸을 더한다.
__shared__ float Bs[TILE][TILE + 1]; // 행 stride = TILE+1 = 17 (홀수)행 stride가 17이면 인접 행 간 주소 간격이 17워드 → 32 뱅크와 서로소가 되어 충돌이 흩어진다(상세 모듈러 산수는 1.7). 패딩 한 칸의 비용은 작은 메모리 낭비뿐이고 충돌 제거 이득이 보통 더 크다.
| 행 | 패딩 없음 (stride 16) | 패딩 +1 (stride 17) |
|---|---|---|
| 행0 | 뱅크 0 | 뱅크 0 |
| 행1 | 뱅크 16 (16칸 점프) | 뱅크 17 |
| 행2 | 뱅크 0 (겹침) | 뱅크 2 (안 겹침) |
패딩이 없으면 짝수 행은 전부 뱅크 0, 홀수 행은 전부 뱅크 16에 몰려 8-way 충돌로 직렬화된다. 패딩 +1(stride 17)이면 행마다 뱅크가 흩어져 겹치지 않는다.
패딩이 항상 필요한 것은 아니다 — 접근 패턴에 달렸다. 사실 2.7 타일 곱의 읽기 $As[ty][k]\cdot Bs[k][tx]$는 고정
k에서 워프가tx(행 방향)로 퍼지는 행 접근이라 16개 뱅크에 고르게 떨어져 충돌이 없다 → 그 패턴엔 패딩이 불필요하다. 패딩이 실질적 효과를 보는 것은 위처럼 열 방향으로 타일을 가르는 패턴(대표: 전치, 연습 3)이다. 프로파일러(Nsight Compute, 2.14)의 bank conflict 지표로 실제 충돌을 확인하고 패딩을 결정하는 것이 정석 — 근거 없는 패딩은 점유율만 낮출 수 있다.
2.6.5 중간 점검 (체크포인트)
- [ ] "naive 행렬곱이 왜 메모리 바운드인가"를 연산-바이트 비율 0.25로 설명한다.
- [ ] 타일링이 글로벌 읽기를 $2N^3$→$2N^3/T$로 줄임을(타일 한 변
T배) 안다(정량은 2.7). - [ ] 두
__syncthreads()의 위치와 이유(적재 후·계산 후)를 설명한다. - [ ] 패딩
[TILE+1]이 무엇을 막는지(뱅크 충돌) 말한다.
2.6.6 흔한 함정
- 동기화 누락 race.
__syncthreads()하나만 빠뜨려도 대부분 우연히 맞고 가끔 틀린다.compute-sanitizer --tool racecheck로 공유 메모리 race를 검출한다. - 조건부
__syncthreads().if(...){ __syncthreads(); }처럼 블록 내 일부 스레드만 도달하는 분기 안에 barrier를 두면 데드락(또는 미정의). barrier는 블록 전원이 도달해야 한다. 경계 처리는 적재 값을 0으로 만들되 barrier 자체는 모든 스레드가 무조건 통과하게. - 뱅크 충돌 방치. 정확성엔 영향 없어 조용히 성능만 저하시킨다. Nsight Compute 충돌 지표 없이는 못 본다(2.14). 패딩은 측정 후.
- 점유율만 보고 타일 맹목적으로 키움.
T=32는 블록 1024 스레드라 더 못 키우고, 공유 메모리·레지스터 압박으로 SM당 블록 수가 떨어져 오히려 느려질 수 있다(2.7).
2.6.7 연습문제
연습 1 (메모리 바운드 산수). naive 행렬곱의 연산-바이트 비율이 0.25인 이유를 $2N^3$ FLOP과 $2N^3\cdot4$ byte로 유도하라.
정답
출력 $N^2$개 각자 A에서 N + B에서 N = 2N 읽기 → 전체 $2N^3$ 원소 = $2N^3\cdot4$ byte(FP32). 연산 = 곱·합 $2N^3$ FLOP. 비율 = $2N^3 / (2N^3\cdot4) = 0.25$ FLOP/byte. 너무 낮아 메모리 바운드(1.10).
연습 2 (동기화 위치). 두 __syncthreads() 중 하나를 지우면 각각 어떤 종류의 race가 나는가?
정답
② 적재 후 동기화 제거 → 다른 스레드가 아직 안 채운 As/Bs 칸을 읽음(read-before-write). ④ 계산 후 동기화 제거 → 빠른 스레드가 다음 단계 적재로 덮어쓰는데 느린 스레드가 아직 읽는 중(write-after-read). 둘 다 간헐적 오답.
연습 3 (전치 뱅크 충돌). $T\times T$ 타일을 공유 메모리에 적재해 행렬을 전치하는 커널에서 패딩이 왜 특히 중요한가?
정답
전치는 tile[ty][tx] 적재와 tile[tx][ty] 읽기가 같은 뱅크를 강하게 친다(전치는 뱅크 충돌의 교과서 예제) → 패딩 [TILE+1]이 전형적으로 큰 이득. (전치 자체는 읽기 병합·쓰기 strided 문제를 공유 메모리로 푸는 것.)
2.6.8 요약 · 다음 장
- naive 행렬곱은 연산-바이트 비율 0.25 FLOP/byte의 메모리 바운드 커널 — 재사용을 글로벌에서 하기 때문.
- 타일링은 글로벌에서 한 번 가져온 타일을 공유 메모리에서 블록 안 여러 스레드가 재사용해 글로벌 읽기를 줄인다.
- 핵심은 두
__syncthreads()(적재 후·계산 후). 하나라도 빠지면 간헐적 race. 패딩[TILE+1]으로 뱅크 충돌을 흩는다(측정 후).
다음 장: 2.7 공유 메모리 II — 타일 행렬곱. naive→tiled GEMM을 직접 빌드·실행·검증하고, 글로벌 트래픽 감소를 수치로 확인하고, 타일 크기 트레이드오프를 측정한다.