在 本系列文章的第一篇 中，我們通過檢查 CUDA C/C++ SAXPY 來研究 CUDA C / C ++的基本元素。在第二篇文章中，我們將討論如何分析這個(gè)和其他 CUDA C / C ++代碼的性能。我們將依賴于這些性能測量技術(shù)在未來的職位，性能優(yōu)化將變得越來越重要。

CUDA 性能度量通常是從主機(jī)代碼中完成的，可以使用 CPU 計(jì)時(shí)器或 CUDA 特定計(jì)時(shí)器來實(shí)現(xiàn)。在討論這些性能度量技術(shù)之前，我們需要討論如何在主機(jī)和設(shè)備之間同步執(zhí)行。

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

讓我們看看數(shù)據(jù)傳輸和來自上一篇文章的 SAXPY 主機(jī)代碼的內(nèi)核啟動(dòng)：

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ī)和設(shè)備之間的數(shù)據(jù)傳輸是synchronous（或blocking）傳輸。同步數(shù)據(jù)傳輸在之前發(fā)出的所有 CUDA 調(diào)用完成之前不會(huì)開始，后續(xù)的 CUDA 調(diào)用在同步傳輸完成之前無法開始。因此，第三行的saxpy內(nèi)核啟動(dòng)在第二行從y到d_y的傳輸完成后才會(huì)發(fā)出。另一方面，內(nèi)核啟動(dòng)是異步的。一旦內(nèi)核在第三行啟動(dòng)，控制權(quán)立即返回到 CPU ，而不是等待內(nèi)核完成。而 MIG ht 似乎為設(shè)備在最后一行主機(jī)數(shù)據(jù)傳輸設(shè)置了一個(gè)競爭條件，數(shù)據(jù)傳輸?shù)淖枞再|(zhì)確保了內(nèi)核在傳輸開始之前完成。

用 CPU 計(jì)時(shí)器計(jì)時(shí)內(nèi)核執(zhí)行

現(xiàn)在讓我們來看看如何使用 CPU 計(jì)時(shí)器為內(nèi)核執(zhí)行計(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);

除了對(duì)通用主機(jī)時(shí)間戳函數(shù)myCPUTimer()的兩次調(diào)用外，我們還使用顯式同步屏障cudaDeviceSynchronize()來阻止 CPU 的執(zhí)行，直到設(shè)備上以前發(fā)出的所有命令都已完成。如果沒有這個(gè)屏障，這段代碼將測量內(nèi)核發(fā)射時(shí)間，而不是內(nèi)核執(zhí)行時(shí)間。

使用 CUDA 事件計(jì)時(shí)

使用主機(jī)設(shè)備同步點(diǎn)（如cudaDeviceSynchronize()的一個(gè)問題是它們會(huì)暫停 GPU 管道。因此， CUDA 通過CUDA 事件 API為 CPU 定時(shí)器提供了一個(gè)相對(duì)輕量級(jí)的替代方案。 CUDA 事件 API 包括在兩個(gè)記錄的事件之間調(diào)用create和破壞事件、record事件和以毫秒為單位計(jì)算已用時(shí)間。

CUDA 事件利用?CUDA?streams. CUDA 流只是按順序在設(shè)備上執(zhí)行的操作序列。在某些情況下[vx3 . 4 可以交叉使用 vx3 . 4]的流。到目前為止， GPU 上的所有操作都發(fā)生在默認(rèn)流或流 0 （也稱為“空流”）中。

在下面的清單中，我們將 CUDA 事件應(yīng)用于 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()將啟動(dòng)和停止事件放入默認(rèn)流 stream 0 。當(dāng)事件到達(dá)流中的事件時(shí)，設(shè)備將記錄事件的時(shí)間戳。函數(shù)cudaEventSynchronize()會(huì)阻止 CPU 的執(zhí)行，直到記錄指定的事件為止。cudaEventElapsedTime()函數(shù)在第一個(gè)參數(shù)中返回錄制start和stop之間經(jīng)過的毫秒數(shù)。該值的分辨率約為半微秒。

內(nèi)存帶寬

現(xiàn)在我們有了一種精確計(jì)時(shí)內(nèi)核執(zhí)行的方法，我們將使用它來計(jì)算帶寬。在評(píng)估帶寬效率時(shí)，我們同時(shí)使用理論峰值帶寬和觀察到的或有效的內(nèi)存帶寬。

理論帶寬

理論帶寬可以使用產(chǎn)品文獻(xiàn)中提供的硬件規(guī)格計(jì)算。例如， NVIDIA Tesla M2050 GPU 使用內(nèi)存時(shí)鐘速率為 1546 MHz 的 DDR （雙數(shù)據(jù)速率） RAM 和 384 位寬的內(nèi)存接口。使用這些數(shù)據(jù)項(xiàng)， NVIDIA Tesla M2050 的峰值理論內(nèi)存帶寬為 148 GB / s ，如下所示。

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

在這個(gè)計(jì)算中，我們將內(nèi)存時(shí)鐘速率轉(zhuǎn)換為赫茲，乘以接口寬度（除以 8 ，將位轉(zhuǎn)換為字節(jié)），再乘以 2 ，這是由于數(shù)據(jù)速率加倍。最后，我們除以 109將結(jié)果轉(zhuǎn)換為 GB / s 。

有效帶寬

我們通過計(jì)時(shí)特定的程序活動(dòng)和了解程序如何訪問數(shù)據(jù)來計(jì)算有效帶寬。我們用下面的等式。

BWEffective=（RB+WB（ VZX50]* 109)

這里，BWEffective有效帶寬，單位為 GB / s ，RB是每個(gè)內(nèi)核讀取的字節(jié)數(shù)，WB是每個(gè)內(nèi)核寫入的字節(jié)數(shù)，t是以秒為單位的運(yùn)行時(shí)間。下面是完整的代碼。

#include



__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: %fn", maxError);

  printf("Effective Bandwidth (GB/s): %fn", N*4*3/milliseconds/1e6);

}

在帶寬計(jì)算中，N*4是每個(gè)數(shù)組讀或?qū)憘鬏數(shù)淖止?jié)數(shù)， 3 的因子表示x的讀取和y的讀寫。經(jīng)過的時(shí)間存儲(chǔ)在變量milliseconds中，以明確單位。請(qǐng)注意，除了添加帶寬計(jì)算所需的功能外，我們還更改了數(shù)組大小和線程塊大小。在 Tesla M2050 上編譯并運(yùn)行此代碼：

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

測量計(jì)算吞吐量

我們剛剛演示了如何測量帶寬，帶寬是數(shù)據(jù)吞吐量的度量。另一個(gè)對(duì)性能非常重要的指標(biāo)是計(jì)算吞吐量。計(jì)算吞吐量的常用度量是 GFLOP / s ，它代表“每秒千兆浮點(diǎn)運(yùn)算”，其中 Giga 是 10 的前綴9. 我們通常測量 SAXPY 的吞吐量，因?yàn)槊恳粋€(gè) SAXPY 運(yùn)算都是有效的

GFLOP/s Effective== 2 N /（ t ：《* 109）

N 是 SAXPY 操作中的元素?cái)?shù)， t 是以秒為單位的運(yùn)行時(shí)間。與理論峰值帶寬一樣，理論峰值 GFLOP / s 可以從產(chǎn)品文獻(xiàn)中獲得（但是計(jì)算它可能有點(diǎn)棘手，因?yàn)樗c體系結(jié)構(gòu)非常相關(guān)）。例如， Tesla M2050 GPU 的單精度浮點(diǎn)吞吐量理論峰值為 1030 GFLOP / s ，雙倍精度的理論峰值吞吐量為 515 GFLOP / s 。

SAXPY 為計(jì)算的每個(gè)元素讀取 12 個(gè)字節(jié)，但是只執(zhí)行一個(gè)乘法加法指令（ 2 個(gè)浮點(diǎn)運(yùn)算），因此很明顯它是帶寬受限的，因此在這種情況下（實(shí)際上在許多情況下），帶寬是衡量和優(yōu)化的最重要的指標(biāo)。在更復(fù)雜的計(jì)算中，在 FLOPs 級(jí)別測量性能可能非常困難。因此，更常見的是使用分析工具來了解計(jì)算吞吐量是否是一個(gè)瓶頸。應(yīng)用程序通常提供特定于問題（而不是特定于體系結(jié)構(gòu)）的吞吐量指標(biāo)，因此對(duì)用戶更有用。例如，天文 n 體問題的“每秒十億次相互作用”，或分子動(dòng)力學(xué)模擬的“每天納秒”。

總結(jié)

這篇文章描述了如何使用 CUDA 事件 API 為內(nèi)核執(zhí)行計(jì)時(shí)。 CUDA 事件使用 GPU 計(jì)時(shí)器，因此避免了與主機(jī)設(shè)備同步相關(guān)的問題。我們提出了有效帶寬和計(jì)算吞吐量性能指標(biāo)，并在 SAXPY 內(nèi)核中實(shí)現(xiàn)了有效帶寬。很大一部分內(nèi)核是內(nèi)存帶寬限制的，因此計(jì)算有效帶寬是性能優(yōu)化的第一步。在以后的文章中，我們將討論如何確定帶寬、指令或延遲是性能的限制因素。

CUDA 事件還可以用于確定主機(jī)和設(shè)備之間的數(shù)據(jù)傳輸速率，方法是在 cudaMemcpy（） 調(diào)用的任一側(cè)記錄事件。

如果你在這個(gè)設(shè)備上運(yùn)行一個(gè)關(guān)于內(nèi)存不足的錯(cuò)誤［ZC9］，你可能會(huì)得到一個(gè)更小的錯(cuò)誤。實(shí)際上，到目前為止，我們的示例代碼還沒有費(fèi)心檢查運(yùn)行時(shí)錯(cuò)誤。在［VZX337］中，我們將學(xué)習(xí)如何在 CUDA C / C ++中執(zhí)行錯(cuò)誤處理以及如何查詢當(dāng)前設(shè)備以確定它們可用的資源，以便我們可以編寫更健壯的代碼。

關(guān)于作者

Mark Harris 是 NVIDIA 杰出的工程師，致力于 RAPIDS 。 Mark 擁有超過 20 年的 GPUs 軟件開發(fā)經(jīng)驗(yàn)，從圖形和游戲到基于物理的模擬，到并行算法和高性能計(jì)算。當(dāng)他還是北卡羅來納大學(xué)的博士生時(shí)，他意識(shí)到了一種新生的趨勢，并為此創(chuàng)造了一個(gè)名字： GPGPU （圖形處理單元上的通用計(jì)算）。

審核編輯：郭婷