OpenCL核心問題

Question

嗨我創建了兩個內核來做一個簡單的匹配deshredder程序與OpenCL和定時運行。這兩個內核做他們應該做的事情，但是由於我無法破譯的原因，其中一個運行速度比另一個慢得多：/唯一真正的區別是我如何存儲發送的數據以及匹配是如何發生的。

__kernel void Horizontal_Match_Orig( 
__global int* allShreds, 
__global int* matchOut, 
const unsigned int shredCount, 
const unsigned int pixelCount) 

{ 
    int match = 0; 
    int GlobalID = get_global_id(0); 
    int currShred = GlobalID/pixelCount; 
    int thisPixel = GlobalID - (currShred * pixelCount); 
    int matchPixel = allShreds[GlobalID];//currShred*pixelCount+thisPixel]; 
    for (int i = 0; i < shredCount; i++) 
    { 

     match = 0; 
     if (matchPixel == allShreds[(i * pixelCount) + thisPixel]) 
     { 
      if (matchPixel == 0) 
      { 
       match = match + 150; 
      } 
      else match = match + 1; 
     } 
     else match = match - 50; 
     atomic_add(&matchOut[(currShred * shredCount) + i], match); 
    } 
} 

此內核水平得到切絲邊緣，所以一個切絲的像素佔用POS 0到n在陣列allShreds再下切絲的像素被從POS n + 1個存儲到M（其中n =像素數量，m =增加的像素數量）。 GPU的每個線程都有一個像素與合作，並與它匹配反對所有其他Shred的對應像素（包括自身）

__kernel void Vertical(
    __global int* allShreds, 
    __global int* matchOut, 
    const int numShreds, 
    const int pixelsPerEdge) 
{ 
    int GlobalID = get_global_id(0); 
    int myMatch = allShreds[GlobalID]; 
    int myShred = GlobalID % numShreds; 
    int thisRow = GlobalID/numShreds; 
    for (int matchShred = 0; matchShred < numShreds; matchShred++) 
    { 
     int match = 0; 
     int matchPixel = allShreds[(thisRow * numShreds) + matchShred]; 
     if (myMatch == matchPixel) 
     { 
      if (myMatch == 0) 
       match = 150; 
      else 
       match = 1; 
     } 
     else match = -50; 
      atomic_add(&matchOut[(myShred * numShreds) + matchShred], match); 
    } 
} 

這個內核垂直得到一絲一毫的邊緣，因此，所有的碎片的第一像素被存儲在pos 0到n中，則所有碎片的第二個像素被存儲在pos n + 1 ot m（其中n =碎片的數量，並且m =被添加到n的碎片的數量）。該過程類似於前一個過程，每個線程獲取一個像素並將其與其他每個細絲的相應像素進行匹配。

兩者都給出了相同的結果針對純序列程序測試的正確結果。從理論上講，它們應該運行在大致相同的時間內，垂直運行的可能性會更快，因爲原子添加對它的影響不應該太大......但是運行速度會更慢......任何想法？

這是我使用啓動它的代碼（我使用它爲C＃包裝）：

theContext.EnqueueNDRangeKernel(1, null, new int[] { minRows * shredcount }, null, out clEvent); 

與全球總工作量等於像素中的每個的總數量（#Shreds X #Pixels一）。

任何幫助，將不勝感激

Answer 1

兩個內核做他們應該做的事情，但一跑遠遠超過了其他的原因，我無法破譯的要慢/ 唯一的區別是如何我存儲正在發送的數據以及匹配是如何發生的。

這讓一切變得不同。這是一個典型的聚結問題。你沒有在你的問題中指定你的GPU模型或供應商，所以我必須保持模糊，因爲實際的數字和行爲完全依賴於硬件，但總的想法是合理的便攜。

GPU中的工作項發出內存請求（讀取和寫入）在一起（通過「warp」/「wavefront」/「sub-group」）到內存引擎。該引擎在事務中提供內存（兩個大小的功耗爲16到128字節的塊）。下面的例子假設大小爲128。

輸入存儲器存取凝聚：如果經的32項工作，閱讀4個字節（int或float）是連續的內存，存儲引擎將問題單筆交易爲所有32個請求。但是對於每個超過128字節的讀取，需要發出另一個事務。在最壞的情況下，這是每個128字節的32個事務，這樣會更昂貴。

---

你的水平內核做以下訪問：

allShreds[(i * pixelCount) + thisPixel] 

的(i * pixelCount)是整個工作項目不變，只是thisPixel變化。給定您的代碼並假設工作項目0具有thisPixel = 0，那麼工作項目1具有thisPixel = 1等等。這意味着您的工作項目正在請求相鄰的讀取，因此您可以獲得完美的合併訪問權限。同樣的電話atomic_add。

在另一方面，你的垂直內核執行以下操作訪問：

allShreds[(thisRow * numShreds) + matchShred] 
// ... 
matchOut[(myShred * numShreds) + matchShred] 

matchShred和numShreds是跨線程不變，僅thisRow和myShred變化。這意味着您請求的讀取距離彼此爲numShreds。這不是順序訪問，因此不能合併。