Shared memory 을 선언한 순간! ( block, warp, bank, thread ) - 사용자 관점 실제 GPU 관점

사용자 관점 (CUDA 코드)

__shared__ float tile[32][32];

• 사용자 관점에서 보는건 block 단위의 2D 배열
• bank가 보이는 것이 아님,

하지만 내부에선 복잡한 하드웨어 동작이 시작됨

 

하드웨어 (SM 내부)에서 실제로 일어나는 일

Block 이 SM에 배정됨 : 이때 shared memory 공간이 block 전용으로 할당됨

shared memory 크기 계산 : tile 크기만큼 메모리 예약

Bank alignment 자동 적용 : 자동으로 32 Bank 로 분리됨 ( 사용자 설정 X )

접근 시 bank index 계산 : (address / 4) % 32 로 자동 계산됨

warp 실행 : warp 가 동시에 접근 - conflict 여부 자동 판단

 

Shared Memory 선언 이후 실제 동작 흐름 (HW 관점)

(1) Block assigned to SM  
    └─ SM 내부 shared memory 공간 확보  
(2) tile[][] 선언 크기만큼 주소 범위 예약  
(3) 주소가 자동으로 32 Bank 규칙에 따라 정렬됨  
(4) warp 실행 시 
      bank = (addr_in_bytes / 4) % 32  
      식으로 bank index 결정  
(5) conflict 발생 여부 판별  
(6) conflict가 없으면 병렬 접근  
    conflict 발생 시 warp-level 직렬화

 

Bank 구조 생성을 사용자가 직접 할 수는 없음

• 사용자는 bank 선언을 할 수 없음 -> HW 가 자동으로 split
• 사용자가 할 수 있는 건 index 패턴 설계 = bank conflict 회피
• 즉 bank 를 만드는 것이 아닌, bank 를 고려해서 접근하는 것이 핵심

 

결론 요약

shared memory 선언 = block 전용 버퍼를 만드는 것

이 순간 내부 GPU 에서는 

32 개의 Bank 구조로 자동 정렬되고, warp 가 접근하면 conflict 여부를 자동 판별한다.

 

개발자는 bank 를 생성하거나 제어할 수 없고, bank conflict 를 발생시키지 않는 접근 패턴만 설계하면 된다.