閱讀93 返回首頁    go 阿裏雲 go 技術社區[雲棲]

CUDA從入門到精通（十）：性能剖析和Visual Profiler

 
 

入門後的進一步學習的內容，就是如何優化自己的代碼。我們前麵的例子沒有考慮任何性能方麵優化，是為了更好地學習基本知識點，而不是其他細節問題。從本節開始，我們要從性能出發考慮問題，不斷優化代碼，使執行速度提高是並行處理的唯一目的。

測試代碼運行速度有很多方法，C語言裏提供了類似於SystemTime()這樣的API獲得係統時間，然後計算兩個事件之間的時長從而完成計時功能。在CUDA中，我們有專門測量設備運行時間的API，下麵一一介紹。

 

翻開編程手冊《CUDA_Toolkit_Reference_Manual》，隨時準備查詢不懂得API。我們在運行核函數前後，做如下操作：

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次計算總時間，這樣估計不容易受隨機擾動影響。我們通過這個例子對比線程並行和塊並行的性能如何。代碼如下：

#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。下麵內容摘自《CUDA_Profiler_Users_Guide》

“Visual Profiler是一個圖形化的剖析工具，可以顯示你的應用程序中CPU和GPU的活動情況，利用分析引擎幫助你尋找優化的機會。”

其實除了可視化的界麵，NVIDIA提供了命令行方式的剖析命令：nvprof。對於初學者，使用圖形化的方式比較容易上手，所以本節使用Visual Profiler。

 

打開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性能提升有了一些想法，好，下一節我們將討論如何優化CUDA程序。

最後更新：2017-04-03 16:48:44

  上一篇：go SQL server 2008 無法連接Local服務器的解決辦法
  下一篇：go 【轉載】[小紅豬]11個物理難題，11種基本粒子