cudaMemcpyToSymbol to a structur element ??

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

When I try to copy copying with cudaMemcpyToSymbol into a structure element, I get an error back.
Is there an alternative here?

typedef struct s_Picture_t
{
	unsigned int	ui_WIDTH;
	unsigned int	ui_HEIGHT;
	unsigned int    ui_SIZE;

	unsigned char * p_Picture;
}s_Picture_t;

typedef struct s_CUDA_CLASS_DEVICE_Holder_t
{
	s_Picture_t s_RGB_Picture;

}s_CUDA_CLASS_DEVICE_Holder_t;

__device__ s_CUDA_CLASS_DEVICE_Holder_t s_CUDA_CLASS_DEVICE_Holder;


ret = cudaMemcpyToSymbol(s_CUDA_CLASS_DEVICE_Holder.s_RGB_Picture.p_Picture, &s_CUDA_CLASS_Holder.p_DEVICE_RGB, sizeof(s_CUDA_CLASS_Holder.p_DEVICE_RGB));

Regards, TinTin

2

At the Moment i do:

typedef struct s_Picture_t
{
	unsigned int	ui_WIDTH;	//4
	unsigned int	ui_HEIGHT;  //4
	unsigned int    ui_SIZE;	//4

	unsigned char * p_Picture;
}s_Picture_t;

typedef struct SIFT_ALIGN(1) s_CUDA_CLASS_DEVICE_Holder_t
{
	s_Picture_t s_RGB_Picture;

}s_CUDA_CLASS_DEVICE_Holder_t;



	ret = cudaMalloc((void **)&s_CUDA_CLASS_Holder.p_DEVICE_RGB, s_CUDA_CLASS_Holder.ui_RGB_SIZE);
	// Copy the pointer 
	ret = cudaGetSymbolAddress((void **)&ptr, s_CUDA_CLASS_DEVICE_Holder);
	ptr += 3 * 4;
	ret = cudaMemcpy(ptr, &s_CUDA_CLASS_Holder.p_DEVICE_RGB, sizeof(s_CUDA_CLASS_Holder.p_DEVICE_RGB), cudaMemcpyHostToDevice);

But this is not really nice…
Regards TinTin

3

The cudaMemcpyToSymbol api has an optional offset parameter, ostensibly for this purpose.

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g9bcf02b53644eee2bef9983d807084c7

ret = cudaMemcpyToSymbol(s_CUDA_CLASS_DEVICE_Holder, &s_CUDA_CLASS_Holder.p_DEVICE_RGB, sizeof(s_CUDA_CLASS_Holder.p_DEVICE_RGB), 3*4);
                                                                                                                                  ^^^

Note that my purpose here is to communicate the availability of the feature in the API, not to suggest that manually computing structure offsets this way is the best way to do it. For example I think it would be better to use offsetof:

http://www.cplusplus.com/reference/cstddef/offsetof/

4

I would simply set the structure values on the CPU and hand it over as additional kernel parameter by value to the CUDA kernel. The structure should be aligned with the resepctive align specifiers so that it is aligned in the same way in CPU Host compiler and GPU NVCC compiler. For that, we employ the ‘HEMI_ALIGN’ macro ´described in Developing Portable CUDA C/C++ Code with Hemi | NVIDIA Technical Blog