---

前言

CUDA（Compute Unified Device Architecture）是NVIDIA推出的一种并行计算平台和编程模型，旨在利用图形处理单元（GPU）的强大计算能力来加速各种应用。在CUDA编程中，异步执行是一项重要特性，它允许CPU和GPU并行工作，提高了整体计算效率和资源利用率。本文将简要介绍CUDA中的GPU异步执行，探讨其工作原理及优势。

---

asyncAPI 异步API

背景

在并行计算中，CPU 和 GPU 是两种不同的处理器。CPU 是中央处理器，擅长执行复杂的计算任务；而 GPU 是图形处理器，擅长处理大量简单的并行任务。通过同时利用这两种处理器，我们可以显著提高计算效率。

异步API的作用

这个示例展示了如何在 CPU 和 GPU 之间实现任务的并行执行。具体来说，通过异步API，CPU 可以在 GPU 进行计算的同时继续执行其他任务，而不会等待 GPU 完成后再开始自己的工作。

工作原理

1. 异步执行: CUDA 流调用是异步的，这意味着当你让 GPU 开始执行某个任务时，CPU 不需要等待这个任务完成，而是可以立即继续执行其他任务。这就像你在家里开了一个洗衣机，然后去做其他事情一样，不需要一直站在洗衣机旁边等着。

2. 事件记录: CUDA 事件类似于一个“标记”或“时间戳”，用来记录某个任务的开始和结束时间。通过这些事件，你可以测量任务执行的时间，或者检查某个任务是否已经完成。

3. CPU 和 GPU 协同工作: 当 GPU 在执行计算任务时，CPU 可以进行其他计算任务或者准备下一批数据。比如，你可以让 GPU 计算图像处理任务的同时，让 CPU 处理用户输入。

4. 查询事件: CPU 可以通过查询 CUDA 事件来检查 GPU 是否完成了任务。如果任务完成，CPU 可以继续进行下一步操作；如果任务尚未完成，CPU 可以选择等待或执行其他任务。

示例说明

假设你有一个计算任务需要在 GPU 上运行，同时你希望在 GPU 计算时，CPU 也能做一些其他工作。这个示例展示了如何实现这种并行执行，并通过 CUDA 事件来协调 CPU 和 GPU 的工作。

1. 启动 GPU 任务: 你让 GPU 开始一个计算任务，例如矩阵乘法。
2. CPU 继续工作: 在 GPU 执行任务时，CPU 继续进行其他计算任务或准备下一批数据。
3. 记录和查询事件: 你在 GPU 任务开始和结束时插入 CUDA 事件，通过这些事件可以知道 GPU 任务何时完成。

简单比喻

就像你在家里同时做多件事情：

• 你启动洗衣机（GPU 开始计算）。
• 在洗衣机工作的时候，你开始做饭（CPU 执行其他任务）。
• 你通过时不时查看洗衣机的状态（查询 CUDA 事件）来判断洗衣机是否完成工作。
• 当洗衣机完成工作后，你可以开始晾衣服（GPU 任务完成后，CPU 执行后续操作）。

通过这种方式，你可以更高效地利用时间，完成更多的任务。同样，通过 CUDA 的异步 API，你可以更高效地利用 CPU 和 GPU，显著提升计算性能。

asyncAPI 的使用

页锁定内存的概念

在计算机内存管理中，页锁定（page-locked）内存是指分配的一块内存区域在操作系统内存管理中被锁定或固定住，不会被交换到硬盘或其他存储设备上。这种内存也称为页锁定内存或锁页内存。

为什么要使用页锁定内存

1. 提高数据传输效率：

   • 页锁定内存的主要优势是它可以提高CPU和GPU之间的数据传输效率。当你使用cudaMallocHost来分配页锁定内存时，这块内存不会被操作系统移动或交换，这样可以加速从主机到设备的数据传输，减少延迟。

2. 防止内存交换：

   • 在标准的内存中，操作系统可能会将一些内存页交换到硬盘中，以释放物理内存。然而，对于CUDA程序来说，这种交换会导致数据传输速度变慢。页锁定内存避免了这种情况，因为它固定在物理内存中，不会被交换到硬盘。

通俗易懂的解释：

分配页锁定内存就像把你的内存预留在内存的“VIP区域”，确保它不会被系统移动或替换。这样，当你需要把数据从计算机的主内存传输到GPU时，可以更快地完成这个操作，因为这块内存已经被锁定在物理内存中，不会受到其他进程的影响。简而言之，页锁定内存可以让CPU和GPU之间的数据交换更高效，从而加快计算速度。

cudaMallocHost 函数介绍

函数原型：

cudaError_t cudaMallocHost(void** ptr, size_t size);

作用：

cudaMallocHost 函数用于在主机内存中分配一块页锁定内存。页锁定内存是指这块内存在操作系统中被固定住，不会被交换到硬盘。此类内存对于高效的数据传输尤其重要，特别是在GPU计算中，因为它可以减少从主机到设备的数据传输延迟。

参数：

1. ptr：一个指向指针的指针，用于返回分配的内存地址。在函数调用后，这个指针会指向分配的页锁定内存。

2. size：要分配的内存大小，以字节为单位。

返回值：

• cudaSuccess：成功分配内存。
• 其他错误码：失败时返回相应的错误码，如 cudaErrorMemoryAllocation（内存分配错误）、cudaErrorInvalidValue（无效的值）等。

详细解释：

• 页锁定内存：页锁定内存是操作系统中不会被换出到硬盘的内存区域。由于它的特殊性质，CPU和GPU之间的数据传输速度较快，适用于高性能计算中的数据传输需求。

• 性能优势：使用页锁定内存可以显著提高数据传输的效率，因为数据在传输过程中不会因为内存被交换到硬盘而导致延迟。对于需要频繁与GPU交换数据的应用程序，使用cudaMallocHost分配的内存非常重要。

同步设备

函数原型：

cudaError_t cudaDeviceSynchronize(void);

作用：

cudaDeviceSynchronize 函数用于在主机端（CPU）等待设备端（GPU）完成所有的CUDA任务和操作。它确保在调用函数时，所有的GPU操作（包括核函数执行、内存传输等）都已经完成，并且没有错误发生。这个函数是同步的，调用它会阻塞主机线程，直到GPU完成所有工作。

详细解释：

1. 确保完成：

   • cudaDeviceSynchronize 保证在函数调用返回之前，所有在设备上提交的CUDA任务（如核函数、内存传输等）都已经完成。这样可以确保主机程序在依赖GPU计算结果时，不会出现数据不一致的问题。

2. 阻塞行为：

   • 函数调用会阻塞CPU线程，直到GPU完成所有之前提交的操作。此时，CPU线程会等待GPU任务的完成，不能继续执行后续代码，直到同步完成。

3. 错误处理：

   • 如果在设备上发生了错误（例如核函数执行错误），cudaDeviceSynchronize 会在返回时报告这些错误。通过检查返回值和错误代码，你可以获取详细的错误信息，并进行调试。

4. 使用场景：

   • 性能测量：在性能测试中，通常会在核函数调用后调用 cudaDeviceSynchronize 来准确测量执行时间，因为它确保所有的GPU操作都完成了。
   • 数据依赖：当主机程序依赖于GPU计算的结果时，可以使用 cudaDeviceSynchronize 确保结果已准备好，以避免读取未完成的结果。

例子：

#include <cuda_runtime.h>
#include <iostream>__global__ void myKernel() {// Kernel code here
}int main() {// Launch a kernelmyKernel<<<1, 1>>>();// Ensure the kernel completescudaError_t err = cudaDeviceSynchronize();if (err != cudaSuccess) {std::cerr << "CUDA error: " << cudaGetErrorString(err) << std::endl;return -1;}std::cout << "Kernel execution completed." << std::endl;return 0;
}

在这个示例中，cudaDeviceSynchronize 被用来确保在检查结果之前，核函数 myKernel 已经完成执行。如果有错误发生，cudaDeviceSynchronize 将返回一个错误码，并输出错误信息。

事件

事件的作用

CUDA的事件（Events）是用来在GPU计算和CPU计算之间同步和测量的工具。简单来说，它们可以帮助你跟踪GPU的工作进度，并且可以用来衡量GPU任务的执行时间。以下是CUDA事件的几个主要用途：

1. 同步GPU和CPU

当你有多个GPU任务需要按顺序执行时，CUDA事件可以确保这些任务的顺序和同步。例如，你可以用事件来确保一个任务在另一个任务开始之前完成。这对于需要在一个GPU计算完成后才开始另一个计算的场景特别有用。

举例：
假设你在GPU上先计算了一个图像的滤镜，然后再计算滤镜后的图像。你可以设置一个事件在滤镜计算完成后触发，然后再开始处理滤镜后的图像。

2. 测量GPU任务的时间

CUDA事件可以用来测量某个任务的执行时间。你可以在任务开始和结束时记录事件，然后计算两个事件之间的时间差，从而知道任务的执行时间。这对于性能优化非常重要，因为你可以通过这种方式找到程序中性能瓶颈的地方。

举例：
如果你有一个复杂的计算任务，你可以用事件记录开始和结束时间，然后计算出这个任务花了多长时间。这样你就知道这个任务的性能如何了。

3. 重叠CPU和GPU的工作

CUDA事件可以帮助你在CPU和GPU之间重叠执行工作。比如，当GPU正在处理数据时，CPU可以同时进行其他计算，避免资源的浪费。这种技术称为异步执行，它可以提高程序的整体效率。

举例：
在训练深度学习模型时，GPU可能需要时间来处理数据，同时你可以让CPU进行数据预处理。通过设置事件来确保GPU的计算和CPU的预处理能够重叠进行，从而提高整体计算效率。

创建一个事件

cudaEventCreate 是 CUDA Runtime API 中的一个函数，用于创建一个 CUDA 事件对象。这个事件对象用于在 GPU 和 CPU 之间进行同步和时间测量。事件可以帮助你跟踪 GPU 的任务执行状态，以及在 CPU 和 GPU 之间实现更高效的工作流。

函数原型

cudaError_t cudaEventCreate(cudaEvent_t *event);

作用

cudaEventCreate 的主要作用是创建一个事件对象，这个对象可以用来：

• 测量 GPU 执行时间：你可以在任务开始和结束时记录事件，通过计算事件之间的时间差来测量任务的执行时间。
• 同步 GPU 和 CPU：事件可以用来确保某个操作在 GPU 上完成后再继续执行 CPU 上的后续操作。
• 检测 GPU 执行状态：你可以使用事件来检查 GPU 是否已经完成某些特定的任务。

参数

• event：一个指向 cudaEvent_t 类型的指针。这个指针将被填充为创建的事件对象的句柄。

返回值

cudaEventCreate 函数返回一个 cudaError_t 类型的错误码，指示函数调用的成功与否。常见的返回值包括：

• cudaSuccess：事件创建成功。
• cudaErrorOutOfMemory：创建事件失败，可能是由于设备内存不足。
• cudaErrorInvalidValue：传递给函数的参数无效。

示例代码

#include <cuda_runtime.h>
#include <iostream>int main() {cudaEvent_t start, stop;// 创建事件cudaError_t err = cudaEventCreate(&start);if (err != cudaSuccess) {std::cerr << "Failed to create start event: " << cudaGetErrorString(err) << std::endl;return 1;}err = cudaEventCreate(&stop);if (err != cudaSuccess) {std::cerr << "Failed to create stop event: " << cudaGetErrorString(err) << std::endl;cudaEventDestroy(start);  // 销毁已创建的事件return 1;}// 使用事件...// 销毁事件cudaEventDestroy(start);cudaEventDestroy(stop);return 0;
}

解释

• 创建事件：cudaEventCreate 函数用于创建一个新的事件对象，并将其句柄返回。事件对象在创建后可以用于记录时间和同步。
• 使用事件：创建的事件可以通过 cudaEventRecord 记录在 GPU 的某个时刻，然后通过 cudaEventElapsedTime 测量两个事件之间的时间差。
• 销毁事件：在不再需要事件时，你应该使用 cudaEventDestroy 销毁事件对象，以释放相关资源。

记录事件

cudaEventRecord 是 CUDA Runtime API 中的一个函数，用于在 CUDA 设备的流中记录一个事件。事件可以用来测量 GPU 操作的时间、同步 GPU 和 CPU 任务等。

函数原型

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream);

作用

cudaEventRecord 的作用是记录一个事件对象到指定的 CUDA 流中。这使得你可以：

• 测量时间：通过在异步操作之前和之后记录事件，可以测量异步操作的执行时间。
• 同步操作：通过记录事件，你可以确保某些 GPU 操作在执行其他操作之前已经完成。
• 检测状态：记录事件后，你可以查询事件来确定 GPU 是否已经完成某些特定的任务。

参数

• event：要记录的事件对象的句柄。你需要先使用 cudaEventCreate 创建这个事件对象。
• stream：事件记录到的 CUDA 流。你可以指定一个 CUDA 流（如 0 表示默认流），也可以指定一个自定义流。

返回值

cudaEventRecord 函数返回一个 cudaError_t 类型的错误码，指示函数调用的成功与否。常见的返回值包括：

• cudaSuccess：事件记录成功。
• cudaErrorInvalidValue：传递给函数的参数无效。
• cudaErrorInvalidResourceHandle：事件或流句柄无效。

为什么在 start 和 stop 之间执行异步函数

在测量异步操作的执行时间时，你需要在异步操作开始之前和结束之后记录事件。这是因为：

1. 精确测量时间：要准确测量一个异步操作的执行时间，你需要记录下操作开始和结束的确切时刻。通过在异步函数调用前后分别记录 start 和 stop 事件，你可以精确计算出异步操作的执行时间。

2. 保证正确的顺序：在异步操作开始之前记录 start 事件，可以确保这个事件的时间戳是在异步操作开始的时刻。而在异步操作结束之后记录 stop 事件，可以确保这个事件的时间戳是在异步操作完成的时刻。这样，你可以确保时间测量涵盖了整个异步操作的执行周期。

异步复制内存

cudaMemcpyAsync 是 CUDA Runtime API 中的一个函数，用于在主机（CPU）和设备（GPU）之间异步地复制数据。与同步的 cudaMemcpy 函数不同，cudaMemcpyAsync 允许数据传输操作在后台进行，CPU 可以在数据传输期间继续执行其他操作，从而提高计算效率。

函数原型

cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count, cudaMemcpyKind kind, cudaStream_t stream);

作用

cudaMemcpyAsync 的主要作用是异步地将数据从主机（CPU）内存复制到设备（GPU）内存，或从设备（GPU）内存复制到主机（CPU）内存。这个函数的异步特性允许在数据传输的同时，CPU 可以执行其他计算任务，从而更有效地利用计算资源。

参数

• dst：目标内存地址，即数据复制到的位置。可以是设备内存或主机内存，取决于 kind 参数。
• src：源内存地址，即数据复制的来源。可以是主机内存或设备内存，取决于 kind 参数。
• count：要复制的数据字节数。
• kind：数据传输的类型，使用 cudaMemcpyKind 枚举值指定。常见的值包括：
  • cudaMemcpyHostToDevice：从主机内存复制到设备内存。
  • cudaMemcpyDeviceToHost：从设备内存复制到主机内存。
  • cudaMemcpyDeviceToDevice：在设备内存之间复制。
  • cudaMemcpyHostToHost：在主机内存之间复制（不常用）。
• stream：指定数据传输操作要使用的 CUDA 流。如果传递 0，表示使用默认流。如果使用非默认流，数据传输和其他操作可以在不同的流中并行执行。

返回值

cudaMemcpyAsync 函数返回一个 cudaError_t 类型的错误码，指示函数调用的成功与否。常见的返回值包括：

• cudaSuccess：数据复制成功。
• cudaErrorInvalidValue：传递给函数的参数无效。
• cudaErrorInvalidMemcpyDirection：指定的复制方向无效。
• cudaErrorInvalidDevicePointer：目标或源地址指针无效。
• cudaErrorInvalidMemcpySize：复制大小无效。

检查事件状态

cudaEventQuery 是 CUDA Runtime API 中的一个函数，用于查询一个 CUDA 事件（event）的状态，检查它是否已经完成。它允许程序在不阻塞的情况下检查事件是否完成，从而可以在程序中实现非阻塞的同步和资源管理。

函数原型

cudaError_t cudaEventQuery(cudaEvent_t event);

作用

cudaEventQuery 的作用是查询指定事件的状态。事件通常用于跟踪 GPU 任务的进度和实现同步。通过调用 cudaEventQuery，可以检查事件是否已经完成，而不会阻塞 CPU 的执行。这在实现非阻塞的同步和异步操作时非常有用。

参数

• event：要查询的 CUDA 事件。该参数是一个 cudaEvent_t 类型的变量，通常通过 cudaEventCreate 创建。

返回值

cudaEventQuery 函数返回一个 cudaError_t 类型的错误码，指示函数调用的成功与否。常见的返回值包括：

• cudaSuccess：事件已经完成。
• cudaErrorNotReady：事件尚未完成。这个返回值表明事件还没有完成，可能需要等待或在后续的代码中再次查询。
• cudaErrorInvalidValue：传递给函数的参数无效。
• cudaErrorInvalidResourceHandle：指定的事件句柄无效。

使用场景

• 非阻塞检查：在异步执行的环境中，cudaEventQuery 允许程序非阻塞地检查任务的完成状态。如果事件尚未完成，程序可以选择继续执行其他任务，而不是等待事件完成。

• 资源管理：在需要同步 CPU 和 GPU 操作的场景中，cudaEventQuery 允许检查 GPU 任务的完成情况，以决定是否可以释放资源或执行后续操作。

获取GPU执行时间

cudaEventElapsedTime 是 CUDA Runtime API 中的一个函数，用于计算两个 CUDA 事件之间的时间间隔。它通常用于测量 GPU 执行的时间，以帮助分析和优化 CUDA 程序的性能。

函数原型

cudaError_t cudaEventElapsedTime(float *ms, cudaEvent_t start, cudaEvent_t stop);

作用

cudaEventElapsedTime 计算从 start 事件记录到 stop 事件记录之间经过的时间，并将结果以毫秒（ms）为单位返回。它提供了一种测量 GPU 执行时间的方式，可以用于性能分析和调优。

参数

• ms：指向 float 类型的指针，用于存储计算出的时间间隔。函数将时间间隔以毫秒为单位写入到这个位置。

• start：表示时间区间的开始事件。通常是一个 cudaEvent_t 类型的事件句柄，通过 cudaEventCreate 创建，并在 GPU 执行开始时记录。

• stop：表示时间区间的结束事件。通常是一个 cudaEvent_t 类型的事件句柄，通过 cudaEventCreate 创建，并在 GPU 执行结束时记录。

返回值

cudaEventElapsedTime 返回一个 cudaError_t 类型的错误码，指示函数调用的成功与否。常见的返回值包括：

• cudaSuccess：函数调用成功，时间间隔计算正常。
• cudaErrorInvalidValue：传递给函数的参数无效。
• cudaErrorInvalidResourceHandle：指定的事件句柄无效。

使用场景

• 性能测量：在 CUDA 编程中，cudaEventElapsedTime 用于测量 GPU 内核函数或其他 CUDA 操作的执行时间，以评估性能和识别瓶颈。

• 性能优化：通过分析时间间隔，可以识别需要优化的代码部分或调整参数设置，以提高计算效率。

示例展示

 /** 此示例说明了 CUDA 事件在 GPU 计时以及重叠的 CPU 和 GPU 执行中的用法。* 事件被插入到 CUDA 调用流中。由于 CUDA 流调用是异步的，CPU 可以在 GPU 执行时执行计算（包括主机和设备之间的 DMA 内存复制）。* CPU 可以查询 CUDA 事件以确定 GPU 是否已完成任务。*/// 包含系统头文件
#include <stdio.h>// 包含 CUDA 运行时头文件
#include <cuda_runtime.h>
#include <cuda_profiler_api.h>// 包含项目头文件
#include <helper_cuda.h>
#include <helper_functions.h>  // 辅助实用函数// 核函数，增加每个元素的值
__global__ void increment_kernel(int* g_data, int inc_value) {int idx = blockIdx.x * blockDim.x + threadIdx.x;  // 计算全局线程索引g_data[idx] = g_data[idx] + inc_value;  // 增加对应索引的值
}// 检查输出结果是否正确
bool correct_output(int* data, const int n, const int x) {for (int i = 0; i < n; i++) {if (data[i] != x) {printf("Error! data[%d] = %d, ref = %d\n", i, data[i], x);return false;}}return true;
}int main(int argc, char* argv[]) {int devID;cudaDeviceProp deviceProps;printf("[%s] - Starting...\n", argv[0]);// 选择最合适的 CUDA 设备devID = findCudaDevice(argc, (const char**)argv);// 获取设备属性checkCudaErrors(cudaGetDeviceProperties(&deviceProps, devID));printf("CUDA device [%s]\n", deviceProps.name);int n = 16 * 1024 * 1024;  // 数组大小int nbytes = n * sizeof(int);  // 数组字节大小int value = 26;  // 增加的值// 分配主机内存int* a = 0;checkCudaErrors(cudaMallocHost((void**)&a, nbytes));  // 分配页锁定主机内存memset(a, 0, nbytes);  // 初始化为 0// 分配设备内存int* d_a = 0;checkCudaErrors(cudaMalloc((void**)&d_a, nbytes));  // 分配设备内存checkCudaErrors(cudaMemset(d_a, 255, nbytes));  // 初始化为 255// 设置核函数启动配置dim3 threads = dim3(512, 1);  // 每个块中的线程数dim3 blocks = dim3(n / threads.x, 1);  // 块的数量// 创建 CUDA 事件句柄cudaEvent_t start, stop;checkCudaErrors(cudaEventCreate(&start));  // 创建开始事件checkCudaErrors(cudaEventCreate(&stop));  // 创建结束事件// 创建和重置计时器StopWatchInterface* timer = NULL;sdkCreateTimer(&timer);sdkResetTimer(&timer);checkCudaErrors(cudaDeviceSynchronize());  // 同步设备，确保之前的所有 CUDA 调用完成float gpu_time = 0.0f;  // 存储 GPU 执行时间// 异步地向 GPU 发送工作（全部发送到流 0）checkCudaErrors(cudaProfilerStart());  // 启动 CUDA ProfilersdkStartTimer(&timer);  // 启动计时器cudaEventRecord(start, 0);  // 记录开始事件cudaMemcpyAsync(d_a, a, nbytes, cudaMemcpyHostToDevice, 0);  // 异步地将数据从主机复制到设备increment_kernel << <blocks, threads, 0, 0 >> > (d_a, value);  // 启动核函数cudaMemcpyAsync(a, d_a, nbytes, cudaMemcpyDeviceToHost, 0);  // 异步地将数据从设备复制到主机cudaEventRecord(stop, 0);  // 记录结束事件sdkStopTimer(&timer);  // 停止计时器checkCudaErrors(cudaProfilerStop());  // 停止 CUDA Profiler// 在等待阶段 1 完成时，让 CPU 执行一些工作unsigned long int counter = 0;// 查询事件，检查 GPU 是否完成while (cudaEventQuery(stop) == cudaErrorNotReady) {counter++;}// 计算 GPU 执行时间checkCudaErrors(cudaEventElapsedTime(&gpu_time, start, stop));// 打印 CPU 和 GPU 的执行时间printf("time spent executing by the GPU: %.2f\n", gpu_time);printf("time spent by CPU in CUDA calls: %.2f\n", sdkGetTimerValue(&timer));printf("CPU executed %lu iterations while waiting for GPU to finish\n", counter);// 检查输出结果是否正确bool bFinalResults = correct_output(a, n, value);// 释放资源checkCudaErrors(cudaEventDestroy(start));checkCudaErrors(cudaEventDestroy(stop));checkCudaErrors(cudaFreeHost(a));checkCudaErrors(cudaFree(d_a));exit(bFinalResults ? EXIT_SUCCESS : EXIT_FAILURE);  // 根据结果返回相应的退出代码
}

---

总结

CUDA的异步执行机制显著提高了计算任务的并行处理能力。通过异步数据传输和核函数执行，CPU和GPU可以同时处理任务，最大限度地利用系统资源。这不仅减少了程序的等待时间，还提高了整体性能。理解和正确应用异步执行，是充分发挥CUDA强大计算能力的关键。通过合理设计和优化，开发者可以在各种高性能计算应用中实现更高的效率和更快的计算速度。