Square Tiling 기초

개요

Square Tiling은 행렬 곱셈(Matrix Multiplication) 연산을 효율적으로 수행하기 위한 기본적인 최적화 기법이다. 큰 행렬을 작은 정사각형 타일(Tile)로 분할하여 캐시 메모리와 레지스터 파일의 locality를 극대화하고, 연산 재사용(Operand Reuse)을 통해 메모리 대역폭 사용량을 줄인다. 이 기법은 GEMM(Generic Matrix Multiply) 커널의 성능을 결정하는 핵심 요소로, GPU, CPU, AI 가속기 등 다양한 하드웨어 플랫폼에서 광범위하게 사용된다.

Square Tiling의 핵심 아이디어는 "작은 타일을 한 번에 처리하면 캐시에 적재된 데이터를 여러 번 재사용할 수 있다"는 것이다. 예를 들어, 1024×1024 행렬 곱셈에서 단순 루프는 O(N³)의 캐시 미스를 발생시키지만, 32×32 타일로 분할하면 캐시 미스가 O(N³/T²)로 감소한다.

핵심 개념

타일링(Tiling)의 원리

행렬 곱셈 C = A × B의 정의에 따르면, 결과 행렬의 각 요소 C[i][j]는 다음과 같이 계산된다:

C[i][j] = Σ (A[i][k] × B[k][j])  (k = 0 ~ N-1)

단순 구현에서는 A의 i번째 행과 B의 j번째 열을 매번 메모리에서 로드하지만, Square Tiling을 적용하면:

1. 행렬 분할: 행렬 A, B를 T×T 크기의 정사각형 타일로 분할
2. 타일 단위 연산: 각 타일 쌍(A 타일, B 타일)에 대해 부분 곱셈 수행
3. 누산: 부분 곱셈 결과를 C 타일에 누산
4. 반복: 모든 타일 쌍에 대해 위 과정을 반복

캐시 레벨과 타일 크기

타일 크기는 하드웨어의 메모리 계층 구조에 따라 결정된다:

최적 타일 크기 공식:

T = √(캐시 크기 / (3 × 요소 크기 × 2))

• 요소 크기: FP16이면 2바이트, FP32이면 4바이트
• 3: A, B, C 세 행렬의 타일이 동시에 캐시에 있어야 함
• 2: A 타일과 B 타일 두 개가 필요

메모리 접근 패턴

Square Tiling은 메모리 접근 패턴을 크게 개선한다:

타일링 전 (단순 루프):
- A 행: O(N)번 로드
- B 열: O(N)번 로드
- 총 캐시 미스: O(N³)

타일링 후:
- A 타일: O(T)번 로드 → O(N/T)번 재사용
- B 타일: O(T)번 로드 → O(N/T)번 재사용
- 총 캐시 미스: O(N³/T²)

Cache-Oblivious 알고리즘

Square Tiling은 하드웨어 캐시 크기를 명시적으로 고려하는 Cache-Aware 접근법이다. 반면, Cache-Oblivious 알고리즘은 캐시 크기 파라미터 없이 재귀적으로 행렬을 분할하여 최적의 캐시 접근 패턴을 달성한다:

• 재귀적 분할: 행렬을 2×2 블록으로 재귀적으로 분할
• 파라미터 불필요: 캐시 크기나 블록 크기를 지정할 필요 없음
• 동적 적응: 멀티프로그래밍 환경에서 캐시 크기가 변해도 최적 성능 유지
• 캐시 미스 복잡도: Θ(N³/(b√M)) (b: 캐시 라인 크기, M: 캐시 크기)

Cache-Oblivious GEMM은 이론적으로 우수하지만, 실제 구현에서는 캐시 라인 단위의 병렬화와 SIMD 벡터화가 어려워 실용성에서는 Cache-Aware tiling이 더 널리 사용된다.

비교/분석

타일링 기법 비교

하드웨어별 타일 크기 예시

연산 재사용 분석

타일 크기 T일 때 연산 재사용률:

재사용률 = N / T

• T = 1 (타일링 없음): 재사용률 = N (최악)
• T = N (전체를 하나의 타일): 재사용률 = 1 (이론적 최적, 메모리 부족)
• T = √N: 재사용률 = √N (균형점)

동작 원리

기본 GEMM 알고리즘 (타일링 없음)

# 단순 GEMM - O(N³) 캐시 미스
for i in range(N):
    for j in range(N):
        C[i][j] = 0
        for k in range(N):
            C[i][j] += A[i][k] * B[k][j]  # 매번 메모리에서 로드

Square Tiling 적용 알고리즘

# Square Tiling GEMM - O(N³/T²) 캐시 미스
T = 32  # 타일 크기
for i in range(0, N, T):
    for j in range(0, N, T):
        C_tile = C[i:i+T, j:j+T]  # C 타일 로드 (T×T)
        for k in range(0, N, T):
            A_tile = A[i:i+T, k:k+T]  # A 타일 로드 (T×T)
            B_tile = B[k:k+T, j:j+T]  # B 타일 로드 (T×T)

            # 타일 곱셈 (이제 캐시에서 연산)
            for ii in range(T):
                for jj in range(T):
                    for kk in range(T):
                        C_tile[ii][jj] += A_tile[ii][kk] * B_tile[kk][jj]

        C[i:i+T, j:j+T] = C_tile  # 결과 저장

GPU 커널 구현

GPU에서는 워프(Warp) 단위로 타일을 처리한다:

// GPU GEMM 커널 (단순화)
__global__ void gemm_tiled(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE][TILE];  // 공유 메모리 A 타일
    __shared__ float Bs[TILE][TILE];  // 공유 메모리 B 타일

    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;

    int row = by * TILE + ty;
    int col = bx * TILE + tx;

    float sum = 0.0f;

    // 타일 루프
    for (int t = 0; t < N; t += TILE) {
        // 타일을 공유 메모리로 로드
        As[ty][tx] = A[row * N + (t + tx)];
        Bs[ty][tx] = B[(t + ty) * N + col];

        __syncthreads();

        // 타일 곱셈
        for (int k = 0; k < TILE; k++) {
            sum += As[ty][k] * Bs[k][tx];
        }

        __syncthreads();
    }

    C[row * N + col] = sum;
}

장단점

장점

1. 캐시 효율성 극대화: 타일 크기만큼의 데이터를 캐시에 적재한 후 여러 번 재사용하여 캐시 미스를 크게 감소시킨다
2. 구현 용이성: 다른 타일링 기법(예: Register Tiling)에 비해 알고리즘이 단순하고 직관적이다
3. 하드웨어 독립성: CPU, GPU, FPGA 등 다양한 플랫폼에서 동일한 원리로 적용 가능하다
4. 선형적 확장성: 타일 크기를 조절하여 다양한 캐시 크기에 대응할 수 있다
5. 텐서 코어 호환성: NVIDIA 텐서 코어의 WMMA/MMA 명령과 자연스럽게 결합된다

단점

1. 파라미터 튜닝 필요: 타일 크기는 하드웨어에 따라 최적 값이 달라지므로 수동 튜닝이 필요하다
2. 비정형 행렬 비효율: 행렬 차원이 타일 크기로 나누어떨어지지 않으면 패딩(padding)이 발생하여 연산 낭비가 있을 수 있다
3. 메모리 오버헤드: 타일을 공유 메모리나 레지스터에 로드하는 과정에서 추가적인 메모리 이동이 발생한다
4. 동기화 비용: GPU 커널에서는 __syncthreads() 같은 동기화 명령이 필요하여 오버헤드가 발생할 수 있다

관련 기술

프레임워크 및 라이브러리

• NVIDIA cuBLAS: Square Tiling을 기본으로 사용하는 GEMM 라이브러리
• NVIDIA CUTLASS: 템플릿 기반 GEMM 커널로 다양한 타일링 전략 지원
• OpenBLAS: CPU 기반 GEMM에서 L1/L2 캐시 타일링 구현
• Intel MKL: x86 프로세서에 최적화된 GEMM 커널
• OpenAI Triton: 타일링을 자동화하는 컴파일러 프레임워크

관련 개념

• GEMV (General Matrix-Vector): 벡터 차원에 대한 locality 최적화
• Strassen 알고리즘: O(N^2.81) 곱셈 횟수, 타일링과 결합 가능. 2×2 행렬 곱셈을 7번의 곱셈으로 줄이고 재귀적으로 적용
• Register Tiling: 레지스터 수준의 미세 타일링
• Loop Unrolling: 타일 내부 루프 최적화
• SIMD Vectorization: 타일 내 연산의 벡터화
• Cannon's Algorithm: 분산 메모리 환경에서 통신을 회피하는 2D 타일링 알고리즘
• AlphaTensor: DeepMind에서 발견한 행렬 곱셈 최적화 알고리즘으로, Strassen보다 더 적은 곱셈으로 n^2.778 복잡도 달성

참고 문헌

• "Optimizing Matrix Multiplication on GPUs" - NVIDIA Developer Blog
• "Anatomy of High-Performance Matrix Multiplication" - Goto & Van De Geijn (2008)
• "CUTLASS: CUDA Templates for Linear Algebra Subroutines" - NVIDIA
• "FlashAttention: Fast and Memory-Efficient Exact Attention" - Dao et al. (2022)
• "Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations" - Tillet et al. (2019)
• "Cache-Oblivious Algorithms" - Harald Prokop (MIT, 1999)
• "AlphaTensor: Discovering novel, efficient and exact matrix multiplication algorithms" - DeepMind (2022)

핵심 정리

1. Square Tiling은 큰 행렬을 작은 정사각형 타일로 분할하여 캐시 locality를 극대화하는 기본 GEMM 최적화 기법이다
2. 타일 크기는 하드웨어의 L1 캐시 크기에 따라 결정되며, 일반적으로 √(캐시 크기/3) 공식으로 계산한다
3. 타일링을 적용하면 캐시 미스가 O(N³)에서 O(N³/T²)로 감소하여, 타일 크기 T만큼의 데이터 재사용이 가능하다
4. GPU에서는 공유 메모리에 타일을 로드하고 워프 단위로 병렬 곱셈을 수행하는 방식으로 구현된다
5. Square Tiling은 구현이 용이하고 하드웨어 독립적이어서, cuBLAS, CUTLASS, Triton 등 다양한 라이브러리의 기본 전략으로 사용된다