0

I understand that when I declare a shared memory array in a kernel, the same sized array is declared by all the threads. A code like

__shared__ int s[5];

will create a 20 byte array in each thread. The way I understand addressing shared memory is that it is universal across all the threads. So, if I address subscript 10 as follows

s[10] = 1900;

it is the exact same memory location across all the threads. It won't be the case that different threads access different shared memory address for subscript 10. Is this correct? The compiler of course throws warnings that the subscript is out of range.

paleonix
  • 2,293
  • 1
  • 13
  • 29
gmemon
  • 2,573
  • 5
  • 32
  • 37
  • 1
    As the name suggests, shared memory arrays are in shared memory, not created in each thread. In other words, there is only one shared array `s`, and that's why there could be bank conflict when accessing the same address in shared memory by different threads. – chaohuang Jul 01 '12 at 01:08
  • 1
    Shared memory are shared by threads in a block, not all the threads of the kernel. It can be treated as a programmable cache. – konjac Jul 02 '12 at 01:46

1 Answers1

4

Actually it will create a 20-byte array per block, not per thread.

Every thread within the block will be able to access these 20 bytes. So if you need to have N bytes per thread, and a block with M threads, you'll need to create a N*M buffer per block.

In your case, if there was 128 threads, you would have had

__shared__ int array[5*128];

And array[10] would have been a valid address for any thread within the block.

paleonix
  • 2,293
  • 1
  • 13
  • 29
cedric
  • 71
  • 2