封面圖片來自互聯(lián)網(wǎng)

本文翻譯自NVIDIA官方博客Parallel Forall顾孽，內(nèi)容僅供參考，如有疑問請訪問原網(wǎng)站:https://devblogs.nvidia.com/parallelforall/how-implement-performance-metrics-cuda-cc/.

在這個系列的第一篇文章中冷冗，我們通過用CUDA C/C++實現(xiàn)SAXPY，學習了CUDA C/C++編程的基本要素惑艇。在這篇文章中蒿辙，我們會學習如何衡量這個程序以及其他CUDAC/C++程序的性能。我們在之后的文章中經(jīng)常用到這種性能度量技術(shù)滨巴，因為程序的性能優(yōu)化將會變得越來越重要思灌。

譯者注：這個系列是指原文的系列，并不是筆者的專欄恭取。

CUDA性能度量通常是在主機端進行的泰偿，我們既可以使用CPU的計時器也可以使用CUDA專門的計時器。在開始學習性能度量技術(shù)之前秽荤，我們需要討論一下如何同步主機和設(shè)備之間的操作甜奄。

主機-設(shè)備同步

讓我們來看一下上一篇博客中SAXPY的數(shù)據(jù)傳輸和核函數(shù)啟動的主機端代碼：

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進行數(shù)據(jù)傳輸?shù)姆绞绞峭絺鬏?或者是阻塞傳輸)方式柠横。同步數(shù)據(jù)傳輸直到前面所有發(fā)布的CUDA調(diào)用全部結(jié)束之后才會開始，而且同步數(shù)據(jù)傳輸結(jié)束之后课兄，隨后的CUDA調(diào)用才會開始牍氛。因此上面第三行的saxpy核函數(shù)只有到第二行的y到d_y的數(shù)據(jù)傳輸結(jié)束之后才會啟動。而在另一方面烟阐，核函數(shù)啟動卻是異步的搬俊。一旦核函數(shù)被啟動，控制權(quán)就立刻返回到CPU蜒茄，并不會等待核函數(shù)執(zhí)行完成唉擂。這樣的話就會對最后一行的設(shè)備到主機數(shù)據(jù)傳輸產(chǎn)生競態(tài)條件(race condition)，但是數(shù)據(jù)傳輸?shù)淖枞匦詴_保核函數(shù)執(zhí)行完成后再開始數(shù)據(jù)傳輸檀葛。

譯者注：這里的競態(tài)條件前面提到過玩祟，簡單說就是前面的數(shù)據(jù)操作還未完成，后面的操作卻又要使用前面的數(shù)據(jù)屿聋，這樣就會導致錯誤的結(jié)果空扎。

使用CPU的計時器來計算核函數(shù)的執(zhí)行時間

現(xiàn)在我們來看一下如何使用CPU的計時器來給核函數(shù)計時。

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

在上面的代碼中润讥，我們除了使用一般的主機時間戳函數(shù)myCPUTimer()转锈，還用到了顯式的同步障礙cudaDeviceSynchronize()來阻塞CPU執(zhí)行，直到設(shè)備上發(fā)布的指令全部執(zhí)行結(jié)束為止楚殿。如果沒有這個同步障礙撮慨，這個代碼測試的就是核函數(shù)的啟動時間而不是執(zhí)行時間。

使用CUDA事件計時

使用類似cudaDeviceSynchronize()函數(shù)的主機設(shè)備同步點的一個問題就是它會拖延GPU管道(stall GPU pipeline)脆粥∑瞿纾基于這個原因，CUDA提供了一個相比CPU計時器更輕量級的選擇冠绢，那就是使用CUDA事件API抚吠。CUDA事件API包括調(diào)用事件創(chuàng)建和銷毀函數(shù)、事件記錄函數(shù)以及以毫秒為單位計算兩個被記錄事件的運行時間的函數(shù)弟胀。

譯者注：這里拖延GPU管道(stall GPU pipeline)的直接結(jié)果就是造成CPU和GPU輪流執(zhí)行，而不再是并行執(zhí)行喊式。于是就使得程序的運行時間等于CPU與GPU時間之和孵户。具體可以參考:https://blogs.msdn.microsoft.com/shawnhar/2008/04/14/stalling-the-pipeline/

CUDA事件使用的是CUDA streams的概念。一個CUDA流只是一系列在設(shè)備上順序執(zhí)行的操作岔留。不同流中的操作可以交替執(zhí)行夏哭，在某些情況下甚至可以交疊執(zhí)行，這個特性可以被用在隱藏主機和設(shè)備間的數(shù)據(jù)傳輸献联。(我們會在之后的文章中討論)竖配。到目前為止何址，我們所有的操作都是在默認的流中進行的，或者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()進行事件的創(chuàng)建和銷毀偎血。在上面的代碼中cudaEventRecord()將事件start和stop放在默認的流中，即0號stream盯漂。函數(shù)cudaEventSynchronize()用來阻塞CPU執(zhí)行直到指定的事件被記錄颇玷。函數(shù)cudaEventElapsedTime()的第一個參數(shù)返回start和stop兩個記錄之間消逝的毫秒時間。這個值的精度大約是0.5ms就缆。

內(nèi)存帶寬

既然我們已經(jīng)可以精確地測量核函數(shù)的運行時間帖渠，那么我們就可以用它來計算帶寬。我們需要使用理論的峰值帶寬和有效內(nèi)存帶寬來評估帶寬效率竭宰。

理論帶寬

理論帶寬可以通過產(chǎn)品資料中的硬件規(guī)格來計算阿弃。例如英偉達Tesla M2050 GPU使用的是時鐘頻率為1546MHz顯存位寬為384-bit的DDR(雙倍數(shù)據(jù)速率)RAM。

使用這些數(shù)據(jù)羞延，我們可以計算出英偉達Tesla M2050的理論峰值帶寬是148 GB/sec：

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

在這個表達式中渣淳，我們將內(nèi)存的時鐘頻率的單位轉(zhuǎn)化為Hz，然后乘以顯存寬度(除以8之后伴箩，單位由比特轉(zhuǎn)化為字節(jié))入愧，又乘以2是因為該顯卡的RAM是DDR(雙倍數(shù)據(jù)速率)。最后我們將結(jié)果除以10^9得到以GB/s的計算結(jié)果嗤谚。

有效帶寬

我們是通過計算特定程序的活動時間和程序如何訪問數(shù)據(jù)來計算機有效帶寬的棺蛛。我們使用下面的公式：

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

這里， $BW_{Effective}$ 是以GB/s的有效帶寬巩步， $R_B$ 是每個核函數(shù)被讀取的字節(jié)數(shù)旁赊， $W_B$ 是每個核函數(shù)被寫入的字節(jié)數(shù)， $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是每次數(shù)組讀或?qū)懙淖止?jié)數(shù)竟闪，因子3的含義是對x的讀以及y的讀和寫共3次讀寫操作离福。程序運行時間被存在變量milliseconds中，把它作為分母即可算出單位時間的帶寬大小炼蛤。注意源程序中除了添加了一些計算帶寬的功能外妖爷，我們也改變了數(shù)組的大小和塊的大小(譯者注：由于該代碼來自之前的博客，所以具體的變化可以對比原來的程序理朋，在這里)絮识。編譯并執(zhí)行上面的代碼绿聘，我們可以得到：

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

測定計算吞吐量

我們剛剛只演示了如何測定帶寬，也叫做數(shù)據(jù)吞吐量次舌。另一種非常重要的性能指標叫做計算吞度量熄攘。一種比較通用的測量計算吞吐量的方法是計算GFLOP/s(Giga-FLoating-point OPerations per second)，代表“每秒10億次的浮點運算數(shù)”垃它，這里的Giga就是千兆鲜屏，即10^9。對于我們的SAXPY計算国拇，測量有效的吞吐量是很簡單的：每個SAXPY元素都會做一次乘法加法操作洛史，因此是典型的2FLOPS，所以我們可以得到：

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

其中酱吝， $N$ 是SAXPY操作的元素個數(shù)也殖， $t$ 是以秒為單位的運行時間。就像理論峰值帶寬一樣务热，理論峰值 $GFLOP/s$ 也可以從產(chǎn)品資料查到(但是計算它卻很難忆嗜，因為它具有架構(gòu)依賴性)。例如崎岂，Tesla M2050 GPU的理論單精度浮點峰值吞吐量是 $1030GFLOP/s$ 捆毫，而雙精度浮點峰值吞吐量是 $515GFLOP/s$ 。SAXPY每次計算讀取12個字節(jié)冲甘，但是僅僅只有一條單獨的乘法加法指令(2 FLOPs)绩卤，所以很明顯這(數(shù)據(jù)吞吐量)就是帶寬限制。而且在這種情況(實際上是大部分情況)下江醇，帶寬是最重要的衡量和優(yōu)化指標濒憋。在更復(fù)雜的計算中，F(xiàn)LOPs級別的性能測定是很困難的陶夜。因此更普遍的方法是使用分析工具來分析計算吞吐量是否是一個瓶頸凛驮。這些應(yīng)用測出的的常常是問題依賴的吞吐量(而不是架構(gòu)依賴的)，這其實對用戶會更有用条辟。例如天文學里每秒百萬次交互作用的N體問題黔夭，或者每天納秒級的分子動態(tài)模擬。

總結(jié)

這篇文章主要介紹了如何用CUDA事件API獲取核函數(shù)的執(zhí)行時間捂贿。CUDA事件使用GPU計時器纠修，因此避免了與主機設(shè)備同步相關(guān)的問題。我們也介紹了有效帶寬和計算吞吐量的性能測定方法厂僧，而且也應(yīng)用這些方法測定了SAXPY例子中核函數(shù)的有效帶寬。另外我們也得出了牛，它的內(nèi)存帶寬占了很大比例颜屠，因此在性能測試中辰妙，計算有效吞吐量是首要的一步。在之后的文章中甫窟，我們會進一步討論在帶寬密浑、指令、或者延遲這些因素中粗井，哪一個是限制程序性能的因素尔破。

CUDA事件也可以用來計算主機和設(shè)備之間數(shù)據(jù)傳輸?shù)乃俾剩椒ê芎唵沃灰獙⒂涗浭录暮瘮?shù)放到cudaMemcpy()調(diào)用的兩邊就可以了浇衬。

如果你在一個很小的GPU上運行文章中的代碼懒构，那么如果你沒有減小數(shù)組的大小，你可能會得到一個關(guān)于不充足設(shè)備內(nèi)存的錯誤消息耘擂。實際上胆剧，我們的實例代碼目前為止還沒有特別檢查運行時錯誤。在下一篇文章中醉冤，我們會學習如何進行錯誤處理以及如何訪問現(xiàn)有設(shè)備來確定已有資源秩霍，這樣的話我們就可以寫出更魯棒的代碼。

[譯]在CUDA C/C++中如何測試代碼性能