5.3 C++ 语言支持¶

本文档为 NVIDIA CUDA Programming Guide 官方文档中文翻译版

原文地址：https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-support.html

本页面是否有帮助？

5.3. C++ 语言支持¶

nvcc 根据以下规范处理 CUDA 和设备代码：

• C++03 (ISO/IEC 14882:2003)，使用 --std=c++03 标志。
• C++11 (ISO/IEC 14882:2011)，使用 --std=c++11 标志。
• C++14 (ISO/IEC 14882:2014)，使用 --std=c++14 标志。
• C++17 (ISO/IEC 14882:2017)，使用 --std=c++17 标志。
• C++20 (ISO/IEC 14882:2020)，使用 --std=c++20 标志。

向 nvcc 传递 -std=c++<version> 标志会启用与指定版本相关的所有 C++ 功能，并使用相应的 C++ 方言选项调用主机预处理器、编译器和链接器。

编译器支持所有受支持标准中的语言功能，但需遵循后续章节中报告的限制。

5.3.1. C++11 语言功能¶

5.3.2. C++14 语言特性¶

5.3.3. C++17 语言特性¶

5.3.4. C++20 语言特性¶

GCC 版本 ≥ 10.0，Clang 版本 ≥ 10.0，Microsoft Visual Studio ≥ 2022，以及 nvc++ 版本 ≥ 20.7。

5.3.5. CUDA C++ 标准库¶

CUDA 提供了一个 C++ 标准库（STL）的实现，称为 libcu++。该库具有以下优点：

• 功能在主机和设备上均可用。
• 与 CUDA 工具包支持的所有 Linux 和 Windows 平台兼容。
• 与最近两个主要版本的 CUDA 工具包支持的所有 GPU 架构兼容。
• 与当前和上一个主要版本的所有 CUDA 工具包兼容。
• 提供新近标准版本（包括 C++20、C++23 和 C++26）中 C++ 标准库功能的 C++17 向后移植。
• 支持扩展数据类型，例如 128 位整数（__int128）、半精度浮点数（__half）、Bfloat16（__nv_bfloat16）和四精度浮点数（__float128）。
• 针对设备代码进行了高度优化。

此外，libcu++ 还提供了 C++ 标准库中未包含的扩展功能，以提高生产力和应用程序性能。这些功能包括数学函数、内存操作、同步原语、容器扩展、CUDA 内部函数的高级抽象、C++ PTX 包装器等。

libcu++ 作为 CUDA 工具包的一部分提供，同时也是开源 CCCL 仓库的一部分。

5.3.6. C 标准库函数¶

5.3.6.1. clock() 和 clock64()¶

__host__ __device__ clock_t   clock();
__device__          long long clock64();

在设备代码中执行时，它返回每个流式多处理器（SM）计数器的值，该计数器在每个时钟周期递增。在内核开始和结束时采样此计数器，将两个值相减，并记录每个线程的结果，可以估算出设备执行该线程所花费的时钟周期数。然而，这个值并不代表设备实际执行该线程指令所花费的时钟周期数。前者大于后者，因为线程是分时执行的。

5.3.6.2. printf()¶

int printf(const char* format[, arg, ...]);

该函数将来自内核的格式化输出打印到主机端的输出流。

内核内的 printf() 函数行为类似于标准 C 库的 printf() 函数。用户应参考其主机系统的手册页以获取 printf() 行为的完整描述。本质上，作为 format 传入的字符串被输出到主机上的一个流。

printf() 命令像任何其他设备端函数一样执行：每个线程执行，并在调用线程的上下文中执行。在多线程内核中，对 printf() 的直接调用将由每个线程使用该线程指定的数据来执行。因此，主机流上会出现多个版本的输出字符串，每个字符串对应一个遇到 printf() 的线程。

与返回打印字符数的 C 标准 printf() 不同，CUDA 的 printf() 返回解析的参数数量。如果格式字符串后没有参数，则返回 0。如果格式字符串为 NULL，则返回 -1。如果发生内部错误，则返回 -2。

在内部，printf() 使用一个共享的数据结构，因此调用 printf() 可能会改变线程的执行顺序。具体来说，调用 printf() 的线程可能比不调用 printf() 的线程执行路径更长，并且该路径的长度取决于 printf() 的参数。但是，请注意，除了在显式的 __syncthreads() 屏障处，CUDA 不保证线程的执行顺序。因此，无法判断执行顺序是否被 printf() 或硬件中的其他调度行为所修改。

格式说明符

与标准 printf() 一样，格式说明符的形式为：%[flags][width][.precision][size]type

支持以下字段。有关所有行为的完整描述，请参阅广泛可用的文档。

• 标志： # , ' ' , 0 , + , -
• 宽度： * , 0-9
• 精度：0-9
• 大小：h、l、ll
• 类型：%cdiouxXpeEfgGaAs

限制

printf() 输出的最终格式化发生在主机系统上。这意味着格式字符串必须能被主机系统的编译器和 C 库理解。尽管已尽最大努力确保 CUDA printf() 函数支持的格式说明符是大多数常见主机编译器所支持格式的通用子集，但其确切行为将取决于主机操作系统。

printf() 接受所有有效的标志和类型组合。这是因为它无法确定在最终输出格式化的主机系统上哪些组合有效、哪些无效。因此，如果程序发出的格式字符串包含无效组合，输出可能是未定义的。

除了格式字符串外，printf() 函数最多可以接受 32 个参数。任何额外的参数将被忽略，格式说明符将按原样输出。

由于 Windows 平台（32 位）和 Linux 平台（64 位）上 long 类型的大小不同，在 Linux 机器上编译然后在 Windows 机器上运行的内核，对于所有包含 %ld 的格式字符串，将产生损坏的输出。为确保安全，建议编译和执行平台保持一致。

主机端缓冲区

printf() 的输出缓冲区在内核启动前被设置为固定大小。该缓冲区是循环的，因此如果内核执行期间产生的输出超过了缓冲区容量，较早的输出将被覆盖。只有在执行以下操作之一时，缓冲区才会被刷新：

• 通过 <<< >>> 或 cuLaunchKernel() 启动内核：在启动开始时，如果 CUDA_LAUNCH_BLOCKING 环境变量设置为 1，则在启动结束时也会刷新；
• 通过 cudaDeviceSynchronize()、cuCtxSynchronize()、cudaStreamSynchronize()、cuStreamSynchronize()、cudaEventSynchronize() 或 cuEventSynchronize() 进行同步；
• 通过任何阻塞版本的 cudaMemcpy*() 或 cuMemcpy*() 进行内存复制；
• 通过 cuModuleLoad() 或 cuModuleUnload() 加载/卸载模块；
• 通过 cudaDeviceReset() 或 cuCtxDestroy() 销毁上下文；
• 在执行由 cudaLaunchHostFunc() 或 cuLaunchHostFunc() 添加的流回调之前。

请注意，程序退出时缓冲区不会自动刷新。

以下 API 函数设置和检索用于将 printf() 参数和内部元数据传输到主机的缓冲区大小。默认大小为 1 兆字节。

• cudaDeviceGetLimit(size_t* size, cudaLimitPrintfFifoSize)
• cudaDeviceSetLimit(cudaLimitPrintfFifoSize, size_t size)

示例

以下代码示例：

#include <stdio.h>

__global__ void helloCUDA(float value) {
    printf("Hello thread %d, value=%f\n", threadIdx.x, value);
}

int main() {
    helloCUDA<<<1, 5>>>(1.2345f);
    cudaDeviceSynchronize();
    return 0;
}

将输出：

Hello thread 2, value=1.2345
Hello thread 1, value=1.2345
Hello thread 4, value=1.2345
Hello thread 0, value=1.2345
Hello thread 3, value=1.2345

查看 Compiler Explorer 上的示例。

以下代码示例：

#include <stdio.h>

__global__ void helloCUDA(float value) {
    if (threadIdx.x == 0)
        printf("Hello thread %d, value=%f\n", threadIdx.x, value);
}

int main() {
    helloCUDA<<<1, 5>>>(1.2345f);
    cudaDeviceSynchronize();
    return 0;
}

将输出：

Hello thread 0, value=1.2345

显然，if() 语句限制了哪些线程调用 printf()，因此只看到一行输出。

查看 Compiler Explorer 上的示例。

5.3.6.3. memcpy() 和 memset()¶

__host__ __device__ void* memcpy(void* dest, const void* src, size_t size);

该函数将 size 字节从 src 指向的内存位置复制到 dest 指向的内存位置。

__host__ __device__ void* memset(void* ptr, int value, size_t size);

该函数将 ptr 指向的内存块的 size 字节设置为 value，value 被解释为 unsigned char。

5.3.6.4. malloc() 和 free()¶

__host__ __device__ void* malloc(size_t size);
// 或者 <cuda/std/cstdlib> 头文件中的 cuda::std::malloc(), cuda::std::calloc()

函数 malloc()（设备端）、cuda::std::malloc() 和 cuda::std::calloc() 从设备堆中分配至少 size 字节，并返回一个指向已分配内存的指针。如果没有足够的内存来满足请求，则返回 NULL。返回的指针保证按 16 字节边界对齐。

__device__ void* __nv_aligned_device_malloc(size_t size, size_t align);
// 或者 <cuda/std/cstdlib> 头文件中的 cuda::std::aligned_alloc()

函数 __nv_aligned_device_malloc() 和 C++ cuda::std::aligned_alloc() 从设备堆中分配至少 size 字节，并返回一个指向已分配内存的指针。如果没有足够的内存来满足请求的大小或对齐要求，则返回 NULL。已分配内存的地址是 align 的倍数。align 必须是非零的 2 的幂。

__host__ __device__ void free(void* ptr);
// 或者 <cuda/std/cstdlib> 头文件中的 cuda::std::free()

设备端函数 free() 和 cuda::std::free() 释放 ptr 指向的内存，该内存必须是由之前对 malloc()、cuda::std::malloc()、cuda::std::calloc()、__nv_aligned_device_malloc() 或 cuda::std::aligned_alloc() 的调用返回的。如果 ptr 是 NULL，则对 free() 或 cuda::std::free() 的调用将被忽略。使用相同的 ptr 重复调用 free() 或 cuda::std::free() 会导致未定义行为。 由给定 CUDA 线程通过 malloc()、cuda::std::malloc()、cuda::std::calloc()、__nv_aligned_device_malloc() 或 cuda::std::aligned_alloc() 分配的内存，将在 CUDA 上下文的整个生命周期内保持分配状态，直到通过调用 free() 或 cuda::std::free() 显式释放。此内存可被其他 CUDA 线程使用，甚至包括后续内核启动的线程。任何 CUDA 线程都可以释放由另一个线程分配的内存；但是，应注意确保同一指针不会被释放超过一次。

堆内存 API

必须在任何在设备代码中分配或释放内存的程序（包括使用 new 和 delete 关键字）之前指定设备内存堆的大小。如果任何程序使用设备内存堆而未显式指定堆大小，则会分配一个默认的 8 兆字节堆。

以下 API 函数用于获取和设置堆大小：

• cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)
• cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

授予的堆大小将至少为 size 字节。cuCtxGetLimit() 和 cudaDeviceGetLimit() 返回当前请求的堆大小。

堆的实际内存分配发生在模块加载到上下文时，无论是通过 CUDA 驱动程序 API（参见模块）显式加载，还是通过 CUDA 运行时 API 隐式加载。如果内存分配失败，模块加载会产生 CUDA_ERROR_SHARED_OBJECT_INIT_FAILED 错误。

堆大小在模块加载后无法更改，并且不会根据需求动态调整。

为设备堆保留的内存是额外的，不包含通过主机端 CUDA API 调用（如 cudaMalloc()）分配的内存。

与主机内存 API 的互操作性

通过设备端函数 malloc()、cuda::std::malloc()、cuda::std::calloc()、__nv_aligned_device_malloc()、cuda::std::aligned_alloc() 或 new 关键字分配的内存，不能通过运行时或驱动程序 API 调用（如 cudaMalloc、cudaMemcpy 或 cudaMemset）来使用或释放。同样，通过主机运行时 API 分配的内存也不能使用设备端函数 free()、cuda::std::free() 或 delete 关键字来释放。

每线程分配示例：

#include <stdlib.h>
#include <stdio.h>

__global__ void single_thread_allocation_kernel() {
    size_t size = 123;
    char*  ptr  = (char*) malloc(size);
    memset(ptr, 0, size);
    printf("Thread %d got pointer: %p\n", threadIdx.x, ptr);
    free(ptr);
}

int main() {
    // 设置堆大小为 128 兆字节。
    // 注意，这必须在任何内核启动之前完成。
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024);
    single_thread_allocation_kernel<<<1, 5>>>();
    cudaDeviceSynchronize();
    return 0;
}

Thread 0 got pointer: 0x20d5ffe20
Thread 1 got pointer: 0x20d5ffec0
Thread 2 got pointer: 0x20d5fff60
Thread 3 got pointer: 0x20d5f97c0
Thread 4 got pointer: 0x20d5f9720

请注意，每个线程都遇到了 malloc() 和 memset() 命令，因此接收并初始化了自己的分配。

在 Compiler Explorer 上查看此示例。

每个线程块分配示例：

#include <stdlib.h>

__global__ void block_level_allocation_kernel() {
    __shared__ int* data;
    // 线程块中的第一个线程执行分配，并通过共享内存将指针
    // 与所有其他线程共享，以便访问可以合并。
    if (threadIdx.x == 0) {
        size_t size = blockDim.x * 64; // 为每个线程分配 64 字节。
        data = (int*) malloc(size);
    }
    __syncthreads();
    // 检查是否分配失败
    if (data == nullptr)
        return;

    // 线程索引到内存中，确保合并访问
    for (int i = 0; i < 64; ++i)
        data[i * blockDim.x + threadIdx.x] = threadIdx.x;
    // 确保所有线程在释放内存前完成操作
    __syncthreads();

    // 只能由一个线程释放内存！
    if (threadIdx.x == 0)
        free(data);
}

int main() {
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128 * 1024 * 1024);
    block_level_allocation_kernel<<<10, 128>>>();
    cudaDeviceSynchronize();
    return 0;
}

在 Compiler Explorer 上查看此示例。

内核启动间持久化分配示例：

#include <stdlib.h>
#include <stdio.h>

const int NUM_BLOCKS = 20;

__device__ int* data_ptrs[NUM_BLOCKS]; // 每个线程块的指针

__global__ void allocate_memory_kernel() {
    // 只有线程块中的第一个线程执行分配，
    // 因为每个线程块只需要一次分配。
    if (threadIdx.x == 0)
        data_ptrs[blockIdx.x] = (int*) malloc(blockDim.x * 4);
    __syncthreads();
    // 检查是否分配失败
    if (data_ptrs[blockIdx.x] == nullptr)
        return;
    // 所有线程并行地将数据清零
    data_ptrs[blockIdx.x][threadIdx.x] = 0;
}

// 简单示例：将线程 ID 存储到每个元素中
__global__ void use_memory_kernel() {
    int* ptr = data_ptrs[blockIdx.x];
    if (ptr != nullptr)
        ptr[threadIdx.x] += threadIdx.x;
}

// 在释放缓冲区之前打印其内容
__global__ void free_memory_kernel() {
    int* ptr = data_ptrs[blockIdx.x];
    if (ptr != nullptr)
        printf("Block %d, Thread %d: final value = %d\n",
            blockIdx.x, threadIdx.x, ptr[threadIdx.x]);
    // 只能由一个线程释放！
    if (threadIdx.x == 0)
        free(ptr);
}

int main() {
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
    // 分配内存
    allocate_memory_kernel<<<NUM_BLOCKS, 10>>>();

    // 使用内存
    use_memory_kernel<<<NUM_BLOCKS, 10>>>();
    use_memory_kernel<<<NUM_BLOCKS, 10>>>();
    use_memory_kernel<<<NUM_BLOCKS, 10>>>();

    // 释放内存
    free_memory_kernel<<<NUM_BLOCKS, 10>>>();
    cudaDeviceSynchronize();
    return 0;
}

5.3.6.5.alloca()¶

__host__ __device__ void* alloca(size_t size);

alloca() 函数在调用者的栈帧内分配 size 字节的内存。返回值是指向所分配内存的指针。当从设备代码调用该函数时，内存的起始地址是 16 字节对齐的。当调用者从 alloca() 返回时，内存会自动释放。

示例：

__device__ void device_function(int num_items) {
    int4* ptr = (int4*) alloca(num_items * sizeof(int4));
    // 使用 ptr
    ...
}

5.3.7. Lambda 表达式¶

编译器通过将 lambda 表达式或闭包类型（C++11）与最内层封闭函数作用域的执行空间相关联，来确定其执行空间。如果没有封闭函数作用域，则执行空间被指定为 __host__。

执行空间也可以使用扩展 lambda 语法显式指定。

示例：

auto global_lambda = [](){ return 0; }; // __host__

void host_function() {
    auto lambda1 = [](){ return 1; };   // __host__
    [](){ return 3; };                  // __host__, 闭包类型（lambda 表达式体）
}

__device__ void device_function() {
    auto lambda2 = [](){ return 2; };   // __device__
}

__global__ void kernel_function(void) {
    auto lambda3 = [](){ return 3; };   // __device__
}

__host__ __device__ void host_device_function() {
    auto lambda4 = [](){ return 4; };   // __host__ __device__
}

using function_ptr_t = int (*)();

__device__ void device_function(float          value,
                                function_ptr_t ptr = [](){ return 4; } /* __host__ */) {}

请参阅 Compiler Explorer 上的示例。

5.3.7.1. Lambda 表达式与 __global__ 函数参数¶

只有当 lambda 表达式或闭包类型的执行空间是 __device__ 或 __host__ __device__ 时，才能用作 __global__ 函数的参数。全局或命名空间作用域的 lambda 表达式不能用作 __global__ 函数的参数。

示例：

template <typename T>
 __global__ void kernel(T input) {}

 __device__ void device_function() {
     // 设备内核调用需要单独编译（-rdc=true 标志）
     kernel<<<1, 1>>>([](){});
     kernel<<<1, 1>>>([] __device__() {});          // 扩展 lambda
     kernel<<<1, 1>>>([] __host__ __device__() {}); // 扩展 lambda
 }

 auto global_lambda = [] __host__ __device__() {};

 void host_function() {
     kernel<<<1, 1>>>([] __device__() {});          // 正确，扩展 lambda
     kernel<<<1, 1>>>([] __host__ __device__() {}); // 正确，扩展 lambda
 //  kernel<<<1, 1>>>([](){});                      // 错误，执行空间为 __host__ 的闭包类型
 //  kernel<<<1, 1>>>(global_lambda);               // 错误，扩展 lambda，但在全局作用域
 }

5.3.7.2. 扩展 Lambda 表达式¶

nvcc 的 --extended-lambda 标志允许在 lambda 表达式中显式标注执行空间。这些标注应出现在 lambda 引导符之后，以及可选的 lambda 声明符之前。当指定 --extended-lambda 标志时，nvcc 会定义宏 __CUDACC_EXTENDED_LAMBDA__。

• 扩展 lambda 定义在 __host__ 或 __host__ __device__ 函数的直接或嵌套块作用域内。
• 扩展设备 lambda 是用 __device__ 关键字标注的 lambda 表达式。
• 扩展主机-设备 lambda 是用 __host__ __device__ 关键字标注的 lambda 表达式。

与标准 lambda 表达式不同，扩展 lambda 可以用作 __global__ 函数中的类型参数。

示例：

void host_function() {
    auto lambda1 = [] {};                      // 不是扩展 lambda：没有显式的执行空间标注
    auto lambda2 = [] __device__ {};           // 扩展 lambda
    auto lambda3 = [] __host__ __device__ {};  // 扩展 lambda
    auto lambda4 = [] __host__ {};             // 不是扩展 lambda
}

__host__ __device__ void host_device_function() {
    auto lambda1 = [] {};                      // 不是扩展 lambda：没有显式的执行空间标注
    auto lambda2 = [] __device__ {};           // 扩展 lambda
    auto lambda3 = [] __host__ __device__ {};  // 扩展 lambda
    auto lambda4 = [] __host__ {};             // 不是扩展 lambda
}

__device__ void device_function() {
    // 此函数内的所有 lambda 都不是扩展 lambda，
    // 因为其外层函数不是 __host__ 或 __host__ __device__ 函数。
    auto lambda1 = [] {};
    auto lambda2 = [] __device__ {};
    auto lambda3 = [] __host__ __device__ {};
    auto lambda4 = [] __host__ {};
}

auto global_lambda = [] __host__ __device__ { }; // 不是扩展 lambda，因为它不是定义在
                                                 // __host__ 或 __host__ __device__ 函数内

5.3.7.3. 扩展 Lambda 类型特征¶

编译器提供了类型特征，用于在编译时检测扩展 lambda 的闭包类型。

bool __nv_is_extended_device_lambda_closure_type(type);

如果 type 是为扩展 __device__ lambda 创建的闭包类，则此函数返回 true，否则返回 false。

bool __nv_is_extended_device_lambda_with_preserved_return_type(type);

如果 type 是为扩展 __device__ lambda 创建的闭包类，并且该 lambda 是使用尾置返回类型定义的，则此函数返回 true，否则返回 false。如果尾置返回类型定义引用了任何 lambda 参数名，则返回类型不被保留。

bool __nv_is_extended_host_device_lambda_closure_type(type);

如果 type 是为扩展 __host__ __device__ lambda 创建的闭包类，则此函数返回 true，否则返回 false。¶

无论是否启用了 lambda 或扩展 lambda，lambda 类型特征都可在所有编译模式下使用。如果扩展 lambda 模式未激活，这些特征将始终返回 false。

示例：

auto lambda0 = [] __host__ __device__ { };

void host_function() {
    auto lambda1 = [] { };
    auto lambda2 = [] __device__ { };
    auto lambda3 = [] __host__ __device__ { };
    auto lambda4 = [] __device__ () -> double { return 3.14; }
    auto lambda5 = [] __device__ (int x) -> decltype(&x) { return 0; }

    using lambda0_t = decltype(lambda0);
    using lambda1_t = decltype(lambda1);
    using lambda2_t = decltype(lambda2);
    using lambda3_t = decltype(lambda3);
    using lambda4_t = decltype(lambda4);
    using lambda5_t = decltype(lambda5);

    // 'lambda0' 不是扩展 lambda，因为它定义在函数作用域之外
    static_assert(!__nv_is_extended_device_lambda_closure_type(lambda0_t));
    static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda0_t));
    static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda0_t));

    // 'lambda1' 不是扩展 lambda，因为它没有执行空间注解
    static_assert(!__nv_is_extended_device_lambda_closure_type(lambda1_t));
    static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda1_t));
    static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda1_t));

    // 'lambda2' 是一个扩展的仅设备 lambda
    static_assert(__nv_is_extended_device_lambda_closure_type(lambda2_t));
    static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda2_t));
    static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda2_t));

    // 'lambda3' 是一个扩展的主机-设备 lambda
    static_assert(!__nv_is_extended_device_lambda_closure_type(lambda3_t));
    static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda3_t));
    static_assert(__nv_is_extended_host_device_lambda_closure_type(lambda3_t));

    // 'lambda4' 是一个具有保留返回类型的扩展仅设备 lambda
    static_assert(__nv_is_extended_device_lambda_closure_type(lambda4_t));
    static_assert(__nv_is_extended_device_lambda_with_preserved_return_type(lambda4_t));
    static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda4_t));

    // 'lambda5' 不是一个具有保留返回类型的扩展仅设备 lambda，
    // 因为它在尾随返回类型中引用了 operator() 的参数类型。
    static_assert(__nv_is_extended_device_lambda_closure_type(lambda5_t));
    static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(lambda5_t));
    static_assert(!__nv_is_extended_host_device_lambda_closure_type(lambda5_t));
}

5.3.7.4. 扩展 Lambda 限制¶

在调用主机编译器之前，CUDA 编译器会将扩展 lambda 表达式替换为在命名空间作用域中定义的占位符类型的实例。该占位符类型的模板参数需要获取包含原始扩展 lambda 表达式的函数的地址。这对于正确执行任何模板参数涉及扩展 lambda 闭包类型的 __global__ 函数模板是必需的。包含函数的计算方式如下。 根据定义，扩展 lambda 表达式存在于 __host__ 或 __host__ __device__ 函数的直接或嵌套块作用域内。

• 如果该函数不是 lambda 表达式的 operator()，则它被视为该扩展 lambda 的封闭函数。
• 否则，该扩展 lambda 定义在一个或多个封闭 lambda 表达式的 operator() 的直接或嵌套块作用域内。如果最外层的 lambda 表达式定义在函数 F 的直接或嵌套块作用域内，则 F 是计算得出的封闭函数。否则，封闭函数不存在。

示例：

void host_function() {
    auto lambda1 = [] __device__ { }; // lambda1 的封闭函数是 "host_function()"
    auto lambda2 = [] {
        auto lambda3 = [] {
            auto lambda4 = [] __host__ __device__ { }; // lambda4 的封闭函数是 "host_function"
        };
    };
}

auto global_lambda = [] {
    auto lambda5 = [] __host__ __device__ { }; // lambda5 的封闭函数不存在
};

扩展 Lambda 限制

1. 扩展 lambda 不能在另一个扩展 lambda 表达式内部定义。 示例：

   void host_function () {
       auto lambda1 = [] __host__ __device__ {
           // 错误，扩展 lambda 定义在另一个扩展 lambda 内部
           auto lambda2 = [] __host__ __device__ { };
       };
   }
   

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

   void host_function () {
       auto lambda1 = [] ( auto ) {
           // 错误，扩展 lambda 定义在泛型 lambda 内部
           auto lambda2 = [] __host__ __device__ { };
       };
   }
   

3. 如果一个扩展 lambda 定义在一个或多个嵌套 lambda 表达式的直接或嵌套块作用域内，那么最外层的 lambda 表达式必须定义在一个函数的直接或嵌套块作用域内。 示例：

   auto lambda1 = [] {
       // 错误，外层封闭 lambda 没有定义在非 lambda-operator() 函数内部
       auto lambda2 = [] __host__ __device__ { };
   };
   

4. 扩展 lambda 的封闭函数必须具有名称，并且其地址必须可访问。如果封闭函数是类成员，则必须满足以下条件：
   • 所有包含该成员函数的类都必须具有名称。
   • 该成员函数在其父类中不能具有 private 或 protected 访问权限。
   • 所有封闭类在其各自的父类中不能具有 private 或 protected 访问权限。 示例：

     void host_function () {
         auto lambda1 = [] __device__ { return 0 ; }; // 正确
         {
             auto lambda2 = [] __device__ { return 0 ; }; // 正确
             auto lambda3 = [] __device__ __host__ { return 0 ; }; // 正确
         }
     }
     struct MyStruct1 {
         MyStruct1 () {
             auto lambda4 = [] __device__ { return 0 ; }; // 错误，封闭函数的地址不可访问
         }
     };
     class MyStruct2 {
         void foo () {
             auto temp1 = [] __device__ { return 10 ; }; // 错误，封闭函数在其父类中具有 private 访问权限
         }
         struct MyStruct3 {
             void foo () {
                 auto temp1 = [] __device__ { return 10 ; }; // 错误，封闭类 MyStruct3 在其父类中具有 private 访问权限
             }
         };
     };
     

5. 在定义扩展 lambda 的位置，必须能够明确地获取其外围函数的地址。然而，这并非总是可行的，例如，当别名声明遮蔽了同名的模板类型参数时。示例：

   template < typename T > struct A {
     using Bar = void ;
     void test ();
   };
   template <> struct A < void > { };
   template < typename Bar >
   void A < Bar >:: test () {
     // 在发送给主机编译器的代码中，nvcc 将在此处注入一个地址表达式，形式如下：
     //   (void (A< Bar> ::*)(void))(&A::test))
     // 然而，类 typedef 'Bar'（指向 void）遮蔽了模板参数 'Bar'，
     // 导致 A<int>::test 中的地址表达式实际上引用的是：
     //    (void (A< void> ::*)(void))(&A::test))
     // 这未能正确获取外围函数 'A<int>::test' 的地址。
     auto lambda1 = [] __host__ __device__ { return 4 ; };
   }
   int main () {
     A < int > var ;
     var . test ();
   }
   

6. 扩展 lambda 不能在函数内部的局部类中定义。示例：

   void host_function () {
     struct MyStruct {
       void bar () {
         // 错误，bar() 是函数内部局部类的成员
         auto lambda2 = [] __host__ __device__ { return 0 ; };
       }
     };
   }
   

7. 扩展 lambda 的外围函数不能具有推导的返回类型。示例：

   auto host_function () {
     // 错误，host_function() 的返回类型是推导得出的
     auto lambda3 = [] __host__ __device__ { return 0 ; };
   }
   

8. 主机-设备扩展 lambda 不能是泛型 lambda，即不能是具有 auto 参数类型的 lambda。示例：

   void host_function () {
     // 错误，__host__ __device__ 扩展 lambda 不能是泛型 lambda
     auto lambda1 = [] __host__ __device__ ( auto i ) { return i ; };
     // 错误，主机-设备扩展 lambda 不能是泛型 lambda
     auto lambda2 = [] __host__ __device__ ( auto ... i ) { return sizeof ...( i ); };
   }
   

9. 如果外围函数是函数模板或成员模板的实例化，或者该函数是类模板的成员，则模板必须满足以下约束：

10. 模板最多只能有一个可变参数，并且必须将其列在模板参数列表的最后。
11. 模板参数必须具有名称。
12. 模板实例化的参数类型不能涉及函数内部的局部类型（扩展 lambda 的闭包类型除外），也不能是私有或受保护的类成员。

示例 1：

template < template < typename ... > class T , typename ... P1 , typename ... P2 >
void bar1 ( const T < P1 ... > , const T < P2 ... > ) {
  // 错误，外围函数具有多个参数包
  auto lambda = [] __device__ { return 10 ; };
}
template < template < typename ... > class T , typename ... P1 , typename T2 >
void bar2 ( const T < P1 ... > , T2 ) {
  // 错误，对于外围函数，参数包未位于模板参数列表的最后
  auto lambda = [] __device__ { return 10 ; };
}
template < typename T , T >
void bar3 () {
  // 错误，对于外围函数，第二个模板参数未命名
  auto lambda = [] __device__ { return 10 ; };
}

示例 2：

template < typename T >
void bar4 () {
  auto lambda1 = [] __device__ { return 10 ; };
}
class MyStruct {
  struct MyNestedStruct {};
  friend int main ();
};
int main () {
  struct MyLocalStruct {};
  // 错误，bar4() 中设备 lambda 的外围函数使用 main 函数的局部类型进行实例化
  bar4 < MyLocalStruct > ();
  // 错误，bar4 中设备 lambda 的外围函数使用类私有成员类型进行实例化
  bar4 < MyStruct :: MyNestedStruct > ();
}

示例：

void host_function () {
    // 正确：仅设备扩展 lambda 允许初始化捕获
    auto lambda1 = [ x = 1 ] __device__ () { return x ; };
    // 错误：主机-设备扩展 lambda 不允许初始化捕获
    auto lambda2 = [ x = 1 ] __host__ __device__ () { return x ; };
    int a = 1 ;
    // 错误：扩展 __device__ lambda 无法通过引用捕获变量
    auto lambda3 = [ & a ] __device__ () { return a ; };
    // 错误：仅设备扩展 lambda 不允许通过引用捕获
    auto lambda4 = [ & x = a ] __device__ () { return x ; };
    struct MyStruct {};
    MyStruct s1 ;
    // 错误：函数的局部类型不能用于捕获变量的类型
    auto lambda6 = [ s1 ] __device__ () { };
    // 错误：初始化捕获不能是 std::initializer_list 类型
    auto lambda7 = [ x = { 11 }] __device__ () { };
    std :: initializer_list < int > b = { 11 , 22 , 33 };
    // 错误：初始化捕获不能是 std::initializer_list 类型
    auto lambda8 = [ x = b ] __device__ () { };
    int var = 4 ;
    auto lambda9 = [ = ] __device__ {
        int result = 0 ;
        if constexpr ( false ) {
            // 错误：仅设备扩展 lambda 不能在 if-constexpr 上下文中首次捕获 'var'
            result += var ;
        }
        return result ;
    };
    auto lambda10 = [ var ] __device__ {
        int result = 0 ;
        if constexpr ( false ) {
            // 正确：'var' 已列在扩展 lambda 的显式捕获列表中
            result += var ;
        }
        return result ;
    };
    auto lambda11 = [ = ] __device__ {
        int result = var ;
        if constexpr ( false ) {
            // 正确：'var' 已在 'if-constexpr' 块外部隐式捕获
            result += var ;
        }
        return result ;
    };
}

template < typename T >
__global__ void kernel ( T in ) { in (); }
__host__ __device__ void host_device_function () {
    // 错误：扩展 lambda 的数量和相对声明顺序依赖于 __CUDA_ARCH__
#if defined(__CUDA_ARCH__)
    auto lambda1 = [] __device__ { return 0 ; };
    auto lambda2 = [] __host__ __device__ { return 10 ; };
#endif
    auto lambda3 = [] __device__ { return 4 ; };
    kernel <<< 1 , 1 >>> ( lambda3 );
}

1. 如上所述，CUDA 编译器将主机函数中定义的设备扩展 lambda 替换为命名空间作用域中定义的占位符类型。除非特征 __nv_is_extended_device_lambda_with_preserved_return_type() 对该扩展 lambda 的闭包类型返回 true，否则该占位符类型不会定义与原始 lambda 声明等效的 operator() 函数。因此，尝试确定此类 lambda 的 operator() 函数的返回类型或参数类型在主机代码中可能无法正常工作，因为主机编译器处理的代码在语义上与 CUDA 编译器处理的输入代码不同。然而，在设备代码内自省 operator() 函数的返回类型或参数类型是可以接受的。请注意，此限制不适用于特征 __nv_is_extended_device_lambda_with_preserved_return_type() 返回 true 的主机或设备扩展 lambda。示例：

   #include <cuda/std/type_traits>
   const char & getRef ( const char * p ) { return * p ; }
   void foo () {
       auto lambda1 = [] __device__ { return "10" ; };
       // 错误：尝试在主机代码中提取设备 lambda 的返回类型
       cuda :: std :: result_of < decltype ( lambda1 )() >:: type xx1 = "abc" ;
       auto lambda2 = [] __host__ __device__ { return "10" ; };
       // 正确：lambda2 表示一个主机-设备扩展 lambda
       cuda :: std :: result_of < decltype ( lambda2 )() >:: type xx2 = "abc" ;
       auto lambda3 = [] __device__ () -> const char * { return "10" ; };
       // 正确：lambda3 表示一个保留了返回类型的设备扩展 lambda
       cuda :: std :: result_of < decltype ( lambda3 )() >:: type xx2 = "abc" ;
       static_assert ( cuda :: std :: is_same_v < cuda :: std :: result_of < decltype ( lambda3 )() >:: type , const char *> );
       auto lambda4 = [] __device__ ( char x ) -> decltype ( getRef ( & x )) { return 0 ; };
       // lambda4 的返回类型未被保留，因为它在尾部返回类型中引用了 operator() 的参数类型。
       static_assert ( ! __nv_is_extended_device_lambda_with_preserved_return_type ( decltype ( lambda4 )));
   }
   

2. 对于仅限设备的扩展 lambda：

   • operator() 参数类型的自省仅在设备代码中受支持。
   • operator() 返回类型的自省仅在设备代码中受支持，除非特征函数 __nv_is_extended_device_lambda_with_preserved_return_type() 返回 true。

3. 如果扩展 lambda 从主机代码传递到设备代码作为 __global__ 函数的参数，那么 lambda 主体中捕获变量的任何表达式都必须保持不变，无论 __CUDA_ARCH__ 宏是否定义及其值如何。此限制源于 lambda 的闭包类布局取决于编译器在处理 lambda 表达式时遇到捕获变量的顺序。如果设备和主机编译的闭包类布局不同，程序可能执行不正确。示例：

   __device__ int result ;
   template < typename T >
   __global__ void kernel ( T in ) {
     result = in ();
   }
   void foo ( void ) {
     int x1 = 1 ;
     // 错误，"x1" 仅在 __CUDA_ARCH__ 定义时被捕获。
     auto lambda1 = [ = ] __host__ __device__ {
   #ifdef __CUDA_ARCH__
       return x1 + 1 ;
   #else
       return 10 ;
   #endif
     };
     kernel <<< 1 , 1 >>> ( lambda1 );
   }
   

4. 如前所述，CUDA 编译器将仅设备的扩展 lambda 表达式替换为发送到主机编译器代码中的占位符类型实例。该占位符类型在主机代码中未定义指向函数的转换运算符；但是，在设备代码中提供了转换运算符。请注意，此限制不适用于主机-设备扩展 lambda。示例：

   template < typename T >
   __global__ void kernel ( T in ) {
     int ( * fp )( double ) = in ;
     fp ( 0 ); // 正确，设备代码中支持转换
     auto lambda1 = []( double ) { return 1 ; };
   }
   void foo () {
     auto lambda_device = [] __device__ ( double ) { return 1 ; };
     auto lambda_host_device = [] __host__ __device__ ( double ) { return 1 ; };
     kernel <<< 1 , 1 >>> ( lambda_device );
     kernel <<< 1 , 1 >>> ( lambda_host_device );
     // 正确，主机代码中支持 __host__ __device__ lambda 的转换
     int ( * fp1 )( double ) = lambda_host_device ;
     // 错误，主机代码中不支持设备 lambda 的转换
     int ( * fp2 )( double ) = lambda_device ;
   }
   

5. 如前所述，CUDA 编译器将仅设备或主机-设备的扩展 lambda 表达式替换为发送到主机编译器代码中的占位符类型实例。此占位符类型可能定义 C++ 特殊成员函数，例如构造函数和析构函数。因此，某些标准 C++ 类型特征在 CUDA 前端编译器中对扩展 lambda 闭包类型的结果可能与主机编译器不同。受影响的类型特征包括：std::is_trivially_copyable、std::is_trivially_constructible、std::is_trivially_copy_constructible、std::is_trivially_move_constructible、std::is_trivially_destructible。必须注意确保这些特征的结果不用于 __global__、__device__、__constant__ 或 __managed__ 函数或变量模板的实例化中。示例：

   #include <cstdio>
   #include <type_traits>
   template < bool b >
   void __global__ kernel () {
     printf ( "hi" );
   }
   template < typename T >
   void kernel_launch () {
     // 错误，此内核启动可能失败，因为 CUDA 前端编译器和主机编译器
     //       可能对扩展 lambda 闭包类型的 std::is_trivially_copyable_v 特征结果不一致
     kernel < std :: is_trivially_copyable_v < T >><<< 1 , 1 >>> ();
     cudaDeviceSynchronize ();
   }
   int main () {
     int x = 0 ;
     auto lambda1 = [ = ] __host__ __device__ () { return x ; };
     kernel_launch < decltype ( lambda1 ) > ();
   }
   

   CUDA 编译器将为 1-12 节中描述的部分情况生成编译器诊断信息；对于 13-17 节的情况，则不会生成诊断信息，但主机编译器可能无法编译生成的代码。

5.3.7.5. 主机-设备 Lambda 优化说明¶

与仅设备 lambda 不同，主机-设备 lambda 可以从主机代码中调用。如前所述，CUDA 编译器将主机代码中定义的扩展 lambda 表达式替换为一个具名占位符类型的实例。对于扩展主机-设备 lambda，其占位符类型通过间接函数调用来调用原始 lambda 的 operator()。如果扩展 lambda 模式未激活，其特性将始终返回 false。

间接函数调用的存在可能导致主机编译器对扩展主机-设备 lambda 的优化程度低于隐式或显式仅为 __host__ 的 lambda。在后一种情况下，主机编译器可以轻松地将 lambda 主体内联到调用上下文中。然而，当遇到扩展主机-设备 lambda 时，主机编译器可能无法轻松地将原始 lambda 主体内联。

5.3.7.6. 按值捕获 *this¶

根据 C++11/C++14 规则，当 lambda 在非 static 类成员函数中定义，且 lambda 主体引用了类成员变量时，必须按值捕获类的 this 指针，而不是引用的成员变量。如果 lambda 是在主机函数中定义并在 GPU 上执行的扩展仅设备或主机-设备 lambda，并且 this 指针指向主机内存，则在 GPU 上访问引用的成员变量将导致运行时错误。

示例：

#include <cstdio>

template <typename T>
__global__ void foo(T in) { printf("value = %d\n", in()); }

struct MyStruct {
    int var;

    __host__ __device__ MyStruct() : var(10) {};

    void run() {
        auto lambda1 = [=] __device__ {
            // 对 "var" 的引用导致按值捕获 'this' 指针 (MyStruct*)
            return var + 1;
        };
        // 内核启动在运行时失败，因为 GPU 无法访问 'this->var'
        foo<<<1, 1>>>(lambda1);
        cudaDeviceSynchronize();
    }
};

int main() {
    MyStruct s1;
    s1.run();
}

C++17 通过引入新的 *this 捕获模式解决了这个问题。在此模式下，编译器复制由 *this 表示的对象，而不是按值捕获 this 指针。*this 捕获模式在 P0018R3 中有更详细的描述。

当使用 --extended-lambda 标志时，CUDA 编译器支持在 __device__ 和 __global__ 函数内定义的 lambda，以及在主机代码中定义的扩展仅设备 lambda 使用 *this 捕获模式。

以下是修改为使用 *this 捕获模式的上述示例：

#include <cstdio>

template <typename T>
__global__ void foo(T in) { printf("\n value = %d", in()); }

struct MyStruct {
    int var;
    __host__ __device__ MyStruct() : var(10) { };

    void run() {
        // 注意 "*this" 捕获说明符
        auto lambda1 = [=, *this] __device__ {
            // 对 "var" 的引用导致按值捕获由 '*this' 表示的对象，
            // GPU 代码将访问 'copy_of_star_this->var'
            return var + 1;
        };
        // 内核启动成功
        foo<<<1, 1>>>(lambda1);
        cudaDeviceSynchronize();
    }
};

int main() {
    MyStruct s1;
    s1.run();
}

struct MyStruct {
    int var;
    __host__ __device__ MyStruct() : var(10) { };

    void host_function() {
        // 正确，在扩展的仅设备 lambda 中使用
        auto lambda1 = [=, *this] __device__ { return var; };

        // 在扩展的主机-设备 lambda 中使用
        // 如果语言方言未启用 *this 捕获，则错误
        auto lambda2 = [=, *this] __host__ __device__ { return var; };

        // 在主机函数的非注解 lambda 中使用
        // 如果语言方言未启用 *this 捕获，则错误
        auto lambda3 = [=, *this]  { return var; };
    }

    __device__ void device_function() {
        // 正确，在仅设备函数中定义的 lambda 中使用
        auto lambda1 = [=, *this] __device__ { return var; };

        // 正确，在仅设备函数中定义的 lambda 中使用
        auto lambda2 = [=, *this] __host__ __device__ { return var; };

        // 正确，在仅设备函数中定义的 lambda 中使用
        auto lambda3 = [=, *this]  { return var; };
    }

    __host__ __device__ void host_device_function() {
        // 正确，在扩展的仅设备 lambda 中使用
        auto lambda1 = [=, *this] __device__ { return var; };

        // 在扩展的主机-设备 lambda 中使用
        // 如果语言方言未启用 *this 捕获，则错误
        auto lambda2 = [=, *this] __host__ __device__ { return var; };

        // 在主机-设备函数的未注解 lambda 中使用
        // 如果语言方言未启用 *this 捕获，则错误
        auto lambda3 = [=, *this]  { return var; };
    }
};

5.3.7.7. 参数依赖查找 (ADL)¶

如前所述，CUDA 编译器在调用主机编译器之前，会用占位符类型替换扩展的 lambda 表达式。该占位符类型的一个模板参数使用了包含原始 lambda 表达式的函数的地址。这可能导致额外的命名空间参与到任何参数类型涉及扩展 lambda 表达式闭包类型的主机函数调用的参数依赖查找 (ADL) 中。因此，主机编译器可能会选择错误的函数。

示例：

namespace N1 {

struct MyStruct {};

template <typename T>
void my_function(T);

}; // namespace N1

namespace N2 {

template <typename T>
int my_function(T);

template <typename T>
void run(T in) { my_function(in); }

} // namespace N2

void bar(N1::MyStruct in) {
    // 对于扩展的仅设备 lambda，发送给主机编译器的代码被替换为
    // 占位符类型实例化表达式
    //    ' __nv_dl_wrapper_t< __nv_dl_tag<void (*)(N1::MyStruct in),(&bar),1> > { }'
    //
    // 因此，命名空间 'N1' 会参与到 N2::run 函数体内对 "my_function()" 调用的 ADL 查找中，导致歧义。
    auto lambda1 = [=] __device__ { };
    N2::run(lambda1);
}

5.3.8. 多态函数包装器¶

nvfunctional 头文件提供了一个多态函数包装器类模板 nvstd::function。该模板的实例可以存储、复制和调用任何可调用目标，例如 lambda 表达式。nvstd::function 可在主机和设备代码中使用。

示例：

#include <nvfunctional>

__host__            int host_function()        { return 1; }
__device__          int device_function()      { return 2; }
__host__ __device__ int host_device_function() { return 3; }

__global__ void kernel(int* result) {
    nvstd::function<int()> fn1 = device_function;
    nvstd::function<int()> fn2 = host_device_function;
    nvstd::function<int()> fn3 = [](){ return 10; };
    *result                    = fn1() + fn2() + fn3();
}

__host__ __device__ void host_device_test(int* result) {
    nvstd::function<int()> fn1 = host_device_function;
    nvstd::function<int()> fn2 = [](){ return 10; };
    *result                    = fn1() + fn2();
}

__host__ void host_test(int* result) {
    nvstd::function<int()> fn1 = host_function;
    nvstd::function<int()> fn2 = host_device_function;
    nvstd::function<int()> fn3 = [](){ return 10; };
    *result                    = fn1() + fn2() + fn3();
}