5.3. C++ 语言支持

5.3.7.1. clock()和clock64()

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

在设备代码中执行时，它返回每个多处理器各自拥有的计数器的值，该计数器在每个时钟周期递增。在核函数开始和结束时对这个计数器采样，将两个值相减，并为每个线程记录结果，即可估算设备执行该线程所花费的时钟周期数。但是，这个值并不代表设备执行该线程指令所花费的实际时钟周期数。由于线程以时间片方式运行，前者大于后者。

5.3.7.2. printf()

__host__ __device__ __tile__ 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()函数在分块代码中的行为与在 SIMT 设备代码中的行为不同。在分块代码中，

• 参数除了可以是标量，还可以是分块。传入分块参数时，分块中的每个元素都按相应的格式说明符打印。

• 返回值始终是所提供的参数个数，即使格式字符串为NULL或发生内部错误也是如此。

• 格式字符串必须是字面量。

• 如果参数个数与格式说明符的个数不匹配，则会报告错误。在 SIMT 设备代码中，这种情况会产生警告。

格式说明符

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

支持以下字段。有关所有行为的完整说明，请参阅广泛可获取的文档。

• 标志：#、' '、0、+、-

• 宽度：*、0-9

• 精度：0-9

• 大小：h、l、ll

• 类型：%cdiouxXpeEfgGaAs

限制

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

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

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

由于long类型在 Windows 平台（32 位）和 Linux 平台（64 位）上的大小不同，在 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()实参和内部元数据传输到主机。默认大小为一兆字节。

• 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

请注意，每个线程都会遇到printf()命令。因此，输出的行数与网格中的线程数相同。

参见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上的示例。

以下代码示例：

#include "cuda_tile.h"
#include <cstdio>

namespace ct = cuda::tiles;

__tile_global__ void kernel() {
  auto ints = ct::iota<ct::tile<int, ct::shape<4, 4>>>();
  printf("%i\n", ints);
}

int main() {
  kernel<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}

将输出：

[[0, 1, 2, 3], [4, 5, 6, 7], [8, 9, 10, 11], [12, 13, 14, 15]]

5.3.7.3. memcpy()和memset()

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

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

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

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

5.3.7.4. malloc()和free()

__host__ __device__ void* malloc(size_t size);
// or cuda::std::malloc(), cuda::std::calloc() in the <cuda/std/cstdlib> header

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

__device__ void* __nv_aligned_device_malloc(size_t size, size_t align);
// or cuda::std::aligned_alloc() in the <cuda/std/cstdlib> header

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

__host__ __device__ void free(void* ptr);
// or cuda::std::free() in the <cuda/std/cstdlib> header

设备端函数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关键字的情况。如果任何程序在未显式指定堆大小的情况下使用设备内存堆，则会分配一个八兆字节的默认堆。

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

• cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)

• cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

实际提供的堆大小将至少为size个字节。cuCtxGetLimit()和cudaDeviceGetLimit()返回当前请求的堆大小。

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

模块加载后，堆大小便无法更改，而且它不会根据需要动态调整大小。

为设备堆保留的内存，是在通过cudaMalloc()等主机端 CUDA API 调用分配的内存之外额外保留的。

与主机内存 API 的互操作性

通过设备端函数malloc()、cuda::std::malloc()、cuda::std::calloc()、__nv_aligned_device_malloc()、cuda::std::aligned_alloc()或new关键字分配的内存，不能通过cudaMalloc、cudaMemcpy或cudaMemset等运行时或驱动程序 API 调用来使用或释放。同样，通过主机运行时 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() {
    // Set a heap size of 128 megabytes.
    // Note that this must be done before any kernel is launched.
    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;
    // The first thread in the block performs the allocation and shares the pointer
    // with all other threads through shared memory, so that access can be coalesced.
    if (threadIdx.x == 0) {
        size_t size = blockDim.x * 64; // 64 bytes per thread are allocated.
        data = (int*) malloc(size);
    }
    __syncthreads();
    // Check for failure
    if (data == nullptr)
        return;

    // Threads index into the memory, ensuring coalescence
    for (int i = 0; i < 64; ++i)
        data[i * blockDim.x + threadIdx.x] = threadIdx.x;
    // Ensure all threads complete before freeing
    __syncthreads();

    // Only one thread may free the memory!
    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]; // Per-block pointer

__global__ void allocate_memory_kernel() {
    // Only the first thread in the block performs the allocation
    // since we need only one allocation per block.
    if (threadIdx.x == 0)
        data_ptrs[blockIdx.x] = (int*) malloc(blockDim.x * 4);
    __syncthreads();
    // Check for failure
    if (data_ptrs[blockIdx.x] == nullptr)
        return;
    // Zero the data with all threads in parallel
    data_ptrs[blockIdx.x][threadIdx.x] = 0;
}

// Simple example: store the thread ID into each element
__global__ void use_memory_kernel() {
    int* ptr = data_ptrs[blockIdx.x];
    if (ptr != nullptr)
        ptr[threadIdx.x] += threadIdx.x;
}

// Print the content of the buffer before freeing it
__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]);
    // Only free from one thread!
    if (threadIdx.x == 0)
        free(ptr);
}

int main() {
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, 128*1024*1024);
    // Allocate memory
    allocate_memory_kernel<<<NUM_BLOCKS, 10>>>();

    // Use memory
    use_memory_kernel<<<NUM_BLOCKS, 10>>>();
    use_memory_kernel<<<NUM_BLOCKS, 10>>>();
    use_memory_kernel<<<NUM_BLOCKS, 10>>>();

    // Free memory
    free_memory_kernel<<<NUM_BLOCKS, 10>>>();
    cudaDeviceSynchronize();
    return 0;
}

参见Compiler Explorer上的示例。

5.3.7.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));
    // use of ptr
    ...
}

5.3.8.1. Lambda 表达式与__global__函数参数

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

示例：

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

 __device__ void device_function() {
     // device kernel call requires separate compilation (-rdc=true flag)
     kernel<<<1, 1>>>([](){});
     kernel<<<1, 1>>>([] __device__() {});          // extended lambda
     kernel<<<1, 1>>>([] __host__ __device__() {}); // extended lambda
 }

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

 void host_function() {
     kernel<<<1, 1>>>([] __device__() {});          // CORRECT, extended lambda
     kernel<<<1, 1>>>([] __host__ __device__() {}); // CORRECT, extended lambda
 //  kernel<<<1, 1>>>([](){});                      // ERROR, closure type with host execution space
 //  kernel<<<1, 1>>>(global_lambda);               // ERROR, extended lambda, but at global scope
 }

参见Compiler Explorer上的示例。

5.3.8.2. 扩展 lambda

• 扩展 lambda定义在__host__或__host__ __device__函数的直接块或嵌套块的作用域内。

• 扩展设备 lambda是使用__device__关键字标注的 lambda 表达式。

• 扩展主机-设备 lambda是使用__host__ __device__关键字标注的 lambda 表达式。

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

示例：

void host_function() {
    auto lambda1 = [] {};                      // NOT an extended lambda: no explicit execution space annotations
    auto lambda2 = [] __device__ {};           // extended lambda
    auto lambda3 = [] __host__ __device__ {};  // extended lambda
    auto lambda4 = [] __host__ {};             // NOT an extended lambda
}

__host__ __device__ void host_device_function() {
    auto lambda1 = [] {};                      // NOT an extended lambda: no explicit execution space annotations
    auto lambda2 = [] __device__ {};           // extended lambda
    auto lambda3 = [] __host__ __device__ {};  // extended lambda
    auto lambda4 = [] __host__ {};             // NOT an extended lambda
}

__device__ void device_function() {
    // none of the lambdas within this function are extended lambdas,
    // because the enclosing function is not a __host__ or __host__ __device__  function.
    auto lambda1 = [] {};
    auto lambda2 = [] __device__ {};
    auto lambda3 = [] __host__ __device__ {};
    auto lambda4 = [] __host__ {};
}

auto global_lambda = [] __host__ __device__ { }; // NOT an extended lambda because it is not defined
                                                 // within a __host__ or __host__ __device__ function

5.3.8.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' is not an extended lambda because it is defined outside function scope
    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' is not an extended lambda because it has no execution space annotations
    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' is an extended device-only 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' is an extended host-device 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' is an extended device-only lambda with preserved return type
    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' is not an extended device-only lambda with preserved return type
    // because it references the operator()'s parameter types in the trailing return type.
    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.8.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__ { }; // enclosing function for lambda1 is "host_function()"
    auto lambda2 = [] {
        auto lambda3 = [] {
            auto lambda4 = [] __host__ __device__ { }; // enclosing function for lambda4 is "host_function"
        };
    };
}

auto global_lambda = [] {
    auto lambda5 = [] __host__ __device__ { }; // enclosing function for lambda5 does not exist
};

扩展 lambda 的限制

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

   void host_function() {
       auto lambda1 = [] __host__ __device__  {
            // ERROR, extended lambda defined within another extended lambda
           auto lambda2 = [] __host__ __device__ { };
       };
   }
   

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

   void host_function() {
       auto lambda1 = [] (auto) {
            // ERROR, extended lambda defined within a generic lambda
           auto lambda2 = [] __host__ __device__ { };
       };
   }
   

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

   auto lambda1 = []  {
       // ERROR, outer enclosing lambda is not defined within a non-lambda-operator() function
       auto lambda2 = [] __host__ __device__ { };
   };
   

4. 扩展 lambda 的外围函数必须具有名称，并且其地址必须可访问。如果外围函数是类成员，则必须满足以下条件：

   • 包含该成员函数的所有外围类都必须具有名称。

   • 该成员函数在其所属类中的访问权限不得为私有或受保护。

   • 所有外围类在各自的外层类中的访问权限都不得为私有或受保护。

   示例：

   void host_function() {
       auto lambda1 = [] __device__ { return 0; }; // OK
       {
           auto lambda2 = [] __device__          { return 0; }; // OK
           auto lambda3 = [] __device__ __host__ { return 0; }; // OK
       }
   }
   
   struct MyStruct1 {
       MyStruct1() {
           auto lambda4 = [] __device__ { return 0; }; // ERROR, address of the enclosing function is not accessible
       }
   };
   
   class MyStruct2 {
       void foo() {
           auto temp1 = [] __device__ { return 10; }; // ERROR, enclosing function has private access in parent class
       }
   
       struct MyStruct3 {
           void foo() {
               auto temp1 = [] __device__ { return 10; };  // ERROR, enclosing class MyStruct3 has private access in its parent class
           }
       };
   };
   

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

   template <typename T>
   struct A {
       using Bar = void;
       void test();
   };
   
   template<>
   struct A<void> { };
   
   template <typename Bar>
   void A<Bar>::test() {
       // In code sent to host compiler, nvcc will inject an address expression here, of the form:
       //   (void (A< Bar> ::*)(void))(&A::test))
       //  However, the class typedef 'Bar' (to void) shadows the template argument 'Bar',
       //  causing the address expression in A<int>::test to actually refer to:
       //    (void (A< void> ::*)(void))(&A::test))
       //  which doesn't take the address of the enclosing routine 'A<int>::test' correctly.
       auto lambda1 = [] __host__ __device__ { return 4; };
   }
   
   int main() {
       A<int> var;
       var.test();
   }
   

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

   void host_function() {
       struct MyStruct {
           void bar() {
               // ERROR, bar() is member of a class that is local to a function
               auto lambda2 = [] __host__ __device__ { return 0; };
           }
       };
   }
   

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

   auto host_function() {
       // ERROR, the return type of host_function() is deduced
       auto lambda3 = [] __host__ __device__ { return 0; };
   }
   

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

   void host_function() {
       // ERROR, __host__ __device__ extended lambdas cannot be a generic lambda
       auto lambda1 = [] __host__ __device__ (auto i) { return i; };
   
       // ERROR, a host-device extended lambda cannot be a generic lambda
       auto lambda2 = [] __host__ __device__ (auto... i) {
           return sizeof...(i);
       };
   }
   

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

   • 模板最多只能具有一个可变参数，并且该参数必须列在模板参数列表的最后。

   • 模板参数必须具有名称。

   • 模板实例化的实参类型不能涉及函数局部的类型（扩展 lambda 的闭包类型除外），也不能涉及作为private或protected类成员的类型。

   示例 1：

   template <template <typename...> class T,
             typename... P1,
             typename... P2>
   void bar1(const T<P1...>, const T<P2...>) {
       // ERROR, enclosing function has multiple parameter packs
       auto lambda = [] __device__ { return 10; };
   }
   
   template <template <typename...> class T,
             typename... P1,
             typename    T2>
   void bar2(const T<P1...>, T2) {
       // ERROR, for enclosing function, the parameter pack is not last in the template parameter list
       auto lambda = [] __device__ { return 10; };
   }
   
   template <typename T, T>
   void bar3() {
       // ERROR, for enclosing function, the second template parameter is not named
       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 {};
       // ERROR, enclosing function for device lambda in bar4() is instantiated with a type local to main
       bar4<MyLocalStruct>();
   
       // ERROR, enclosing function for device lambda in bar4 is instantiated with a type
       //        that is a private member of a class
       bar4<MyStruct::MyNestedStruct>();
   }
   

10. 使用 Microsoft Visual Studio 主机编译器时，外围函数必须具有外部链接。存在此限制是因为主机编译器不支持将非外部链接函数的地址用作模板实参。CUDA 编译器的转换需要这些地址才能支持扩展 lambda。

11. 使用 Microsoft Visual Studio 主机编译器时，不得在if constexpr块的主体内定义扩展 lambda。

12. 扩展 lambda 对捕获的变量有以下限制：

    • 在发送给主机编译器的代码中，变量可能会先按值传递给一系列辅助函数，然后才用于直接初始化表示该扩展 lambda 闭包类型的类类型的字段。然而，C++ 标准规定，应使用被捕获的变量直接初始化闭包类型的字段。

    • 变量只能按值捕获。

    • 如果数组的维数大于 7，则无法捕获该数组类型的变量。

    • 对于数组类型的变量，在发送给主机编译器的代码中，首先对闭包类型的数组字段进行默认初始化，然后从被捕获数组变量的相应元素向每个数组元素进行复制赋值。因此，在主机代码中，数组元素类型必须既可默认构造，又可复制赋值。

    • 不能捕获作为可变实参包元素的函数参数。

    • 被捕获变量的类型不能是函数局部的类型（扩展 lambda 闭包类型除外），也不能是private或protected类成员。

    • 主机-设备扩展 lambda 不支持初始化捕获。不过，设备扩展 lambda 支持初始化捕获，但初始化器为数组或类型为std::initializer_list的情况除外。

    • 扩展 lambda 的函数调用运算符不是constexpr。扩展 lambda 的闭包类型不是字面类型。声明扩展 lambda 时不能使用constexpr和consteval说明符。

    • 在词法上嵌套于扩展 lambda 内部的if-constexpr块内，不能隐式捕获变量，除非该变量已在if-constexpr块之外被隐式捕获，或出现在该扩展 lambda 的显式捕获列表中。

    示例：

    void host_function() {
        // CORRECT, an init-capture is allowed for an extended device-only lambda
        auto lambda1 = [x = 1] __device__ () { return x; };
    
        // ERROR, an init-capture is not allowed for an extended host-device lambda
        auto lambda2 = [x = 1] __host__ __device__ () { return x; };
    
        int a = 1;
        // ERROR, an extended __device__ lambda cannot capture variables by reference
        auto lambda3 = [&a] __device__ () { return a; };
    
        // ERROR, by-reference capture is not allowed for an extended device-only lambda
        auto lambda4 = [&x = a] __device__ () { return x; };
    
        struct MyStruct {};
        MyStruct s1;
        // ERROR, a type local to a function cannot be used in the type of a captured variable
        auto lambda6 = [s1] __device__ () { };
    
        // ERROR, an init-capture cannot be of type std::initializer_list
        auto lambda7 = [x = {11}] __device__ () { };
    
        std::initializer_list<int> b = {11,22,33};
        // ERROR, an init-capture cannot be of type std::initializer_list
        auto lambda8 = [x = b] __device__ () { };
    
        int  var     = 4;
        auto lambda9 = [=] __device__ {
            int result = 0;
            if constexpr(false) {
                //ERROR, An extended device-only lambda cannot first-capture 'var' in if-constexpr context
                result += var;
            }
            return result;
        };
    
        auto lambda10 = [var] __device__ {
            int result = 0;
            if constexpr(false) {
                // CORRECT, 'var' already listed in explicit capture list for the extended lambda
                result += var;
            }
            return result;
        };
    
        auto lambda11 = [=] __device__ {
            int result = var;
            if constexpr(false) {
                // CORRECT, 'var' already implicit captured outside the 'if-constexpr' block
                result += var;
            }
            return result;
        };
    }
    

13. 解析函数时，CUDA 编译器会为函数中的每个扩展 lambda 分配一个计数器值。该计数器值用于传递给主机编译器的替换后的具名类型中。因此，函数内是否存在扩展 lambda 不应依赖于__CUDA_ARCH__的特定值，也不应依赖于__CUDA_ARCH__未被定义。示例：

    template <typename T>
    __global__ void kernel(T in) { in(); }
    
    __host__ __device__ void host_device_function() {
        // ERROR, the number and relative declaration order of
        //        extended lambdas depend on __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);
    }
    

14. 如上所述，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"; };
    
        // ERROR, attempt to extract the return type of a device lambda in host code
        cuda::std::result_of<decltype(lambda1)()>::type xx1 = "abc";
    
        auto lambda2 = [] __host__ __device__ { return "10"; };
    
        // CORRECT, lambda2 represents a host-device extended lambda
        cuda::std::result_of<decltype(lambda2)()>::type xx2 = "abc";
    
        auto lambda3 = [] __device__ () -> const char* { return "10"; };
    
        // CORRECT, lambda3 represents a device extended lambda with preserved return type
        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's return type is not preserved because it references the operator()'s
        // parameter types in the trailing return type.
        static_assert(!__nv_is_extended_device_lambda_with_preserved_return_type(decltype(lambda4)));
    }
    

15. 对于仅设备扩展 lambda：

    • 仅在设备代码中支持对operator()的参数类型进行自省。

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

16. 如果扩展 lambda 从主机代码传递到设备代码，例如作为__global__函数的实参传递，那么无论是否定义了__CUDA_ARCH__宏，也无论其值为何，lambda 体中捕获变量的任何表达式都必须保持不变。之所以有此限制，是因为 lambda 的闭包类布局取决于编译器处理 lambda 表达式时遇到被捕获变量的顺序。如果设备编译与主机编译中的闭包类布局不同，程序可能会执行错误。示例：

    __device__ int result;
    
    template <typename T>
    __global__ void kernel(T in) { result = in(); }
    
    void foo(void) {
        int x1 = 1;
        // ERROR, "x1" is only captured when __CUDA_ARCH__ is defined.
        auto lambda1 = [=] __host__ __device__ {
    #ifdef __CUDA_ARCH__
            return x1 + 1;
    #else
            return 10;
    #endif
        };
        kernel<<<1, 1>>>(lambda1);
    }
    

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

    template <typename T>
    __global__ void kernel(T in) {
        int (*fp)(double) = in;
        fp(0); // CORRECT, conversion in device code is supported
        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);
    
        // CORRECT, conversion for a __host__ __device__ lambda is supported in host code
        int (*fp1)(double) = lambda_host_device;
    
        // ERROR, conversion for a device lambda is not supported in host code
        int (*fp2)(double) = lambda_device;
    }
    

18. 如前所述，在发送给主机编译器的代码中，CUDA 编译器会将仅设备扩展 lambda 表达式或主机-设备扩展 lambda 表达式替换为占位类型实例。此占位类型可能定义 C++ 特殊成员函数，例如构造函数和析构函数。因此，对于扩展 lambda 的闭包类型，某些标准 C++ 类型特征在 CUDA 前端编译器中可能产生与主机编译器中不同的结果。受影响的类型特征如下：：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() {
        // ERROR, this kernel launch may fail, because CUDA frontend compiler and host compiler
        //        may disagree on the result of std::is_trivially_copyable_v trait on the
        //        closure type of the extended lambda
        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.8.5. 主机-设备 lambda 优化注意事项

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

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

5.3.8.6. *this按值捕获

根据 C++11/C++14 规则，当在非static类成员函数中定义 lambda，且 lambda 体引用了类成员变量时，必须按值捕获该类的this指针，而不是所引用的成员变量。如果该 lambda 是在主机函数中定义并在 GPU 上执行的仅设备扩展 lambda 或主机-设备扩展 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__ {
            // reference to "var" causes the 'this' pointer (MyStruct*) to be captured by value
            return var + 1;
        };
        // Kernel launch fails at run time because 'this->var' is not accessible from the GPU
        foo<<<1, 1>>>(lambda1);
        cudaDeviceSynchronize();
    }
};

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

C++17 通过引入新的*this捕获模式解决了此问题。在此模式下，编译器会复制*this所指代的对象，而不是按值捕获this指针。P0018R3中更详细地介绍了*this捕获模式。

使用--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() {
        // note the "*this" capture specification
        auto lambda1 = [=, *this] __device__ {
            // reference to "var" causes the object denoted by '*this' to be captured by
            // value, and the GPU code will access 'copy_of_star_this->var'
            return var + 1;
        };
        // Kernel launch succeeds
        foo<<<1, 1>>>(lambda1);
        cudaDeviceSynchronize();
    }
};

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

对于主机代码中定义的未加注解的 lambda 或主机-设备扩展 lambda，不允许使用*this捕获模式，除非所选语言方言启用了*this捕获。下面是受支持和不受支持的用法示例：

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

    void host_function() {
        // CORRECT, use in an extended device-only lambda
        auto lambda1 = [=, *this] __device__ { return var; };

        // Use in an extended host-device lambda
        // Error if *this capture not enabled by language dialect
        auto lambda2 = [=, *this] __host__ __device__ { return var; };

        // Use in an non-annotated lambda in host function
        // Error if *this capture not enabled by language dialect
        auto lambda3 = [=, *this]  { return var; };
    }

    __device__ void device_function() {
        // CORRECT, use in a lambda defined in a device-only function
        auto lambda1 = [=, *this] __device__ { return var; };

        // CORRECT, use in a lambda defined in a device-only function
        auto lambda2 = [=, *this] __host__ __device__ { return var; };

        // CORRECT, use in a lambda defined in a device-only function
        auto lambda3 = [=, *this]  { return var; };
    }

    __host__ __device__ void host_device_function() {
        // CORRECT, use in an extended device-only lambda
        auto lambda1 = [=, *this] __device__ { return var; };

        // Use in an extended host-device lambda
        // Error if *this capture not enabled by language dialect
        auto lambda2 = [=, *this] __host__ __device__ { return var; };

        // Use in an unannotated lambda in a host-device function
        // Error if *this capture not enabled by language dialect
        auto lambda3 = [=, *this]  { return var; };
    }
};