⚡ AI Optimization

뱅크 충돌 방지

개요

뱅크 충돌(Bank Conflict)은 GPU 공유 메모리(Shared Memory)에서 동일한 뱅크에 있는 여러 주소에 대한 동시 접근이 발생할 때 발생하는 성능 병목 현상이다. NVIDIA GPU의 공유 메모리는 32개의 메모리 뱅크로 구성되어 있으며, 인접한 32비트 단어가 순차적으로 다른 뱅크에 할당된다. 동일한 뱅크에 여러 스레드가 동시에 접근하면 serialization이 발생하여, N-way 충돌 시 N배의 지연 시간이 추가된다.

뱅크 충돌은 GPU 커널 성능에 직접적인 영향을 미치며, 특히 GEMM, 컨볼루션, 어텐션 등 메모리 집약적 워크로드에서 두드러진다. NVIDIA A100/H100 아키텍처에서는 공유 메모리 대역폭이 ~19 TB/s에 달하지만, 뱅크 충돌이 발생하면 실질 대역폭이 수십 배까지 저하될 수 있다.

핵심 개념

공유 메모리 뱅크 아키텍처

NVIDIA GPU의 공유 메모리는 32개의 뱅크(bank)로 분할되어 있다:

항목	설명
뱅크 수	32개 (Compute Capability 7.0 이상)
뱅크당 버스 폭	4 Bytes (32 bits)
동시 접근	뱅크당 1 사이클에 1개 접근
스트라이드 매핑	주소 % 32 = 뱅크 번호

접근 패턴 예시:

스레드 0: 주소 0  → 뱅크 0
스레드 1: 주소 1  → 뱅크 1
스레드 2: 주소 2  → 뱅크 2
...
스레드 31: 주소 31 → 뱅크 31
스레드 32: 주소 32 → 뱅크 0 (스트라이드 패턴)

뱅크 충돌 유형

충돌 유형	설명	사이클 수	예시
No Conflict	모든 접근이 다른 뱅크에 분산	1 사이클	스트라이드=1
2-way Conflict	동일 뱅크에 2개 접근	2 사이클	스트라이드=2
N-way Conflict	동일 뱅크에 N개 접근	N 사이클	스트라이드=N
Broadcast	동일 주소에 대한 동시 읽기	1 사이클	모든 스레드 같은 주소

스트라이드 기반 충돌 분석

공유 메모리 접근 패턴에 따른 뱅크 충돌:

패턴 1: 스트라이드 1 (No Conflict)
Thread[i] → smem[i]
뱅크: 0, 1, 2, ..., 31 → 모두 다른 뱅크

패턴 2: 스트라이드 2 (2-way Conflict)
Thread[i] → smem[i * 2]
뱅크: 0, 2, 4, ..., 30, 0, 2, ... → 짝수 뱅크만 사용

패턴 3: 스트라이드 32 (No Conflict)
Thread[i] → smem[i * 32]
뱅크: 0, 0, 0, ..., 0 → 모든 스레드가 뱅크 0에 접근 (32-way 충돌!)

비교/분석

방지 기법 비교

기법	원리	오버헤드	유연성	범용성
패딩 (+4)	행 끝에 더미 요소 추가	SMEM 12.5% 낭비	낮음	높음
스위즐 (Swizzle)	비트 연산으로 뱅크 매핑 변경	연산 오버헤드	높음	중간
프로그래밍 가능한 스위즐링	하드웨어 기반 주소 변환	하드웨어 지원	높음	낮음
전치 (Transpose)	메모리 재배치	별도 전치 커널	중간	중간

패딩 기법 상세

가장 단순하고 널리 사용되는 기법으로, 각 행 끝에 더미 요소를 추가하여 다음 행의 시작 뱅크를 이동시킨다:

// 패딩 전: 뱅크 충돌 발생
__shared__ float smem[32][32];  // 32×32, 4바이트 단위

// 패딩 후: 뱅크 충돌 회피
__shared__ float smem[32][33];  // 32×33, 각 행 끝에 +1 요소

패딩 효과:
- 행 0: 뱅크 0~31 사용
- 행 1: 뱅크 1~32 사용 (뱅크 31 다음은 뱅크 0)
- 행 2: 뱅크 2~33 사용
- 모든 행이 서로 다른 뱅크 오프셋으로 시작

SMEM 사용량 변화:
- 패딩 전: 32×32×4 = 4,096 Bytes
- 패딩 후: 32×33×4 = 4,224 Bytes (+3.125%)

스위즐 기법 상세

스위즐(Swizzle)은 비트 연산을 통해 주소를 변환하여 뱅크 매핑을 변경하는 기법이다:

// 기본 스위즐: XOR 연산
int bank = (row ^ col) % 32;

// 또는 비트 시프트 스위즐
int bank = ((row << 2) ^ col) % 32;

스위즐 장점:
- SMEM 낭비 없음
- 복잡한 접근 패턴에서도 뱅크 충돌 회피 가능
- 하드웨어 지원 시 오버헤드 최소화

프로그래밍 가능한 스위즐링 (Hopper+)

NVIDIA Hopper 아키텍처부터 도입된 기능으로, 런타임에 스위즐 패턴을 조정할 수 있다:

// Hopper H100에서 지원
__shared__ uint32_t smem[32][32];
// LDMATRIX 명령과 결합하여 자동 스위즐 적용

동작 원리

GEMM 커널에서의 뱅크 충돌 회피

전형적인 GEMM 커널에서 공유 메모리 타일 로딩 시 뱅크 충돌 발생 패턴과 회피 방법:

// 뱅크 충돌 발생 예시
__global__ void gemm_conflict(float* A, float* B, float* C) {
    __shared__ float smem_A[32][32];  // 32×32 타일
    __shared__ float smem_B[32][32];

    int tid = threadIdx.x;

    // 타일 로딩: 뱅크 충돌 발생
    for (int i = tid; i < 32*32; i += blockDim.x) {
        int row = i / 32;
        int col = i % 32;
        smem_A[row][col] = ...;  // 같은 열에 대한 접근 시 뱅크 충돌
    }
}

// 뱅크 충돌 회피 패딩 적용
__global__ void gemm_padded(float* A, float* B, float* C) {
    __shared__ float smem_A[32][33];  // +1 패딩
    __shared__ float smem_B[32][33];

    int tid = threadIdx.x;

    // 타일 로딩: 뱅크 충돌 회피
    for (int i = tid; i < 32*32; i += blockDim.x) {
        int row = i / 32;
        int col = i % 32;
        smem_A[row][col] = ...;  // 패딩으로 뱅크 충돌 회피
    }
}

컨볼루션 커널에서의 뱅크 충돌

컨볼루션 연산에서 필터와 입력 데이터의 접근 패턴이 뱅크 충돌을 유발하는 경우:

// Winograd 컨볼루션에서의 뱅크 충돌 회피
__shared__ float tile[4][4][4];  // 3D 타일

// 패딩 적용: 각 차원에 +1
__shared__ float tile_padded[4][4][5];  // 마지막 차원 +1

어텐션 메커니즘에서의 뱅크 충돌

FlashAttention 등 어텐션 커널에서 Q/K/V 행렬의 공유 메모리 접근 시:

// Q 행렬 로딩 시 뱅크 충돌 회피
__shared__ float Q[TILE][TILE + 1];  // 행 패딩

// K/V 전치 접근 시 뱅크 충돌 회피
__shared__ float K[TILE + 1][TILE];  // 열 패딩

장단점

장점

성능 향상: 뱅크 충돌을 제거하면 공유 메모리 접근이 매 사이클 1회로 최적화되어, 이론적 대역폭에 근접하는 성능 달성
일관된 지연: 뱅크 충돌로 인한 불확실한 지연 시간이 제거되어, 커널 실행 시간의 예측 가능성 향상
점유율 유지: 뱅크 충돌 회피를 위한 추가 리소스 없이 기존 점유율을 유지하면서 성능 개선
호환성: 패딩 기법은 모든 CUDA 아키텍처에서 적용 가능하며, 하드웨어 의존성이 낮음
낮은 구현 복잡도: 패딩은 선언 시 +1만 추가하면 되며, 스위즐은 XOR 연산 몇 줄로 구현 가능

단점

SMEM 공간 낭비: 패딩 기법은 공유 메모리 사용량을 3~12% 증가시켜, 점유율(Occupancy)에 영향을 미칠 수 있음
복잡한 패턴 회피 어려움: 단순 스트라이드 패턴은 쉽게 회피할 수 있지만, 복잡한 2D/3D 접근 패턴은 추가적인 분석이 필요
하드웨어 의존적 최적화: 스위즐 기법은 GPU 아키텍처에 따라 최적 패턴이 달라짐
전치 오버헤드: 메모리 전치 기법은 별도의 커널 실행이 필요하여 전체 파이프라인 지연 증가
디버깅 어려움: 뱅크 충돌은 Nsight Compute로만 확인 가능하며, 소스 코드 수준에서 직관적으로 파악하기 어려움