• cuda
• c
• c++

Fighting_Bird 2016年11月21日發(fā)布

• 
  

• 1 推薦
• 
  
  • 4 收藏，667 瀏覽

  本文翻譯自NVIDIA官方博客Parallel Forall，內(nèi)容僅供參考，如有疑問請訪問原網(wǎng)站:https://devblogs.nvidia.com/p....

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

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

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

  主機(jī)-設(shè)備同步

  讓我們來看一下上一篇博客中SAXPY的數(shù)據(jù)傳輸和核函數(shù)啟動的主機(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é)束之后才會開始，而且同步數(shù)據(jù)傳輸結(jié)束之后，隨后的CUDA調(diào)用才會開始。因此上面第三行的saxpy核函數(shù)只有到第二行的y到d_y的數(shù)據(jù)傳輸結(jié)束之后才會啟動。而在另一方面，核函數(shù)啟動卻是異步的。一旦核函數(shù)被啟動，控制權(quán)就立刻返回到CPU，并不會等待核函數(shù)執(zhí)行完成。這樣的話就會對最后一行的設(shè)備到主機(jī)數(shù)據(jù)傳輸產(chǎn)生競態(tài)條件(race condition)，但是數(shù)據(jù)傳輸?shù)淖枞匦詴_保核函數(shù)執(zhí)行完成后再開始數(shù)據(jù)傳輸。

  譯者注：這里的競態(tài)條件前面提到過，簡單說就是前面的數(shù)據(jù)操作還未完成，后面的操作卻又要使用前面的數(shù)據(jù)，這樣就會導(dǎo)致錯誤的結(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);

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

  使用CUDA事件計時

  使用類似cudaDeviceSynchronize()函數(shù)的主機(jī)設(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....

  CUDA事件使用的是CUDA streams的概念。一個CUDA流只是一系列在設(shè)備上順序執(zhí)行的操作。不同流中的操作可以交替執(zhí)行，在某些情況下甚至可以交疊執(zhí)行，這個特性可以被用在隱藏主機(jī)和設(shè)備間的數(shù)據(jù)傳輸。(我們會在之后的文章中討論)。到目前為止，我們所有的操作都是在默認(rèn)的流中進(jìn)行的，或者0號流(也叫做空流)。

  下面的代碼中，我們使用了CUDA事件API來對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類型，通過cudaEventCreate()和cudaEventDestroy()進(jìn)行事件的創(chuàng)建和銷毀。在上面的代碼中cudaEventRecord()將事件start和stop放在默認(rèn)的流中，即0號stream。函數(shù)cudaEventSynchronize()用來阻塞CPU執(zhí)行直到指定的事件被記錄。函數(shù) cudaEventElapsedTime()的第一個參數(shù)返回start和stop兩個記錄之間消逝的毫秒時間。這個值的精度大約是0.5ms。

  內(nèi)存帶寬

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

  理論帶寬

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

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

  BWTheoretical=1546?106?(384/8)?2/109=148GB/s

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

  有效帶寬

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

  BWEffective=(RB+WB)/(t?109)

  這里，BWEffective是以GB/s的有效帶寬，RB是每個核函數(shù)被讀取的字節(jié)數(shù)，WB是每個核函數(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 elementssaxpy<<<(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); }

  在上面的帶寬計算(譯者注：即表達(dá)式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ù)吞吐量。另一種非常重要的性能指標(biāo)叫做計算吞度量。一種比較通用的測量計算吞吐量的方法是計算GFLOP/s(Giga-FLoating-point OPerations per second)，代表“每秒10億次的浮點運算數(shù)”，這里的Giga就是千兆，即10^9。對于我們的SAXPY計算，測量有效的吞吐量是很簡單的：每個SAXPY元素都會做一次乘法加法操作，因此是典型的2FLOPS，所以我們可以得到：

  GFLOP/sEffective=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)化指標(biāo)。在更復(fù)雜的計算中，FLOPs級別的性能測定是很困難的。因此更普遍的方法是使用分析工具來分析計算吞吐量是否是一個瓶頸。這些應(yīng)用測出的的常常是問題依賴的吞吐量(而不是架構(gòu)依賴的)，這其實對用戶會更有用。例如天文學(xué)里每秒百萬次交互作用的N體問題，或者每天納秒級的分子動態(tài)模擬。

  總結(jié)

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

  CUDA事件也可以用來計算主機(jī)和設(shè)備之間數(shù)據(jù)傳輸?shù)乃俾?#xff0c;方法很簡單只要將記錄事件的函數(shù)放到cudaMemcpy()調(diào)用的兩邊就可以了。

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

總結(jié)

以上是生活随笔為你收集整理的[译]在CUDA C/C++中如何衡量代码性能的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。