我們在編寫完CUDA程序后，?還要從性能出發考慮問題，不斷優化代碼，使執行速度提高是并行處理的唯一目的。
?????測試代碼運行速度有很多方法，C語言里提供了類似于SystemTime()這樣的API獲得系統時間，然后計算兩個事件之間的時長從而完成計時功能。在CUDA中，我們有專門測量設備運行時間的API，下面一一介紹。
翻開編程手冊《CUDA_Toolkit_Reference_Manual》，隨時準備查詢不懂得API。我們在運行核函數前后，做如下操作：

1 2 3 4 5 6 7 8 9 10

cudaEvent_t?start,stop;//事件對象 cudaEventCreate(&start);//創建事件 cudaEventCreate(&stop);//創建事件 cudaEventRecord(start,stream);//記錄開始 myKernel<<<dimg,dimb,size_smem,stream>>>(parameter?list);//執行核函數 cudaEventRecord(stop,stream);//記錄結束事件 cudaEventSynchronize(stop);//事件同步，等待結束事件之前的設備操作均已完成 float?elapsedTime; cudaEventElapsedTime(&elapsedTime,start,stop);//計算兩個事件之間時長（單位為ms）

核函數執行時間將被保存在變量elapsedTime中。通過這個值我們可以評估算法的性能。下面給一個例子，來看怎么使用計時功能。
上面的例子規模很小，只有5個元素，處理量太小不足以計時。將規模擴大為1024，此外將反復運行1000次計算總時間，這樣估計不容易受隨機擾動影響。我們通過這個例子對比線程并行和塊并行的性能如何。代碼如下：

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 123 124 125 126 127 128 129 130 131 132 133 134 135 136 137 138 139 140

#include?"cuda_runtime.h" #include?"device_launch_parameters.h" #include?<stdio.h> cudaError_t?addWithCuda(int?*c,?const?int?*a,?const?int?*b,?size_t?size); __global__?void?addKernel_blk(int?*c,?const?int?*a,?const?int?*b) { ????int?i?=?blockIdx.x; ????c[i]?=?a[i]+?b[i]; } __global__?void?addKernel_thd(int?*c,?const?int?*a,?const?int?*b) { ????int?i?=?threadIdx.x; ????c[i]?=?a[i]+?b[i]; } int?main() { ????const?int?arraySize?=?1024; ????int?a[arraySize]?=?{0}; ????int?b[arraySize]?=?{0}; ????for(int?i?=?0;i<arraySize;i++) ????{ ????????a[i]?=?i; ????????b[i]?=?arraySize-i; ????} ????int?c[arraySize]?=?{0}; ????//?Add?vectors?in?parallel. ????cudaError_t?cudaStatus; ????int?num?=?0; ????cudaDeviceProp?prop; ????cudaStatus?=?cudaGetDeviceCount(&num); ????for(int?i?=?0;i<num;i++) ????{ ????????cudaGetDeviceProperties(&prop,i); ????} ????cudaStatus?=?addWithCuda(c,?a,?b,?arraySize); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"addWithCuda?failed!"); ????????return?1; ????} ????//?cudaThreadExit?must?be?called?before?exiting?in?order?for?profiling?and ????//?tracing?tools?such?as?Nsight?and?Visual?Profiler?to?show?complete?traces. ????cudaStatus?=?cudaThreadExit(); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaThreadExit?failed!"); ????????return?1; ????} ????for(int?i?=?0;i<arraySize;i++) ????{ ????????if(c[i]?!=?(a[i]+b[i])) ????????{ ????????????printf("Error?in?%d\n",i); ????????} ????} ????return?0; } //?Helper?function?for?using?CUDA?to?add?vectors?in?parallel. cudaError_t?addWithCuda(int?*c,?const?int?*a,?const?int?*b,?size_t?size) { ????int?*dev_a?=?0; ????int?*dev_b?=?0; ????int?*dev_c?=?0; ????cudaError_t?cudaStatus; ????//?Choose?which?GPU?to?run?on,?change?this?on?a?multi-GPU?system. ????cudaStatus?=?cudaSetDevice(0); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaSetDevice?failed!??Do?you?have?a?CUDA-capable?GPU?installed?"); ????????goto?Error; ????} ????//?Allocate?GPU?buffers?for?three?vectors?(two?input,?one?output)????. ????cudaStatus?=?cudaMalloc((void**)&dev_c,?size?*?sizeof(int)); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaMalloc?failed!"); ????????goto?Error; ????} ????cudaStatus?=?cudaMalloc((void**)&dev_a,?size?*?sizeof(int)); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaMalloc?failed!"); ????????goto?Error; ????} ????cudaStatus?=?cudaMalloc((void**)&dev_b,?size?*?sizeof(int)); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaMalloc?failed!"); ????????goto?Error; ????} ????//?Copy?input?vectors?from?host?memory?to?GPU?buffers. ????cudaStatus?=?cudaMemcpy(dev_a,?a,?size?*?sizeof(int),?cudaMemcpyHostToDevice); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaMemcpy?failed!"); ????????goto?Error; ????} ????cudaStatus?=?cudaMemcpy(dev_b,?b,?size?*?sizeof(int),?cudaMemcpyHostToDevice); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaMemcpy?failed!"); ????????goto?Error; ????} ????cudaEvent_t?start,stop; ????cudaEventCreate(&start); ????cudaEventCreate(&stop); ????cudaEventRecord(start,0); ????for(int?i?=?0;i<1000;i++) ????{ //??????? addKernel_blk<<<size,1>>>(dev_c,?dev_a,?dev_b); ????????addKernel_thd<<<1,size>>>(dev_c,?dev_a,?dev_b); ????} ????cudaEventRecord(stop,0); ????cudaEventSynchronize(stop); ????float?tm; ????cudaEventElapsedTime(&tm,start,stop); ????printf("GPU?Elapsed?time:%.6f?ms.\n",tm); ????//?cudaThreadSynchronize?waits?for?the?kernel?to?finish,?and?returns ????//?any?errors?encountered?during?the?launch. ????cudaStatus?=?cudaThreadSynchronize(); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaThreadSynchronize?returned?error?code?%d?after?launching?addKernel!\n",?cudaStatus); ????????goto?Error; ????} ????//?Copy?output?vector?from?GPU?buffer?to?host?memory. ????cudaStatus?=?cudaMemcpy(c,?dev_c,?size?*?sizeof(int),?cudaMemcpyDeviceToHost); ????if?(cudaStatus?!=?cudaSuccess)? ????{ ????????fprintf(stderr,?"cudaMemcpy?failed!"); ????????goto?Error; ????} Error: ????cudaFree(dev_c); ????cudaFree(dev_a); ????cudaFree(dev_b);???? ????return?cudaStatus; }

addKernel_blk是采用塊并行實現的向量相加操作，而addKernel_thd是采用線程并行實現的向量相加操作。分別運行，得到的結果如下圖所示：
線程并行：

線程塊并行：

可見性能竟然相差近16倍！因此選擇并行處理方法時，如果問題規模不是很大，那么采用線程并行是比較合適的，而大問題分多個線程塊處理時，每個塊內線程數不要太少，像本文中的只有1個線程，這是對硬件資源的極大浪費。一個理想的方案是，分N個線程塊，每個線程塊包含512個線程，將問題分解處理，效率往往比單一的線程并行處理或單一塊并行處理高很多。這也是CUDA編程的精髓。
這種分析程序性能的方式比較粗糙，只知道大概運行時間長度，對于設備程序各部分代碼執行時間沒有一個深入的認識，這樣我們就有個問題，如果對代碼進行優化，那么優化哪一部分呢？是將線程數調節呢，還是改用共享內存？這個問題最好的解決方案就是利用Visual?Profiler。
Visual?Profiler是一個圖形化的剖析工具，可以顯示你的應用程序中CPU和GPU的活動情況，利用分析引擎幫助你尋找優化的機會。
其實除了可視化的界面，NVIDIA提供了命令行方式的剖析命令：nvprof。
打開Visual?Profiler，可以從CUDA?Toolkit安裝菜單處找到。主界面如下：

我們點擊File->New?Session，彈出新建會話對話框，如下圖所示：

其中File一欄填入我們需要進行剖析的應用程序exe文件，后面可以都不填（如果需要命令行參數，可以在第三行填入），直接Next，見下圖：

第一行為應用程序執行超時時間設定，可不填；后面三個單選框都勾上，這樣我們分別使能了剖析，使能了并發核函數剖析，然后運行分析器。
點Finish，開始運行我們的應用程序并進行剖析、分析性能。

上圖中，CPU和GPU部分顯示了硬件和執行內容信息，點某一項則將時間條對應的部分高亮，便于觀察，同時右邊詳細信息會顯示運行時間信息。從時間條上看出，cudaMalloc占用了很大一部分時間。下面分析器給出了一些性能提升的關鍵點，包括：低計算利用率（計算時間只占總時間的1.8%，也難怪，加法計算復雜度本來就很低呀！）；低內存拷貝/計算交疊率（一點都沒有交疊，完全是拷貝——計算——拷貝）；低存儲拷貝尺寸（輸入數據量太小了，相當于你淘寶買了個日記本，運費比實物價格還高！）；低存儲拷貝吞吐率（只有1.55GB/s）。這些對我們進一步優化程序是非常有幫助的
我們點一下Details，就在Analysis窗口旁邊。得到結果如下所示：

通過這個窗口可以看到每個核函數執行時間，以及線程格、線程塊尺寸，占用寄存器個數，靜態共享內存、動態共享內存大小等參數，以及內存拷貝函數的執行情況。這個提供了比前面cudaEvent函數測時間更精確的方式，直接看到每一步的執行時間，精確到ns。
在Details后面還有一個Console，點一下看看。

個其實就是命令行窗口，顯示運行輸出。看到加入了Profiler信息后，總執行時間變長了（原來線程并行版本的程序運行時間只需4ms左右）。這也是“測不準定理”決定的，如果我們希望測量更細微的時間，那么總時間肯定是不準的；如果我們希望測量總時間，那么細微的時間就被忽略掉了。
后面Settings就是我們建立會話時的參數配置，不再詳述。
小伙伴們，可以試一下！~~

總結

以上是生活随笔為你收集整理的cuda profiler使用的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。