CUDA学习（六十九）

设备内存空间说明符：
__device__ ， __shared__和__constant__内存空间说明符不允许用于：

• 类，结构和联合数据成员，
• 形参
• 在主机上执行的函数中的局部变量。

__shared__和__constant__变量隐含了静态存储。
__device__和__constant__变量定义只能在命名空间范围（包括全局命名空间范围）中使用。
__device__，__constant__和__shared__在命名空间范围内定义的类变量类型的变量不能有非空的构造函数或非空的析构函数。 类型的构造函数在翻译单元中的某一点被认为是空的，如果它是一个简单的构造函数或它满足以下所有条件：

• 构造函数已被定义。
• 构造函数没有参数，初始化列表是空的，函数体是空的复合语句。
• 它的类没有虚函数，也没有虚拟基类。
• 它的类的所有基类的默认构造函数可以被认为是空的。
• 对于类的所有非静态数据成员，它们是类类型（或其数组），默认构造函数可以被认为是空的。

如果一个类的析构函数是一个微不足道的析构函数，或者它满足以下所有条件，则该类的析构函数在翻译单元中的某个点上被认为是空的：

• 析构函数已被定义。
• 析构函数体是一个空的复合语句。
• 它的类没有虚函数，也没有虚拟基类。
• 它的类的所有基类的析构函数可以被认为是空的。
• 对于类的所有类非静态数据成员（或其数组），析构函数可以被认为是空的。

在整个程序编译模式下编译时（有关此模式的说明，请参阅nvcc用户手册），不能使用extern关键字将__device__，__shared__和__constant__变量定义为外部变量。 唯一的例外是在__shared__中描述的动态分配的__shared__变量。
在单独的编译模式下编译时（有关此模式的说明，请参阅nvcc用户手册），可以使用extern关键字将__device__，__shared__和__constant__变量定义为外部变量。 如果nvlink无法找到外部变量的定义（除非它是一个动态分配的__shared__变量），它会产生一个错误。
__managed__存储空间说明符
标有__managed__内存空间说明符的变量（“受管理”变量）具有以下限制：

• 管理变量的地址不是常量表达式。
• 托管变量（managed variable）不应具有const限定类型

• 当CUDA运行时可能不处于有效状态时，不应使用受管变量的地址或值，包括以下情况：

• • 在静态或动态初始化或销毁具有静态或线程本地存储持续时间的对象时。
  • 在调用exit（）后执行的代码中（例如，标有gcc的“__attribute __（（destructor））”）的函数。
  • 在CUDA运行时可能未初始化时执行的代码中（例如，使用gcc的“__attribute __（（构造函数）”）标记的函数）。
• 托管变量不能用作decltype（）表达式的未隐含的id表达式参数。
• 托管变量具有与动态分配的托管内存相同的一致性和一致性行为。
• 当包含托管变量的CUDA程序在具有多个GPU的执行平台上运行时，变量只会分配一次，而不是每个GPU。

__device__ __managed__ int xxx = 10; // OK
int *ptr = &xxx; // error: use of managed variable
                 // (xxx) in static initialization
struct S1_t {
    int field;
    S1_t(void) : field(xxx) { };
};
struct S2_t {
    ~S2_t(void) { xxx = 10; }
};
S1_t temp1; // error: use of managed variable
            // (xxx) in dynamic initialization
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
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
}