# constant and shared memoty

**URL:** <https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806>\
**Category:** CUDA Programming and Performance\
**Created:** [December 12, 2008, 6:20am UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806 "2008-12-12T06:20:16Z")\
**Posts on this page:** 9\
**Page:** 1

<div class="post-metadata">

**Author:** ![ykhong](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@ykhong](https://forums.developer.nvidia.com/u/ykhong)\
**Post date:** [December 12, 2008, 6:20am UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/1 "2008-12-12T06:20:16Z")

</div>

Hello

I have two problems about constant and shared memory on CUDA

(1) Can I copy a float “array” on host memory to device

```
  constant memory? If it can, how to do that? Can anyone give

   me a simple example copy a float array from host memory 

   to device constant memory?

```

(2) I want to create shared memory dynamically. Suppose I want to

```
  create two 2D float array on shared memory depend on a variable 

  like __shared__ float array1;

        __shared__ float array2;

  I know it has to use the dynamical creation from calling the kernel function,

  but I encountered some problems on compiling.

  Can anyone give me a very simple example to do that?

```

Thank you very much.

---

<div class="post-metadata">

**Author:** ![SPWorley](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@SPWorley](https://forums.developer.nvidia.com/u/SPWorley)\
**Post date:** [December 12, 2008, 7:19am UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/2 "2008-12-12T07:19:29Z")

</div>

Declaration:

**constant** float dFloats[1024];

in main():

float hFloats[1024]={123.4, 1.0} ; // fill hFloats on host…

CUDA\_SAFE\_CALL(cudaMemcpyToSymbol(dFloats, hFloats, 1024\*sizeof(float), 0, cudaMemcpyHostToDevice));

Your kernel can then read (but not write) from the dFloats array.

#2

No, you can’t dynamically allocate shared memory.  
But you could statically allocate a large chunk and just use whatever portion you like of it. If the dimensions keep changing dynamically, you could change to a 1D array and do your own array multiplies, ie array1[x+y\*xsize] .

---

<div class="post-metadata">

**Author:** ![Gimurk](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Gimurk](https://forums.developer.nvidia.com/u/Gimurk)\
**Post date:** [December 12, 2008, 7:30am UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/3 "2008-12-12T07:30:11Z")

</div>

(1) constant memory has nothing with the type of data element, just make sure that its declaration is placed outside any function, and it seem can’t be dynamically allocated, copying data between constant mem and other mem region should use

cudaMemcpyToSymbol(). for example:

[codebox] **constant** float cons[256];

unsigned float \*vals;

cudaMemcpyToSymbol(cons, vals, sizeof(float) \* 256);[/codebox]

(2) dynamically allocated array in shared memory by execution configuration should be one dimension, i suggest that you access it as the way to one-dimension array, not the way to 2D

---

<div class="post-metadata">

**Author:** ![ykhong](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@ykhong](https://forums.developer.nvidia.com/u/ykhong)\
**Post date:** [December 12, 2008, 8:43am UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/4 "2008-12-12T08:43:07Z")

</div>

I try the following code

but the answer is very strange

1.000000000

2.000000000

3.000000000

4.000000000

5.000000000

6.000000000

7.000000000

8.000000000

9.000000000

10.000000000

Does anyone have any idea about it?

[codebox] **global** void kernel(float \*devicearray, const float \*arrayD)

{

for(int loop=0; loop\<10; loop++)

\*(devicearray+loop) = \*(arrayD+loop)\*10;

return;

};

_constant_\_ float arrayD[100];

int main(void)

{

cudaError error;

float arrayH[10];

float \*devicearray;

for(int loop=0; loop\<10; loop++)

\*(arrayH+loop) = loop+1;

cudaMemcpyToSymbol(arrayD, arrayH, sizeof(float)\*10, cudaMemcpyHostToDevice);

cudaMalloc((void \*\*)&devicearray, sizeof(float)\*10);

kernel\<\<\<1,1\>\>\>(devicearray, arrayD);

error = cudaGetLastError();

if(error!=cudaSuccess)

printf(“%s\n”,cudaGetErrorString(error));

cudaMemcpy(arrayH, devicearray, sizeof(float)\*10, cudaMemcpyDeviceToHost);

for(int loop=0; loop\<10; loop++)

printf(“%10.9f\n”, \*(arrayH+loop));

cudaFree(devicearray);

return 0;

};[/codebox]

---

<div class="post-metadata">

**Author:** ![Reimar](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Reimar](https://forums.developer.nvidia.com/u/Reimar)\
**Post date:** [December 12, 2008, 3:03pm UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/5 "2008-12-12T15:03:46Z")

</div>

> [@](#):
>
> [codebox] **global** void kernel(float \*devicearray, const float \*arrayD)
> 
> …
> 
> _constant_\_ float arrayD[100];
> 
> …
> 
> kernel\<\<\<1,1\>\>\>(devicearray, arrayD);
> 
> [/codebox]

Can’t work, the compiler can’t know that the second function argument is in constant memory.

You have to use the “arrayD” global variable directly in your kernel code.

---

<div class="post-metadata">

**Author:** ![ykhong](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@ykhong](https://forums.developer.nvidia.com/u/ykhong)\
**Post date:** [December 12, 2008, 3:21pm UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/6 "2008-12-12T15:21:20Z")

</div>

> [@](#):
>
> Can’t work, the compiler can’t know that the second function argument is in constant memory.
> 
> You have to use the “arrayD” global variable directly in your kernel code.

I also use the globle variable directly

but it also gives the same result.

Can anyont help me for this problem?

I will deeply appricate fot you. Thanks.

---

<div class="post-metadata">

**Author:** ![frea1](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@frea1](https://forums.developer.nvidia.com/u/frea1)\
**Post date:** [December 12, 2008, 3:30pm UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/7 "2008-12-12T15:30:54Z")

</div>

> [@](#):
>
> No, you can’t dynamically allocate shared memory.
> 
> But you could statically allocate a large chunk and just use whatever portion you like of it. If the dimensions keep changing dynamically, you could change to a 1D array and do your own array multiplies, ie array1[x+y\*xsize] .

You can allocate ‘dynamic’ amount of used shared memory, before kernel execution. In other words array sizes don’t have to be hardcoded.

Examples are in the SDK, search for file sharedmem.cuh.

---

<div class="post-metadata">

**Author:** ![SPWorley](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@SPWorley](https://forums.developer.nvidia.com/u/SPWorley)\
**Post date:** [December 12, 2008, 4:02pm UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/8 "2008-12-12T16:02:31Z")

</div>

You’re not using the right arguments for the copy to symbol function. It specifies an offset (probably 0). You’re giving only 4 arguments to a 5 argument function.

See the CUDA reference manual for the function definition.

---

<div class="post-metadata">

**Author:** ![gshi](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@gshi](https://forums.developer.nvidia.com/u/gshi)\
**Post date:** [December 12, 2008, 5:06pm UTC](https://forums.developer.nvidia.com/t/constant-and-shared-memoty/6806/9 "2008-12-12T17:06:20Z")

</div>

You made a couple of mistakes

here is a modified version of your code that works

[codebox]\_\_global\_\_ void kernel(float \*devicearray, const float \*arrayD)

{

for(int loop=0; loop\<10; loop++)

\*(devicearray+loop) = \*(arrayD+loop)\*10;

return;

};

_constant_\_ float arrayD[100];

int main(void)

{

cudaError error;

float arrayH[10];

float \*devicearray;

for(int loop=0; loop\<10; loop++)

\*(arrayH+loop) = loop+1;

cudaMemcpyToSymbol(arrayD, arrayH, sizeof(float)\*10, cudaMemcpyHostToDevice);

cudaMalloc((void \*\*)&devicearray, sizeof(float)\*10);

kernel\<\<\<1,1\>\>\>(devicearray, arrayD);

error = cudaGetLastError();

if(error!=cudaSuccess)

printf(“%s\n”,cudaGetErrorString(error));

cudaMemcpy(arrayH, devicearray, sizeof(float)\*10, cudaMemcpyDeviceToHost);

for(int loop=0; loop\<10; loop++)

printf(“%10.9f\n”, \*(arrayH+loop));

cudaFree(devicearray);

return 0;

};[/codebox]

[/quote]