如何执行cuda c/c++中的性能指标

原文链接

https://developer.nvidia.com/blog/how-implement-performance-metrics-cuda-cc/

　　在上一篇文章中，我们通过cuda c实现SAXPY来了解了cuda c的一些基本知识。在这篇文章中我们将讨论如何分析此代码和其他cuda c代码的性能。在未来的文章中，在越来越重要的性能优化领域，我们将依靠着写性能测量技术。

　　cuda性能测量通常是在host端代码上完成度，可以通过cpu计时器或者cuda特定计时器来实现。在我们了解性能测试技术之前，我们需要讨论如何在host端和device端实现同步。

host端-device端同步

　　先来看一下上一篇文章中host端和device端的数据传输以及核函数启动：

cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice);

// perform SAXPY in 1M elements
int threads=256;
int blocks=(N+threads-1)/threads;   //  上取整
saxpy<<<blocks,threads>>>(N,2.0f,d_x,d_y);
    
cudaMemcpy(y,d_y,N*sizeof(float),cudaMemcpyDeviceToHost);

　　使用cudaMemcpy()在host端和device端传输数据是同步传输(阻塞传输)。在所有先前发出的cuda调用完成之前，同步数据传输不会开始，并且在同步传输完成之前无法开始后续的cuda调用。因此，在第二行y到d_y的传输完成之前，第三行上的saxpy核函数不会启动。另一方面，核函数启动是异步的。第三行的核函数一旦启动，控制权会立刻返回给cpu，不会等待核函数执行完成（自己跑自己的，不管你了）。这也许会为最后一行代码设备向主机传送数据设置了竞争条件，数据传输的阻塞性质确保了核函数向host端传输数据开始之前已经运行完成。

使用cpu定时器计时核函数执行

　　让我们来看看书和使用cpu计时器为核函数执行计时。（伪代码？，t1，t2没声明，myCPUTimer好像也没有这个东西）

cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice);

// perform SAXPY in 1M elements
int threads=256;
int blocks=(N+threads-1)/threads;   //  上取整

t1=myCPUTimer();
saxpy<<<blocks,threads>>>(N,2.0f,d_x,d_y);
cudaDeviceSynchronize();
t2=myCPUTimer();
    
cudaMemcpy(y,d_y,N*sizeof(float),cudaMemcpyDeviceToHost);

 　　除了对主机通用时间戳函数myCPUTimer()的两次调用，我们还使用我们还是用显示同步阻塞cudaDeviceSynchronize()来阻塞cpu执行，直到device端所有之前发出的命令都执行完成。如果不加上这个阻塞，这段代码的量的将是核函数的启动时间而不是核函数的执行时间。

使用cuda event计时

　　使用device-host同步点会有一个问题，synchronize()会停止gpu的pipline。因此，cuda通过cuda event api(https://docs.nvidia.com/cuda/cuda-runtime-api/index.html#group__CUDART__EVENT) 为cpu计时器提供了一个相对轻量级的替代方案,cuda event api包括创建和销毁事件以及计算两个事件之间经过的时间的调用(以毫秒为单位)。

　　cuda event利用率cuda流的概念，cuda流是在device上按顺序执行的一系列操作，不同流中的操作可以交错没在某些情况下可以重叠----该属性可以用于隐藏device和host之间的数据传输（https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#asynchronous-concurrent-execution）（后面再讨论https://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/）。到目前为止，gpu上天的所有操作都发生在默认流或流0(也称作null stream)中。

　　下面我们将在saxpy代码中使用cuda event

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

cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice);
cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice);

// perform SAXPY in 1M elements
int threads=256;
int blocks=(N+threads-1)/threads;   //  上取整

cudaEventRecord(start);
saxpy<<<blocks,threads>>>(N,2.0f,d_x,d_y);
cudaEventRecord(stop);

cudaMemcpy(y,d_y,N*sizeof(float),cudaMemcpyDeviceToHost);
cudaEventSynchronize(stop);
float milliseconds=0;
cudaEventElapsedTime(&milliseconds,start,stop);

　　cuda event的数据类型是cudaEvent_t，他的创建和销毁通过cudaEventCreate()和cudaEventDestory()。在上面的代码中，cudaEventRecord()将开始和停止event放入默认流中，即stream 0.当device在流中达到该event时，该device将记录该event的时间戳。函数cudaEventSynchronize()阻止cpu执行，直到指定的event被记录。cudaEventElapsedTime()函数在第一个参数中返回开始和停止记录之间经过的毫秒数。该值具有大约1/2微妙的分辨率。

内存带宽

　　现在我们有了一种准确计时核函数执行的方法，我们将使用它来计算带宽。在评估带宽效率是，我们同时使用理论峰值带宽和观察到的或有效的内存带宽

理论带宽

　　理论带宽可以使用产品资料中提供的硬件规格计算得出。例如NVIDIA Tesla M2050gpu使用内存时钟频率为1564MHz的DDR（双倍数据速率）RAM和384位宽的内存接口。使用这些数据项，该卡内存带宽的理论值为148GB/s，计算方式如下

　　

　　在计算中，我们将内存时钟频率转换为Hz，乘以接口宽度（除以8，将会单位转化为字节）并乘以2，这是由于双倍数据速率。最终，我们将结果除以10的9次幂，将结果转化为GB/s

有效带宽

　　我们通过对特定程序活动进行计时和了解我们的程序如何访问数据来计算有效带宽。我们可以通过下面公式计算

　　

　　公式中的BWEffective是以GB/s为单位的有效带宽，R_B是每个内核读取的字节数，W_B是每个内核写入的字节数，t是以秒为单位经过的时间。我们可以修改我们的saxpy代码来计算有效带宽

　　

#include <stdio.h>

__global__ void saxpy(int n,float a,float *x,float *y)
{
    int i=blockIdx.x*blockDim.x+threadIdx.x;
    if(i<n)
    {
        y[i]=a*x[i]+y[i];
    }
}

int main(void)
{
    int N=1<<20;    // 1左移20
    float *x,*y,*d_x,*d_y;
    x=(float*)malloc(N*sizeof(float));
    y=(float*)malloc(N*sizeof(float));

    cudaMalloc(&d_x,N*sizeof(float));
    cudaMalloc(&d_y,N*sizeof(float));

    for(int i=0;i<N;++i)
    {
        x[i]=1.0f;
        y[i]=2.0f;
    }

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

    cudaMemcpy(d_x,x,N*sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpy(d_y,y,N*sizeof(float),cudaMemcpyHostToDevice);

    // perform SAXPY in 1M elements
    int threads=512;
    int blocks=(N+threads-1)/threads;   //  上取整

    cudaEventRecord(start);
    saxpy<<<blocks,threads>>>(N,2.0f,d_x,d_y);
    cudaEventRecord(stop);
    
    cudaMemcpy(y,d_y,N*sizeof(float),cudaMemcpyDeviceToHost);
    cudaEventSynchronize(stop);
    float milliseconds=0;
    cudaEventElapsedTime(&milliseconds,start,stop);

    float maxError=0.0f;
    for(int i=0;i<N;++i)
    {
        maxError=max(maxError,abs(y[i]-4.0f));
    }
    printf("Max error: %f\n",maxError);
    printf("Effective Bandwidth(GB/s):%f\n",N*4*3/milliseconds/1e6);

    cudaFree(d_x);
    cudaFree(d_y);
    free(x);
    free(y);

    return 0;
}

　　在带宽的计算中，N*4是每次读取或写入数组时传输的字节数，而3则表示读取x（2）和写入y（1），消耗时间存储在变量milliseconds让他的单位更明确。请注意，除了添加带宽计算所需功能外，我们还更改了数组大小和线程块大小（256=>512）,编译运行，得到结果

测量计算吞吐量（throughput）

　　我们杠杆演示了如何衡量带宽，这是衡量数据吞吐量的一个指标。另一个非常重要的性能指标是计算吞吐量。计算吞吐量的常用度量是GFLOP/s，它代表每秒千兆浮点操作数（Giga-FLoating-point OPerations per second），Giga是10e9，对于saxpy计算，衡量有效吞吐量很简单，每个saxpy元素执行乘加运算，通常以两个FLOP衡量（2*N），所以

 　　N是saxpy操作中元素数量，t是经过时间（以秒为单位）。与理论峰值带宽一样，理论峰值GFLOP/s可以从产品文献中收集（但计算它可能有一点棘手，因为它非常依赖架构）。例如Tesla M2050gpu单精度浮点兔兔量理论峰值为1030GFLOP/s，双精度峰值吞吐量为515GFLOP/s

　　saxpy为每个计算元素读取12个字节，但是只执行一条乘加指令（2FLOP），所以很明显她将受到带宽限制（我认为是读取多，计算少，所以性能影响首先受到带宽限制），因此在这种情况下（实际上在许多情况下），带宽是最重要的衡量和优化指标。在更复杂的计算中，在FLOP级别测量性能可能非常困难。因此，更常见的是使用分析工具来了解计算吞吐量是否到达瓶颈。应用程序通常提供特定于问题的吞吐量指标（而不是特定于架构），因此对用户更有用。For example, “Billion Interactions per Second” for astronomical n-body problems, or “nanoseconds per day” for molecular dynamic simulations.（看不懂）

总结

　　这篇文章描述了如何使用cuda event api为核函数计时。cuda event使用gpu计时器，因此避免了device-host端同步问题。我们提出了有效带宽和计算吞吐量性能指标，并在saxpy核函数中实现了有效带宽。大部分内核都受到内存带宽限制，因此计算有效带宽是性能优化的良好开端，在以后的文章中，我们将讨论如何确定哪个因素是性能的限制因素（带宽，指令或延迟）。

　　cuda event还可以用于确定device端host端的数据传输速率，通过记录event在任何一端的cudaMemcpy()调用。

　　如果你在较小的gpu上运行本文中的代码，除非减小数组大小，否则可能会受到设备内存不足的错误消息。事实上，到目前为止我们的示例代码并没有费心检查运行时的错误。在下篇文章中，我们将学习如何在cuda c/c++中执行错误处理以及如何查询当前设备已确定他们的可用资源，以便我们可以写出更健壮的代码。