Shared memory bank ( warp 와 bank 의 갯수는 하드웨어-레벨에서 설계적으로 32개로 맞춘 것 )

SM(Streaming Multiprocessor) Shared Memory 는 32 개의 bank 로 나뉘어 있음

한 warp 가 동시에 shared memory 를 읽을 때, 

각 스레드가 서로 다른 bank 를 읽으면 - conflict 없음

여러 스레드가 같은 bank 를 서로 다른 주소로 읽으면 bank conflict

(여러 스레드가 같은 주소를 읽으면 broadcast, 예외적으로 빠름)

 

Shared Memory 내부 구조

Shared Memory 는 단일 덩어리가 아님

Shared Memory
 ├─ Bank 0
 ├─ Bank 1
 ├─ Bank 2
 ├─ ...
 ├─ Bank 31   ← 보통 32개 bank (warp 크기와 맞춤)

각 bank 는 한 번에 하나의 주소 요청만 처리 가능

 

같은 bank 에 여러 스레드가 동시에 접근

t0 → bank 0
t1 → bank 1
t2 → bank 2
...
t31 → bank 31   → 32개가 모두 다른 bank → 완전 병렬 OK

//////////////

t0 → bank 0
t1 → bank 0
t2 → bank 0
...
t31 → bank 0   → 전부 bank 0으로 몰림

 

 

bank 결정 방식은 보통

bank_id = (address / 4 bytes) % 32

float32 기준 : 주소가 4 바이트씩 증가하니 stride 패턴에 따라 같은 bank 로 모일 수 있다.

 

예를 들어 stride = 32 float 배열 접근

A[threadIdx.x * 32]

모든 thread 가 index * 32 위치 접근

모듈로 32 하면 항상 0

전부 bank 0

 

 

shared memory = 32 개의 통로, banks 로 구성

warp = 32 개의 스레드가 동시에 접근하는 팀

 

32개의 스레드가 같은 cycle 에 shared memory 에 접근한다.

 

이 32 스레드가 모두 다른 bank 로 들어가면 병렬로 단번에 요청 처리 가능

여러 스레드가 같은 bank 로 몰릴 시, 해당 bank 는 한 번에 한 요청만 수행, 직렬화됨

 

메모리 주소가 bank 를 결정한다

각 스레드가 읽는 shared memory 주소에 따라 어느 bank로 가는지가 결정된다.

대부분 GPU 에서 주소 흐름이 꺾여서 모이는 패턴이면, 여러 스레드가 같은 bank 에 몰린다.