優化Cuda內核時間執行

Question

我是一名學習Cuda的學生，我想優化內核函數的執行時間。結果，我意識到計算兩張圖片之間差異的簡短程序。所以，我比較在C經典的CPU執行，並在CUDA中C的GPU執行之間的執行時間

在這裏你可以找到我談論代碼：

int *imgresult_data = (int *) malloc(width*height*sizeof(int)); 
int size = width*height; 

switch(computing_type) 
{ 

    case GPU: 

    HANDLE_ERROR(cudaMalloc((void**)&dev_data1, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data2, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data_res, size*sizeof(int))); 

    HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice)); 

    float time; 
    cudaEvent_t start, stop; 

    HANDLE_ERROR(cudaEventCreate(&start)); 
    HANDLE_ERROR(cudaEventCreate(&stop)); 
    HANDLE_ERROR(cudaEventRecord(start, 0)); 

    for(int m = 0; m < nb_loops ; m++) 
    { 
     diff<<<height, width>>>(dev_data1, dev_data2, dev_data_res); 
    } 

    HANDLE_ERROR(cudaEventRecord(stop, 0)); 
    HANDLE_ERROR(cudaEventSynchronize(stop)); 
    HANDLE_ERROR(cudaEventElapsedTime(&time, start, stop)); 

    HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost)); 

    printf("Time to generate: %4.4f ms \n", time/nb_loops); 

    break; 

    case CPU: 

    clock_t begin = clock(), diff; 

    for (int z=0; z<nb_loops; z++) 
    { 
     // Apply the difference between 2 images 
     for (int i = 0; i < height; i++) 
     { 
      tmp = i*imgresult_pitch; 
      for (int j = 0; j < width; j++) 
      { 
       imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp]; 
      } 
     } 
    } 
    diff = clock() - begin; 

    float msec = diff*1000/CLOCKS_PER_SEC; 
    msec = msec/nb_loops; 
    printf("Time taken %4.4f milliseconds", msec); 

    break; 
} 

這裏是我的內核功能：

__global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res) 
{ 
    int row = blockIdx.x; 
    int col = threadIdx.x; 
    int v = col + row*blockDim.x; 

    if (row < MAX_H && col < MAX_W) 
    { 
     data_res[v] = (int) data2[v] - (int) data1[v]; 
    } 
} 

我獲得的這些執行時間爲每一個

• CPU：1,3210ms
• GPU：0,3229ms

我不知道爲什麼GPU結果不低，因爲它應該是。我是Cuda的初學者，所以如果有一些經典錯誤，請全面。編號1： 謝謝您的反饋。我試圖從內核中刪除'if'條件，但它並沒有改變我的程序執行時間。

但是，在安裝Cuda分析器後，它告訴我我的線程沒有併發運行。我不明白爲什麼我會有這種信息，但看起來確實如此，因爲我的GPU只比CPU有5到6倍的速度。這個比例應該更大，因爲每個線程都應該同時處理一個像素到所有其他線程。如果你對我做錯了什麼有所瞭解，那將是有益的...

流量。

Answer 1

可能還有其他代碼問題，但這裏是我所看到的。在__global__ void diff以下行被認爲不是最佳的：

if (row < MAX_H && col < MAX_W) 
{ 
    data_res[v] = (int) data2[v] - (int) data1[v]; 
} 

條件運算內核結果經線發散內部。這意味着在經紗內部的if和else部件按順序執行，而不是並行執行。另外，正如您可能已經意識到的，if僅在邊界處評估爲false。爲了避免分歧和不必要的計算，分成兩個部分圖像：

1. 中央部分，其中row < MAX_H && col < MAX_W總是true。爲這個區域創建一個額外的內核。這裏不需要if。

2. 將使用您的diff內核的邊界區域。

很明顯，您將修改調用內核的代碼。

---

而且在一個單獨的說明：

1. GPU具有面向吞吐量的體系結構，而不是延遲導向爲CPU。這意味着在處理少量數據時，CPU可能比CUDA更快。您是否嘗試過使用大型數據集？

2. CUDA Profiler是一個非常方便的工具，它會告訴你在代碼中不是最優的。

Answer 2

我不認爲你正確測量時間，記憶體複製是一個耗時的GPU步驟，你應該在測量時間時考慮到這一點。

我看到一些細節，你可以測試：

1. 我想你正在使用MAX_H和MAX_H爲常數，你可能會考慮這樣做使用cudaMemcpyToSymbol（）。

2. 記得使用__syncthreads（）同步你的線程，所以你不會在每次循環迭代之間得到問題。

3. CUDA與warps一起工作，所以每個block的線程數和線程數可以更好地工作8倍，但不超過每塊512個線程，除非您的硬件支持它。這是一個使用每塊128個線程的示例：< < <（cols * rows + 127）/ 128,128 >>>。

4. 還記得在GPU中釋放您分配的內存並摧毀您創建的時間事件。

5. 在你的內核函數中，你可以有一個變量int v = threadIdx.x + blockIdx.x * blockDim.x。

6. 除了執行時間之外，您是否測試過您的結果是正確的？我認爲在使用數組時由於填充而應該使用cudaMallocPitch（）和cudaMemcpy2D（）。