CUDA学习（七十四）

Constexpr变量：
'V'表示命名空间范围变量或已标记为constexpr且没有执行空间注释的类静态成员变量（例如__device__，__constant__，__shared__）。 V被认为是主机代码变量。
如果V的标量类型不是long double，并且该类型不是volatile限定的，则可以直接在设备代码中使用V的值。 另外，如果V是非标量类型，那么如果对函数的调用是常量表达式，则可以在constexpr __device__或__host__ __device__函数内使用V的标量元素。 器件源代码不能包含对V的引用或取V的地址。
如果V的标量类型为14而不是long double，并且该类型不是volatile限定的，则可以直接在设备代码中使用V的值。 另外，如果V是非标量类型，那么如果函数的调用是一个常量表达式，则在constexpr __device__或__host__ __device__函数内可以使用V的标量元素。15.设备源代码不能包含对V 或采取V的地址。
例子：

constexpr int xxx = 10;
constexpr int yyy = xxx + 4;
struct S1_t { static constexpr int qqq = 100; };
constexpr int host_arr[] = { 1, 2, 3 };
constexpr __device__ int get(int idx) { return host_arr[idx]; }
__device__ int foo(int idx) {
    int v1 = xxx + yyy + S1_t::qqq; // OK
    const int &v2 = xxx; // error: reference to host constexpr
                         // variable
    const int *v3 = &xxx; // error: address of host constexpr
                          // variable
    const int &v4 = S1_t::qqq; // error: reference to host constexpr
                               // variable
    const int *v5 = &S1_t::qqq; // error: address of host constexpr
                                // variable
    v1 += get(2); // OK: 'get(2)' is a constant
                  // expression.
    v1 += get(idx); // error: 'get(idx)' is not a constant
                    // expression
    v1 += host_arr[2]; // error: 'host_arr' does not have
                       // scalar type.
    return v1;
}

内联命名空间：
对于输入CUDA转换单元，CUDA编译器可以调用主机编译器来编译转换单元中的主机代码。 在传递给主机编译器的代码中，如果输入CUDA转换单元包含任何以下实体的定义，CUDA编译器将注入其他编译器生成的代码：

• __global__函数或函数模板实例化
• __device__, __constant__
• 具有表面或纹理类型的变量

编译器生成的代码包含对已定义实体的引用。 如果实体在内联名称空间内定义，并且在封闭名称空间中定义了另一个具有相同名称和类型签名的实体，则主引擎编译器可能认为此引用不明确，并且主机编译将失败。
通过在内联命名空间中定义的这些实体使用唯一名称可以避免此限制。
例子：

__device__ int Gvar;
inline namespace N1 {
    __device__ int Gvar;
}
// <-- CUDA compiler inserts a reference to "Gvar" at this point in the
// translation unit. This reference will be considered ambiguous by the
// host compiler and compilation will fail.

inline namespace N1 {
    namespace N2 {
        __device__ int Gvar;
    }
}
namespace N2 {
    __device__ int Gvar;
}
// <-- CUDA compiler inserts reference to "::N2::Gvar" at this point in
// the translation unit. This reference will be considered ambiguous by
// the host compiler and compilation will fail.

内联未命名的名称空间：
以下实体不能在内联未命名名称空间内的名称空间范围内声明：

• __device__，__shared__和__constant__变量
• __global__函数和函数模板
• 具有表面或纹理类型的变量

例子：

inline namespace {
    namespace N2 {
        template <typename T>
        __global__ void foo(void); // error
        __global__ void bar(void) { } // error
        template <>
        __global__ void foo<int>(void) { } // error
        __device__ int x1b; // error
        __constant__ int x2b; // error
        __shared__ int x3b; // error
        texture<int> q2; // error
        surface<int> s2; // error
    }
};