使用cudaMemcpy/cudaMemcpyToSymbol出现invalid argument错误如何解决

• 你混淆了设备端符号d_pa本身的存储地址和它指向的设备内存地址：d_pa是一个存放在设备全局内存区的指针变量，你在kernel中调用malloc给它赋值，只是让这个指针指向了设备堆上的一块内存，而cudaMemcpyToSymbol、cudaGetSymbolAddress两个接口操作的是指针变量本身的设备地址，不是它指向的内存地址。
• 版本1错误：cudaMemcpyToSymbol(d_pa,&pa,sizeof(params))是直接把结构体内容覆写到d_pa这个指针变量的内存空间，不仅类型不匹配，结构体大小也大于64位平台下指针的8字节长度，属于越界写入，直接触发非法参数报错。
• 版本2错误：cudaGetSymbolAddress拿到的是d_pa这个指针本身的设备地址，你往这个地址写入结构体，本质还是覆盖指针本身的值，根本没有写入你在kernel里malloc出来的结构体内存区域。

方案1：直接用设备端结构体（最简洁，推荐） #

不需要用指针，直接将d_pa定义为设备端结构体变量，省略kernel里的malloc操作即可：

// stackoverflow.h 修改
extern __device__ params d_pa; // 去掉*，定义为结构体而非指针

// stackoverflow.cu 修改
__device__ params d_pa; // 同步修改定义
params pa;

// 删掉paramsDeviceInit核函数及对应的调用、cudaDeviceSynchronize代码

void setParams() {
    pa.a=1;
    pa.b=2;
    // 直接拷贝到设备端结构体符号即可
    errCheck(cudaMemcpyToSymbol(d_pa,&pa,sizeof(params)));
    printf("Host: a=%d, b=%f\n",pa.a,pa.b);
}

方案2：保留指针写法的正确实现 #

如果业务逻辑必须用指针，正确操作流程是先在主机端分配设备内存、拷贝结构体内容，再把设备内存地址赋值给设备端的d_pa指针：

// 头文件无需修改，保留__device__ params *d_pa;
// 删掉paramsDeviceInit核函数，不需要在kernel里做malloc操作

void setParams() {
    pa.a=1;
    pa.b=2;

    // 1. 主机端先分配设备端结构体内存
    params *dev_pa;
    errCheck(cudaMalloc((void**)&dev_pa, sizeof(params)));
    // 2. 把主机结构体拷贝到刚分配的设备内存
    errCheck(cudaMemcpy(dev_pa, &pa, sizeof(params), cudaMemcpyHostToDevice));
    // 3. 把设备内存的地址拷贝到设备端的d_pa指针符号里
    errCheck(cudaMemcpyToSymbol(d_pa, &dev_pa, sizeof(params*)));
    
    printf("Host: a=%d, b=%f\n",pa.a,pa.b);
}

注意：kernel内调用的malloc是设备端堆分配，分配出来的地址无法直接在主机端的cudaMemcpy接口中使用，这种方式性能差还容易出错，不推荐用于参数传递场景。

内容的提问来源于stack exchange，提问作者Gregor Freund