CUDA学习（七十七）

扩展的Lambda类型特征：
编译器提供类型特征来在编译时检测扩展lambda表达式的闭包类型（closure types）：
__nv_is_extended_device_lambda_closure_type（type）：如果'type'是为扩展的__device__ lambda创建的闭包类，那么该特征为true，否则为false。
__nv_is_extended_host_device_lambda_closure_type（type）：如果'type'是为扩展的__host__ __device__ lambda创建的闭包类，那么该特征为true，否则为false。
这些特征可以在所有编译模式下使用，而不管是否启用lambda或lambda扩展:
例子：

#define IS_D_LAMBDA(X) __nv_is_extended_device_lambda_closure_type(X)
#define IS_HD_LAMBDA(X) __nv_is_extended_host_device_lambda_closure_type(X)
auto lam0 = [] __host__ __device__{};
void foo(void) {
    auto lam1 = [] {};
    auto lam2 = [] __device__{};
    auto lam3 = [] __host__ __device__{};
    // lam0 is not an extended lambda (since defined outside function scope)
    static_assert(!IS_D_LAMBDA(decltype(lam0)), "");
    static_assert(!IS_HD_LAMBDA(decltype(lam0)), "");
    // lam1 is not an extended lambda (since no execution space annotations)
    static_assert(!IS_D_LAMBDA(decltype(lam1)), "");
    static_assert(!IS_HD_LAMBDA(decltype(lam1)), "");
    // lam2 is an extended __device__ lambda
    static_assert(IS_D_LAMBDA(decltype(lam2)), "");
    static_assert(!IS_HD_LAMBDA(decltype(lam2)), "");
    // lam3 is an extended __host__ __device__ lambda
    static_assert(!IS_D_LAMBDA(decltype(lam3)), "");
    static_assert(IS_HD_LAMBDA(decltype(lam3)), "");
}

扩展的Lambda限制：
在调用主机编译器之前，CUDA编译器将用名称空间范围中定义的占位符类型实例替换扩展的lambda表达式。 占位符类型的模板参数需要获取封装原始扩展lambda表达式的函数的地址。 这对于正确执行其模板参数涉及扩展lambda的闭包类型的__global__函数模板是必需的。 关闭功能计算如下。
根据定义，扩展lambda存在于__host__或__host__ __device__函数的直接或嵌套块范围内。 如果这个函数不是lambda表达式的operator（），那么它被认为是扩展lambda的封闭函数。 否则，扩展lambda在一个或多个封闭lambda表达式的operator（）的直接或嵌套块范围内定义。 如果最外面的这样的lambda表达式是在函数F的立即或嵌套块范围中定义的，那么F是计算出来的封闭函数，否则封闭函数不存在。
例子：

void foo(void) {
    // enclosing function for lam1 is "foo"
    auto lam1 = [] __device__{};
    auto lam2 = [] {
        auto lam3 = [] {
            // enclosing function for lam4 is "foo"
            auto lam4 = [] __host__ __device__{};
        };
    };
}
auto lam6 = [] {
    // enclosing function for lam7 does not exist
    auto lam7 = [] __host__ __device__{};
};

以下是对扩展lambda表达式的限制：
1.扩展的lambda不能在另一个扩展的lambda表达式中定义。 例：

void foo(void) {
    auto lam1 = [] __host__ __device__{
        // error: extended lambda defined within another extended lambda
        auto lam2 = [] __host__ __device__{};
    };

2.扩展lambda不能在泛型lambda表达式中定义。 例：

void foo(void) {
    auto lam1 = [](auto) {
        // error: extended lambda defined within a generic lambda
        auto lam2 = [] __host__ __device__{};
    }

3.如果在一个或多个嵌套lambda表达式的直接或嵌套块范围内定义扩展lambda，则最外面的这种lambda表达式必须在函数的直接或嵌套块范围内定义。 例：

auto lam1 = [] {
    // error: outer enclosing lambda is not defined within a
    // non-lambda-operator() function.
    auto lam2 = [] __host__ __device__{};
};

4.扩展lambda的封闭函数必须被命名并且它的地址可以被采用。 如果封闭函数是类成员，则必须满足以下条件：

• 包含成员函数的所有类都必须有一个名称。
• 成员函数不能在其父类中拥有私有或受保护的访问权限。
• 所有封闭类不得在其各自的父类中拥有私有或受保护的访问权限。

例子：

void foo(void) {
    // OK
    auto lam1 = [] __device__{ return 0; };
    {
        // OK
        auto lam2 = [] __device__{ return 0; };
        // OK
        auto lam3 = [] __device__ __host__{ return 0; };
    }
}
struct S1_t {
    S1_t(void) {
        // Error: cannot take address of enclosing function
        auto lam4 = [] __device__{ return 0; };
    }
};
class C0_t {
    void foo(void) {
        // Error: enclosing function has private access in parent class
        auto temp1 = [] __device__{ return 10; };
    }
    struct S2_t {
        void foo(void) {
            // Error: enclosing class S2_t has private access in its
            // parent class
            auto temp1 = [] __device__{ return 10; };
        }
    };
};