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将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__.

这里的文档支持这一点,但受到各种限制:

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

如果V在使用点之前使用常量表达式初始化,则V的值可以直接在设备代码中使用,并且它具有以下类型之一:

• 内置浮点类型,除非Microsoft编译器用作主机编译器,
• 内置整体式.

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

在其他情况下,一些可能的选择是:

1. 使用编译器宏定义的常量:

   #define COMMON 1.0
   

2. 如果变量的选择范围是离散且有限的,则使用模板.

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