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

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

#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

这是什么问题，我怎样才能得到这两个设备和主机内存常数初始化同步？

来源

2012-02-26 davinchi

由于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 
}

这有相同的结果天鹅，但它是在真正的代码的脸更灵活：它可以让你在不断阵列具有运行时定义的值，例如，并允许您在__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(); 
}

来源

2012-02-28 06:02:01 harrism

我自己找到了解决方案，并且完全匹配您的解决方案。谢谢！ – davinchi 2012-02-28 15:40:20

是的，这是更强大，我同意。但它也是更多类型;） – CygnusX1 2012-02-28 16:28:21

使用__constant__限定符显式分配设备上的内存。无法从主机访问该内存，即使使用新的CUDA Unified Addressing（仅适用于分配给cudaMalloc（）及其朋友的内存）。用const限定变量只是说“这是一个指向（...）”的常量指针。

正确的做法是确实有两个阵列：一个在主机上，另一个在设备上。初始化您的主机阵列，然后使用cudaMemcpyToSymbol（）在运行时将数据复制到设备阵列。有关如何做到这一点的详细信息，请参阅本主题：http://forums.nvidia.com/index.php?showtopic=69724

来源

2012-02-26 22:40:38 mtrberzi

这个答案是几乎没有...但@davinchi希望有一台主机/设备功能使用他的常数表。为此，他应该使用'#ifdef __CUDA_ARCH__'并在其中访问__costant__数组，并在'＃else'中访问该数组的主机副本。 – harrism 2012-02-27 01:24:33

我觉得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()等功能的设备阵列可能会更困难。

来源

2012-02-27 20:55:36 CygnusX1

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

回答

相关问题

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