# Am I using the constant memory properly?

**URL:** <https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414>\
**Category:** CUDA Programming and Performance\
**Created:** [May 14, 2013, 1:16pm UTC](https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414 "2013-05-14T13:16:12Z")\
**Posts on this page:** 5\
**Page:** 1

<div class="post-metadata">

**Author:** ![Florimond](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Florimond](https://forums.developer.nvidia.com/u/Florimond)\
**Post date:** [May 14, 2013, 1:16pm UTC](https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414/1 "2013-05-14T13:16:12Z")

</div>

I read several old topics pretending one can’t dynamically allocate constant memory, even from the host, as this guy wants to do the same thing that I do → [http://stackoverflow.com/questions/271273/dynamic-allocation-of-constant-memory-in-cuda](http://stackoverflow.com/questions/271273/dynamic-allocation-of-constant-memory-in-cuda)

The typical answer is that one should texture memory instead. Anyway, I ran the following test:

```auto
__constant__ int* c;
__global__ void kernel()
{
	/* UPDATE: ok, the following works, suggesting it's not constant memory... */
	c[4] = 666;

	for (int i = 0; i < 10; ++i)
	{
		printf("%d\n", c[i]);
	}
}

int main()
{
	int host_c[10] = {0,1,2,3,4,5,6,7,8,9};
	int* dev_c;

	cudaMalloc(&dev_c, 10 * sizeof(int));
	cudaMemcpy(dev_c, host_c, 10 * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(c, &dev_c, sizeof(int*));
	
	kernel<<<1,1>>>();
}

```

This seems to work, as I get no error and my 10 numbers are correctly printed.

I know from the book “Programming massively parallel processors” that the constant memory is just global memory that is subsequently cached.

My questions are:

- Just because the "c" pointer is in constant memory, is the area of global memory that c end up pointing automatically considered as "constant memory" by the device? Or am I naive and doing it wrong?

**Update:** Sorry but I just thought now about this basic test answering this first question: I can modify the data pointed by c. This suggest these data are not considered as constant memory! Ok but my next question still worth answering…

- Is it really the same to use texture memory instead of constant memory? If it is, why did they invent this "constant memory"?

I need to load data from a file. Those data may use more that 65535 bytes, so I actually need an dynamically allocated vector of dynamically allocated constant memory banks to hold those data. But once those data are on the device, they won’t change, ever…

**Update:** Sorry but I just thought now about this basic test answering my first question: I can modify the data pointed by c. This suggest these data are not considered as constant memory! Ok but my second question still worth answering…

---

<div class="post-metadata">

**Author:** ![wlangdon](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@wlangdon](https://forums.developer.nvidia.com/u/wlangdon)\
**Post date:** [May 14, 2013, 2:10pm UTC](https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414/2 "2013-05-14T14:10:48Z")

</div>

May be have a look at previous discussions about “constant” memory

> **[function parameter vs constant memory](https://forums.developer.nvidia.com/default/topic/481176/function-parameter-vs-constant-memory/)**
>
> I have some kernels,they have same parameter like width,height,pitch,addr make parameter list very long. to reduce kernel function parameter count,can i put them to constant memory.how to do it? thanks to any advice.

Bill

---

<div class="post-metadata">

**Author:** ![Florimond](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Florimond](https://forums.developer.nvidia.com/u/Florimond)\
**Post date:** [May 14, 2013, 2:45pm UTC](https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414/3 "2013-05-14T14:45:51Z")

</div>

Thank you for your answer, Bill.  
I guess you suggested this discussion mainly because of its last answer:

> [@](#):
>
> There are definite problems with constant memory. For example, it really does matter how many threads in a warp try to read different words of the constant memory simultaneously. Essentially if its more than one, the threads are forced to run sequentially, whereas if they all read the same word they all do it in parallel.

This is definitely an important issue to me, as each of my threads has to access a different float from a vector…

---

<div class="post-metadata">

**Author:** ![Florimond](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Florimond](https://forums.developer.nvidia.com/u/Florimond)\
**Post date:** [May 14, 2013, 3:08pm UTC](https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414/4 "2013-05-14T15:08:10Z")

</div>

To be precise, each of my thread just transfer a float from a vector currently held in global memory (and which will never be modified) to another vector also in global memory, without doing any computation! It’s just an initalization operation of a vector for further computations by other threads. I was wondering if I could speed up at least the reading by using constant memory. Apparently not.  
I definitely need more information about concurent reading and writing in global memory.  
I just begun to read this → [url][https://developer.nvidia.com/content/how-access-global-memory-efficiently-cuda-cc-kernels[/url]](https://developer.nvidia.com/content/how-access-global-memory-efficiently-cuda-cc-kernels%5B/url%5D)

Update: I just checked, I could even skip the writing to initialize the other vector in global memory, but I definitely need the fastest concurent reading achievable.  
Sorry, I didn’t write the original non parallel code…

Anyway, my second question still worth answering: why did they invent constant memory if you can do at least the same with texture memory?

---

<div class="post-metadata">

**Author:** ![cbuchner1](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/cbuchner1/32/14168_2.png) [@cbuchner1](https://forums.developer.nvidia.com/u/cbuchner1)\
**Post date:** [May 15, 2013, 9:40am UTC](https://forums.developer.nvidia.com/t/am-i-using-the-constant-memory-properly/29414/5 "2013-05-15T09:40:49Z")

</div>

Constant memory is a convenience for the programmer. In the simplest case it’s just adding **constant** and one cudaMemcpyToSymbol() call.

Textures usually require more setup, but they are not limited to 64kb page sizes and they also require reads through tex…1D/2D/3D() calls. Also they support interpolation, coordinate clamping and other goodies.