本文翻譯自NVIDIA官方博客Parallel Forall，內(nèi)容僅供參考，如有疑問(wèn)請(qǐng)?jiān)L問(wèn)原網(wǎng)站:https://devblogs.nvidia.com/parallelforall/how-implement-performance-metrics-cuda-cc/.
在這個(gè)系列的第一篇文章中，我們通過(guò)用CUDA C/C++實(shí)現(xiàn)SAXPY，學(xué)習(xí)了CUDA C/C++編程的基本要素。在這篇文章中，我們會(huì)學(xué)習(xí)如何衡量這個(gè)程序以及其他CUDAC/C++程序的性能。我們?cè)谥蟮奈恼轮薪?jīng)常用到這種性能度量技術(shù)，因?yàn)槌绦虻男阅軆?yōu)化將會(huì)變得越來(lái)越重要。
譯者注：這個(gè)系列是指原文的系列，并不是筆者的專(zhuān)欄。
CUDA性能度量通常是在主機(jī)端進(jìn)行的，我們既可以使用CPU的計(jì)時(shí)器也可以使用CUDA專(zhuān)門(mén)的計(jì)時(shí)器。在開(kāi)始學(xué)習(xí)性能度量技術(shù)之前，我們需要討論一下如何同步主機(jī)和設(shè)備之間的操作。
主機(jī)-設(shè)備同步
讓我們來(lái)看一下上一篇博客中SAXPY的數(shù)據(jù)傳輸和核函數(shù)啟動(dòng)的主機(jī)端代碼：
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進(jìn)行數(shù)據(jù)傳輸?shù)姆绞绞峭絺鬏?或者是阻塞傳輸)方式。同步數(shù)據(jù)傳輸直到前面所有發(fā)布的CUDA調(diào)用全部結(jié)束之后才會(huì)開(kāi)始，而且同步數(shù)據(jù)傳輸結(jié)束之后，隨后的CUDA調(diào)用才會(huì)開(kāi)始。因此上面第三行的saxpy核函數(shù)只有到第二行的y到d_y的數(shù)據(jù)傳輸結(jié)束之后才會(huì)啟動(dòng)。而在另一方面，核函數(shù)啟動(dòng)卻是異步的。一旦核函數(shù)被啟動(dòng)，控制權(quán)就立刻返回到CPU，并不會(huì)等待核函數(shù)執(zhí)行完成。這樣的話就會(huì)對(duì)最后一行的設(shè)備到主機(jī)數(shù)據(jù)傳輸產(chǎn)生競(jìng)態(tài)條件(race condition)，但是數(shù)據(jù)傳輸?shù)淖枞匦詴?huì)確保核函數(shù)執(zhí)行完成后再開(kāi)始數(shù)據(jù)傳輸。
譯者注：這里的競(jìng)態(tài)條件前面提到過(guò)，簡(jiǎn)單說(shuō)就是前面的數(shù)據(jù)操作還未完成，后面的操作卻又要使用前面的數(shù)據(jù)，這樣就會(huì)導(dǎo)致錯(cuò)誤的結(jié)果。
使用CPU的計(jì)時(shí)器來(lái)計(jì)算核函數(shù)的執(zhí)行時(shí)間
現(xiàn)在我們來(lái)看一下如何使用CPU的計(jì)時(shí)器來(lái)給核函數(shù)計(jì)時(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);
在上面的代碼中，我們除了使用一般的主機(jī)時(shí)間戳函數(shù)myCPUTimer()，還用到了顯式的同步障礙cudaDeviceSynchronize()來(lái)阻塞CPU執(zhí)行，直到設(shè)備上發(fā)布的指令全部執(zhí)行結(jié)束為止。如果沒(méi)有這個(gè)同步障礙，這個(gè)代碼測(cè)試的就是核函數(shù)的啟動(dòng)時(shí)間而不是執(zhí)行時(shí)間。
使用CUDA事件計(jì)時(shí)
使用類(lèi)似cudaDeviceSynchronize()函數(shù)的主機(jī)設(shè)備同步點(diǎn)的一個(gè)問(wèn)題就是它會(huì)拖延GPU管道(stall GPU pipeline)。基于這個(gè)原因，CUDA提供了一個(gè)相比CPU計(jì)時(shí)器更輕量級(jí)的選擇，那就是使用CUDA事件API。CUDA事件API包括調(diào)用事件創(chuàng)建和銷(xiāo)毀函數(shù)、事件記錄函數(shù)以及以毫秒為單位計(jì)算兩個(gè)被記錄事件的運(yùn)行時(shí)間的函數(shù)。
譯者注：這里拖延GPU管道(stall GPU pipeline)的直接結(jié)果就是造成CPU和GPU輪流執(zhí)行，而不再是并行執(zhí)行。于是就使得程序的運(yùn)行時(shí)間等于CPU與GPU時(shí)間之和。具體可以參考:https://blogs.msdn.microsoft.com/shawnhar/2008/04/14/stalling-the-pipeline/
CUDA事件使用的是CUDA streams的概念。一個(gè)CUDA流只是一系列在設(shè)備上順序執(zhí)行的操作。不同流中的操作可以交替執(zhí)行，在某些情況下甚至可以交疊執(zhí)行，這個(gè)特性可以被用在隱藏主機(jī)和設(shè)備間的數(shù)據(jù)傳輸。(我們會(huì)在之后的文章中討論)。到目前為止，我們所有的操作都是在默認(rèn)的流中進(jìn)行的，或者0號(hào)流(也叫做空流)。
下面的代碼中，我們使用了CUDA事件API來(lái)對(duì)SAXPY代碼進(jìn)行性能度量。
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類(lèi)型，通過(guò)cudaEventCreate()和cudaEventDestroy()進(jìn)行事件的創(chuàng)建和銷(xiāo)毀。在上面的代碼中cudaEventRecord()將事件start和stop放在默認(rèn)的流中，即0號(hào)stream。函數(shù)cudaEventSynchronize()用來(lái)阻塞CPU執(zhí)行直到指定的事件被記錄。函數(shù)cudaEventElapsedTime()的第一個(gè)參數(shù)返回start和stop兩個(gè)記錄之間消逝的毫秒時(shí)間。這個(gè)值的精度大約是0.5ms。
內(nèi)存帶寬
既然我們已經(jīng)可以精確地測(cè)量核函數(shù)的運(yùn)行時(shí)間，那么我們就可以用它來(lái)計(jì)算帶寬。我們需要使用理論的峰值帶寬和有效內(nèi)存帶寬來(lái)評(píng)估帶寬效率。
理論帶寬
理論帶寬可以通過(guò)產(chǎn)品資料中的硬件規(guī)格來(lái)計(jì)算。例如英偉達(dá)Tesla M2050 GPU使用的是時(shí)鐘頻率為1546MHz顯存位寬為384-bit的DDR(雙倍數(shù)據(jù)速率)RAM。
使用這些數(shù)據(jù)，我們可以計(jì)算出英偉達(dá)Tesla M2050的理論峰值帶寬是148 GB/sec：
在這個(gè)表達(dá)式中，我們將內(nèi)存的時(shí)鐘頻率的單位轉(zhuǎn)化為Hz，然后乘以顯存寬度(除以8之后，單位由比特轉(zhuǎn)化為字節(jié))，又乘以2是因?yàn)樵擄@卡的RAM是DDR(雙倍數(shù)據(jù)速率)。最后我們將結(jié)果除以10^9得到以GB/s的計(jì)算結(jié)果。
有效帶寬
我們是通過(guò)計(jì)算特定程序的活動(dòng)時(shí)間和程序如何訪問(wèn)數(shù)據(jù)來(lái)計(jì)算機(jī)有效帶寬的。我們使用下面的公式：
這里，是以GB/s的有效帶寬，
是每個(gè)核函數(shù)被讀取的字節(jié)數(shù)，
是每個(gè)核函數(shù)被寫(xiě)入的字節(jié)數(shù)，
是以秒為單位的運(yùn)行時(shí)間。我們可以修改SAXPY例子來(lái)計(jì)算有效帶寬，下面是完整的代碼：
#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);
}
在上面的帶寬計(jì)算(譯者注：即表達(dá)式N*4*3/milliseconds/1e6)中，N*4是每次數(shù)組讀或?qū)懙淖止?jié)數(shù)，因子3的含義是對(duì)x的讀以及y的讀和寫(xiě)共3次讀寫(xiě)操作。程序運(yùn)行時(shí)間被存在變量milliseconds中，把它作為分母即可算出單位時(shí)間的帶寬大小。注意源程序中除了添加了一些計(jì)算帶寬的功能外，我們也改變了數(shù)組的大小和塊的大小(譯者注：由于該代碼來(lái)自之前的博客，所以具體的變化可以對(duì)比原來(lái)的程序，在這里)。編譯并執(zhí)行上面的代碼，我們可以得到：
$ ./saxpy
Max error: 0.000000
Effective Bandwidth (GB/s): 110.374872
測(cè)定計(jì)算吞吐量
我們剛剛只演示了如何測(cè)定帶寬，也叫做數(shù)據(jù)吞吐量。另一種非常重要的性能指標(biāo)叫做計(jì)算吞度量。一種比較通用的測(cè)量計(jì)算吞吐量的方法是計(jì)算GFLOP/s(Giga-FLoating-point OPerations per second)，代表“每秒10億次的浮點(diǎn)運(yùn)算數(shù)”，這里的Giga就是千兆，即10^9。對(duì)于我們的SAXPY計(jì)算，測(cè)量有效的吞吐量是很簡(jiǎn)單的：每個(gè)SAXPY元素都會(huì)做一次乘法加法操作，因此是典型的2FLOPS，所以我們可以得到：
其中，是SAXPY操作的元素個(gè)數(shù)，
是以秒為單位的運(yùn)行時(shí)間。就像理論峰值帶寬一樣，理論峰值
也可以從產(chǎn)品資料查到(但是計(jì)算它卻很難，因?yàn)樗哂屑軜?gòu)依賴(lài)性)。例如，Tesla M2050 GPU的理論單精度浮點(diǎn)峰值吞吐量是
，而雙精度浮點(diǎn)峰值吞吐量是
。SAXPY每次計(jì)算讀取12個(gè)字節(jié)，但是僅僅只有一條單獨(dú)的乘法加法指令(2 FLOPs)，所以很明顯這(數(shù)據(jù)吞吐量)就是帶寬限制。而且在這種情況(實(shí)際上是大部分情況)下，帶寬是最重要的衡量和優(yōu)化指標(biāo)。在更復(fù)雜的計(jì)算中，F(xiàn)LOPs級(jí)別的性能測(cè)定是很困難的。因此更普遍的方法是使用分析工具來(lái)分析計(jì)算吞吐量是否是一個(gè)瓶頸。這些應(yīng)用測(cè)出的的常常是問(wèn)題依賴(lài)的吞吐量(而不是架構(gòu)依賴(lài)的)，這其實(shí)對(duì)用戶(hù)會(huì)更有用。例如天文學(xué)里每秒百萬(wàn)次交互作用的N體問(wèn)題，或者每天納秒級(jí)的分子動(dòng)態(tài)模擬。
總結(jié)
這篇文章主要介紹了如何用CUDA事件API獲取核函數(shù)的執(zhí)行時(shí)間。CUDA事件使用GPU計(jì)時(shí)器，因此避免了與主機(jī)設(shè)備同步相關(guān)的問(wèn)題。我們也介紹了有效帶寬和計(jì)算吞吐量的性能測(cè)定方法，而且也應(yīng)用這些方法測(cè)定了SAXPY例子中核函數(shù)的有效帶寬。另外我們也得出，它的內(nèi)存帶寬占了很大比例，因此在性能測(cè)試中，計(jì)算有效吞吐量是首要的一步。在之后的文章中，我們會(huì)進(jìn)一步討論在帶寬、指令、或者延遲這些因素中，哪一個(gè)是限制程序性能的因素。
CUDA事件也可以用來(lái)計(jì)算主機(jī)和設(shè)備之間數(shù)據(jù)傳輸?shù)乃俾剩椒ê芎?jiǎn)單只要將記錄事件的函數(shù)放到cudaMemcpy()調(diào)用的兩邊就可以了。
如果你在一個(gè)很小的GPU上運(yùn)行文章中的代碼，那么如果你沒(méi)有減小數(shù)組的大小，你可能會(huì)得到一個(gè)關(guān)于不充足設(shè)備內(nèi)存的錯(cuò)誤消息。實(shí)際上，我們的實(shí)例代碼目前為止還沒(méi)有特別檢查運(yùn)行時(shí)錯(cuò)誤。在下一篇文章中，我們會(huì)學(xué)習(xí)如何進(jìn)行錯(cuò)誤處理以及如何訪問(wèn)現(xiàn)有設(shè)備來(lái)確定已有資源，這樣的話我們就可以寫(xiě)出更魯棒的代碼。
