제2부 CUDA 프로그래밍

2.7 공유 메모리 II — 타일 행렬곱

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

2.6에서 타일링의 개념과 동기화 흐름을 봤다. 이 장은 그것을 직접 코드로 — naive에서 출발해 타일 GEMM을 빌드·실행·검증하고, 글로벌 트래픽 감소를 수치로 확인하고, 타일 크기 트레이드오프를 측정한다.

배우는 것:

• 복붙해서 빌드·실행 가능한 tiled GEMM 커널을 작성하고 정확성을 검증한다.
• 글로벌 읽기가 $2N^3$→$2N^3/T$로 줄어드는 것을 수치로 확인한다.
• 타일 크기 선택의 트레이드오프(재사용↑ vs 공유 메모리·점유율↓)를 측정으로 판단한다.

선수 지식: 2.6(타일링 메커니즘·동기화·뱅크 충돌), 2.1~2.2(빌드·실행·검증).

코드는 CUDA Toolkit 12.x nvcc 가정(시점 의존). sm_XX는 보유 GPU에 맞춤(2.12).

2.7.2 워크스루 — naive → tiled GEMM

2.7.2.1 1단계 — naive 커널 (기준선)

// gemm.cu (공통 헤더·CHECK 매크로)
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <cuda_runtime.h>
#define CHECK(x) do{cudaError_t e=(x); if(e!=cudaSuccess){ \
  printf("CUDA error %s:%d: %s\n",__FILE__,__LINE__,cudaGetErrorString(e)); exit(1);} }while(0)

// C = A * B, 모두 N x N (row-major)
__global__ void gemm_naive(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) sum += A[row*N + k] * B[k*N + col];
        C[row*N + col] = sum;
    }
}

2.6.2처럼 글로벌에서 $2N^3$ 원소를 읽는다. 기준선으로만.

2.7.2.2 2단계 — 타일 커널

#define TILE 16
__global__ void gemm_tiled(const float* A, const float* B, float* C, int N) {
    __shared__ float As[TILE][TILE + 1];   // 패딩(+1): 뱅크 충돌 대비(방어적 — 이 곱은 행 접근이라 사실 충돌 없음, 2.6.4)
    __shared__ float Bs[TILE][TILE + 1];
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * TILE + ty;      // 이 스레드가 맡은 C의 행
    int col = blockIdx.x * TILE + tx;      // 이 스레드가 맡은 C의 열
    float sum = 0.0f;

    for (int p = 0; p < (N + TILE - 1) / TILE; ++p) {     // K 방향 단계
        // ① 협력 적재 (경계 밖은 0으로)
        int aCol = p * TILE + tx;
        int bRow = p * TILE + ty;
        As[ty][tx] = (row < N && aCol < N) ? A[row*N + aCol] : 0.0f;
        Bs[ty][tx] = (bRow < N && col < N) ? B[bRow*N + col] : 0.0f;
        __syncthreads();                   // ② 타일 적재 완료 보장

        // ③ 공유 메모리에서 T번 곱-합
        #pragma unroll
        for (int k = 0; k < TILE; ++k) sum += As[ty][k] * Bs[k][tx];
        __syncthreads();                   // ④ 덮어쓰기 전 계산 완료 보장
    }
    if (row < N && col < N) C[row*N + col] = sum;
}

2.7.2.3 호스트 드라이버 — 빌드·실행·검증

static void gemm_cpu(const float* A, const float* B, float* C, int N) {
    for (int i = 0; i < N; ++i)
        for (int j = 0; j < N; ++j) {
            float s = 0.0f;
            for (int k = 0; k < N; ++k) s += A[i*N+k]*B[k*N+j];
            C[i*N+j] = s;
        }
}
int main() {
    const int N = 512;                       // TILE 배수가 아니어도 동작(경계 처리)
    size_t bytes = (size_t)N*N*sizeof(float);
    float *hA=(float*)malloc(bytes), *hB=(float*)malloc(bytes);
    float *hC=(float*)malloc(bytes), *hRef=(float*)malloc(bytes);
    for (int i=0;i<N*N;++i){ hA[i]=(float)(rand()%7-3); hB[i]=(float)(rand()%7-3); }

    float *dA,*dB,*dC;
    CHECK(cudaMalloc(&dA,bytes)); CHECK(cudaMalloc(&dB,bytes)); CHECK(cudaMalloc(&dC,bytes));
    CHECK(cudaMemcpy(dA,hA,bytes,cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(dB,hB,bytes,cudaMemcpyHostToDevice));

    dim3 block(TILE, TILE);
    dim3 grid((N+TILE-1)/TILE, (N+TILE-1)/TILE);
    cudaEvent_t t0,t1; CHECK(cudaEventCreate(&t0)); CHECK(cudaEventCreate(&t1));
    CHECK(cudaEventRecord(t0));
    gemm_tiled<<<grid, block>>>(dA, dB, dC, N);   // ← naive로 바꿔 비교
    CHECK(cudaEventRecord(t1));
    CHECK(cudaGetLastError()); CHECK(cudaEventSynchronize(t1));
    float ms=0; CHECK(cudaEventElapsedTime(&ms,t0,t1));

    CHECK(cudaMemcpy(hC,dC,bytes,cudaMemcpyDeviceToHost));
    gemm_cpu(hA,hB,hRef,N);
    double maxerr=0;
    for (int i=0;i<N*N;++i) maxerr=fmax(maxerr, fabs(hC[i]-hRef[i]));
    double gflop = 2.0*(double)N*N*N/1e9;
    printf("N=%d  kernel=%.3f ms  ~%.1f GFLOP/s  maxerr=%.3g\n", N, ms, gflop/(ms/1e3), maxerr);
    return 0;
}

nvcc -O3 -arch=sm_80 gemm.cu -o gemm   # sm_XX는 보유 GPU에 맞춤
./gemm
# 예상(개략, 환경 의존):  N=512  kernel=0.18 ms  ~1491 GFLOP/s  maxerr=0

gemm_tiled를 gemm_naive로 바꿔 같은 N으로 실행하면 같은 maxerr(정확성 동일)에 시간이 더 길게(GFLOP/s 더 낮게) 나온다. 정확한 배수는 GPU·N·캐시 거동에 달려 단정하지 않는다(2.7.4). 핵심은 정확성은 같고 처리량이 오른다는 점이다.

실전 주의: 이 커널은 교육용이다. 양산 성능은 레지스터 블로킹·이중 버퍼링·벡터화·텐서코어를 더한 cuBLAS·CUTLASS가 압도한다 — 3.1·3.4에 맡긴다.

2.7.3 수치로 — 글로벌 읽기가 얼마나 줄어드나

$N\times N$ 행렬곱, 타일 $T\times T$:

• naive: 출력 $N^2$개, 각자 2N 읽기 → $2N^3$ 글로벌 읽기.
• tiled: 출력 타일 $(N/T)^2$개. 한 타일은 K 방향 N/T 단계, 각 단계 A타일($T^2$)+B타일($T^2$)=$2T^2$를 글로벌에서 한 번 읽음 → 한 타일당 $(N/T)\cdot2T^2 = 2NT$. 전체 = $(N/T)^2\cdot2NT = 2N^3/T$.

감소 배수 = $2N^3 / (2N^3/T) = T$. 타일 한 변 크기 T만큼 글로벌 읽기가 준다. 요지: 글로벌에서 공유로 한 번 적재한 원소 하나를 블록 안에서 T번 재사용하기 때문.

구체 수치(N=1024, FP32):

(비율 = $2N^3 / (2N^3\cdot4/T) = T/4$ FLOP/byte → T=16이면 4, T=32이면 8. 핵심은 T에 비례해 비율이 오른다는 것.)

2.7.4 트레이드오프 — 타일 크기 선택

타일 크기 T는 추가 비용 없이 키울 수 없다. 세 힘이 맞선다:

1. 재사용(↑ 클수록 좋음): 글로벌 트래픽 ∝ 1/T.
2. 공유 메모리 용량(↓ 제약): 블록당 공유 메모리 = $2\cdot T^2\cdot4$ 바이트(As+Bs, FP32). SM당 공유 메모리는 한정(세대 의존). 타일이 크면 블록 하나가 공유 메모리를 많이 먹어 SM에 동시에 상주시킬 블록 수(점유율) 가 줄고 지연 은닉이 약해진다(1.6).
3. 블록 크기·레지스터(↓ 제약): $T\times T$ 스레드. T=32면 1024 스레드(블록 최대)라 더 못 키운다. 레지스터 압박도 커진다.

견해: T=16 또는 T=32로 시작해 프로파일러(2.14)로 점유율·뱅크 충돌·메모리 처리량을 보고 정한다. 본격 성능은 한 스레드가 출력 여러 개를 맡는 레지스터 블로킹·이중 버퍼링으로 더 높이지만, 그것은 이미 라이브러리(cuBLAS·CUTLASS, 3.1·3.4)가 극한까지 한 영역이다.

비엔비디아 정직: 타일링/blocking은 엔비디아 고유가 아니라 캐시 블로킹의 일반 원리다(CPU BLAS·rocBLAS·SYCL 모두 같은 아이디어). CUDA의 차별점은 공유 메모리를 명시 제어한다는 점과 그 위에 텐서코어가 더해진다는 점.

naive 대비 실측 배수를 단정하지 않는 이유: naive 커널도 L2 캐시가 재사용을 일부 흡수해 "이론상 T배"가 그대로 시간 단축으로 안 나타난다. 작은 N은 전부 L2/L1에 들어가 차이가 작고, 큰 N·캐시 초과 영역에서 타일링 이득이 두드러진다. 수치는 직접 측정하라.

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

• [ ] "naive 대비 왜 빨라지는가"를 수치로 설명한다(글로벌 읽기 T배 감소, 비율 0.25→T/4).
• [ ] tiled 커널을 빌드·실행하고 maxerr=0(또는 미세 FP 오차)으로 정확성을 확인한다.
• [ ] 타일 크기를 바꿔 시간을 측정하고 트레이드오프를 해석한다.

출력이 틀렸다면 의심 순서:

1. maxerr가 실행마다 다르게 0이 아니다 → 동기화 누락 race(2.6). 가장 흔하다.
2. maxerr가 항상 같은 큰 값 → 인덱싱 버그(As[ty][k] vs As[k][ty], row/col 뒤바뀜, aCol/bRow).
3. N이 TILE 배수가 아닐 때만 가장자리가 틀림 → 타일 경계 처리 누락(적재 시 0 패딩, 저장 시 row<N && col<N).
4. 정확하지만 naive보다 안 빠름 → N이 너무 작아 전부 캐시에 들어감, 또는 점유율 저하. 프로파일러로 확인(2.14).

2.7.6 흔한 함정

• 동기화 누락 race (대표). compute-sanitizer --tool racecheck ./gemm로 공유 메모리 race를 검출한다. 작은 N으로 반복 실행해 maxerr 변동을 본다.
• 타일 경계(N이 TILE 배수 아님). 적재 시 경계 밖을 안 막으면 잘못된 글로벌을 읽어 결과 오염. 삼항 가드가 처리.
• 뱅크 충돌 방치. 정확성엔 영향 없어 조용히 성능만 저하시킨다(2.6). 패딩은 측정 후.
• 점유율만 보고 타일 맹목적으로 키움. T=32는 블록 1024라 더 못 키우고, 공유 메모리·레지스터 압박으로 SM당 블록 수가 떨어져 오히려 느려질 수 있다.

2.7.7 연습문제

연습 1 (감소 배수). naive와 tiled(T=16)의 글로벌 읽기 원소 수를 N=2048에 대해 계산하라. 감소 배수는?

연습 2 (동기화 race 재현). 둘째 __syncthreads()(④)를 지우고 N=512로 20회 반복해 maxerr를 관찰하라. racecheck로 확인하라.

연습 3 (타일 크기 측정). TILE을 8·16·32로 바꿔 N=2048에서 시간을 측정하고 naive와 비교해 표로. 어떤 T가 가장 빠른가? 예상과 다르면 왜?

연습 4 (직사각 일반화). $A:M\times K$, $B:K\times N$으로 일반화하려면 커널의 어디를 바꾸나? 단계 수·경계 조건은?

2.7.8 요약 · 다음 장

• 타일링은 글로벌 읽기를 $2N^3$→$2N^3/T$로 타일 한 변 T배 줄인다. 정확성은 naive와 같고 처리량이 오른다.
• 핵심은 두 __syncthreads()(적재 후·계산 후) + 패딩(뱅크 충돌). 타일 크기는 재사용↑ vs 공유 메모리·점유율↓ 트레이드오프 — 프로파일러로 정한다.
• 실측 배수는 L2 캐시 거동 때문에 "이론상 T배"와 달라 직접 측정해야 한다. 본격 성능은 cuBLAS(3.1)·CUTLASS(3.4)의 영역이며, 이 장은 그 라이브러리가 안에서 무엇을 하는지를 실습으로 익히는 단계다.

다음 장: 2.8 동기화 I — barrier·atomic. __syncthreads()(블록 barrier)와 원자연산(atomic)으로 스레드들이 안전하게 협력·집계하는 법을 본다.