cudaMemcpyToSymbol()的invalid device symbol问题

这个问题调试了好久，也从网上查了不少资料，但一直没找到解决方法，直到最近修改完成。因为查找解决方案花费很大功夫，为了方便，特此记录。

下面是问题描述。

我写一个cuda程序需要用到__constant__常量，而在host端给__constant__常量赋值的函数是cudaMemcpyToSymbol()函数。我在文件中声明了2个__constant__常量和3个__constant__二维数组，在host函数中把数据从host端传递到device端。为了更好说明，下面附上关键代码：

// ...// 声明__constant__ double cd;__constant__ int ci;__constant__ double ca1[4][4];__constant__ double ca2[4][4];__constant__ double ca3[4][4];void init(){    // 声明host端与__constant__对应变量并赋值    // 数据传递    size_t sized = sizeof(double);    sizei = sizeof(int);    cudaMemcpyToSymbol(&cd, &hd, sized, 0, cudaMemcpyHostToDevice);    cudaMemcpyToSymbol(&ci, &hi, sizei, 0, cudaMemcpyHostToDevice);    cudaMemcpyToSymbol(ca1, ha1, sized*4*4, 0, cudaMemcpyHostToDevice);    // ca2和ca3的数据传递...}

之后在kernel函数中使用ca1，ca2，ca3（二维数组）的时候都没问题，但等到使用cd和ci的时候，取值都是0。经过检查，是使用cudaMemcpyToSymbol函数的时候出了问题。

// ...cudaError_t err;err = cudaMemcpyToSymbol(&cd, &hd, sized, 0, cudaMemcpyHostToDevice);printf("%d: %s.\n", err, cudaGetErrorString(err));err = cudaMemcpyToSymbol(&ci, &hi, sizei, 0, cudaMemcpyHostToDevice);printf("%d: %s.\n", err, cudaGetErrorString(err));// ...

打印结果：

13: invalid device symbol（无效的device symbol）

13: invalid device symbol
对于cudaMemcpyToSymbol函数，这种错误一般出现的原因是函数的第一个参数（设备端）数据类型有问题。我试过了更改计算能力，把&改成""，都失败了。直到把cd和ci改成指针类型：

// 声明__constant__ double *cd;__constant__ int *ci;// 赋值void init(){    // ...    // 分配空间    cudaMalloc((void**)&cd, sized);    cudaMalloc((void**)&ci, sizei);    // ...    // 数据传输    cudaError_t err;    err = cudaMemcpyToSymbol(cd, &hd, sized, 0, cudaMemcpyHostToDevice);    printf("%d: %s.\n", err, cudaGetErrorString(err));    err = cudaMemcpyToSymbol(ci, &hi, sizei, 0, cudaMemcpyHostToDevice);    printf("%d: %s.\n", err, cudaGetErrorString(err));    // ...}// kernel函数__global__ void kernel(...){    // ...    if(tid == 0 && bid == 0) printf("cd:%lf, ci:%d.\n", cd, ci);    // ...}

打印结果：

0: no error

0: no error

cd:cd的值, ci:ci的值

看起来没毛病，但实际上使用的时候会在调试过程中报错：

error: expression must have arithmetic or enum type（表达式必须是算数或枚举类型）

在之后才查到原来__constant__类型可以是变量、数组，但不能是指针。

所以最后还是要从cudaMemcpyToSymbol这个函数本身入手。CUDA Runtime API中关于此函数是这么说的：

__host__ ​cudaError_t cudaMemcpyToSymbol ( const void* symbol, const void* src, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyHostToDevice )

Copies data to the given symbol on the device.

Parameters
symbol
- Device symbol address
src
- Source memory address
count
- Size in bytes to copy
offset
- Offset from start of symbol in bytes
kind
- Type of transfer

单看这个看不出什么来。然而我在cuda手册的3.2.2 Device Memory部分找到了如下三个例子：

__constant__ float constData[256];float data[256];cudaMemcpyToSymbol(constData, data, sizeof(data));cudaMemcpyFromSymbol(data, constData, sizeof(data));__device__ float devData; float value = 3.14f; cudaMemcpyToSymbol(devData, &value, sizeof(float));__device__ float* devPointer; float* ptr;cudaMalloc(&ptr, 256 * sizeof(float));cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

发现cudaMemcpyToSymbol函数的第一个参数都没有用&。

于是我抱着尝试的心态修改了代码：

// 声明__constant__ double cd;__constant__ int ci;// 赋值void init(){    // ...    // 分配空间    cudaMalloc((void**)&cd, sized);    cudaMalloc((void**)&ci, sizei);    // ...    // 数据传输    cudaError_t err;    err = cudaMemcpyToSymbol(cd, &hd, sized, 0, cudaMemcpyHostToDevice);    printf("%d: %s.\n", err, cudaGetErrorString(err));    err = cudaMemcpyToSymbol(ci, &hi, sizei, 0, cudaMemcpyHostToDevice);    printf("%d: %s.\n", err, cudaGetErrorString(err));    // ...}

打印结果：

0: no error
0: no error

检查kernel函数中的值也没问题，也能正确使用。到此问题解决。

关于问题的出现以及这样解决的原因，到现在我也没有弄明白，如果找到了原因，再进行补充。