Structure into __constant__memory?

Hello all,

is it possible to store a structure in constant memory?

Let’s say I have the following structure:

struct my_struct {
int a;
double b;
string c;
} ;

Is it possible to store such a structure in constant memory? And how?

Thanks!

constant struct my_struct c_mystruct;

but you can’t use std::string in there (maybe use a fixed size char array instead). It’s best if your struct is a POD struct (plain old data).

Also note you have to assign values to a struct living on the host, then use cudaMemcpyToSymbol to send the data to GPU, like

struct mystruct h_struct_on_host = { /* assign as needed */ };
cudaMemcpyToSymbol(c_mystruct, &h_struct_on_host, sizeof(struct my_struct), 0, cudaMemcpyHostToDevice);

(that extra 0 given is a byte offset relative to c_mystruct which we have to set to 0)

The simples way IMO is to hand over the struct as a parameter (per value) to the kernel call. Then you don’t have declare a constant variable, invoke the cudaMemcpyToSymbol call etc…

Add a proper alignment macro (to multiples of 8 or 16 bytes) to the struct, so that it is aligned the same way on CPU on GPU. We use the HEMI_ALIGN macro (with value 8) from https://github.com/harrism/hemi/blob/master/hemi/hemi.h
See also https://devblogs.nvidia.com/developing-portable-cuda-cc-code-hemi/

As already mentioned, the struct should be a POD type (no constructors, virtual fns. etc.)
The data layout of elements in a struct is implementation-specific (so can differ between host and device compiler), but usually it is in the way described in https://stackoverflow.com/questions/2748995/c-struct-memory-layout