07.给核函数计时

原创

已于 2024-08-25 16:04:23 修改 · 836 阅读

29 ·

CC 4.0 BY-SA版权

文章标签：

#机器学习 #语言模型 #算法 #c语言

于 2024-07-18 19:38:45 首次发布

在内核的性能转换过程中，了解核函数的执行需要多长时间是很有帮助并且十分关键的。衡量核函数性能的方法有很多。最简单的方法是在主机端使用一个CPU或GPU计时器来计算内核的执行时间。

1.用CPU计时器计时

计时我们使用gettimeofday() 函数，是linux下的一个库函数，创建一个cpu计时器，从1970年1月1日0点以来到现在的秒数，需要头文件sys/time.h。

完整代码：https://github.com/dive-into-cuda

#include <sys/time.h>
double cpuSecond()
{
   
   
  struct timeval tp;
  gettimeofday(&tp,NULL);
  return((double)tp.tv_sec+(double)tp.tv_usec*1e-6);
}

接着使用cpuSecond函数来测试你的核函数，由于核函数调用与主机端程序是异步的，你需要用cudaDeviceSynchronize函数来等待所有的GPU线程运行结束。为了更好的测试，需要修改一下之前的程序。

设置数据集大小来对一个有16M个元素的大向量进行测试
```
int nElem = 1 <<24;
```

由于GPU的可扩展性，你需要借助块和线程的索引来计算一个按行优先的数组索引i，并对核函数进行修改，添加限定条件（i＜N）来检验索引值是否越界

__global__ void sumArraysOnGPU(float *A, float *B, float *C, const int N)
{
     
     
    int i = blockId.x * blockDim.x + threadIdx.x;
    if(i < N) 
    {
     
     
        C[i] = A[i] + B[i];
    }  
}

在开始写计时代码之前，我们先来写一个头文件，放我们整个项目代码中常用的函数声明、宏定义、结构体定义等。文件名：include/utils.h

#ifndef _UTILS_H
#define _UTILS_H
#define CHECK(call)                                                             \
{
     
                                                                                    \
    const cudaError_t error = call;                                             \
    if (error != cudaSuccess)                                                   \
    {
     
                                                                                \
        printf("Error: %s:%d, ", __FILE__, __LINE__);                           \
        printf("code: %d, reason: %s\n", error, cudaGetErrorString(error));     \
        exit(1);                                                                \
    }                                                                           \