cudaMemcpyToSymbol using or not using string

Question

I was trying to copy a structure to constant memory in this way:

struct Foo {
    int a, b, c;
};

__constant__ Foo cData;

int main() {
    Foo hData = {1, 2, 3};
    cudaMemcpyToSymbol(cData, &hData, sizeof(Foo));
    // ...
}

And this worked fine, in my kernel I could access the constant data directly:

__global__ void kernel() {
    printf("Data is: %d %d %d\n", cData.a, cData.b, cData.c); // 1 2 3
}

But then I tried to use a const char * as symbol name, and things stopped working:

cudaMemcpyToSymbol("cData", &hData, sizeof(Foo)); // prints 0 0 0

I thought both versions were similar, but it seems I was wrong.

What is happening?

EDIT: I'd like to report this same behavior with cudaGetSymbolAddress, which works for me if no const char * is used:

__constant__ int someData[10];
__constant__ int *ptrToData;

int *dataPosition;
cudaGetSymbolAddress((void **)&dataPosition, someData); // Works
// cudaGetSymbolAddress((void **)&dataPosition, "someData"); // Do not work
cudaMemcpyToSymbol(ptrToData, &dataPosition, sizeof(int *));

Answer 1

As of CUDA 5, using a string for symbol names is no longer supported. This is covered in the CUDA 5 release notes here

•The use of a character string to indicate a device symbol, which was possible with certain API functions, is no longer supported. Instead, the symbol should be used directly.

One of the reasons for this has to do with enabling of a true device linker, which is new functionality in CUDA 5.

Answer 2

Because of getting the same error again and again, I want to share this sample code that shows nearly all of the example cases for this problem (so I may refer here later when I make same mistakes again).

//file: main.cu
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>

__constant__ float constData[256];
__device__ float devData;
__device__ float* devPointer;

int main(int argc, char **argv)
{
  cudaFree(0);

  float data[256];
  cudaError_t err = cudaMemcpyToSymbol(constData, data, sizeof(data));
  printf("Err id: %d, str: %s\n", err, cudaGetErrorString(err));

  float value = 3.14f;
  err = cudaMemcpyToSymbol(devData, &value, sizeof(float));
  printf("Err id: %d, str: %s\n", err, cudaGetErrorString(err));

  float* ptr;
  cudaMalloc(&ptr, 256 * sizeof(float));
  err = cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));
  printf("Err id: %d, str: %s\n", err, cudaGetErrorString(err));
  cudaFree(ptr);

  return EXIT_SUCCESS;
}

I was getting "invalid device symbol" and many others which are related to _constant_ _device_ memory usage. This code gives no such errors at runtime.