CUDA cudaMemcpyFromSymbol“无效的设备符号”错误？

Question

所以我看到一个parent 的问题，关于如何使用cudaMemcpyToSymbol 从主机复制到 GPU 上的常量内存。

我的问题是如何做相反的事情，使用cudaMemcpyFromSymbol从设备常量内存复制到主机。

在下面的最小可重现示例中，我得到了

• 1) invalid device symbol 使用 cudaMemcpyFromSymbol(const_d_a, b, size); 时出错，或
• 2) 如果我使用cudaMemcpyFromSymbol(&b, const_d_a, size, cudaMemcpyDeviceToHost)，则得到segmentation fault。

我已经咨询了manual，它建议我按照 1) 进行编码，this SO question 建议我按照 2) 进行编码。他们都不在这里工作。

有人可以帮忙提出解决方法吗？我一定是理解不正确...谢谢！

代码如下：

// a basic CUDA function to test working with device constant memory
#include <stdio.h>
#include <cuda.h>

const unsigned int N = 10;    // size of vectors

__constant__ float const_d_a[N * sizeof(float)];  

int main()
{
    float * a, * b;  // a and b are vectors. c is the result
    a = (float *)calloc(N, sizeof(float));
    b = (float *)calloc(N, sizeof(float));

    /**************************** Exp 1: sequential ***************************/
    int i;
    int size = N * sizeof(float);
    for (i = 0; i < N; i++){
        a[i] = (float)i / 0.23 + 1;
    }


    // 1. copy a to constant memory
    cudaError_t err = cudaMemcpyToSymbol(const_d_a, a, size);
    if (err != cudaSuccess){
        printf("%s in %s at line %d\n", cudaGetErrorString(err), __FILE__, __LINE__);
        exit(EXIT_FAILURE);
    }

    cudaError_t err2 = cudaMemcpyFromSymbol(const_d_a, b, size);
    if (err2 != cudaSuccess){
        printf("%s in %s at line %d\n", cudaGetErrorString(err2), __FILE__, __LINE__);
        exit(EXIT_FAILURE);
    }

    double checksum0, checksum1;
    for (i = 0; i < N; i++){
        checksum0 += a[i];
        checksum1 += b[i];
    }

    printf("Checksum for elements in host memory is %f\n.", checksum0);
    printf("Checksum for elements in constant memory is %f\n.", checksum1);

    return 0;
}

Answer 1

在 CUDA 中，各种 cudaMemcpy* 操作是在 C 标准库 memcpy 例程之后建模的。在该函数中，第一个指针总是是目标指针，第二个指针是总是源指针。对于所有 cudaMemcpy* 函数也是如此。

因此，如果你想做cudaMemcpyToSymbol，符号最好是传递给函数的第一个（目标）参数（第二个参数是主机指针）。如果你想做cudaMemcpyFromSymbol，符号需要是第二个参数（源位置），主机指针是第一个参数。这不是你在这里所拥有的：

cudaError_t err2 = cudaMemcpyFromSymbol(const_d_a, b, size);
                                          ^        ^
                                          |       This should be the symbol.
                                          |  
                                   This is supposed to be the host destination pointer.

您可以通过查看API documentation 来发现这一点。

如果我们在该行代码中颠倒这两个参数的顺序：

cudaError_t err2 = cudaMemcpyFromSymbol(b, const_d_a, size);

您的代码将毫无错误地运行，并且打印的最终结果将匹配。

在这些函数中，没有必要在a 或b 指针中使用与号。 a 和 b 已经是指针。在您链接的示例中，pi_gpu_h 不是指针。它是一个普通的变量。要使用cudaMemcpyFromSymbol 向其复制某些内容，必须获取该普通变量的地址，因为该函数需要一个（目标）指针。

顺便说一句，这看起来不对：

__constant__ float const_d_a[N * sizeof(float)];  

这实际上是一个静态数组声明，除了 __constant__ 装饰器之外，它应该等同于在 C 或 C++ 中的实现方式。如果您想要存储N float 数量，则无需在此处将N 乘以sizeof(float)。只需 N 本身就可以做到这一点：

__constant__ float const_d_a[N];

但是，保持原样不会对您发布的代码造成问题。