13_Shared_Memory_Maximizing_CUDA

Shared Memory

 

• Memory Bank
  • 데이터의 접근 통로
  • Shared Memory는 32개의 bank로 구성되어 있음
  • 구성단위: 4 byte or 8 byte (compute capability 3.x ~이상)
  • bank는 서로 독립적이며, 서로 다른 bank는 동시 접근 가능
• Bank Conflict
  • 하나의 bank에 여러 thread가 접근하는 경우 하나의 thread씩 처리할 수 있다.
  • 즉, thread의 직렬화 발생! → 성능 저하

 

 

---

No Bank Conflict

성능 저하

CUDA를 사용할때 Bank Conflict를 조심해야 한다.

 

 

 

 

---

Example

Bank Conflict 발생

이전글에서

row = threadIdx.x 일 경우 Bank Conflict 발생!

이를 해결해주기 위해 Global Memory에서 Shared Memory로 옮길때

Transpose를 해주면서 넘겨주면 해결이 됩니다.

 

 

 

 

 

출처:

https://www.youtube.com/watch?v=hyCEFNbumOo&list=PLBrGAFAIyf5pp3QNigbh2hRU5EUD0crgI&index=25