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变量的这种行为吗？如果不能，还有其他选择吗？

更新 顺便说一句，上述行为在Windows上无法重现。

- wotopul

Answer 1

对于普通的浮点数或整型，只需在全局范围内将变量标记为const即可：

const double common = 1.0;

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

这在文档这里有支持，但需要遵守各种限制：

让 'V' 表示一个具有const修饰类型且不带执行空间注释（例如__device__、__constant__、__shared__）的命名空间作用域变量或类静态成员变量。V 被视为主机代码变量。

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

• 内置浮点类型，除非使用 Microsoft 编译器作为主机编译器，
• 内置整数类型。

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

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

1. Use a compiler macro defined constant:

   #define COMMON 1.0
   

2. Use templating, if the range of choices on the variable is discrete and limited.

3. For other options/cases, it may be necessary to manage explicit host and device copies of the variable, e.g. using __constant__ memory on the device, and a corresponding copy on the host. Host and device paths within the __host__ __device__ function that accesses the variable could then differentiate behavior based on a nvcc compiler macro (e.g. #ifdef __CUDA_ARCH__ ...

- Robert Crovella