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

问题出在哪里，如何才能同时初始化设备和主机内存常量？

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相同的结果，但是面对实际代码则更加灵活:例如，它使您可以在常量数组中具有运行时定义的值，并允许您使用CUDA API函数(如cudaMemcpyToSymbol/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();
}