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

【问题标题】：CUDA host and device using same __constant__ memoryCUDA 主机和设备使用相同的 __constant__ 内存
【发布时间】：2012-03-16 11:40:25
【问题描述】：

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

#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

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

【问题讨论】：

标签： cuda

【解决方案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 的结果相同，但在面对实际代码时更灵活：例如，它允许您在常量数组中拥有运行时定义的值，并允许您使用 CUDA API 函数，如 @ 987654322@/cudsaMemcpyFromSymbol 在__constant__ 数组上。

一个更真实的完整例子：

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

【讨论】：

我自己找到了解决方案，它与您的完全匹配。谢谢！
是的，这更强大，我同意。但它也需要更多的输入；）

【解决方案2】：

使用 __constant__ 限定符显式分配设备上的内存。没有办法从主机访问该内存——即使使用新的 CUDA 统一寻址东西（仅适用于使用 cudaMalloc() 及其朋友分配的内存）。用 const 限定变量只是说“这是一个指向 (...) 的常量指针”。

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

【讨论】：

这个答案差不多了...但是@davinchi 想让主机/设备函数使用他的常量表。为此，他应该使用#ifdef __CUDA_ARCH__ 并在其中访问costant 数组，并在#else 中访问该数组的主机副本。

【解决方案3】：

我认为 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() 之类的函数的设备数组可能会更加困难。

【讨论】：

实际上cudaMemcpyToSymbol 可以很好地配合这个技巧。尽管两个数组具有相同的名称，但似乎 Visual Studio 至少足够聪明地知道哪个是哪个，所以 cudaMemcpyToSymbol(CFG, CFG, sizeof(CFG)) 行确实有效！

【解决方案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();
}

【讨论】：