如何確定CUDA GPU的性能？

Question

我正在寫一個cuda程序，用於匹配分辨率〜180X180的每個輸入圖像，以及約10,000個分辨率爲〜128 * 128的模板圖像。目標是實現實時性能，即在1秒內對25〜30個輸入圖像進行模板匹配（每個圖像都有10,000個模板）。

目前我使用以下方法

1. 預裝的所有模板上的GPU全局內存保存運行時的I/O操作。
2. 創建單個內核以將一個源圖像與所有模板圖像進行匹配，併爲所有正匹配返回一個數組。
3. 在時域中執行所有操作（不使用FFT）。原因是，我嘗試過基-4實現，但它需要大量的中間全局讀取和寫入，最終花費更多時間。

到目前爲止1個輸入圖像到10,000個模板，它需要大約2秒。

我的問題是：

1. 是否有辦法，以確定此任務是achieveable實時或不？我的意思是在最大FLOPS和I/O帶寬限制e.t.c的幫助下。
2. 如何計算GPU是否最大程度地被充分利用？
3. 提高性能的可能方法？

機規格：i7-4770，8GB，GTX-680]

闡釋目前的內核代碼：

1. 所有模板圖像[尺寸爲RGB約128X128]每加載在GPU內存上。想法是在運行時操作中保存I/O。
2. 每個輸入圖像都加載到紋理內存上，原因在於紋理對於2D尋址來說是很好的選擇。
3. 每個「塊」有1024個線程。
4. 每個線程計算每個輸出像素的值，輸出大小爲[31X31 = 961像素]。
5. 啓動的塊數等於匹配的模板圖像的數量。

內核代碼：

__global__ void cudaMatchTemplate(TemplateArray *templates, uchar *Match) 
{ 
    int global = blockIdx.x*blockDim.x + threadIdx.x; 

    __shared__ int idx[TEMPLATE_MATCH_DIM]; 
    __shared__ float out_shared[TEMPLATE_MATCH_DIM]; 

    //halving the template size.... 
    int rows = (templates[blockIdx.x].nHeight)/2; 
    int cols = (templates[blockIdx.x].nWidth)/2; 

    int fullCol = templates[blockIdx.x].nWidth; 

    int x = templates[blockIdx.x].nMatchLeft; 
    int y = templates[blockIdx.x].nMatchTop; 

    int offset_y = (threadIdx.x/TEMPLATE_MATCH_SIZE); 
    int offset_x = (threadIdx.x - offset_y*TEMPLATE_MATCH_SIZE); 

    // *************** Performing match in time domain *****************************// 
    int sum = 0; 
    float temp; 
    int idxXFactor = 3*(2*(offset_x) + x); 
    int idxYFactor = 2*(offset_y) + y ; 

    for (int i = 0; i < rows; i++) 
    { 
     int I=3*i*fullCol; 
     int sourceIdxY = idxYFactor + 2*i; 
     for (int j = 0; j < cols; j++) 
     { 
      int J=3*j; 
      int sourceIdxX = idxXFactor + 2*J;   
      int templateIdx = 2*I+2*J; 
      //**** R *****// 
      temp = float(tex2D(SourceImgColorTex,sourceIdxX,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx]); 
      sum = sum + temp*temp; 
      //**** G *****// 
      temp = float(tex2D(SourceImgColorTex,sourceIdxX+1,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +1]); 
      sum = sum + temp*temp; 
      //**** B *****// 
      temp = float(tex2D(SourceImgColorTex,sourceIdxX+2,sourceIdxY)) - float(templates[blockIdx.x].pRAWPixels_gpu[templateIdx +2]); 
      sum = sum + temp*temp; 
     } 
    } 

    __syncthreads(); 

//placing all values in shared memory for comparison. 
    if(threadIdx.x < TEMPLATE_MATCH_DIM) 
    { 
     idx[threadIdx.x] = threadIdx.x; 
     out_shared[threadIdx.x] = sum; 
    } 
    __syncthreads(); 


// //computing the Min location.....// 

#pragma unroll 
    for(int s=512; s>0; s>>=1) 
    { 
     if ((threadIdx.x < s) &&((threadIdx.x + s)<TEMPLATE_MATCH_DIM)) 
     { 
      idx[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? idx[threadIdx.x] : idx[threadIdx.x + s]; 
      out_shared[threadIdx.x] = out_shared[threadIdx.x] < out_shared[threadIdx.x + s] ? out_shared[threadIdx.x] : out_shared[threadIdx.x + s];   
     } 

    } 

    __syncthreads(); 

    if(threadIdx.x <1) 
    { 
     int half_Margin = MARGIN_FOR_TEMPLATE_MATCH/2; 
     int matchY = idx[0]/TEMPLATE_MATCH_SIZE ; 
     int matchX = idx[0] - matchY * TEMPLATE_MATCH_SIZE; 

     int diff = absolute(half_Margin - matchX) + absolute(half_Margin - matchY); 
     if(diff < THRESHOLD) 
     { 
      Match[blockIdx.x] = 1; 
     } 
     else 
      Match[blockIdx.x] = 0; 

    } 
} 

Answer 1

我會盡力回答您的大多數問題

是否有辦法，以確定此任務是achieveable實時或不？我的意思是在最大FLOPS和I/O帶寬限制e.t.c的幫助下。

我不知道如何確定內核是否是實時的實現，您可以使用CUDA Occupancy Calculator， 您可以考慮使用紋理，表面內存，常量內存最大化CUDA內核，固定主機的內存和更多這些都取決於你的算法實現。

如何計算GPU是否最大程度地完全使用？

您可以使用CUDA佔用率計算器和CUDA可視化分析器。 我強烈建議使用視覺分析器，它會引導您瞭解CUDA。

提高性能的可能方法？

有一些有趣的方法，這樣做， 首先你可以用上面的方法， 最大限度地提高您的內核調用。如果這還不夠，請按順序嘗試在同一複製數據和計算的工作使用流對象實現管道時間。

如果不能解決問題，嘗試使用延遲，同時操作多個線程訪問GPU，因爲CC 3.5 CUDA啓動了HyperQ，這可能會幫助您並行完成多個調用。

如果這不起作用，請考慮使用多個GPU設備。

請保持我們的發佈