# cudaMalloc and cudaMemcpy for pointer in struct

**URL:** <https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717>\
**Category:** CUDA Programming and Performance\
**Created:** [November 22, 2014, 9:37pm UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717 "2014-11-22T21:37:56Z")\
**Posts on this page:** 8\
**Page:** 1

<div class="post-metadata">

**Author:** ![akoutmos](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@akoutmos](https://forums.developer.nvidia.com/u/akoutmos)\
**Post date:** [November 22, 2014, 9:37pm UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/1 "2014-11-22T21:37:56Z")

</div>

I am working on creating a ray tracer that supports multiple triangular meshes and I am having issues with getting some of my data structures moved over to the GPU. Right now I am using tinyobjloader to load OBJ files with multiple meshes into my rendering program. Everything looks good CPU side, but when I try to move everything over to the GPU and run my kernel, NSight shows that I have quite a few memory access violations. Here is the offending data structure and variables of interest (sorry for being over verbose. I figure more details up front makes things clearer):

```auto
__device__ Sphere* dev_spheres;
__device__ Plane* dev_planes;
__device__ BoundingBox* dev_boundingVolumes;
__device__ Triangle** dev_triangles; //A sparse 2d array containing all the triangles that make up each mesh in the OBJ scene {ex. {Array of triangles for mesh 1}, {Array of triangles for mesh 2}, {Array of triangles for mesh 3}, ...}

__device__ RenderData* dev_renderInfo;

__device__ Camera dev_cam;
__device__ LightSource dev_light;

struct RenderData
{
	int sphereCount; 
	int planeCount;  
	int triangularMeshes; //How many triangular meshes are in the scene as described in the OBJ file
	int* trianglesPerMesh; //An array of size 'triangularMeshes' listing how many triangles are in each mesh (ex. {996, 1024, 673...})
};

```

Here are some of the memory allocations that I am performing:

```auto
//----These work and I am able to view the struct vars with NSight when a breakpoint is triggered:----
HANDLE_ERROR( cudaMemcpyToSymbol( dev_light, temp_l, sizeof(LightSource) ) );
HANDLE_ERROR( cudaMemcpyToSymbol( dev_cam, temp_c, sizeof(Camera) ) );

//----These are the ones giving me problems:----
//Malloc and memcpy spheres
HANDLE_ERROR( cudaMalloc((void**)&dev_spheres, sizeof(Sphere) * SPHERE_COUNT) );
HANDLE_ERROR( cudaMemcpy(dev_spheres, temp_s, sizeof(Sphere) * SPHERE_COUNT, cudaMemcpyHostToDevice) );
free( temp_s );

//Malloc and memcpy planes
HANDLE_ERROR( cudaMalloc((void**)&dev_planes, sizeof(Plane) * PLANE_COUNT) );
HANDLE_ERROR( cudaMemcpy(dev_planes, temp_p, sizeof(Plane) * PLANE_COUNT, cudaMemcpyHostToDevice) );
free( temp_p );

//Malloc and memcpy triangle
HANDLE_ERROR( cudaMalloc((void**) &dev_triangles, sizeof(Triangle*) * temp_data->triangularMeshes) );
for(int i = 0; i < temp_data->triangularMeshes; i++)
{
	char* offset = (((char*)dev_triangles) + (sizeof(Triangle*) * i));
	HANDLE_ERROR( cudaMalloc((void**) &offset, sizeof(Triangle) * temp_data->trianglesPerMesh[i]) );
	HANDLE_ERROR( cudaMemcpy(offset, temp_t[i], sizeof(Triangle) * temp_data->trianglesPerMesh[i], cudaMemcpyHostToDevice) );

	free(temp_t[i]);
}
free(temp_t);

	
//Malloc and memcpy render data
HANDLE_ERROR( cudaMalloc((void**) &dev_renderInfo, sizeof(RenderData)) );
char* planeOffset = (((char*)dev_renderInfo) + offsetof(RenderData, planeCount));
HANDLE_ERROR( cudaMemcpy(planeOffset, &temp_data->planeCount, sizeof(int), cudaMemcpyHostToDevice) );

char* sphereOffset = (((char*)dev_renderInfo) + offsetof(RenderData, sphereCount));
HANDLE_ERROR( cudaMemcpy(sphereOffset, &temp_data->sphereCount, sizeof(int), cudaMemcpyHostToDevice) );

char* triOffset = (((char*)dev_renderInfo) + offsetof(RenderData, triangularMeshes));
HANDLE_ERROR( cudaMemcpy(triOffset, &temp_data->triangularMeshes, sizeof(int), cudaMemcpyHostToDevice) );
	
char* triCountOffset = (((char*)dev_renderInfo) + offsetof(RenderData, trianglesPerMesh));
HANDLE_ERROR( cudaMalloc((void**) &(triCountOffset), sizeof(int) * temp_data->triangularMeshes) );
HANDLE_ERROR( cudaMemcpy(triCountOffset, temp_data->trianglesPerMesh, sizeof(int) * temp_data->triangularMeshes, cudaMemcpyHostToDevice) );

free(temp_data->trianglesPerMesh);
free(temp_data);

```

Some questions about my code:

I think my main problem is a lack of fulling understanding the cudaMemcpy and cudaMalloc calls and misusing them. I don’t think offseting the malloc call for the pointer in the struct is the right way of doing it, but I am not really sure the best way to do this. Also is **device** really needed for specifying device allocated variables? As a side note I am on using CUDA runtime 6.0. Should i be using unified memory instead of explicitly making all the memory allocations myself?

Thank you in advance for the help. It is much appreciated :).

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [November 22, 2014, 9:55pm UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/2 "2014-11-22T21:55:41Z")

</div>

Yes, you are misusing cudaMalloc. You cannot use cudaMalloc directly on a **device** variable:

```auto
HANDLE_ERROR( cudaMalloc((void**)&dev_spheres, sizeof(Sphere) * SPHERE_COUNT) );

```

It is illegal to take the address of a device variable (&dev\_spheres) in host code. Furthermore, the pointer that you pass to cudaMalloc must point to a location on the host. You could do something like this if you really wanted to:

```auto
Sphere *h_dev_spheres;
HANDLE_ERROR( cudaMalloc((void**)&h_dev_spheres, sizeof(Sphere) * SPHERE_COUNT) );
HANDLE_ERROR( cudaMemcpyToSymbo(dev_spheres, &h_dev_spheres, sizeof(Sphere *) );

```

but I wouldn’t recommend it for multiple reasons relating to unnecessary complexity. For example, I cannot conveniently use cudaMemcpy or even cudaMemcpyToSymbol to copy to the intended allocation at that point. Instead I would have to use cudaMemcpy to h\_dev\_spheres, knowing that the data will end up in the same place as is pointed to by dev\_spheres.

Use of **device** variables as storage for pointers is therefore in my opinion not that sensible. Instead I would convert all of those cases to ordinary dynamically allocated variables using cudaMalloc and cudaMemcpy, and pass those variables to kernels that need to refer to them as kernel parameters.

---

<div class="post-metadata">

**Author:** ![akoutmos](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@akoutmos](https://forums.developer.nvidia.com/u/akoutmos)\
**Post date:** [November 23, 2014, 8:35am UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/3 "2014-11-23T08:35:39Z")

</div>

Thanks for the response! That definitely makes things quite a bit clearer. I’ll change up my mem allocations and see if that fixes things.

---

<div class="post-metadata">

**Author:** ![akoutmos](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@akoutmos](https://forums.developer.nvidia.com/u/akoutmos)\
**Post date:** [November 23, 2014, 10:30am UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/4 "2014-11-23T10:30:10Z")

</div>

So I was able to fix all the other allocations, but I am still having issues creating the sparse 2d array.

Currently I have:

```auto
Triangle** h_temp_tri;
	HANDLE_ERROR( cudaMalloc((void**) &h_temp_tri, sizeof(Triangle*) * temp_data->triangularMeshes) );
	HANDLE_ERROR( cudaMemcpyToSymbol(dev_triangles, &h_temp_tri, sizeof(Triangle*) * temp_data->triangularMeshes) );

```

I have not yet allocated space for each individual mesh, but shouldn’t that allocate space for the given number of Triangle\* pointers?

I am currently getting an invalid argument error at the memCpyToSymbol line.

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [November 23, 2014, 2:01pm UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/5 "2014-11-23T14:01:08Z")

</div>

Yes, it’s observably incorrect because this:

&h\_temp\_tri

produces a triple level of indirection:

* * *

but you are casting it to a double pointer:

(void \*\*)

So although that isn’t the crux of the issue, it indicates that the methodology is flawed because there should never be a need to do this with cudaMalloc.

In general, 2D arrays (i.e. arrays that could be accessed via double C subscripts on the host) can present some challenges. There are a variety of ways to handle it, but the simplest approach is usually to “flatten” the array to 1D and then, if desired, simulate 2D access where needed via pointer or index arithmetic.

---

<div class="post-metadata">

**Author:** ![akoutmos](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@akoutmos](https://forums.developer.nvidia.com/u/akoutmos)\
**Post date:** [November 25, 2014, 5:43am UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/6 "2014-11-25T05:43:08Z")

</div>

I ended up flattening the array which made the whole thing much easier to work with. I only have one more issue which I can’t seem to work out. I am currently passing in the cudaMalloced triangles pointer, but would prefer to copy it to the global symbol. It makes it easier seeing as I use it in several places and would rather not pass it from function to function.

When uncommented, the cudaMemcpyToSymbol fails and the application crashes. Does the source arg have to be a double pointer? I figured it would just copy the address to the symbol.

```auto
__device__ Triangle* dev_triangles;

Triangle* d_temp_tri;
HANDLE_ERROR( cudaMalloc((void**) &d_temp_tri, sizeof(Triangle) * temp_data->totalTriangles) );	
//HANDLE_ERROR( cudaMemcpyToSymbol(dev_triangles, d_temp_tri, sizeof(Triangle*)) );
HANDLE_ERROR( cudaMemcpy(d_temp_tri, temp_t, sizeof(Triangle) * temp_data->totalTriangles, cudaMemcpyHostToDevice) );

```

Thanks for helping me with getting this up and running :)!  
[External Media](http://krcomp.com/pictures/Capture.PNG "Media hosted on another site. Click to open in a new tab.")

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [November 25, 2014, 9:26am UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/7 "2014-11-25T09:26:03Z")

</div>

Change this:

```auto
HANDLE_ERROR( cudaMemcpyToSymbol(dev_triangles, d_temp_tri, sizeof(Triangle*)) );

```

to this:

```auto
HANDLE_ERROR( cudaMemcpyToSymbol(dev_triangles, &d_temp_tri, sizeof(Triangle*)) );

```

Your intention in this case is to pass the pointer value itself, instead of whatever the pointer points to. So you must pass a pointer-to-pointer to get cudaMemcpy to copy the proper thing.

---

<div class="post-metadata">

**Author:** ![akoutmos](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@akoutmos](https://forums.developer.nvidia.com/u/akoutmos)\
**Post date:** [November 26, 2014, 5:32am UTC](https://forums.developer.nvidia.com/t/cudamalloc-and-cudamemcpy-for-pointer-in-struct/35717/8 "2014-11-26T05:32:39Z")

</div>

Worked! Thanks :)!