Memory Banks Conflict Question

Accelerated Computing CUDA CUDA Programming and Performance
#1 
__shared__ char shared[32];

char data = shared[BaseIndex + tid];

Above code is the example that is given in the CUDA documentation, however, I’m not clear on why is:

shared[1], shared[2], shared[3], shared[4];

Shares the same memory banks… Well, it says that’s because if “s” is odd, then it will conflict, but I’d like to see some type of proof…? I can’t see this from all of the examples given…

Thanks,

Grady

#2

Banks are 32-bit, so four chars will be packed into a single bank.

#3

On the same lines, would it be correct to say that a “double” would occupy 2 banks?

#4

Thanks! I got it… :)…

#5

Yes.

#6

That means that each char is 8 bits, so that’s why it’s being multiplied by 4 to get 32 for each of the char, am I correct?

#7

yes, chars are 1byte = 8bit values

#8

Thanks!