3

How to allocate shared variables in CUDA? I have a kernel where data needs to be shared across threads belonging to a particular block. I need two shared variables named sid and eid. I use it like this:

extern __shared__ int sid, eid

but it is giving me an error that __shared__ variables cannot have external linkage.

paleonix
  • 2,293
  • 1
  • 13
  • 29
username_4567
  • 4,737
  • 12
  • 56
  • 92
  • possible duplicate of [allocating shared memory](http://stackoverflow.com/questions/5531247/allocating-shared-memory) – talonmies Aug 22 '12 at 05:38

1 Answers1

9

There are two ways to allocate shared memory : static and dynamic

1、static

  __shared__ int Var1[10]

2、dynamic : should add "extern" keyword

extern __shared__ int Var1[]

If you use dynamic way to allocate shared memory, you should set the shared memory size when you call the function. for example:
testKernel <<< grid, threads, size>>>(...)
the third parameter is the size of shared memory. In this way, all the shared memories start from the same address. If you want to define several shared memory variables, you should write code like following.

__global__ void func(...)
{
    extern __shared__ char array[];
    short * array0 = (short*)array;
    float * array1 = (float*)(&array0[128]);
}
Ricky Dev
  • 45
  • 1
  • 6
Samuel
  • 5,977
  • 14
  • 55
  • 77
  • Now if this function is __device__ function and i'm calling it from some __global__ function then how to achieve this thing? – username_4567 Aug 22 '12 at 06:32
  • you can pass the shared memory adress to the device function – Samuel Aug 22 '12 at 06:33
  • @Samuel if you declare your shared memory outside the global function, it should be visible in by other __device__ functions in the file, no ? – Seltymar Oct 10 '12 at 07:28
  • I think the ansower is yes.You can write a testbed to test it. I think it's not good to declare shared memory outside the global function – Samuel Oct 11 '12 at 02:01
  • How is this possible to have ```float** array``` instead of ```float* array0``` and ```float* array1```? In this case the storage of ```float**``` should also be assigned dynamically. – Mohammad Jul 26 '17 at 17:36