Shared memory types

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I am currently doing the following to access my shared memory:
extern shared in sdata;
extern shared unsigned int usdata;

Now I have two declarations because I want to access fields as ints in one situation and unsigned int in another. Does this setup work without problems? Are the two declarations treated as pointers to the beginning of the shared memory? Is there a better way?

2

I never tried to have more than one extern and I think it is potentially dangerous as I don’t know how to tell the runtime how to bind the extern to the shared mem argument of the kernel launch. It is done somehow, so you better stick to the solution in the manual, that is do one extern and cast appropriate pointers.

extern __shared__ char sharedArray[];

int* sdata = (int*)sharedArray;

unsigned int* usdata = (unsigned int*)sharedArray;

Or make usdata point to some offset in the array if you don’t want it to overlap with sdata.

One final hint: always check the .ptx to make sure the compiler doesn’t drop the shared attribute on the int pointers.

Peter