CUDA编程【3】计时

CPU计时

下面详细介绍三种常用的测量CUDA核函数执行时间的方法：CUDA Events、CPU计时器 和 Nsight工具。

---

1. 使用CUDA Events

CUDA Events是专门用于测量GPU操作时间的高精度工具。它通过在GPU上插入事件来记录时间戳，适合测量核函数执行时间。

实现步骤：

1. 创建两个CUDA事件（cudaEvent_t）。
2. 在核函数执行前记录开始事件。
3. 在核函数执行后记录结束事件。
4. 同步结束事件以确保记录完成。
5. 计算两个事件之间的时间差。
6. 销毁事件以释放资源。

示例代码：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void myKernel() {
    // 核函数代码
    for (int i = 0; i < 1000000; i++); // 模拟一些计算
}

int main() {
    // 创建CUDA事件
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // 记录开始事件
    cudaEventRecord(start);

    // 执行核函数
    myKernel<<<1, 1>>>();

    // 记录结束事件
    cudaEventRecord(stop);

    // 同步结束事件
    cudaEventSynchronize(stop);

    // 计算时间差（单位：毫秒）
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("Kernel execution time: %f ms\n", milliseconds);

    // 销毁事件
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    return 0;
}

优点：

• 高精度，直接测量GPU时间。
• 简单易用，适合大多数场景。

---

2. 使用CPU计时器

通过CPU计时器（如std::chrono或clock()）测量核函数执行前后的时间差。这种方法依赖于CPU时间，适合粗略测量。

实现步骤：

1. 在核函数执行前获取CPU时间。
2. 执行核函数。
3. 在核函数执行后获取CPU时间。
4. 计算时间差。

示例代码：

#include <cuda_runtime.h>
#include <chrono>
#include <stdio.h>

__global__ void myKernel() {
    // 核函数代码
    for (int i = 0; i < 1000000; i++); // 模拟一些计算
}

int main() {
    // 获取开始时间
    auto start = std::chrono::high_resolution_clock::now();

    // 执行核函数
    myKernel<<<1, 1>>>();
    cudaDeviceSynchronize(); // 确保核函数执行完成

    // 获取结束时间
    auto end = std::chrono::high_resolution_clock::now();

    // 计算时间差（单位：毫秒）
    std::chrono::duration<double, std::milli> elapsed = end - start;
    printf("Kernel execution time: %f ms\n", elapsed.count());

    return 0;
}

优点：

• 不需要额外的CUDA API。
• 适合快速粗略测量。

缺点：

• 精度较低，受CPU和GPU同步影响。
• 不适合测量非常短的时间。

---

3. 使用Nsight工具

Nsight Compute和Nsight Systems是NVIDIA提供的性能分析工具，可以详细测量核函数的执行时间以及其他性能指标。

Nsight Systems 是一个系统级的性能分析工具，可以捕获 GPU、CPU 和内存等硬件资源的性能数据。
sight Compute 是一个核函数级别的性能分析工具，专注于 CUDA 核函数的性能分析。

实现步骤：

1. 安装Nsight Compute或Nsight Systems。
2. 编译CUDA程序时添加调试信息（如-lineinfo）。
3. 使用Nsight工具运行程序并捕获性能数据。
4. 查看核函数的执行时间和其他性能指标。

示例步骤：

安装Nsight

• 从NVIDIA官网下载并安装Nsight Compute或Nsight Systems。

编译CUDA程序

nvcc -o my_program my_program.cu -lineinfo

使用Nsight Compute

1. 打开Nsight Compute。
2. 选择要分析的可执行文件（my_program）。
3. 点击“Start”运行程序并捕获性能数据。
4. 在分析结果中查看核函数的执行时间。

使用Nsight Systems

1. 打开Nsight Systems。
2. 选择要分析的可执行文件（my_program）。
3. 点击“Start”运行程序并捕获性能数据。
4. 在时间线视图中查看核函数的执行时间。

优点：

• 提供详细的性能分析，包括执行时间、内存带宽、寄存器使用等。
• 适合深入优化和分析。

---

补充：cuda events

1. 创建和销毁 CUDA Events

cudaEventCreate

cudaError_t cudaEventCreate(cudaEvent_t* event);

• 功能：创建一个 CUDA 事件。
• 参数：
  • event：指向 cudaEvent_t 类型的指针，用于存储创建的事件。
• 返回值：返回 cudaSuccess 表示成功，否则返回错误代码。
• 示例：

  cudaEvent_t start;
  cudaEventCreate(&start);
  

cudaEventCreateWithFlags

cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);

• 功能：创建一个带有特定标志的 CUDA 事件。
• 参数：
  • event：指向 cudaEvent_t 类型的指针，用于存储创建的事件。
  • flags：指定事件行为的标志，常用的标志包括：
    • cudaEventDefault：默认行为。
    • cudaEventBlockingSync：事件同步时阻塞主机线程。
    • cudaEventDisableTiming：禁用事件的时间记录功能（用于纯同步）。
    • cudaEventInterprocess：允许事件在进程间共享（需要 CUDA 11.0 及以上）。
• 返回值：返回 cudaSuccess 表示成功，否则返回错误代码。
• 示例：

  cudaEvent_t stop;
  cudaEventCreateWithFlags(&stop, cudaEventBlockingSync);
  

cudaEventDestroy

cudaError_t cudaEventDestroy(cudaEvent_t event);

• 功能：销毁一个 CUDA 事件。
• 参数：
  • event：要销毁的事件。
• 返回值：返回 cudaSuccess 表示成功，否则返回错误代码。
• 示例：

  cudaEventDestroy(start);
  

---

2. 记录 CUDA Events

cudaEventRecord

cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);

• 功能：在指定的 CUDA 流中记录一个事件。
• 参数：
  • event：要记录的事件。
  • stream：指定事件所属的 CUDA 流，默认为 0（默认流）。
• 返回值：返回 cudaSuccess 表示成功，否则返回错误代码。
• 示例：

  cudaEventRecord(start, 0); // 在默认流中记录事件
  

---

3. 同步 CUDA Events

cudaEventSynchronize

cudaError_t cudaEventSynchronize(cudaEvent_t event);

• 功能：阻塞主机线程，直到指定的事件完成。
• 参数：
  • event：要同步的事件。
• 返回值：返回 cudaSuccess 表示成功，否则返回错误代码。
• 示例：

  cudaEventSynchronize(stop); // 等待事件完成
  

cudaEventQuery

cudaError_t cudaEventQuery(cudaEvent_t event);

• 功能：查询事件是否完成，而不阻塞主机线程。
• 参数：
  • event：要查询的事件。
• 返回值：
  • cudaSuccess：事件已完成。
  • cudaErrorNotReady：事件未完成。
  • 其他错误代码：表示发生了错误。
• 示例：

  if (cudaEventQuery(stop) == cudaSuccess) {
      printf("Event completed.\n");
  } else {
      printf("Event not ready.\n");
  }
  

---

4. 计算时间间隔

cudaEventElapsedTime

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

• 功能：计算两个事件之间的时间间隔（以毫秒为单位）。
• 参数：
  • ms：指向存储时间间隔的浮点变量的指针。
  • start：起始事件。
  • stop：结束事件。
• 返回值：返回 cudaSuccess 表示成功，否则返回错误代码。
• 注意：两个事件必须在同一个 CUDA 流中记录。
• 示例：

  float milliseconds;
  cudaEventElapsedTime(&milliseconds, start, stop);
  printf("Elapsed time: %f ms\n", milliseconds);
  

---

5. 高级用法

跨流同步

CUDA Events 可以用于跨流同步。例如，可以在一个流中记录事件，然后在另一个流中等待该事件完成：

cudaEvent_t event;
cudaEventCreate(&event);

// 在流 1 中记录事件
cudaEventRecord(event, stream1);

// 在流 2 中等待事件完成
cudaStreamWaitEvent(stream2, event, 0);

禁用时间记录

如果仅需要事件用于同步而不需要时间记录，可以使用 cudaEventDisableTiming 标志创建事件，以减少开销：

cudaEvent_t syncEvent;
cudaEventCreateWithFlags(&syncEvent, cudaEventDisableTiming);

---

6. 完整示例

以下是一个完整的示例，展示如何使用 CUDA Events 测量核函数的执行时间：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void kernel() {
    // 简单的核函数
    for (int i = 0; i < 1000000; i++);
}

int main() {
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    // 记录开始事件
    cudaEventRecord(start, 0);

    // 启动核函数
    kernel<<<1, 1>>>();

    // 记录结束事件
    cudaEventRecord(stop, 0);

    // 等待结束事件完成
    cudaEventSynchronize(stop);

    // 计算时间间隔
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

    printf("Kernel execution time: %f ms\n", milliseconds);

    // 销毁事件
    cudaEventDestroy(start);
    cudaEventDestroy(stop);

    return 0;
}

---

7. 总结

CUDA Events 的主要函数包括：

• cudaEventCreate / cudaEventCreateWithFlags：创建事件。
• cudaEventDestroy：销毁事件。
• cudaEventRecord：在流中记录事件。
• cudaEventSynchronize：同步事件。
• cudaEventQuery：查询事件状态。
• cudaEventElapsedTime：计算时间间隔。