由于对结构体进行cudaMemcpyToSymbol，导致出现cudaInvalidSymbol错误

Question

这是一个最小化的示例，使用 Cuda 编译工具版本 6.5.12 编译，在 sync_and_check("cudaMemcpyToSymbol") 处抛出了 cudaInvalidSymbol 异常：

#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");
}

- Leo

---

如果符号不是结构体，它能工作吗？如果不能，请尝试在这里的代码片段并查看是否有效... https://dev59.com/8GQo5IYBdhLWcg3whPrI - Christian Sarofeen

Answer 1

问题在于引用。正确的调用方式是：

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

我感到困惑，因为手册上写着：

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

无论他们所说的“符号”是什么 :) 复制到结构元素中

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

它的功能非常出色。

- Leo