Shared Memory Bank Conflict test

세 가지 패턴 별 접근

• broadcast
  • Nsight Compute 상에서 shared load 관련 warning 없음
  • hardward broadcast 로 처리 - conflict free
• stride - 1 ( 완전 conflict-free)
  • shared memory 를 고루 배치
  • 가장 이상적 패턴
• stride - 32 ( 의도적 conflict)
  • uncoalesced shared accesses 로 분류

 

실전 GEMM / Conv에서의 의미

• shared tiling 설계 시:
  • thread → smem index 매핑이 bank conflict를 만들지 않도록
  • laneId, warpId 기반으로 index를 재배치해야 한다.
• warp-level tiling / Tensor Core mma를 쓸 때도
  • shared layout을 잘못 잡으면
  • cp.async / ld.shared 가 bank conflict 때문에 병목이 된다.
• Nsight Compute는
  • Memory Workload Analysis Tables
  • L1 Wavefronts Shared Excessive
  • Source Counters (uncoalesced shared accesses)
  • 에서 이런 문제를 정확히 잡아주므로, bank conflict 튜닝 때 필수 도구가 된다