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

1 2 3 4 5 6 7 8 9 10

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

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

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是采用塊并行實現(xiàn)的向量相加操作，而addKernel_thd是采用線程并行實現(xiàn)的向量相加操作。分別運行，得到的結(jié)果如下圖所示：
線程并行：

線程塊并行：

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

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

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

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

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

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

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

總結(jié)

以上是生活随笔為你收集整理的cuda性能测试的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

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