CUDA学习（七十五）

thread_local：
设备代码中不允许使用thread_local存储说明符。
__global__函数和函数模板:
如果在__global__函数模板实例的模板参数中使用与lambda表达式关联的闭包类型，那么lambda表达式必须在__device__或__global__函数的直接或嵌套块作用域中定义，或者必须是扩展lambda。
例子：

template <typename T>
__global__ void kernel(T in) { }
__device__ void foo_device(void)
{
    // All kernel instantiations in this function
    // are valid, since the lambdas are defined inside
    // a __device__ function.
    kernel << <1, 1 >> >([] __device__{});
    kernel << <1, 1 >> >([] __host__ __device__{});
    kernel << <1, 1 >> >([] {});
}
auto lam1 = [] {};
auto lam2 = [] __host__ __device__{};
void foo_host(void)
{
    // OK: instantiated with closure type of an extended __device__ lambda
    kernel << <1, 1 >> >([] __device__{});
    // OK: instantiated with closure type of an extended __host__ __device__
    // lambda
    kernel << <1, 1 >> >([] __host__ __device__{});
    // error: unsupported: instantiated with closure type of a lambda
    // that is not an extended lambda
    kernel << <1, 1 >> >([] {});
    // error: unsupported: instantiated with closure type of a lambda
    // that is not an extended lambda
    kernel << <1, 1 >> >(lam1);
    // error: unsupported: instantiated with closure type of a lambda
    // that is not an extended lambda
    kernel << <1, 1 >> >(lam2);
}

__global__函数或函数模板不能声明为constexpr，并且不能有尾随返回类型。
以下类型不适用于__global__函数或函数模板的参数：

• 右值引用类型
• std::initializer_list
• 限制合格的参考类型

可变参数__global__函数模板具有以下限制：

• 只允许一个包参数。
• 包参数必须在模板参数列表中最后列出。

例子：

template <template <typename...> class Wrapper, typename... Pack>
__global__ void foo1(Wrapper<Pack...>);
// error: pack parameter is not last in parameter list
template <typename... Pack, template <typename...> class Wrapper>
__global__ void foo2(Wrapper<Pack...>);
// error: multiple parameter packs
template <typename... Pack1, int...Pack2, template<typename...> class Wrapper1,
    template<int...> class Wrapper2>
__global__ void foo3(Wrapper1<Pack1...>, Wrapper2<Pack2...>);

__device __ / __ constant __ / __ shared__变量:
__device__，__constant__和__shared__变量不能用关键字constexpr标记，并且不能有右值引用类型。
默认功能:
CUDA编译器忽略默认功能上的执行空间说明符。
例子：

struct S1 {
    // warning: __host__ annotation on a defaulted function is ignored
    __host__ S1() = default;
};
struct S2 {
    // warning: __device__ annotation on a defaulted function is ignored
    __device__ ~S2() = default;
};

C ++ 14的特性：
主机编译器默认启用的C ++ 14功能也受nvcc支持。 传递nvcc -std = c ++ 14标志将打开所有C ++ 14功能，并使用相应的C ++ 14方言选项16调用主机预处理器，编译器和链接器。本节介绍对受支持的C ++ 14个功能。
具有推导返回类型的函数：
如果__device__函数推导出返回类型，则CUDA前端编译器将在调用主机编译器之前将函数声明更改为具有void返回类型。 这可能会导致在主机代码中反演__device__函数的推导返回类型的问题。 因此，除非在__CUDA_ARCH__未定义时引用不存在，否则CUDA编译器将发出编译时错误以引用设备函数体外的此类推导返回类型。
例子：

__device__ auto fn1(int x) {
    return x;
}
__device__ decltype(auto) fn2(int x) {
    return x;
}
__device__ void device_fn1() {
    // OK
    int(*p1)(int) = fn1;
}
// error: referenced outside device function bodies
decltype(fn1(10)) g1;
void host_fn1() {
    // error: referenced outside device function bodies
    int(*p1)(int) = fn1;
    struct S_local_t {
        // error: referenced outside device function bodies
        decltype(fn2(10)) m1;
        S_local_t() : m1(10) { }
    };
}
// error: referenced outside device function bodies
template <typename T = decltype(fn2)>
void host_fn2() { }
template<typename T> struct S1_t { };
// error: referenced outside device function bodies
struct S1_derived_t : S1_t<decltype(fn1)> { };

变量模板：
__device __ / __ constant__变量模板在Windows上不能具有const限定类型。
例子：

// error: a __device__ variable template cannot
// have a const qualified type on Windows
template <typename T>
__device__ const T d1(2);
int *const x = nullptr;
// error: a __device__ variable template cannot
// have a const qualified type on Windows
template <typename T>
__device__ T *const d2(x);
// OK
template <typename T>
__device__ const T *d3;
__device__ void fn() {
    int t1 = d1<int>;
    int *const t2 = d2<int>;
    const int *t3 = d3<int>;
}

[[deprecated]] 属性：
CUDA前端编译器接受[[deprecated]]属性，并在发送给主机编译器的代码中重新生成它。 但是，CUDA前端编译器不会针对此属性的使用发出任何警告。