CUDA __host__ __device__变量

Question

在CUDA函数类型限定符__device__和__host__中可以一起使用，在这种情况下，为主机和设备编译函数。这可以消除复制粘贴。但是，不存在__host__ __device__变量。我正在寻找一个优雅的方式做这样的事情：

__host__ __device__ const double common = 1.0; 

__host__ __device__ void foo() { 
    ... access common 
} 

__host__ __device__ void bar() { 
    ... access common 
} 

我发现下面的代码相符，无差错运行。 （在Ubuntu 14.04就获得CUDA 7.5和gcc 4.8.4为主机编译器的结果）

#include <iostream> 

__device__ const double off = 1.0; 

__host__ __device__ double sum(int a, int b) { 
    return a + b + off; 
} 

int main() { 
    double res = sum(1, 2); 
    std::cout << res << std::endl; 
    cudaDeviceReset(); 
    return 0; 
} 

$ nvcc main.cu -o main && ./main 
4 

其实 ，nvcc --cuda main.cu转换铜文件到这一点：

... 
static const double off = (1.0); 
# 5 "main.cu" 
double sum(int a, int b) { 
# 6 "main.cu" 
return (a + b) + off; 
# 7 "main.cu" 
} 
# 9 "main.cu" 
int main() { 
# 10 "main.cu" 
double res = sum(1, 2); 
# 11 "main.cu" 
(((std::cout << res)) << (std::endl)); 
# 12 "main.cu" 
cudaDeviceReset(); 
# 13 "main.cu" 
return 0; 
# 14 "main.cu" 
} 
... 

但是，没有惊讶的是，如果变量off而不const预选赛（__device__ double off = 1.0）我得到以下输出声明：

$ nvcc main.cu -o main && ./main 
main.cu(7): warning: a __device__ variable "off" cannot be directly read in a host function 

3 

因此，回到原来的问题，我可以依靠这种行为与全球__device__ const变量？如果不是，还有什么其他选择？

UPD顺便说一下，上述行为不会在Windows上重现。

Answer 1

对于普通浮点或整数类型应该是足够简单地在全局范围内来标记该变量作为const：

const double common = 1.0; 

那么它应该是可用在任何后续功能，无论主机，__host__，__device__，或__global__。

此文档here，受到各种限制在被支撑：

让“V”表示一个命名空间范围变量或具有常量限定类型和不具有执行一类静态成员变量空间注释（例如，__device__,__constant__,__shared__）。 V被认为是主机代码变量。

v的值可以在设备代码可以直接使用，如果 V具有已与使用点之前一个常量表达式初始化， 它具有以下类型之一：

• 内置浮动除了使用Microsoft编译器作为主机编译器时，除了内置点类型外，还可以使用内建整数类型。

设备源代码不能包含对V的引用或取V的地址。

在其他情况下，一些可能的选项是：

1. 使用定义的常数一个编译器宏：

   #define COMMON 1.0 
   

2. 使用模板，如果选择在可变的范围是离散并受到限制。

3. 对于其他选项/情况，可能需要管理变量的显式主机和设备副本，例如，在设备上使用__constant__内存，并在主机上使用相应的副本。访问变量的函数中的主机和设备路径然后可以基于nvcc编译器宏来区分行为（例如，#ifdef __CUDA_ARCH__ ...