OpenCL內核在速度更快的GPU上執行速度更慢

我對OpenCL相當陌生，並且試圖實現DSP算法以比較其與不同GPU上的性能相比，標準CPU實現。儘管我已經獲得了巨大的性能提升，但是我發現奇怪的是，我在GT240上獲得了與GTX 480相同的增益。我的程序執行了兩個內核，而另一個在GTX 480上加速慢下來。OpenCL內核在速度更快的GPU上執行速度更慢

GT240：內核1：226us，內核2：103us。

GTX 480：內核1：35us，內核2：293us。

這些數字是使用Visual Profiler獲得的。下面是內核2的代碼，在較大的卡上它幾乎慢了3倍。該內核需要一個iTotalBins x iNumAngles大的內存塊，並計算iNumAngles長度的每一行的最大值，並將曲線擬合到3個相鄰值。

__kernel void max_curve_fit_gpu (__global float* fCorrelationResult, 
          const int iNumAngles, 
          const int iTotalBins, 
          __global float* fResult){ 

// Get the thread ID which is used as the index the bin the direction is being calculated for 
const int iBinNum = get_global_id(0); 
const int iOffset = iBinNum*iNumAngles; 

// Find the max for this bin 
float fMax = 0; 
int iMaxIndex = 0; 
for (int iAngle=0; iAngle<iNumAngles; iAngle++) 
{ 
    if (fMax < fCorrelationResult[iOffset + iAngle]) 
    { 
     fMax = fCorrelationResult[iOffset + iAngle]; 
     iMaxIndex = iAngle; 
    } 
} 

// Do the curve fit 
float fPrev, fNext, fA, fB, fAxis; 
fPrev = fCorrelationResult[iOffset + (iMaxIndex + iNumAngles - 1) % iNumAngles]; 
fNext = fCorrelationResult[iOffset + (iMaxIndex + iNumAngles + 1) % iNumAngles]; 

fB = (fPrev - fNext)*0.5f; 
fA = (fNext + fPrev) - fMax*2.0f; 
fAxis = fB/fA; 

    // Store the result 
fResult[iBinNum] = iMaxIndex + fAxis; }

視覺探查還表明，有對內核2. 135％全局內存指令重播我有一個版本的最大搜索不使用的if-else statment，但它更慢運行在GPU的。

任何幫助將不勝感激。

來源

2012-01-25 Almero

使用OpenCL，當內核在特定GPU上運行時，是否有可能進入較低級別並找出有關注冊和共享內存使用情況的內容？

從我對NVIDIA CUDA的有限接觸中，利用率可能是關鍵。 GT240的計算能力爲1.2，GTX480爲2.0，所以後者有2個寄存器和3個共享內存。我的猜測是，OpenCL爲第二個內核生成的代碼未能在480上使用這些資源。例如，可能存在共享內存銀行衝突。

來源

2012-01-25 15:38:52

在你的代碼中，線程T將訪問fCorrelationResult[T*iNumAngles+iAngle]，這意味着你沒有合併訪問，也可能是內存組衝突。銀行衝突可能會解釋您觀察到的現象。

您應該轉置您的矩陣，並改爲訪問fCorrelationResult[T+iAngle*iNumBins]。你一定會得到一個不錯的加速，並且可能會在兩個GPU之間更經常的基準測試。

來源

2012-01-25 22:59:08

OpenCL內核在速度更快的GPU上執行速度更慢

回答

相關問題