2

featWidth * featHeight * 31キューブを別のキューブと「畳み込み」しようとしていmodelWidth * modelHeight * 31ます。問題は、このカーネルが非常に遅いことです (まあ、私はシーケンシャルな CPU コードよりも高速にできていますが、OpenMP バージョンと同じくらい遅いです)。私は Quadro FX 1800 を使用しています (ええ、64 個の CUDA コア...)。

__constant__ float d_model[31*22*22];
#define IMUL(a,b) ( __mul24((a), (b)) )
#define IMAD(a,b,c) ( __mul24((a), (b)) + (c) )
__global__ void dMatch(float *score, const int featWidth, const int featHeight, const int modelWidth, const int modelHeight, const int scoreWidth, const int scoreHeight)
{
  const int x = IMAD(blockIdx.x, blockDim.x, threadIdx.x);
  const int y = IMAD(blockIdx.y, blockDim.y, threadIdx.y);
  if(x < scoreWidth && y < scoreHeight)
  {
   const int scoreIdx = IMAD(x, scoreHeight, y);
   score[scoreIdx] = 0.f;
   const int baseFeatIdx = IMUL(x,scoreHeight) + IMAD(modelHeight-1, x, y);
   for(int z = 0; z < 31; ++z)
   {
     // Index positionning
     int featIdx =  IMAD(z, IMUL(featWidth,featHeight), baseFeatIdx);
     int modelIdx = IMUL(z, IMUL(modelWidth,modelHeight));

     float value = 0.f;

     // filter
     for(int xx=0; xx<modelWidth; xx++)
     {
       const int xxmodelIdx = IMAD(xx, modelHeight, modelIdx);
       const int xxfeatIdx = IMAD(xx, featHeight, featIdx);
       for(int yy=0; yy<modelHeight; yy++)
       {
         value += d_model[xxmodelIdx+yy] * tex1Dfetch(texFeatures,xxfeatIdx+yy);
       }
     }
     score[scoreIdx] += value;
  }
 }
}

とにかく、8*8ブロック内のスレッドとグリッド サイズ(scoreWidth/8)*(scoreHeight/8)(scoreWidth と scoreHeight は結果のマトリックス サイズ) でこのカーネルを起動します。私のコードで何が間違っているのか、何がかなり遅いのか、何か手がかりがあるかどうか知りたいです。

編集：

tera のおかげで、はるかに高速なバージョン (480 ミリ秒のプロセスで 150 ミリ秒のドロップ!):

__global__ void dMatch(float *score, const int featWidth, const int featHeight, const int modelWidth, const int modelHeight, const int scoreWidth, const int scoreHeight)
{
    const int y = IMUL(4,IMAD(blockIdx.x, blockDim.x, threadIdx.x));
    const int x = IMAD(blockIdx.y, blockDim.y, threadIdx.y);
    if(x < scoreWidth && y < scoreHeight)
    {
    const int scoreIdx = IMAD(x, scoreHeight, y);
    const int baseFeatIdx = IMUL(x,scoreHeight) + IMAD(modelHeight-1, x, y);
    float value=0.f, value1 = 0.f, value2 = 0.f, value3 = 0.f;
    float feat,feat1,feat2,feat3;

    // Index positionning
    int featIdx =  0;
    int modelIdx = 0;
    int xxmodelIdx;
    int xxfeatIdx; 
    float val;
    for(int z = 0; z < 31; ++z)
    {
        featIdx = IMAD(z,IMUL(featWidth,featHeight),baseFeatIdx);
        modelIdx = IMUL(z,IMUL(modelWidth,modelHeight));

        // filter
        for(int xx=0; xx<modelWidth; xx++)
        {
            xxmodelIdx  = IMAD(xx, modelHeight, modelIdx);
            xxfeatIdx = IMAD(xx, featHeight, featIdx);
            feat=tex1Dfetch(texFeatures,xxfeatIdx+0);
            feat1=tex1Dfetch(texFeatures,xxfeatIdx+1);
            feat2=tex1Dfetch(texFeatures,xxfeatIdx+2);
            feat3=tex1Dfetch(texFeatures,xxfeatIdx+3);
            for(int yy=0; yy<modelHeight; yy++)
            {
                val = d_model[xxmodelIdx+yy];
                value += val * feat;
                value1 += val * feat1;
                value2 += val * feat2;
                value3 += val * feat3;
                feat = feat1;
                feat1 = feat2;
                feat2 = feat3;
                feat3 = tex1Dfetch(texFeatures,xxfeatIdx+yy+4);
            }
        }
    }
    score[scoreIdx] = value;
    if(y+1 < scoreHeight)
        score[scoreIdx+1] = value1;
    if(y+2 < scoreHeight)
        score[scoreIdx+2] = value2;
    if(y+3 < scoreHeight)
        score[scoreIdx+3] = value3;
}

これで起動しましたdim3 threads(16,16); dim3 grid(divup(scoreHeight,64), divup(scoreWidth,16));。

4

1 に答える 1

1

プロファイラーは何と言っていますか？NVidia NSight (Windows の Visual Studio およびLinux のEclipse用のプラグイン) を使用すると、ストールがどこにあるかを確認し、パフォーマンスを最適化するためのさまざまなヒントを得ることができます。

私の推測では (プロファイラーを見ずに)、あなたが持っているブロックが小さすぎるということです。基本的なスケジューリング要素である warp 内には 32 個のスレッドがあります。NVIDIA GPU は、現在のスレッドが前の命令を実行している間に他のスレッドで動作することによってレイテンシを隠すことができるため、高速化できます。SM ごとに 8 ブロック (Tesla と Fermi の場合) または 16 (Kepler の場合) のブロックを持つことは可能ですが、ピークで 16 ～ 32 のワープがあり、これは非常に小さい可能性があります (私が間違っている可能性がありますが、ブロックの起動には一定の遅延があります)。もっと大きなブロックの使用を検討します。

コードを正しく理解していれば、テクスチャ フェッチは最適modelHeight - 1ではありbaseFeatIdません。したがって、テクスチャ フェッチは完全にランダムであり、データの局所性を利用しません。逆にすると、より効率的になります。featIdxxxfeatIdxxy

ただし、プロファイラーで確認することをお勧めします。問題が GPU に依存する計算である場合は、計算側に集中する必要があります。問題がメモリ バウンドである場合は、メモリ アクセス パターンを調べる必要があります。最適化のスポットのように見える他のいくつかの部分があるかもしれませんが、ボトルネックが何であるかを見るまではわかりません. それがわかったら、ベスト プラクティス ガイドの特定の章を読みたいと思うかもしれません。

于 2013-04-19T11:51:54.370 に答える