CUDA学习（七十一）

函数参数：
__global__函数参数通过常量内存传递给设备，限制为4 KB。
__global__函数不能有可变数量的参数。
函数中的静态变量:
在__device__或__global__函数的主体中，只有__shared__变量或没有任何设备内存空间说明符的变量可以用静态存储类声明。 在__device__` host函数的主体中，只有未注释的静态变量（即没有设备内存空间说明符）可以用静态存储类声明。 未注释的函数范围静态变量与名称空间范围中定义的__device__`变量具有相同的限制。 如果它们是类类型的，则它们不能具有非空的构造函数或非空的析构函数（请参阅设备内存空间说明符）。
下面显示了函数范围静态变量的合法和非法使用示例:

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 int i2 = 11; // OK
    static S1_t i3; // OK
    static S1_t i4 = { 22 }; // OK
    static __shared__ int i5; // OK
    int x = 33;
    static int i6 = x; // error: dynamic initialization is not allowed
    static S1_t i7 = { x }; // error: dynamic initialization is not allowed
    static S2_t i8; // error: dynamic initialization is not allowed
    static S3_t i9(44); // error: dynamic initialization is not allowed
}
__host__ __device__ void f2() {
    static int i1; // OK
    static __shared__ int i2; // error: __shared__ variable inside
                              // a host function
}

函数指针：
主机代码中使用的__global__函数的地址不能用于设备代码（例如，启动内核）。 同样，在设备(计算能力大于3.5)代码中采用的__global__函数的地址不能用于主机代码。
不允许在主机代码中使用__device__函数的地址。
函数递归:
__global__函数不支持递归。
友函数：
__global__函数或函数模板不能在友函数声明中定义。
例子：

struct S1_t {
    friend __global__
        void foo1(void); // OK: not a definition
    template<typename T>
    friend __global__
        void foo2(void); // OK: not a definition
    friend __global__
        void foo3(void) { } // error: definition in friend declaration
    template<typename T>
    friend __global__
        void foo4(void) { } // error: definition in friend declaration
};

类：
数据成员：
不支持静态数据成员，除了那些也是const限定的成员（请参阅Const限定变量）。
函数成员：
静态成员函数不能是__global__函数。
虚函数：
当派生类中的函数覆盖基类中的虚函数时，覆盖函数和覆盖函数上的执行空间说明符（即__host__，__device__）必须匹配。
不允许将具有虚函数的类的对象作为参数传递给__global__函数
编译器将虚函数表放置在全局或常量内存中。
虚基类：
不允许将__global__函数作为参数传递给从虚基类派生的类的对象。
Anonymous Unions：
名称空间作用域匿名联合的成员变量不能在__global__或__device__函数中引用。
Windows的具体：
CUDA编译器遵循IA64 ABI进行类布局，而Microsoft主编译器不支持。 这可能会导致CUDA编译器计算与Microsoft主编译器不同的类布局和大小，对于满足以下任何条件的类类型“T”或对于任何具有T作为字段类型或基类类型 班级类型（直接或间接）：

• T有虚函数。
• T有一个虚拟基类。
• T具有多个直接或间接空基类的继承。
• T的所有直接和间接基类（'B'）都是空的，并且T（'F'）的第一个字段的类型在其定义中使用B，使得B在F的定义中的偏移量0处布局。

只要受影响的类别类型专门用于主机或设备代码，程序应该正常工作; 请勿在主机和设备代码之间传递此类类型的对象（例如，作为__global__函数的参数或通过cudaMemcpy *（）调用）