1

What is a bank conflict in devices with 2.x devices? As I understand the CUDA C programming guide, in 2.x devices, if two threads access the same 32 bit word in the same shared memory bank, it does not cause a bank conflict. Instead, the word is broadcasted. When the two threads write the same 32 bit word in the same shared memory bank, then only one thread succeeds.

Since on-chip memory is 64 KB (48 KB for shared memory and 16 KB for L1, or vice versa), and it is organized in 32 banks, I am assuming that each bank consists of 2 KB. So I think that bank conflicts will arise if two threads access two different 32 bit words in the same shared memory bank. Is this correct?

paleonix
  • 2,293
  • 1
  • 13
  • 29
gmemon
  • 2,573
  • 5
  • 32
  • 37
  • 1
    That's correct. I kind of felt like dinging your question with a -1 because bank conflicts are explained so thoroughly in the CUDA C Programming Guide. – Roger Dahl Jun 30 '12 at 17:10

1 Answers1

4

Your description is correct. There are many access patterns that can generate bank conflicts, but here's a simple and common example: strided access.

__shared__ int smem[512];

int tid = threadIdx.x;

x = smem[tid * 2]; // 2-way bank conflicts
y = smem[tid * 4]; // 4-way bank conflicts
z = smem[tid * 8]; // 8-way bank conflicts
// etc.

Bank ID = index % 32, so if you look at the pattern of addresses in the x, y, and z accesses, you can see that in each warp of 32 threads, for x, 2 threads will access each bank, for y, 4 threads will access each bank, and for z, 8 threads will access each bank.

harrism
  • 26,505
  • 2
  • 57
  • 88