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__。

这在文档 here 中得到支持，但受到各种限制：

Let 'V' denote a namespace scope variable or a class static member variable that has const qualified type and does not have execution space annotations (e.g., __device__, __constant__, __shared__). V is considered to be a host code variable.

The value of V may be directly used in device code, if V has been initialized with a constant expression before the point of use, and it has one of the following types:

builtin floating point type except when the Microsoft compiler is used as the host compiler,

builtin integral type.

Device source code cannot contain a reference to V or take the address of V.

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

使用编译器宏定义常量：
```
#define COMMON 1.0
```
如果变量的选择范围是离散且有限的，则使用模板。
对于其他 options/cases，可能需要管理变量的显式主机和设备副本，例如使用设备上的 __constant__ 内存，以及主机上的相应副本。 __host__ __device__ 函数中访问变量的主机和设备路径然后可以根据 nvcc 编译器宏区分行为（例如 #ifdef __CUDA_ARCH__ ...

CUDA host device 变量

CUDA host device variables

c++

cuda

gpgpu

nvcc

CUDA __host__ __device__ 变量

CUDA __host__ __device__ variables

c++

cuda

gpgpu

nvcc

CUDA host device 变量

CUDA host device variables