Shared Memory - Bank Conflict

Shaered Memory 의 구조

Bank 개수 : 32 개, 각 warp 의 thread 수와 동일

Bank 너비 : 4bytes(float 기준), 주소 % 32 - bank 결정

Access 단위 : Warp 단위 ( 32 threads), conflict / broadcast 판단은 warp 내부에서만 발생

 

Bank 매핑 규칙

bank = (address_in_bytes / 4) % 32 // float 기준

• thread 는 thread 단위로 동작하는 것처럼 보이지만, shared memory 접근은 warp 단위로 동시에 발생한다.
• 따라서 bank conflict 여부는 thread 가 아니라 warp 내부 패턴으로 결정된다.

 

Access 패턴 시나리오

Broadcast (Conflict 없음)

__shared_- float sh[1024];
float x = sh[100];

thraed_id (tid) [0~31], addr - 100, bank - 4, 결과 - Broadcast

같은 주소 - GPU가 1번만 읽고 warp 전체에 값 전달, conflict X (오히려 최적)

 

Bank Conflict (32-way)

int td = threadIdx.x;
float x = sh[tid * 32];

서로 다른 주소 + 같은 bank -> warp 가 직렬적으로 32번 처리

32-way bank conflict 

 

Bank Conflict 발생 조건

같은 bank + 다른 주소 (in one warp) → bank conflict   
같은 bank + 같은 주소 → broadcast        (OK)