CUDA计时函数：精确测量GPU代码执行时间

在GPU编程中，精确测量代码执行时间是性能优化的关键步骤。CUDA提供了专门的计时工具来帮助开发者准确获取核函数（Kernel）、内存拷贝等操作的耗时。本文将详细介绍CUDA计时函数的使用方法，并通过实例代码演示如何高效测量GPU代码的执行时间。

---

为什么需要CUDA计时函数？

在CPU和GPU异构计算中，CPU和GPU的工作是异步的。若使用传统的CPU计时方法（如clock()或std::chrono），可能无法准确测量GPU代码的执行时间。CUDA的事件（Event）机制能够直接在GPU硬件层面记录时间戳，避免了CPU-GPU同步带来的误差。

---

一、CUDA事件（Event）计时原理

CUDA事件是基于GPU内部时钟的轻量级计时工具，原理如下：

1. 事件记录：在代码中插入事件标记，记录GPU执行到该点的时间戳。

2. 时间差计算：通过两个事件的时间戳差值计算代码段的执行时间。

---

二、CUDA计时函数的使用步骤

1. 创建事件对象

使用cudaEventCreate创建事件对象：

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

2. 记录事件时间戳

在需要计时的代码段前后插入事件记录：

cudaEventRecord(start);  // 记录起始时间戳
// 执行需要计时的代码（例如核函数）
myKernel<<<grid, block>>>(...);
cudaEventRecord(stop);   // 记录结束时间戳

3. 同步事件

由于GPU执行是异步的，需等待事件完成：

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

4. 计算时间差（毫秒）

使用cudaEventElapsedTime计算时间差：

float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);
printf("执行时间: %.3f ms\n", milliseconds);

5. 销毁事件对象

释放事件资源：

cudaEventDestroy(start);
cudaEventDestroy(stop);

---

三、完整示例代码

以下代码演示了如何测量一个核函数的执行时间：

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

// 简单的核函数：向量加法
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    const int N = 1 << 20;  // 1M元素
    int *a, *b, *c;
    int *d_a, *d_b, *d_c;

    // 分配主机内存
    a = (int*)malloc(N * sizeof(int));
    b = (int*)malloc(N * sizeof(int));
    c = (int*)malloc(N * sizeof(int));

    // 分配设备内存
    cudaMalloc(&d_a, N * sizeof(int));
    cudaMalloc(&d_b, N * sizeof(int));
    cudaMalloc(&d_c, N * sizeof(int));

    // 初始化数据
    for (int i = 0; i < N; i++) {
        a[i] = i;
        b[i] = i * 2;
    }

    // 拷贝数据到设备
    cudaMemcpy(d_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, N * sizeof(int), cudaMemcpyHostToDevice);

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

    // 定义线程块和网格大小
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;

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

    // 启动核函数
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, N);

    // 记录结束事件
    cudaEventRecord(stop);
    cudaEventSynchronize(stop);  // 等待核函数执行完成

    // 计算时间差
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("核函数执行时间: %.3f ms\n", milliseconds);

    // 拷贝结果回主机
    cudaMemcpy(c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

    // 验证结果（可选）
    bool success = true;
    for (int i = 0; i < N; i++) {
        if (c[i] != a[i] + b[i]) {
            success = false;
            break;
        }
    }
    if (success) {
        printf("结果验证成功！\n");
    } else {
        printf("结果验证失败！\n");
    }

    // 释放资源
    cudaEventDestroy(start);
    cudaEventDestroy(stop);
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(a);
    free(b);
    free(c);

    return 0;
}

四、注意事项

1. 同步的必要性：

   如果未调用cudaEventSynchronize，时间差的计算可能不准确。
   核函数启动后，CPU会继续执行后续代码，必须同步等待GPU完成。

2. 事件的开销：

   CUDA事件的创建和销毁有一定开销，避免在频繁调用的代码段中使用。

3. 多流（Stream）环境：

   若使用多流并行，需为每个流单独创建事件，并通过cudaEventRecord指定流。

4. 时间单位：

   cudaEventElapsedTime返回的时间单位为毫秒（ms），分辨率约为0.5微秒。