Shared Memory Bank Conflict Test - padding 을 통해 bank conflict 회피 가능

kernel 에 임의 패딩을 포함 시키는 형태로  bank conflict 회피 하는 형태

__shared__ float sh[32][32];

int lane = threadIdx.x % 32;
// ...
float v = sh[lane][0];

/////////////////////////////

__shared__ float sh[32][33];

int lane = threadIdx.x % 32;
float v = sh[lane][0];

 

커널 시간을 측정하기 위한 ncu 명령어 사용

[conflict] kernel time: 34.195 ms  // 이 값은 "프로파일링 오버헤드 포함" 시간

l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum = 19,840
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum = 19,840
...
warp_issue_stalled_* 관련 비율은 0%


/////////////////////////////////////////////////////////////////////////////////

[padded]  kernel time: 31.299 ms  // 역시 프로파일링 오버헤드 포함

l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum = 0
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum = 0
...
warp_issue_stalled_* 관련 비율은 0%

임의의 padding 통해 bank conflict 0 달성 성공

 

하지만 실행 시간 측면에서 거의 동일했음

테스트 환경 및 세팅 규모가 충분히 크지 않아 confllict 차이에 따른 실행 시간 차이를 관측할 수 없었음