带扩展Lambda的CUDA代码拷贝构造函数歧义及删除函数问题咨询

问题代码 #

template< bool b >
struct Foo {
  Foo() = default;
  __host__   Foo( const Foo & ) requires(  b ) {}
  __device__ Foo( const Foo & ) requires( !b ) {}
};

template< typename Lambda >
__global__
void kernel( Lambda ) {}

int main() {
  Foo< true > foo;
  auto la = [foo] __device__ (){ };
  kernel<<< 1, 1 >>>( la );
}

编译环境与命令 #

• NVCC版本：12.1
• 主机编译器：GCC 11.3
• 编译命令：nvcc main.cu -std=c++20 --expt-extended-lambda

编译错误信息 #

1. 拷贝构造歧义错误：

copy constructor for class "Foo<true>" is ambiguous
    auto la = [foo] __attribute__((device)) (){ };
               ^

2. 删除函数引用错误：

function "lambda []()->void::(const lambda []()->void &)"
(declared implicitly) cannot be referenced -- it is a deleted function
    kernel<<< 1, 1 >>>( la );
                        ^

疑问 #

• 为何Foo<true>的拷贝构造函数会存在歧义？
• 哪个函数被隐式删除？编译器在何处尝试调用该函数？

解答 #

1. Foo<true>拷贝构造函数歧义的原因 #

当实例化Foo<true>时，按C++20规则，只有__host__修饰的拷贝构造函数满足requires(b)约束（b=true），__device__修饰的拷贝构造函数因requires(!b)不满足（!b=false），应该被SFINAE规则排除在重载决议的候选函数集之外。

但NVCC在处理CUDA设备属性（__host__/__device__）与C++20 requires约束的交互时存在异常：它没有正确移除不满足requires条件的__device__拷贝构造函数。这导致主机上下文（main函数中构造lambda并捕获foo时），编译器看到两个签名完全相同的拷贝构造函数（仅属性和约束不同），触发重载决议歧义。

2. 被隐式删除的函数及调用场景 #

被隐式删除的是lambda的拷贝构造函数。

原因是：lambda捕获了Foo<true>实例，但Foo<true>的拷贝构造存在歧义，编译器无法为lambda生成合法的隐式拷贝构造函数，因此将其标记为删除。

当执行kernel<<<1,1>>>(la);时，kernel参数按值传递，需要拷贝lambda实例（用于设备端执行），此时编译器尝试调用lambda的拷贝构造函数，但该函数已被删除，因此触发第二个错误。

内容的提问来源于stack exchange，提问作者tommsch