[譯]在CUDA C/C++中如何衡量代碼性能

本文翻譯自NVIDIA官方博客Parallel Forall，內容僅供參考，若有疑問請訪問原網站:https://devblogs.nvidia.com/p....

在這個系列的第一篇文章中，咱們經過用CUDA C/C++實現SAXPY，學習了CUDA C/C++編程的基本要素。在這篇文章中，咱們會學習如何衡量這個程序以及其餘CUDAC/C++程序的性能。咱們在以後的文章中常常用到這種性能度量技術，由於程序的性能優化將會變得愈來愈重要。

譯者注：這個系列是指原文的系列，並非筆者的專欄。

CUDA性能度量一般是在主機端進行的，咱們既可使用CPU的計時器也可使用CUDA專門的計時器。在開始學習性能度量技術以前，咱們須要討論一下如何同步主機和設備之間的操做。

主機-設備同步

讓咱們來看一下上一篇博客中SAXPY的數據傳輸和核函數啓動的主機端代碼：

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

saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);

cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

這裏使用cudaMemcpy進行數據傳輸的方式是同步傳輸(或者是阻塞傳輸)方式。同步數據傳輸直到前面全部發布的CUDA調用所有結束以後纔會開始，並且同步數據傳輸結束以後，隨後的CUDA調用纔會開始。所以上面第三行的saxpy核函數只有到第二行的y到d_y的數據傳輸結束以後纔會啓動。而在另外一方面，核函數啓動倒是異步的。一旦核函數被啓動，控制權就馬上返回到CPU，並不會等待覈函數執行完成。這樣的話就會對最後一行的設備到主機數據傳輸產生競態條件(race condition)，可是數據傳輸的阻塞特性會確保核函數執行完成後再開始數據傳輸。

譯者注：這裏的競態條件前面提到過，簡單說就是前面的數據操做還未完成，後面的操做卻又要使用前面的數據，這樣就會致使錯誤的結果。

使用CPU的計時器來計算核函數的執行時間

如今咱們來看一下如何使用CPU的計時器來給核函數計時。

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

t1 = myCPUTimer();
saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);
cudaDeviceSynchronize();
t2 = myCPUTimer();

cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

在上面的代碼中，咱們除了使用通常的主機時間戳函數myCPUTimer()，還用到了顯式的同步障礙 cudaDeviceSynchronize()來阻塞CPU執行，直到設備上發佈的指令所有執行結束爲止。若是沒有這個同步障礙，這個代碼測試的就是核函數的啓動時間而不是執行時間。

使用CUDA事件計時

使用相似cudaDeviceSynchronize()函數的主機設備同步點的一個問題就是它會拖延GPU管道(stall GPU pipeline)。基於這個緣由，CUDA提供了一個相比CPU計時器更輕量級的選擇，那就是使用CUDA事件API。CUDA事件API包括調用事件建立和銷燬函數、事件記錄函數以及以毫秒爲單位計算兩個被記錄事件的運行時間的函數。

譯者注：這裏拖延GPU管道(stall GPU pipeline)的直接結果就是形成CPU和GPU輪流執行，而再也不是並行執行。因而就使得程序的運行時間等於CPU與GPU時間之和。具體能夠參考:https://blogs.msdn.microsoft....

CUDA事件使用的是CUDA streams的概念。一個CUDA流只是一系列在設備上順序執行的操做。不一樣流中的操做能夠交替執行，在某些狀況下甚至能夠交疊執行，這個特性能夠被用在隱藏主機和設備間的數據傳輸。(咱們會在以後的文章中討論)。到目前爲止，咱們全部的操做都是在默認的流中進行的，或者0號流(也叫作空流)。

下面的代碼中，咱們使用了CUDA事件API來對SAXPY代碼進行性能度量。

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

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

cudaEventRecord(start);
saxpy<<<(N+255)/256, 256>>>(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事件是cudaEvent_t類型，經過cudaEventCreate()和cudaEventDestroy()進行事件的建立和銷燬。在上面的代碼中cudaEventRecord()將事件start和stop放在默認的流中，即0號stream。函數cudaEventSynchronize()用來阻塞CPU執行直到指定的事件被記錄。函數 cudaEventElapsedTime()的第一個參數返回start和stop兩個記錄之間消逝的毫秒時間。這個值的精度大約是0.5ms。

內存帶寬

既然咱們已經能夠精確地測量核函數的運行時間，那麼咱們就能夠用它來計算帶寬。咱們須要使用理論的峯值帶寬和有效內存帶寬來評估帶寬效率。

理論帶寬

理論帶寬能夠經過產品資料中的硬件規格來計算。例如英偉達Tesla M2050 GPU使用的是時鐘頻率爲1546MHz顯存位寬爲384-bit的DDR(雙倍數據速率)RAM。

使用這些數據，咱們能夠計算出英偉達Tesla M2050的理論峯值帶寬是148 GB/sec：

$$BW_{Theoretical}=1546 * 106 * (384/8) * 2 / 109 = 148 GB/s $$

在這個表達式中，咱們將內存的時鐘頻率的單位轉化爲Hz，而後乘以顯存寬度(除以8以後，單位由比特轉化爲字節)，又乘以2是由於該顯卡的RAM是DDR(雙倍數據速率)。最後咱們將結果除以10^9獲得以GB/s的計算結果。

有效帶寬

咱們是經過計算特定程序的活動時間和程序如何訪問數據來計算機有效帶寬的。咱們使用下面的公式：

$$BW_{Effective} = (R_B + W_B) / (t * 109)$$

這裏，$BW_{Effective}$是以GB/s的有效帶寬，$R_B$是每一個核函數被讀取的字節數，$W_B$是每一個核函數被寫入的字節數，$t$是以秒爲單位的運行時間。咱們能夠修改SAXPY例子來計算有效帶寬，下面是完整的代碼：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#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 = 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);

    cudaEventRecord(start);

    // Perform SAXPY on 1M elements
    saxpy<<<(N+511)/512, 512>>>(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);
}

在上面的帶寬計算(譯者注：即表達式N*4*3/milliseconds/1e6)中，N*4是每次數組讀或寫的字節數，因子3的含義是對x的讀以及y的讀和寫共3次讀寫操做。程序運行時間被存在變量milliseconds中，把它做爲分母便可算出單位時間的帶寬大小。注意源程序中除了添加了一些計算帶寬的功能外，咱們也改變了數組的大小和塊的大小(譯者注：因爲該代碼來自以前的博客，因此具體的變化能夠對比原來的程序，在這裏)。編譯並執行上面的代碼，咱們能夠獲得：

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

測定計算吞吐量

咱們剛剛只演示瞭如何測定帶寬，也叫作數據吞吐量。另外一種很是重要的性能指標叫作計算吞度量。一種比較通用的測量計算吞吐量的方法是計算GFLOP/s(Giga-FLoating-point OPerations per second)，表明「每秒10億次的浮點運算數」，這裏的Giga就是千兆，即10^9。對於咱們的SAXPY計算，測量有效的吞吐量是很簡單的：每一個SAXPY元素都會作一次乘法加法操做，所以是典型的2FLOPS，因此咱們能夠獲得：

$$GFLOP/{s_{Effective}} = 2N / (t * 109)$$

其中，$N$是SAXPY操做的元素個數，$t$是以秒爲單位的運行時間。就像理論峯值帶寬同樣，理論峯值$GFLOP/s$也能夠從產品資料查到(可是計算它卻很難，由於它具備架構依賴性)。例如，Tesla M2050 GPU的理論單精度浮點峯值吞吐量是$1030GFLOP/s$，而雙精度浮點峯值吞吐量是$515GFLOP/s$。SAXPY每次計算讀取12個字節，可是僅僅只有一條單獨的乘法加法指令(2 FLOPs)，因此很明顯這(數據吞吐量)就是帶寬限制。並且在這種狀況(其實是大部分狀況)下，帶寬是最重要的衡量和優化指標。在更復雜的計算中，FLOPs級別的性能測定是很困難的。所以更廣泛的方法是使用分析工具來分析計算吞吐量是不是一個瓶頸。這些應用測出的的經常是問題依賴的吞吐量(而不是架構依賴的)，這其實對用戶會更有用。例如天文學裏每秒百萬次交互做用的N體問題，或者天天納秒級的分子動態模擬。

總結

這篇文章主要介紹瞭如何用CUDA事件API獲取核函數的執行時間。CUDA事件使用GPU計時器，所以避免了與主機設備同步相關的問題。咱們也介紹了有效帶寬和計算吞吐量的性能測定方法，並且也應用這些方法測定了SAXPY例子中核函數的有效帶寬。另外咱們也得出，它的內存帶寬佔了很大比例，所以在性能測試中，計算有效吞吐量是首要的一步。在以後的文章中，咱們會進一步討論在帶寬、指令、或者延遲這些因素中，哪個是限制程序性能的因素。

CUDA事件也能夠用來計算主機和設備之間數據傳輸的速率，方法很簡單隻要將記錄事件的函數放到cudaMemcpy()調用的兩邊就能夠了。

若是你在一個很小的GPU上運行文章中的代碼，那麼若是你沒有減少數組的大小，你可能會獲得一個關於不充足設備內存的錯誤消息。實際上，咱們的實例代碼目前爲止尚未特別檢查運行時錯誤。在下一篇文章中，咱們會學習如何進行錯誤處理以及如何訪問現有設備來肯定已有資源，這樣的話咱們就能夠寫出更魯棒的代碼。