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