# Dynamic Shared memory Allocation

**URL:** https://forums.developer.nvidia.com/t/dynamic-shared-memory-allocation/21671
**Category:** CUDA Programming and Performance
**Created:** [March 23, 2011, 3:08pm UTC](https://forums.developer.nvidia.com/t/dynamic-shared-memory-allocation/21671 "2011-03-23T15:08:18Z")
**Posts on this page:** 4
**Page:** 1

<div class="post-metadata">

### Author: ![gzaki](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@gzaki](https://forums.developer.nvidia.com/u/gzaki)
#### Post date: [March 23, 2011, 3:08pm UTC](https://forums.developer.nvidia.com/t/dynamic-shared-memory-allocation/21671/1 "2011-03-23T15:08:18Z")

</div>

Dear All,

I am trying to dynamically allocate multiple arrays that will reside on the shared memory. Here is a code snippet:

**global** void kernel\_function(char \*Oarray1, float \*Oarray3, short \*Oarray2, int a) {

```
 extern __shared__ char array[];

 char *array1 = (char*) array;
 short *array2 = (short*) ALIGN((char*) (&array1[a]), sizeof(short));
 float *array3 = (float*) ALIGN((char*) (&array2[a]), sizeof(float));

```

Where ALIGN is a device function that will take care of the alignment. I launch the kernel using:

kernel\_function\<\<\<dimGrid, dimBlock, (size\_t) shared\>\>\>(ddata1, ddata3, ddata2, a);

I get the following error:

ptxas /tmp/tmpxft\_00006b0d\_00000000-2\_dynamic-shared.ptx, line 0; fatal : (C6017) Unaligned access for SMEM (unknown symbol) in entr  
y \_Z15kernel\_functionPcPfPsi; the offset should be 4-byte aligned.

But if I replace a with a constant, the program works fine.

Any ideas or code example that shows how can I perform these steps where the size of the shared memory array is not know till runtime?

Thanks.

---

<div class="post-metadata">

### Author: ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)
#### Post date: [March 23, 2011, 8:20pm UTC](https://forums.developer.nvidia.com/t/dynamic-shared-memory-allocation/21671/2 "2011-03-23T20:20:37Z")

</div>

I have the feeling that the problem might somehow be related to the way how [font=“Courier New”]ALIGN[/font] takes care of the alignment. Can you post it’s definition?

Also, does the calculation of [font=“Courier New”]shared[/font] take alignment into account?

---

<div class="post-metadata">

### Author: ![tera](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@tera](https://forums.developer.nvidia.com/u/tera)
#### Post date: [March 23, 2011, 8:31pm UTC](https://forums.developer.nvidia.com/t/dynamic-shared-memory-allocation/21671/3 "2011-03-23T20:31:08Z")

</div>

By the way: You can get rid of ALIGN if you just order the arrays to have decreasing natural alignment:

```auto
extern __shared__ char array[];

float *array3 = (float*) array;

short *array2 = (short*) &array3[a];

char *array1 = (char*) &array2[a];

```

---

<div class="post-metadata">

### Author: ![gzaki](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@gzaki](https://forums.developer.nvidia.com/u/gzaki)
#### Post date: [March 25, 2011, 2:58pm UTC](https://forums.developer.nvidia.com/t/dynamic-shared-memory-allocation/21671/4 "2011-03-25T14:58:48Z")

</div>

Ordering the array works.

Thanks.