SEGV during cudaMemcpyToSymbol()

const __constant__ SomeType some_instance = {0};

…successfully compiles, but fails at runtime on the emu target. Backtrace:

#0  0x0000003331c7b391 in memcpy () from /lib64/libc.so.6

#1  0x00002aaaaaae3e88 in cudaMemcpyToSymbol () from /usr/local/cuda/lib/libcudart.so

#2  0x000000000040c5df in cudaMemcpyToSymbol<SomeType> (symbol=@0x61d960, src=0x7fffeef0a100, count=136, offset=0,

    kind=cudaMemcpyHostToDevice) at /usr/local/cuda/bin/../include/cuda_runtime.h:148

Relevant portion of symbol table:

$ objdump -t emu/blah |grep some_instance

000000000040f140 l     O .rodata        0000000000000088              some_instance

000000000061d960 l     O .bss   0000000000000088              __shadow_some_instance

So the SEGV is because I shouldn’t be writing to a read-only page, right? Ok, I can live w/that. But if I do this:

__constant__ SomeType some_instance_non_const = {0};

…then I get this compiler error:

./foo.cu:145:1: error: pasting "{" and "__constant__" does not give a valid preprocessing token

Also,

const __constant__ SomeType no_init;

…fails to compile, with:

error: const variable "no_init"

          requires an initializer -- class "SomeType" has no explicitly

          declared default constructor

  const __attribute__((__constant__)) SomeType no_init;

SomeType is declared like:

typedef struct 

{

    int some_member;

    float some_other_member;

} SomeType;

I’m using “–host-compilation=c++”, btw.

I have also tried declaring SomeType as:

struct SomeType

{

    int some_member;

    float some_other_member;

};

Unfortunately, I fear that my case is more complex than I present here, since I’m unable to reproduce it with a simpler example. Can anyone suggest any hints on where to look next?

Which of the above examples is more appropriate for declaring constant GPU memory?