CUDA中的性能

Question

我似乎無法弄清楚影響內核性能的基本因素。我實現了兩個簡單的內核，一個加載兩個圖像並逐個像素地添加它們，另一個加載兩個圖像，並按位進行加載。現在，我對它們進行了模板化，以便內核可以拍攝8位和32位圖像，以及1,3或4通道圖像。

所以，起初我有兩個內核加載全局內存爲uchar3和float3，與uchar4等我也不太清楚有關使用三元組，不過，由於合併一起，所以我想我給它分析運行。我認爲，由於操作與通道號無關，因此我可以像讀取圖像一樣，將圖像看作是三倍寬度的1通道圖像，而不是其實際的圖像。

實際上，uchar3全球負荷很大，太多比uchar負荷慢。我的努力得到了證明。但是，唉，這隻發生在算術內核上。按位與運算顯示了完全相反的結果！

現在，我知道我可以將圖像數據加載爲uint s而不是uchar s，用於位運算，它應該完美地處理合並。但讓我們假設我只是想學習和理解正在發生的事情。

讓我們忘記float3 s和float4 s等。我的問題是與uchar版本的內核。所以，簡而言之，爲什麼uchar負載有時比負載快，有時不是？

我正在使用GTX 470計算能力2.0。

PS。根據CUDA編程指南，邏輯操作和添加操作具有相同的吞吐量。 （我的內核實際上必須首先將uchar s轉換爲uint s，但這應該在兩個內核中發生。）所以執行長度應該與我收集的內容大致相同。

算術添加內核（uchar版）：

__global__ void add_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep) 
{ 
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x; 
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y; 

    if (yCoordinate >= height) 
     return; 

#pragma unroll IMAGE_MULTIPLIER 
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i) 
    { 
     // Load memory. 
     uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 
     uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 

     // Write output. 
     *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] + inputElementTwo[0]; 
    } 
} 

位與內核：

__global__ void and_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep) 
{ 
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x; 
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y; 

    if (yCoordinate >= height) 
     return; 

#pragma unroll IMAGE_MULTIPLIER 
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i) 
    { 
     // Load memory. 
     uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 
     uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)); 

     // Write output. 
     *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] & inputElementTwo[0]; 
    } 
} 

的uchar3版本除了加載/存儲線如下內容相同：

 // Load memory. 
    uchar3 inputElementOne = *reinterpret_cast<uchar3*>(inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3); 
    uchar3 inputElementTwo = *reinterpret_cast<uchar3*>(inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3); 

    // Write output. 
    *reinterpret_cast<uchar3*>(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3) 
     = make_uchar3(inputElementOne.x + inputElementTwo.x, inputElementOne.y + inputElementTwo.y, inputElementOne.z + inputElementTwo.z); 

與AND內核類似。 （我不確定我完全記得內核，說實話...明天我會證實這一點）。

Answer 1

uchar3由於SM指令集中沒有24位加載，因此編譯器會將負載分解爲單獨的負載。因此，他們從不合並。在一定程度上，緩存將緩解這種情況。

但是，根據確切的執行配置，每個線程可能只有大約10.7個字節的緩存（您的示例可能會接近該值，因爲內核很簡單，所以很多線程可以同時在一個SM上運行）。由於緩存沒有完全關聯，所以在發生顛簸之前，每個線程的可用字節數可能會小很多。何時發生這種情況取決於許多因素，包括指令的確切調度，即使對於具有相同記錄吞吐量的指令，這些因素也可能不同。

您可以比較兩個版本的cuobjdump -sass可執行文件的輸出以查看編譯器的靜態調度是否相同。然而，運行時的動態調度如何運行基本上是不可觀測的。

正如您已經注意到的，圖像的所有通道都以相同的方式進行處理，所以在線程之間分配它們並不重要。最好的選擇是使用uchar4而不是uchar3或uchar，它（假設圖像的合適對齊）將使您可以獨立於緩存進行合併訪問。這應該導致更短和更一致的執行時間。