# Shared Memory initialization **URL:** \ **Category:** CUDA Programming and Performance\ **Created:** [February 20, 2007, 6:49am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244 "2007-02-20T06:49:24Z")\ **Posts on this page:** 20\ **Page:** 1
**Author:** ![shyam.pm](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@shyam.pm](https://forums.developer.nvidia.com/u/shyam.pm)\ **Post date:** [February 20, 2007, 6:49am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/1 "2007-02-20T06:49:24Z")
Hi, I am working with CUDA on a windows platform. I had a query regarding initialization of an array on shared memory. If the array is initialized as follows :\_\_shared float Array[16 \* 16] it works fine. But when I try to pass in the size of the array as an argument to the kernel function (for example \_\_shared float Array [m] where m is passed in as an argument) it results in a build error: expression must have a constant value. Is it not possible to declare the size of the array on the fly? I tried using the extern \_\_shared float declaration but I am using quite a few variables and am facing memory overlap problems. I would greatly appreciate any help. Thanks, Shyam ---
**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\ **Post date:** [February 20, 2007, 10:59am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/2 "2007-02-20T10:59:57Z")
Hi Shyam, This is explained in the programming guide. Because it’s not possible to allocate memory on the fly inside **global** or **device** functions, you have to either allocate **shared** arrays statically, or pass the size through the execution configuration (the \<\<\< \>\>\> parameters). For the same reason, you can allocate at most one dynamic **shared** array per kernel invocation. Here’s some example code: ```auto __global__ void myKernel(float *deviceData) {    extern __shared__ float myData[];    // kernel code... } void main() {    // initialization, data download, etc.    int sharedMemSize = calcSharedMemSize(); // depends on the problem    myKernel<<>>(data);    // shutdown } ``` The key here is that you declared the **shared** array extern to tell the compiler it would be allocated elsewhere. Then at runtime the cuda runtime API allocates the shared memory based on the third parameter in the execution configuration. Because of this, only one dynamically-sized shared memory array per kernel is supported. Mark ---
**Author:** ![shyam.pm](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@shyam.pm](https://forums.developer.nvidia.com/u/shyam.pm)\ **Post date:** [February 20, 2007, 1:48pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/3 "2007-02-20T13:48:38Z")
> [@](#): > > Hi Shyam, > > This is explained in the programming guide.  Because it’s not possible to allocate memory on the fly inside **global** or **device** functions, you have to either allocate **shared** arrays statically, or pass the size through the execution configuration (the \<\<\< \>\>\> parameters).  For the same reason, you can allocate at most one dynamic **shared** array per kernel invocation.  Here’s some example code: > > ```auto > __global__ void myKernel(float *deviceData) > > { > >    extern __shared__ float myData[]; > >    // kernel code... > > } > > void main() > > { > >    // initialization, data download, etc. > >    int sharedMemSize = calcSharedMemSize(); // depends on the problem > >    myKernel<<>>(data); > >    // shutdown > > } > > ``` > > The key here is that you declared the **shared** array extern to tell the compiler it would be allocated elsewhere.  Then at runtime the cuda runtime API allocates the shared memory based on the third parameter in the execution configuration.  Because of this, only one dynamically-sized shared memory array per kernel is supported. > > Mark > > [snapback]161767[/snapback] Thanks Mark. I did try that too. Infact I tried giving the offsets myself as mentioned in the guide for dynamically allocated arrays but with the number of variables I am using, it wasn’t feasible. So I would have to statically allocate the size of the arrays. Thanks again Shyam ---
**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\ **Post date:** [February 20, 2007, 2:41pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/4 "2007-02-20T14:41:18Z")
OK, without any actual code to look at it’s hard to say what the problem is. Mark ---
**Author:** ![shyam.pm](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@shyam.pm](https://forums.developer.nvidia.com/u/shyam.pm)\ **Post date:** [February 20, 2007, 3:47pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/5 "2007-02-20T15:47:36Z")
ok this might give you a better idea… extern “C” void runTest(float \*data, float \*estimates, int n ) { // copy host memory to device …Memcpy… //allocate device memory to store results …cudaMalloc… // setup execution parameters dim3 grid(1,1,1); dim3 threads(3,n,1); // execute the kernel kernel\<\<\< grid, threads\>\>\>(d\_C, d\_A, d\_B,d\_RHS, n) ; n is the height of the array which is computed in my main cpp file by using the sizeof() function and passed into **global** void kernel( float\* C, float\* A, float\* B, float\* D, unsigned int n) { // access thread id const unsigned int tid = threadIdx.x; // 0-3 // warp id const unsigned int wid = threadIdx.y; // 0-n const unsigned int ltid = (wid \* blockDim.x) + tid; // access number of threads in this block const unsigned int num\_threads = blockDim.x \* blockDim.y; ``` __shared__ float As[183]; __shared__ float Bs[3]; __shared__ float Cs[183]; __shared__ float Ds[61]; __shared__ float CTs[183]; __shared__ float Ps[3 * 3]; __shared__ float Rs[3]; __shared__ float Ls[3*3]; __shared__ float Us[3*3]; __shared__ float Ys[3]; __shared__ float Xs[3]; ``` I cannot use the extern like you mentioned but if i try to allocate these variables as ``` __shared__ float As[n*3].....ans so on ``` I get the errors i mentioned earlier… Thanks Shyam ---
**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\ **Post date:** [February 21, 2007, 10:01am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/6 "2007-02-21T10:01:23Z")
(Note, please wrap code in code blocks using the buttons on the edit pane so that the forum will not reformat your code.) I see. As I said, there is no dynamic memory allocation inside CUDA kernel code. What you could do is declare one array and define offsets and macros to access “virtual” arrays. ```auto // define some macros to access your "virtual" arrays #define As(i) data[AsIndex + (i)] #define Bs(i) data[BsIndex + (i)] #define Cs(i) data[CsIndex + (i)] #define Ds(i) data[DsIndex + (i)] // etc... __global__ void myFunc(float *deviceDataA, float *deviceDataB); {    extern __shared__ float data[];    unsigned AsIndex = 0;    unsigned BsIndex = n    unsigned BsIndex = n+3;    unsigned CsIndex = 2*n+3;    unsigned DsIndex = 4*n+3;    // etc...    // now, when you need to access the "virtual" arrays, just use the macros:    As(threadIdx.x) = deviceDataA[threadIdx.x];    Bs(threadIdx.x) = deviceDataB[threadIdx.x];    // etc... } ``` Now you just allocate “data” at runtime by using the execution configuration. Mark ---
**Author:** ![shyam.pm](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@shyam.pm](https://forums.developer.nvidia.com/u/shyam.pm)\ **Post date:** [February 21, 2007, 10:02pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/7 "2007-02-21T22:02:55Z")
Hi, Thanks a ton. It worked perfectly. But there is a limit on the amount of shared memory I can allocate from the execution configuration right?? Is there a way I could get around this as I am trying to develop a generic program that could do the computation on any number of floating points. In my other post I have mentioned block sizes which I think would be the best solution for this. But I am confused about how I could use blocks with the number of variables I am using. Also there is a sequence of operations I am trying to perform on the data. Thanks Again, Shyam ---
**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\ **Post date:** [February 22, 2007, 1:08pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/8 "2007-02-22T13:08:54Z")
There is a physical limitation of 16KB per thread block on G80. You have to work within that physical limit. You can always re-use memory space later by re-using temporary storage for intermediates and loading more data later in the kernel to replace data you are finished with. To extend the earlier example: ```auto // define some macros to access your "virtual" arrays #define As(i) data[AsIndex + (i)] #define Bs(i) data[BsIndex + (i)] #define Cs(i) data[CsIndex + (i)] #define Ds(i) data[DsIndex + (i)] #define Temp(i) data[(i)] // etc... __global__ void myFunc(float *deviceDataA, float *deviceDataB,                                    float * deviceDataC, float * deviceDataD,                                    float *deviceDataOut); {   extern __shared__ float data[];   unsigned temp = 0;   unsigned AsIndex = n;   unsigned BsIndex = 2*n   // note that we restart Cs at n.  This means we can't use Cs and Ds at the same   // time as As, and Bs   unsigned CsIndex = n;       unsigned DsIndex = 2*n;   // etc...   int tid = threadIdx.x;   // now, when you need to access the "virtual" arrays, just use the macros:   As(tid) = deviceDataA[tid];   Bs(tid) = deviceDataB[tid];   __syncthreads();   // compute using As and Bs and store in temp   Temp(tid) = As(tid) + Bs(tid);   // now load in C and D:   Cs(tid) = deviceDataC[tid];   Ds(tid) = deviceDataD[tid];   __syncthreads();   // compute using temp, As, and Bs and write to output   deviceDataOut[threadIdx.x] = Cs(tid) * (Ds(tid) + Temp(tid)); } ``` Also, for constant data, consider using constant memory or texture memory. Mark ---
**Author:** ![xsyourpal](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@xsyourpal](https://forums.developer.nvidia.com/u/xsyourpal)\ **Post date:** [February 26, 2007, 12:38am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/9 "2007-02-26T00:38:12Z")
> [@](#): > > Hi, > > Thanks a ton. It worked perfectly. But  there is a limit on the amount of shared memory I can allocate from the execution configuration right?? Is there a way I could get around this as I am trying to develop a generic program that could do the computation on any number of floating points. In my other post I have mentioned  block sizes  which I think would be the best solution for this. But I am confused about how I could use blocks with the number of variables I am using. Also there is a sequence of operations I am trying to perform on the data. > > Thanks Again, > > Shyam > > [snapback]162501[/snapback] ---
**Author:** ![JohnH](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@JohnH](https://forums.developer.nvidia.com/u/JohnH)\ **Post date:** [February 26, 2007, 4:21pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/10 "2007-02-26T16:21:42Z")
I use a similar method that you guys might like: ```auto extern __shared__ float data[]; float *As = data; float *Bs = data+183; // 183 is size of As // etc ``` Then you can use “As”, etc as you’d expect. [External Image](http://hqnveipbwb20/public/style_emoticons/<#EMO_DIR#>/homestar.gif "Image hosted on another site. Click to open in a new tab.") ---
**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\ **Post date:** [February 27, 2007, 11:02am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/11 "2007-02-27T11:02:37Z")
Generally this should work, but I’ve had problems in the past with shared memory pointers. Perhaps those issues are fixed in the current beta release. Mark ---
**Author:** ![Cyril\_Zeller](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Cyril\_Zeller](https://forums.developer.nvidia.com/u/Cyril_Zeller)\ **Post date:** [February 28, 2007, 10:27am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/12 "2007-02-28T10:27:47Z")
Yes, this works in the current Beta release. Either this: ```auto extern __shared__ float data[]; __device__ func()     // some device or __global__ function {    float *As = data;    float *Bs = data+183; // 183 is size of As } ``` of this: ```auto extern __shared__ float data[]; __shared__ float *As = data; __shared__ float *Bs = data+183; // 183 is size of As ``` The programming guide is outdated in this regard. Section 4.2.2.4 should say: [i]Pointers in code that is executed on the device are supported as long as the compiler is able to resolve whether they point to either the shared memory space or the global memory space, otherwise they are restricted to only point to memory allocated or declared in the global memory space. Dereferencing a pointer either to global or shared memory in code that is executed on the host or to host memory in code that is executed on the device results in an undefined behavior, most often in a segmentation fault and application termination.[/i] ---
**Author:** ![paulius](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@paulius](https://forums.developer.nvidia.com/u/paulius)\ **Post date:** [March 4, 2007, 5:22pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/13 "2007-03-04T17:22:34Z")
> [@](#): > > Hi, > > I am working with CUDA on a windows platform. I had a query regarding initialization of an array on shared memory. If the array is initialized as follows :\_\_shared float Array[16 \* 16] it works fine. But when I try to pass in the size of the array as an argument to the kernel function (for example \_\_shared float Array [m] where m is passed in as an argument) it results in a build error: expression must have a constant value. Is it not possible to declare the size of the array on the fly? I tried using the extern \_\_shared float declaration but I am using quite a few variables and  am facing memory overlap problems.  I would greatly appreciate any help. > > Thanks, > > Shyam > > [snapback]161717[/snapback] One additional note, Shyam. Inability to allocate arrays on the fly as you tried (float Array[m]) is not unique to CUDA. You cannot do this in C or C++ since, as the compiler complained, array dimension in the declaration must be a constant. ---
**Author:** ![shyam.pm](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@shyam.pm](https://forums.developer.nvidia.com/u/shyam.pm)\ **Post date:** [March 6, 2007, 6:40am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/14 "2007-03-06T06:40:52Z")
> [@](#): > > One additional note, Shyam.  Inability to allocate arrays on the fly as you tried (float Array[m]) is not unique to CUDA.  You cannot do this in C or C++ since, as the compiler complained, array dimension in the declaration must be a constant. > > [snapback]166998[/snapback] Hi Paulius, Thanks for the note. I had never tried that in C/C++ but assumed it would work. I am not much of a programmer. :) Shyam ---
**Author:** ![seibert](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/seibert/32/520688_2.png) [@seibert](https://forums.developer.nvidia.com/u/seibert)\ **Post date:** [March 6, 2007, 3:17pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/15 "2007-03-06T15:17:15Z")
> [@](#): > > Hi Paulius, > > Thanks for the note. I had never tried that in C/C++ but assumed it would work. I am not much of a programmer. :) > > Shyam > > [snapback]167805[/snapback] The g++ compiler will let you declare arrays with sizes that are variables, even though it is not part of the C++ standard. This has confused a number of people I work with when they moved to a Solaris compiler which did not support that extension. ---
**Author:** ![jhanweck](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jhanweck](https://forums.developer.nvidia.com/u/jhanweck)\ **Post date:** [March 9, 2007, 5:56pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/16 "2007-03-09T17:56:13Z")
> [@](#): > > Yes, this works in the current Beta release. > > Either this: > > ```auto > extern __shared__ float data[]; > > __device__ func()     // some device or __global__ function > > { > >    float *As = data; > >    float *Bs = data+183; // 183 is size of As > > } > > ``` > > of this: > > ```auto > extern __shared__ float data[]; > > __shared__ float *As = data; > > __shared__ float *Bs = data+183; // 183 is size of As > > ``` Should this work for different data types? For instance. I’m having trouble with this technique, something like: ```auto extern __shared__ float shmem[]; __shared__ unsigned int *int_data= (unsigned int *) shmem + 0; __shared__ float *float_data= shmem + num_threads; ``` This works in the emulator, but on the card (GeForce 8800 GTS), the float data doesn’t reference properly. ---
**Author:** ![jhanweck](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jhanweck](https://forums.developer.nvidia.com/u/jhanweck)\ **Post date:** [March 22, 2007, 10:23pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/17 "2007-03-22T22:23:35Z")
When assigning a constant value to shared memory, is it better to let one thread do the work, or is there an optimization that occurs? For example: ```auto __shared__ float x; x = -1.0f;      // is this better? if (tid == 0) x = -1.0f;    // or this? ``` The ptx code for the straight assignment is just two instructions, while the if(tid) approach is 4 instructions. Will the straight assignment will get serialized, with each thread repeating the assignment in turn? Or is there some kind of special treatment for a constant assignment to shared memory? ---
**Author:** ![prkipfer](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@prkipfer](https://forums.developer.nvidia.com/u/prkipfer)\ **Post date:** [March 23, 2007, 10:13am UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/18 "2007-03-23T10:13:12Z")
The manual says that the result of concurrent writes to the same memory location is undefined. As long as all threads write the same value, you can probably get away with it, but you better stick with the second version as the threads might not be in sync and so a second thread might write the -1 while the first has already moved on modifying x further. If x isn’t modified further down, you should make it **constant**. Peter ---
**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\ **Post date:** [March 26, 2007, 5:26pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/19 "2007-03-26T17:26:36Z")
My hunch is that the second is better, but you need to follow it with a \_\_syncthreads() to avoid the issue Peter mentions. Also, you may be able to do this: ```auto if (tid < 32) x = -1.0f;     __syncthreads(); ``` This would mean the entire first warp writes the value, so that the first warp is non-divergent, unlike the (tid == 0) case. However, divergence in only one warp is probably not much faster overall than divergence in zero warps, unless you have very few warps per block (which is usually bad for efficiency anyway). Mark [/quote] ---
**Author:** ![jhanweck](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jhanweck](https://forums.developer.nvidia.com/u/jhanweck)\ **Post date:** [March 26, 2007, 8:04pm UTC](https://forums.developer.nvidia.com/t/shared-memory-initialization/244/20 "2007-03-26T20:04:10Z")
Thanks Mark and Peter. My question was specifically targeting whether there was some subtle optimization for initializing shared memory to a constant. Apparently there’s not. I suspect, then that the method where ALL threads write the (same) constant to one shared memory location would cause bank conflicts? I didn’t detect any, but it’s a pretty small part of my code. If it did, that would settle the issue. :P I left out the \_\_syncthreads() explicitly because I was initializing other shared variables. But point well taken!