0

CUDA SDK (2.3) の matrixMultiply カーネルの一部を次に示します。

for (int a = aBegin, b = bBegin;
         a <= aEnd;
         a += aStep, b += bStep) {

    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

    int XI=wA * ty + tx;
    int XII=wB * ty + tx;
    ////////////////////
    // PREFETCH BLOCK //
    ////////////////////
    AS(ty, tx) = A[a + XI];
    BS(ty, tx) = B[b + XII];

    __syncthreads();

    for (int k = 0; k < BLOCK_SIZE; ++k)
        Csub += AS(ty, k) * BS(k, tx);

    __syncthreads();
}

このバージョンの行列乗算は、タイルを共有メモリに取り込み、共有メモリ帯域幅で計算を実行します。次の反復のデータを L1 キャッシュにプリフェッチして、パフォーマンスを向上させたい。ここで提案されているようにプリフェッチ組み込みを使用し、PREFETCH BLOCK上記に次のコマンドを挿入しました。

    long long int k,kk;
    k=((long long int)A+aStep); if(k<=aEnd) prefetch_l1(k+XI);
    kk=((long long int)B+bStep); if(kk<=aEnd) prefetch_l1(kk+XII);

テスト後、2 つのバージョン (プリフェッチありまたはなし) のパフォーマンスは非常に似ています (3 回の実行の平均)。

プリフェッチなし: 6434.866211 (ミリ秒)

プリフェッチあり: 6480.041016 (ミリ秒)

質問：

プリフェッチによって速度が向上することを期待していますが、結果に混乱しています。これら2つの実装が非常に近いパフォーマンスを発揮する理由を正当化する団体はありますか? 間違ったプリフェッチを実行している可能性があります。

前もって感謝します。

詳細情報:

GPU: テスラ C2050

CUDA バージョン: 4.0

inline __device__ void prefetch_l1 (unsigned int addr)
{

  asm(" prefetch.global.L1 [ %1 ];": "=r"(addr) : "r"(addr));
}
4

1 に答える 1

1

(任意のアーキテクチャでの) プリフェッチは、次の場合にのみ有効です。

  • メモリ帯域幅に余裕があり、
  • つまり、データが実際に必要になる前に十分に前もって、予備のメモリ帯域幅が利用できるときに、プリフェッチを開始できます。

これらの基準を満たすことができない場合、プリフェッチは役に立たず、むしろ害を及ぼす可能性があります。

于 2012-12-02T22:05:10.727 に答える