cudaInvalidSymbol due to cudaMemcpyToSymbol to a struct

Question

Here is the minimum example that throws cudaInvalidSymbol at sync_and_check("cudaMemcpyToSymbol"), compiled using Cuda compilation tools, release 6.5, V6.5.12:

#include<stdio.h>

__inline __host__ void gpuAssert(cudaError_t code, const char* command, const char *file, int line, 
                 bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"Error while calling %s in %s line %d: %s\n", command,
          file, line, cudaGetErrorString(code));
      if (abort) exit(code);
   }
}

#define sync_and_check(command) { gpuAssert(cudaGetLastError(), command, __FILE__, __LINE__); }

struct S
{
    float one;
};

__device__ __constant__ S d_s;

__global__ void kernel(float* f)
{
    (*f) = d_s.one;
}

int main()
{
    cudaGetLastError();

    S s;
    s.one = 1.f;
    cudaMemcpyToSymbol(&d_s, &s, sizeof(S));
    sync_and_check("cudaMemcpyToSymbol");

    float* d_f;
    cudaMalloc(&d_f, sizeof(float));
    sync_and_check("cudaMalloc");

    dim3 dimGrid(1), dimBlock(32);
    kernel<<<dimGrid, dimBlock>>>(d_f);
    sync_and_check("kernel");

    cudaFree(d_f);
    sync_and_check("cudaFree");
}

Answer 1

The problem lies in referencing. The correct invocation is

cudaMemcpyToSymbol(d_s, &s, sizeof(S));

I was confused because the man page says

cudaError_t cudaMemcpyToSymbol (const void * symbol, const void * src,
   size_t count, size_t offset = 0, enum cudaMemcpyKind kind =
   cudaMemcpyHostToDevice)
   Copies count bytes from the memory area pointed to by src to the memory
   area pointed to by offset bytes from the start of symbol symbol. The
   memory areas may not overlap. symbol is a variable that resides in
   global or constant memory space. kind can be either
   cudaMemcpyHostToDevice or cudaMemcpyDeviceToDevice.

   Parameters:
       symbol - Device symbol address

Whatever they mean by "symbol" :) Copying to struct elements

    cudaMemcpyToSymbol(d_s.one, &(s.one), sizeof(S));

works just as well.