2012-01-25 10 views
0

私はOpenCLをかなり新しくしており、標準的なCPU実装と比較して異なるGPUでその性能を比較するDSPアルゴリズム を実装しようとしています。 私は巨大なパフォーマンスの向上を達成しましたが、奇妙なのは、はるかに高速なGTX 480とほぼ同じゲインをGT240で得ることです。私のプログラムは2つのカーネルを実行し、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; } 

のVisual Profilerはまた、私はのif-elseなステートメントを使用していない最大の検索のバージョンを使用しているカーネルの2の135パーセントのグローバルメモリ命令のリプレイがあることを示しているが、それは、両方のGPUの上でも、実行速度が遅いです。

ご協力いただきますようお願い申し上げます。

答えて

0

OpenCLを使用すると、特定のGPUでカーネルを実行しているときに、下位レベルに進み、レジスタと共有メモリの使用状況を知ることは可能ですか?

私のNVIDIA CUDAへの露出が限られているので、ここでは使用率が重要です。 GT240は計算能力1.2でGTX480は2.0なので、後者は2倍のレジスタと3倍の共有メモリを持っています。私の推測では、2番目のカーネルのためにOpenCLによって生成されたコードは、480上でこれらのリソースを利用することができません。例えば、共有メモリバンク競合があるかもしれません。

3

コードでは、スレッドTfCorrelationResult[T*iNumAngles+iAngle]にアクセスします。つまり、アクセスが合併していない可能性があり、メモリバンクの競合が発生している可能性があります。銀行の争いは、あなたが観察する現象を説明するかもしれません。

マトリックスを転置して、代わりにfCorrelationResult[T+iAngle*iNumBins]にアクセスしてください。確かに、2つのGPUの間で素早いスピードアップと、おそらくより一般的なベンチマークを得るでしょう。

関連する問題