CUDA中如何获取内核函数相对于程序启动的启动时间

• 缺少程序启动的基准时间戳
  gettimeofday返回的是系统自1970年1月1日0点以来的绝对时间值，你没有在程序刚启动时记录基准时间做差值计算，直接输出的是绝对时间的毫秒值，自然会出现1e12量级的异常结果。
• 计时逻辑不符合CUDA异步执行特性
  CUDA内核启动接口<<<>>>是异步调用，CPU侧调用完接口会立刻返回执行后续代码，你在内核启动前调用gettimeofday拿到的是CPU发起内核调用的时间，不是内核实际在GPU侧启动执行的时间，存在明显误差。
• 数据类型精度不足
  float类型的有效位数只有6~7位十进制数，存储1e12量级的时间戳时会出现严重的精度丢失，应该用double类型存储时间计算结果。

方案1：CPU侧相对计时（简单实现，误差可接受的场景） #

#include <sys/time.h>

// 内核函数定义
template<int BLOCK_SIZE> __global__ void MatrixMulCUDA(float *C, float *A, float *B, int wA, int wB) {
    // 原有内核逻辑
}

int main(){
    // 程序启动立刻记录基准时间（作为零点）
    struct timeval base_time;
    gettimeofday(&base_time, NULL);

    // 原有memcpy、初始化等逻辑
    // ...

    struct timeval kernel_call_time;
    gettimeofday(&kernel_call_time, NULL);
    MatrixMulCUDA<32><<<matrixMulgrid, matrixMulthreads>>>(C, A, B, dimsA.x, dimsB.x);
    cudaDeviceSynchronize();

    // 计算相对于程序启动的时间，单位为毫秒
    double relative_start_time = (1000000.0 * (kernel_call_time.tv_sec - base_time.tv_sec) 
                                + (kernel_call_time.tv_usec - base_time.tv_usec)) / 1000.0;

    printf("kernel start time = %f ms\n", relative_start_time);
    return 0;
}

方案2：CUDA事件计时（精确获取GPU侧内核启动时间） #

如果需要精准获取内核在GPU上的实际启动时间，应该使用CUDA官方提供的事件接口，避免异步调度带来的误差：

#include <sys/time.h>
#include <cuda_runtime.h>

template<int BLOCK_SIZE> __global__ void MatrixMulCUDA(float *C, float *A, float *B, int wA, int wB) {
    // 原有内核逻辑
}

int main(){
    // 记录程序启动基准时间
    struct timeval base_time;
    gettimeofday(&base_time, NULL);
    double base_ms = (1000000.0 * base_time.tv_sec + base_time.tv_usec) / 1000.0;

    // 初始化CUDA事件
    cudaEvent_t kernel_start;
    cudaEventCreate(&kernel_start);
    // 若要对齐CPU时间，可在CUDA上下文初始化时额外记录基准事件ctx_start
    // cudaEvent_t ctx_start;
    // cudaEventCreate(&ctx_start);
    // cudaEventRecord(ctx_start, 0);

    // 原有memcpy、初始化等逻辑
    // ...

    // 在内核前插入启动事件
    cudaEventRecord(kernel_start, 0);
    MatrixMulCUDA<32><<<matrixMulgrid, matrixMulthreads>>>(C, A, B, dimsA.x, dimsB.x);
    // 等待事件记录完成
    cudaEventSynchronize(kernel_start);

    // 输出相对于CUDA上下文启动的内核启动时间，单位为毫秒
    float kernel_start_ms;
    // cudaEventElapsedTime(&kernel_start_ms, ctx_start, kernel_start);
    // printf("kernel start time = %f ms\n", kernel_start_ms);

    cudaEventDestroy(kernel_start);
    // cudaEventDestroy(ctx_start);
    return 0;
}

内容的提问来源于stack exchange，提问作者mehran