CUDA主机和设备使用相同的__constant__内存

Question

我有一个使用常量内存的设备/主机功能。它在设备上运行正常，但在主机上似乎这个内存没有初始化。

#include <iostream>
#include <stdio.h>


const __constant__ double vals[2] = { 0.0, 1000.0 };

__device__ __host__ double f(size_t i)
{
    return vals[i];
}

__global__ void kern()
{
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]);
}

int main() {
    std::cerr << f(0) << " " << f(1) << std::endl;
    kern<<<1, 2>>>();
    cudaThreadSynchronize();
}

这将打印文本（需要CC 2.0或以上版本）

0 0
vals[0] = 0.000000
vals[1] = 1000.000000

问题是什么，如何同时初始化设备和主机内存常量？

- davinchi

Answer 1

由于 CygnusX1 在 MurphEngineer 的答案评论中误解了我的意思，也许我应该发表自己的回答。 我的意思是这样的：我想说的是：

__constant__ double dc_vals[2] = { 0.0, 1000.0 };
       const double hc_vals[2] = { 0.0, 1000.0 };

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

这与Cygnus的结果相同，但在面对实际代码时更加灵活:例如，它允许您在常量数组中具有运行时定义的值，并允许您在 __constant__ 数组上使用 CUDA API 函数，如 cudaMemcpyToSymbol / cudsaMemcpyFromSymbol。

一个更现实的完整示例:

#include <iostream>
#include <stdio.h>

__constant__ double dc_vals[2];
       const double hc_vals[2];

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

__global__ void kern()
{
    printf("vals[%d] = %lf\n", threadIdx.x, vals[threadIdx.x]);
}

int main() {
    hc_vals[0] = 0.0;
    hc_vals[1] = 1000.0;

    cudaMemcpyToSymbol(dc_vals, hc_vals, 2 * sizeof(double), 0, cudaMemcpyHostToDevice);

    std::cerr << f(0) << " " << f(1) << std::endl;
    kern<<<1, 2>>>();
    cudaThreadSynchronize();
}

- harrism

---

我自己找到了解决方案，它完全与你的匹配。谢谢！ - davinchi

是的，我同意这更健壮。但是需要输入的内容也更多 ;) - CygnusX1

Answer 2

我认为MurphEngineer解释了为什么它不起作用。

为了快速解决这个问题，您可以按照harrism的想法做一些类似于这样的事情：

#ifdef __CUDA_ARCH__
#define CONSTANT __constant__
#else
#define CONSTANT
#endif

const CONSTANT double vals[2] = { 0.0, 1000.0 };

使用这种方法，主机编译将创建一个普通的主机常量数组，而设备编译将创建一个设备__constant__编译。

请注意，使用这种技巧，如果您决定这样做，可能较难使用CUDA API访问该设备数组，例如cudaMemcpyToSymbol()函数。

- CygnusX1

---

实际上，cudaMemcpyToSymbol 在使用这个技巧时可以正常工作。即使两个数组具有相同的名称，至少在 Visual Studio 中似乎聪明到知道哪个是哪个，因此像 cudaMemcpyToSymbol(CFG, CFG, sizeof(CFG)) 这样的语句实际上是可以正常工作的！ - poby

Answer 3

使用 __constant__ 限定符可以在设备上显式分配内存。没有办法从主机上访问该内存，即使使用新的 CUDA 统一寻址功能也不行（它只适用于使用 cudaMalloc() 和其相关函数分配的内存）。使用 const 修饰变量仅表示“这是一个指向常量的指针(...)”。

正确的做法是确实需要有两个数组：一个在主机上，一个在设备上。初始化主机数组，然后使用 cudaMemcpyToSymbol() 在运行时将数据复制到设备数组中。有关如何执行此操作的更多信息，请参见此线程：http://forums.nvidia.com/index.php?showtopic=69724。

- mtrberzi

---

这个答案接近正确……但@davinchi希望使用他的常量表作为主机/设备函数。为了做到这一点，他应该使用#ifdef __CUDA_ARCH__，在其中访问__constant__数组，在#else中访问数组的主机副本。 - harrism

Answer 4

绝对棒极了。我曾经遇到同样的问题，这提供了一个解决方案。然而，harrism 提出的代码在编译时会出现错误。这里是修正后的代码，使用 nvcc 可以正确编译：

#include <iostream>
#include <stdio.h>

__constant__ double dc_vals[2];
       const double hc_vals[2] = {0.0, 1000.0};

__device__ __host__ double f(size_t i)
{
#ifdef __CUDA_ARCH__
    return dc_vals[i];
#else
    return hc_vals[i];
#endif
}

__global__ void kern()
{

    printf("Device: vals[%d] = %lf\n", threadIdx.x, f(threadIdx.x));
}

int main() {

    cudaMemcpyToSymbol(dc_vals, hc_vals, 2 * sizeof(double), 0, cudaMemcpyHostToDevice);

    std::cerr << "Host: " << f(0) << " " << f(1) << std::endl;
    kern<<<1, 2>>>();
    cudaThreadSynchronize();
}

- Rodrigo