CUDA中關於C++特性的限制
CUDA官方文檔中對C++語言的支持和限制，懶得每次看英文文檔，自己嘗試翻譯一下（沒有放lambda表達式的相關內容，太過於復雜，我選擇不用）。官方文檔https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#c-cplusplus-language-support
__CUDA_ARCH__
在如下3種情況下，開發者不應該依賴於__CUDA_ARCH__宏是否定義或__CUDA_ARCH__的具體值來確定代碼：
typedef
typedef用於下列幾種情況時：
__global__函數及其函數模版__device__和__constant__變量- 紋理和surfaces
如下所示：
#if !defined(__CUDA_ARCH__)
typedef int mytype;
#else
typedef double mytype;
#endif
__device__ mytype xxx; // error: xxx's type depends on __CUDA_ARCH__
__global__ void foo(mytype in, // error: foo's type depends on __CUDA_ARCH__
mytype *ptr) {
*ptr = in;
}
函數模版實例化
__global__函數實例化：
__device__ int result;
template <typename T>
__global__ void kern(T in) {
result = in;
}
__host__ __device__ void foo(void) {
#if !defined(__CUDA_ARCH__)
kern<<<1,1>>>(1); // error: "kern<int>" 的實例化必須依賴於__CUDA_ARCH__未定義
}
單獨編譯
extern的函數或變量定義：
#if !defined(__CUDA_ARCH__)
void foo(void) {/*...*/ } // error: foo函數實現依賴於__CUDA_ARCH__未定義
#endif
特別是若頭文件中定義的函數實現依賴於__CUDA_ARCH__，而用到該函數的多個對象在不同arch下編譯會導致該函數的行為不一致，從而引起函數沖突。
CUDA還規定，若設備端代碼調用了一個被聲明為extern的函數，則該函數和設備端代碼必須在同一arch下編譯。
空間描述符
內存空間描述符
__device__、__shared__、__managed__、__constant__內存空間描述符不可用於如下情況：
class、struct、union數據成員- 形式參數
- 在主機端執行的函數中不可聲明非extern變量
- 在設備端執行的函數中，除
__shared__外另外三個描述符描述的變量不可聲明為既非extern亦非static
此外，若有此四類描述符的變量的類型為自定義類（class或struct），則該類必須擁有空構造函數和空析構函數。
關於C++類在什么情況下會擁有空構造函數和空析構函數可參考《深度探索C++對象模型》
在全編譯模式下，此四類描述符變量不能被定義為extern變量，無論主機端設備端（當然，動態共享內存聲明不包含在內），但在單獨編譯模式下可以。
__shared__變量不能將初始化作為其聲明的一部分。
托管變量
__managed__托管內存變量具有如下限制：
- 托管變量的地址非常量
- 托管變量不應有
const限定 - 托管變量不應為引用類型
- 當CUDA運行時可能不處於有效狀態時，不得使用托管變量的地址或值：
- 在對象的靜態/動態初始化或析構過程中，而該對象屬於靜態或本地作用域
- 在exit后被調用的代碼中（如gcc的
__attribute__((destructor))屬性） - 在CUDA運行時被初始化前的代碼中（如gcc的
__attribute__((constructor))屬性）
- 托管變量不可作為
decltype()表達式的unparenthesized參數 - 托管變量和動態分配的托管內存一樣具有連貫性和一致性行為（coherence and consistency）
- 當含托管變量的CUDA程序運行在多GPU設備中時，該托管變量只會被分配一次，而不是每個GPU分配一個
- 在主機端執行的函數中不可使用非外部鏈接(
extern）的托管變量 - 在設備端執行的函數中不可使用非外部鏈接(
extern）或靜態鏈接（static）的托管變量
示例：
__device__ __managed__ int xxx = 10; // OK
struct S1_t {
int field;
S1_t(void) : field(xxx) { };
};
S1_t temp1; // error: use of managed variable
// (xxx) in dynamic initialization
struct S2_t {
~S2_t(void) { xxx = 10; }
};
S2_t temp2; // error: use of managed variable
// (xxx) in the destructor of
// object with static storage
// duration
__device__ __managed__ const int yyy = 10; // error: const qualified type
__device__ __managed__ int &zzz = xxx; // error: reference type
int *ptr = &xxx; // error: use of managed variable
// (xxx) in static initialization
template <int *addr> struct S3_t { };
S3_t<&xxx> temp; // error: address of managed
// variable(xxx) not a
// constant expression
__global__ void kern(int *ptr) {
assert(ptr == &xxx); // OK
xxx = 20; // OK
}
int main(void) {
int *ptr = &xxx; // OK
kern<<<1,1>>>(ptr);
cudaDeviceSynchronize();
xxx++; // OK
decltype(xxx) qqq; // error: managed variable(xxx) used
// as unparenthized argument to
// decltype
decltype((xxx)) zzz = yyy; // OK
}
const常量
__device__, __constant__ 和__shared__變量不允許被聲明為constexpr的。
若常量V是一個被const或constexpr 限定符修飾的變量或類的靜態變量，且V沒有被內存空間描述符（__device__、__shared__、__constant__）修飾，則常量V是一個主機端常量。但常量V仍然可以被設備端代碼直接訪問，只要V滿足如下條件：
V在使用點前用一個常量表達式初始化V沒有被volatile限定符修飾V是一個內置整型（int)或內置浮點型（float），但constexpr比const寬松，只要是非long double的標量即可。
對於constexpr常量V，若函數F是一個__device__ constexpr或__host__ __device__ constexpr函數，且該函數被常量表達式調用，則即使常量V是一個非標量類型，也可以被函數F直接使用。
顯然，設備端代碼不能引用V或取V的地址。
默認情況下，一個constexpr函數不能被執行空間不兼容的另一個函數調用，但可以通過nvcc選項--expt-relaxed-constexpr移除這個限制，從而可以在主機端調用__device__ constexpr函數，反之亦然，也可以在設備端調用__host__ constexpr函數。開發者可以通過__CUDACC_RELAXED_CONSTEXPR__宏是否定義來判斷編譯器是否開啟這個選項。
需要注意即使模版函數被constexpr關鍵字標記，但該模版函數的實例化函數不一定就是constexpr函數。
示例：
const int xxx = 10;
struct S1_t { static const int yyy = 20; };
constexpr int host_arr[] = { 1, 2, 3};
constexpr __device__ int get(int idx) { return host_arr[idx]; }
extern const int zzz;
const float www = 5.0;
__device__ void foo(void) {
int local1[xxx]; // OK
int local2[S1_t::yyy]; // OK
int val1 = xxx; // OK
const float val5 = www; // OK
int val2 = S1_t::yyy; // OK
int val3 = zzz; // error: zzz在使用點前沒有被常量表達式初始化
const int &val3 = xxx; // error: __device__不能引用一個主機端常量
const int *val4 = &xxx; // error: __device__不能取一個主機端常量的地址
int v1 = xxx + 4 + S1_t::yyy; // OK
v1 += get(2); // OK
v1 += get(idx); // get(idx)不是一個常量表達式
v1 += host_arr[2]; // host_arr[2]不是一個標量
}
const int zzz = 20; // error: 注意zzz是在使用點后被初始化的
函數和類
__global__ 函數
__global__ 函數傳參是通過常量內存傳入設備端的，且規定參數大小不得大於4KB。此外__global__ 函數不支持可變長參數。另外，開發者不能將一個操作符函數（如operator+、operator-等等）聲明為__global__的，目前__global__函數尚不支持遞歸，不支持作為類的靜態成員函數，支持類的友元聲明但不支持在友元聲明同時進行定義，例如：
class S1_t {
friend __global__
void foo1(void); // OK: 友元聲明但未定義
template<typename T>
friend __global__
void foo2(void); // OK: __global__函數模版也是一樣
friend __global__
void foo3(void) { } // error: 友元聲明的同時進行定義
template<typename T>
friend __global__
void foo4(void) { } // error: __global__函數模版也是一樣
};
我們可以取到函數的函數指針，但主機端代碼獲取到的__global__函數指針不可用於設備端代碼，反之亦然。顯然，主機端代碼不能獲取device函數的函數指針，設備端代碼也不能獲取__host__函數的函數指針。類似的，不允許在設備端調用創建於主機端的對象的虛函數，反之亦然。
__global__函數或模版不能被聲明為constexpr的，其參數不能為std::initializer_list或va_list類型，不能為右值引用類型。
不允許將帶有虛函數的類的對象作為參數傳遞給__global__函數，同樣的，不允許將虛繼承的派生類的對象作為參數傳遞給__global__函數。
不能在__global__函數實例化或__device__、__constant__變量實例化的類型模板參數、無類型模板參數、模板模板參數中使用如下類型或模版：
- 定義為
__host__、__host__ __device__的類型或模版 - 類型或模版是某個類的
private或protected成員且該類的父類（若有）沒有定義在__device__或__global__函數中 - 匿名類型
- 上述任何類型的復合
例：
template <typename T>
__global__ void myKernel(void) { } //__global__模版函數
class myClass {
private:
struct inner_t { };
public:
static void launch(void) {
myKernel<inner_t><<<1,1>>>(); // error: inner_t類是private成員
}
};
template <typename T> __device__ T d1; //__device__模版變量
template <typename T1, typename T2> __device__ T1 d2;
void fn() {
struct S1_t { };
d1<S1_t> = {}; // error: S1_t是__host__的
auto lam1 = [] { };
d2<int, decltype(lam1)> = 10; // error: lam1是一個匿名類型
}
__global__支持可變參數模版，但只允許一個pack參數，且該pack參數必須置於模版參數最后。
static修飾符
CUDA尚不支持類的static靜態數據成員，除非同時被static const限定符修飾。而在__device__、__global__執行空間的函數中，僅允許普通變量（無任何內存空間描述符）和__shared__變量使用static修飾符，而在__device__ __host__函數中則只允許普通變量使用static修飾符。此外，需要注意若static變量類型為自定義類class，則該類必須擁有空構造函數和空析構函數。此外，static變量不允許動態初始化，例如：
struct S1_t {
int x;
};
struct S2_t {
int x;
__device__ S2_t(void) { x = 10; } //非空構造函數
};
struct S3_t {
int x;
__device__ S3_t(int p) : x(p) { } //非空構造函數
};
__device__ void f1() {
static int i1; // OK
static S1_t i3; // OK，空構造函數
static int i2 = 11; // OK，靜態初始化
static S1_t i4 = {22}; // OK，靜態初始化
static __shared__ int i5; // OK，__device__函數__shared__變量可用static修飾
int x = 33;
static int i6 = x; // error: 動態初始化
static S1_t i7 = {x}; // error: 動態初始化
static S2_t i8; // error: 非空構造函數
static S3_t i9(44); // error: 非空構造函數
}
函數的執行空間
若函數F在首次聲明時被顯式或隱式聲明為默認函數（如果首次顯式聲明為默認函數時有指定執行空間，指定的執行空間會被忽略，但如果不是在首次聲明時顯式默認，則執行空間為指定的執行空間，下述規則忽略），則函數F的執行空間描述符（__host__、__device__）為所有調用函數F的函數的執行空間描述符的集合（__global__也視為__device__）。例如：
class Base {
int x;
public:
__host__ __device__ Base(void) : x(10) {}
};
class Derived : public Base { //隱式聲明了一個默認構造函數
int y;
};
class Other: public Base { //隱式聲明了一個默認構造函數
int z;
};
__device__ void foo(void) {
Derived D1; //Derived的默認構造函數僅被foo函數調用，因此
//Derived::Derived()的執行空間為__device__
Other D2;
}
__host__ void bar(void) {
Other D3; //Other的默認構造函數被foo和bar調用，因此
//Other::Other()的執行空間為__host__ __device__
}
需要注意，當派生類重寫基類的虛函數時，必須保證執行空間與基類虛函數的執行空間一致。但若函數F、D為虛析構函數，且D沒有隱式定義或在非首次聲明的聲明中顯式默認，則F覆蓋的每個虛析構函數D的執行空間的集合既為F的執行空間。例如：
struct Base1 { virtual __host__ __device__ ~Base1() {/*...*/} }; //~Base1()顯式定義非默認
struct Derived1 : Base1 { }; // ~Derived1()的執行空間為 __host__ __device__
struct Base2 { virtual __device__ ~Base2(); }; //首次聲明非顯式默認
__device__ Base2::~Base2() = default; //非首次聲明的聲明中顯示默認
struct Derived2 : Base2 { }; // ~Derived2()的執行空間為__device__
命名空間范圍內匿名union的成員變量不能被__global__或__device__函數引用。
封閉類成員函數的執行空間與定義該封閉類的的最內層的那個指定了執行空間的函數的執行空間相同，若所有嵌套函數都沒有指定執行空間，或該類不是定義在函數中，則其成員函數的執行空間為__host__。
C++特性
CUDA不允許對內置變量（gridDim、blockIdx、blockDim、threadIdx、warpSize）賦值。
CUDA在設備端同樣支持volatile限定符。
默認情況下，std::move和std::forward函數的執行空間為__host__ __device__，因此開發者同樣可以在設備端調用這兩個函數。
CUDA新增__int128_、_Complex、__float128類型，但這些類型只能在主機端使用，且__float128類型只支持64位Linux平台，同時注意編譯器可能會以精度較低的浮點數類型處理__float128類型的常量表達式。
CUDA設備端代碼不支持long double類型，不支持thread_local限定符。
CUDA還支持gcc等編譯器的deprecated 屬性，nvcc選項-Wno-deprecated-declarations將禁用所有棄用警告，而-Werror=deprecated-declarations選項會將所有棄用警告轉換為error。
C++ RTTI（運行時類型識別）特性（typeid 運算符、std::type_info、dynamic_cast運算符）僅支持主機端代碼，不支持設備端代碼。類似的，C++異常僅支持主機端代碼，不支持設備端代碼（包括__global__函數）。
CUDA目前尚不支持設備端的STL。
默認情況下，std::initializer_list的成員函數默認為__host__ __device__執行空間，因此開發者可以在主機端和設備端都能調用std::initializer_list的成員函數，示例：
#include <initializer_list>
__device__ int foo(std::initializer_list<int> in);
__device__ void bar(void) {
foo({4,5,6}); // (a) OK
int i = 4;
foo({i,5,6}); // (b) OK
}