在 本系列文章的第一篇 中，我們通過檢查 CUDA C/C++ SAXPY 來研究 CUDA C / C ++的基本元素。在第二篇文章中，我們將討論如何分析這個和其他 CUDA C / 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()在主機和設備之間的數據傳輸是synchronous（或blocking）傳輸。同步數據傳輸在之前發出的所有 CUDA 調用完成之前不會開始，后續的 CUDA 調用在同步傳輸完成之前無法開始。因此，第三行的saxpy內核啟動在第二行從y到d_y的傳輸完成后才會發出。另一方面，內核啟動是異步的。一旦內核在第三行啟動，控制權立即返回到 CPU ，而不是等待內核完成。而 MIG ht 似乎為設備在最后一行主機數據傳輸設置了一個競爭條件，數據傳輸的阻塞性質確保了內核在傳輸開始之前完成。

用 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 管道。因此， CUDA 通過CUDA 事件 API為 CPU 定時器提供了一個相對輕量級的替代方案。 CUDA 事件 API 包括在兩個記錄的事件之間調用create和破壞事件、record事件和以毫秒為單位計算已用時間。

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

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

內存帶寬

現在我們有了一種精確計時內核執行的方法，我們將使用它來計算帶寬。在評估帶寬效率時，我們同時使用理論峰值帶寬和觀察到的或有效的內存帶寬。

理論帶寬

理論帶寬可以使用產品文獻中提供的硬件規格計算。例如， NVIDIA Tesla M2050 GPU 使用內存時鐘速率為 1546 MHz 的 DDR （雙數據速率） RAM 和 384 位寬的內存接口。使用這些數據項， NVIDIA Tesla M2050 的峰值理論內存帶寬為 148 GB / s ，如下所示。

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

在這個計算中，我們將內存時鐘速率轉換為赫茲，乘以接口寬度（除以 8 ，將位轉換為字節），再乘以 2 ，這是由于數據速率加倍。最后，我們除以 109將結果轉換為 GB / s 。

有效帶寬

我們通過計時特定的程序活動和了解程序如何訪問數據來計算有效帶寬。我們用下面的等式。

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

這里，BWEffective有效帶寬，單位為 GB / s ，RB是每個內核讀取的字節數，WB是每個內核寫入的字節數，t是以秒為單位的運行時間。下面是完整的代碼。

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

}

在帶寬計算中，N*4是每個數組讀或寫傳輸的字節數， 3 的因子表示x的讀取和y的讀寫。經過的時間存儲在變量milliseconds中，以明確單位。請注意，除了添加帶寬計算所需的功能外，我們還更改了數組大小和線程塊大小。在 Tesla M2050 上編譯并運行此代碼：

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

測量計算吞吐量

我們剛剛演示了如何測量帶寬，帶寬是數據吞吐量的度量。另一個對性能非常重要的指標是計算吞吐量。計算吞吐量的常用度量是 GFLOP / s ，它代表“每秒千兆浮點運算”，其中 Giga 是 10 的前綴9. 我們通常測量 SAXPY 的吞吐量，因為每一個 SAXPY 運算都是有效的

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

N 是 SAXPY 操作中的元素數， t 是以秒為單位的運行時間。與理論峰值帶寬一樣，理論峰值 GFLOP / s 可以從產品文獻中獲得（但是計算它可能有點棘手，因為它與體系結構非常相關）。例如， Tesla M2050 GPU 的單精度浮點吞吐量理論峰值為 1030 GFLOP / s ，雙倍精度的理論峰值吞吐量為 515 GFLOP / s 。

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

總結

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

CUDA 事件還可以用于確定主機和設備之間的數據傳輸速率，方法是在 cudaMemcpy（） 調用的任一側記錄事件。

如果你在這個設備上運行一個關于內存不足的錯誤［ZC9］，你可能會得到一個更小的錯誤。實際上，到目前為止，我們的示例代碼還沒有費心檢查運行時錯誤。在［VZX337］中，我們將學習如何在 CUDA C / C ++中執行錯誤處理以及如何查詢當前設備以確定它們可用的資源，以便我們可以編寫更健壯的代碼。

關于作者

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

審核編輯：郭婷