SHARED MEMORY BANK CONFLICT 퀴즈

🧩 Q1. bank id 계산

Warp thread 0~31이 shared memory에서 다음 index를 읽는다:

idx = tid * 4
bank = ?

❓ Q1-1. 각 tid가 접근하는 bank 번호를 0~31까지 적어라.

❓ Q1-2. conflict는 몇-way인가?

 

 

🧩 Q2. stride=8 bank mapping

다음 index 패턴이 있다고 하자:

idx = tid * 8

❓ Q2-1. bank = (tid * 8) % 32 일 때,

bank 값은 어떤 패턴으로 반복되는가?
(0, 8, 16, 24, 0, 8, 16, 24, … 이런 식으로 적으라는 뜻)

❓ Q2-2. 한 bank 당 몇 thread가 몰리는가?

❓ Q2-3. conflict는 몇-way인가?

 

🧩 Q3. 최악의 stride 구하기

Warp가 최악의 32-way bank conflict를 만드는 stride는 무엇인가?

❓ Q3. stride = ?

(“tid 0~31 모두 bank 0으로 가도록 하는 stride”를 구하라)

 

🧩 Q4. padding이 해결하는 이유

다음 shared memory 선언과 접근을 사용한다고 하자:

__shared__ float sh[32 * 33];

idx = tid * 33
bank = (idx % 32)

❓ Q4-1. tid=0~31이 각각 어떤 bank로 매핑되는가?

❓ Q4-2. conflict가 사라지는 이유를 한 줄로 설명하라.

 

 

🧩 Q5. 실험 결과 해석 (너의 출력 기반 문제)

너의 실행 결과는:

stride=1  → 6.23 ms
stride=2  → 6.24 ms
stride=4  → 9.64 ms
stride=8  → 17.15 ms
stride=16 → 29.85 ms
stride=32 → 57.92 ms

❓ Q5-1. 왜 stride=2는 stride=1과 거의 같았는가?

❓ Q5-2. 왜 stride=32는 가장 느렸는가?

❓ Q5-3. 왜 stride=8 → stride=16 → stride=32 로 갈수록 거의 “2배 가까이” 증가하는가?