CUDA性能測量與錯誤處理

• 討論如何通過CUDA事件來測量它的性能
• 如何通過CUDA代碼進行調試

1.測量CUDA程序的性能

1.1 CUDA事件

• CPU端的計時器可能無法給出正確的內核執行時間
• CUDA事件等于是在你的CUDA應用運行的特定時刻被記錄的時間戳，通過使用CUDA事件API，由GPU來記錄這個時間戳
• 使用CUDA測量時間需要兩個步驟：創建事件和記錄事件，記錄事件（開始時間與結束時間）
• 代碼如下：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in Array
#define N	50000
//Defining Kernel function for vector addition
__global__ void gpuAdd(int* d_a, int* d_b, int* d_c) {//Getting Thread index of current kernelint tid = threadIdx.x + blockIdx.x * blockDim.x;while (tid < N){d_c[tid] = d_a[tid] + d_b[tid];tid += blockDim.x * gridDim.x;}}int main(void) {//Defining host arraysint h_a[N], h_b[N], h_c[N];//Defining device pointersint* d_a, * d_b, * d_c;//----------創建事件記錄起止時間---------------------cudaEvent_t e_start, e_stop;cudaEventCreate(&e_start);cudaEventCreate(&e_stop);//第一次記錄時間戳cudaEventRecord(e_start, 0);// allocate the memorycudaMalloc((void**)&d_a, N * sizeof(int));cudaMalloc((void**)&d_b, N * sizeof(int));cudaMalloc((void**)&d_c, N * sizeof(int));//Initializing Arraysfor (int i = 0; i < N; i++) {h_a[i] = 2 * i * i;h_b[i] = i;}// Copy input arrays from host to device memorycudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);//Calling kernels passing device pointers as parametersgpuAdd << <512, 512 >> > (d_a, d_b, d_c);//Copy result back to host memory from device memorycudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);cudaDeviceSynchronize();//再次記錄時間戳cudaEventRecord(e_stop, 0);//等待所有GPU工作都完成cudaEventSynchronize(e_stop);float elapsedTime;//計算時間插值cudaEventElapsedTime(&elapsedTime, e_start, e_stop);printf("Time to add %d numbers: %3.1f ms\n", N, elapsedTime);int Correct = 1;printf("Vector addition on GPU \n");//Printing result on consolefor (int i = 0; i < N; i++) {if ((h_a[i] + h_b[i] != h_c[i])){Correct = 0;}}if (Correct == 1){printf("GPU has computed Sum Correctly\n");}else{printf("There is an Error in GPU Computation\n");}//Free up memorycudaFree(d_a);cudaFree(d_b);cudaFree(d_c);return 0;
}

1.2 NVIDIA Visual Profiler

• 如果你在程序中使用了CUDA，代碼的性能并未提升，在這種情況下，能夠可視化地查看代碼的哪些部分花費了最長的時間完成將非常有用，這叫剖析內核執行代碼
• 英偉達提供了以上用途的工具 nvvp ，就在標準的CUDA安裝包里，在電腦的如下路徑可以被找到：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.7\libnvvp：
  
• 執行它需要安裝java環境，即安裝jdk8即可，可以去官網下載，也可以從我的鏈接 jdk8下載，然后需要配置環境變量C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.7\extras\CUPTI\lib64 C:\Program Files\Java\jdk-1.8\bin
  
• 打開nvvp 會出現如下窗口，此工具會分析你的代碼執行過程，采集GPU上的性能數據，運行結束后會給你一個詳細的報告，包括每個內核的執行時間，代碼中每個詳細操作的時間戳，以及代碼存儲器的使用情況
  
• 想要得到詳細報告，可依次點擊 File -> New Session，然后在彈出的對話框中選擇程序的.exe文件
  
• Profiler 是分析內核執行情況的重要工具，它也可以用來比較兩個內核的性能。它會告訴你就是是代碼里的何種操作拉低了性能

2. CUDA中的錯誤處理

• 如果系統中沒有可用的GPU設備怎么辦？顯存不足怎么辦？
• 學會在CUDA程序里邊添加錯誤處理代碼很有好處

#include "cuda_runtime.h"
#include "device_launch_parameters.h"#include <stdio.h>__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {*d_c = *d_a + *d_b;
}
int main()
{//Defining host variablesint h_a, h_b, h_c;//Defining Device Pointersint *d_a, *d_b, *d_c;//Initializing host variablesh_a = 1;h_b = 4;//定義錯誤結果變量cudaError_t cudaStatus;// Allocate GPU buffers for three vectors (two input, one output)    .cudaStatus = cudaMalloc((void**)&d_c, sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}cudaStatus = cudaMalloc((void**)&d_a, sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}cudaStatus = cudaMalloc((void**)&d_b, sizeof(int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!");goto Error;}// Copy input vectors from host memory to GPU buffers.cudaStatus = cudaMemcpy(d_a,&h_a, sizeof(int), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}cudaStatus = cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}// Launch a kernel on the GPU with one thread for each element.gpuAdd<<<1, 1>>>(d_a, d_b, d_c);// Check for any errors launching the kernelcudaStatus = cudaGetLastError();if (cudaStatus != cudaSuccess) {fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));goto Error;}// Copy output vector from GPU buffer to host memory.cudaStatus = cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}printf("Passing Parameter by Reference Output: %d + %d = %d\n", h_a, h_b, h_c);
Error:cudaFree(d_c);cudaFree(d_a);cudaFree(d_b);return 0;
}

• -----------------------END----------------------------