CUDA進階第三篇：CUDA計時方式

cuda_study發表於2018-03-06
寫在前面內容分為兩部分，第一部分為翻譯《Professional CUDA C Programming》 Section 2. CUDA Programming Model中的TIMING YOUR KERNEL；第二部分為自己的經驗。經驗不足，歡迎各位大大補充。
寫CUDA，追求的就是加速比，想要得到準確的時間，計時函式就是必不可少
計時通常分為兩種情況，（1）直接得到介面函式的時間，一般用於得到加速比；（2）獲得介面函式核心函式、記憶體拷貝函式等所耗時間，一般用於優化程式碼時。
情況（1）方法有兩種，CPU計時函式和GPU計時函式。
情況（2）有三種工具nsight，nvvp，nvprof
本文會詳細介紹情況（1）的兩種方法；情況（2），nsight不會用，簡單介紹一下nvvp和nvprof的用法。
CPU計時函式
在利用CPU計時函式時，要考慮的一個問題是：核函式的執行是非同步執行的，所以必須加上核函式同步函式，才能得到準確的時間。
示例程式碼如下：
  1. double cpuSecond() {
  2.     struct timeval tp;
  3.     gettimeofday(&tp,NULL);
  4.     return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6);
  5. }
  6. double iStart = cpuSecond();
  7. function(argument list);
  8. cudaDeviceSynchronize();  // 同步函式
  9. double iElaps = cpuSecond() – iStart;

複製程式碼

GPU計時函式
GPU計時函式就不需要考慮同步問題，直接用計時事件函式就可以了，示例程式碼如下：

  1. cudaEvent_t start, stop;
  2. float elapsedTime = 0.0;
  4. cudaEventCreate(&start);
  5. cudaEventCreate(&stop);
  6. cudaEventRecord(start, 0);
  8. function(argument list);;
  10. cudaEventRecord(stop, 0);
  11. cudaEventSynchronize(stop);
  13. cudaEventElapsedTime(&elapsedTime, start, stop);
  15. cout << elapsedTime << endl;
  17. cudaEventDestroy(start);
  18. cudaEventDestroy(stop);

複製程式碼

如何獲得精確的計時正常情況下，第一次執行核函式的時間會比第二次慢一些。這是因為GPU在第一次計算時需要warmup。所以想要第一次核函式的執行時間是不精確的。獲得精確的計時我總結為三種，如下

  • 迴圈執行一百次所需要計時的部分，求平均值，將第一次的誤差縮小100倍。這種方法的優點是簡單粗暴。但缺點也很明顯：(1)程式的執行時間大大增長，特別是比較大的程式(2)要考慮記憶體溢位問題，C++的記憶體需要程式猿自己手動管理。寫出執行一次不出記憶體溢位問題的程式很容易，但是寫出迴圈執行一百次而不出記憶體溢位問題的程式碼就有一定難度了(3)計時不是特別準確，雖然誤差已經被縮小了100倍。
  • 在計時之前先執行一個warmup函式，warmup函式隨便寫，比如我從cuda sample裡的vectoradd作為warmup函式。這種方法的優點是程式執行時間縮短；缺點是需要在程式中新增一個函式，而且因為GPU亂序並行的執行方式，核函式的兩次執行時間並不能完全保持一樣。所以推薦使用方法3.
  • 先執行warmup函式，在迴圈10遍計時部分。

nvvp和nvprof的用法
nvprof是自cuda5.0開始存在的一個命令列Profiler，你可以只用nvprof來你程式碼的一些執行細節。簡單用法如下：

  1. $ nvprof ./sumArraysOnGPU-timer

複製程式碼

你就可以得到如下：

  1. ./sumArraysOnGPU-timer Starting…
  2. Using Device 0: Tesla M2070
  3. ==17770== NVPROF is profiling process 17770, command: ./sumArraysOnGPU-timer
  4. Vector size 16777216
  5. sumArraysOnGPU <<<16384, 1024>>> Time elapsed 0.003266 sec
  6. Arrays match.
  7. ……

複製程式碼

關於nvprof的更多引數資訊，可以使用幫助命令：

  1. $ nvprof –help

複製程式碼

The NVIDIA Visual Profiler（nvvp）是一款圖形化介面的Profiler，也是我一直在用的Profiler。

本文來自大光叔叔的原創，http://blog.csdn.net/litdaguang/article/details/50520549

原文釋出時間為：2016-7-8 10:35:49

原文由：十四王爺釋出，版權歸屬於原作者 

本文來自雲棲社群合作伙伴NVIDIA，瞭解相關資訊可以關注NVIDIA官方網站


相關文章