CUDA host device variables

本文关键字：variables device host CUDA 更新时间：2023-10-16

在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-file翻译为:

...
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上重现

对于普通的浮点型或整型，只要在全局范围内将变量标记为const就足够了:

const double common = 1.0;

它应该可以在任何后续函数中使用，无论是host、__host__、__device__还是__global__。

这在这里的文档中得到支持，但有各种限制:

设'V'表示命名空间作用域变量或类静态成员变量，该变量具有const限定类型并且没有执行空间注释(例如，__device__, __constant__, __shared__)。V被认为是一个主机代码变量。
V的值可以直接用在设备代码中，如果V在使用点之前已经用常量表达式初始化，并且它有以下类型之一:
内置浮点类型，除非Microsoft编译器用作主机编译器，
内置积分型。
设备源代码不能包含对V的引用或V的地址

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

使用编译器宏定义的常量:
```
#define COMMON 1.0
```
使用模板，如果变量的选择范围是离散的和有限的
对于其他选项/情况，可能需要管理变量的显式主机和设备副本，例如在设备上使用__constant__内存，并在主机上使用相应的副本。访问变量的__host__ __device__函数中的主机和设备路径可以根据nvcc编译器宏(例如#ifdef __CUDA_ARCH__ ...

CUDA __host__ __device__ variables