2

私のカーネルはfloat、以下のランダム アクセス パターンでサイズ 8 x 8 の配列を使用します。

// inds - big array of indices in range 0,...,7
// flts - 8 by 8 array of floats

// kernel essentially processes large 2D array by looping through slow coordinate
// and having block/thread parallelization of fast coordinate.

__global__ void kernel (int const* inds, float const* flt, ...)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;  // Global fast coordinate
    int idy;                                          // Global slow coordinate
    int sx = gridDim.x * blockDim.x;                  // Stride of fast coordinate

    for ( idy = 0; idy < 10000; idy++ )       // Slow coordinate loop
    {
        int id = idx + idy * sx;              // Global coordinate in 2D array

        int ind = inds[id];                   // Index of random access to small array

        float f0 = flt[ind * 8 + 0];
        ...
        float f7 = flt[ind * 8 + 7];

        NEXT I HAVE SOME ALGEBRAIC FORMULAS THAT USE f0, ..., f7
    }
}

flt配列にアクセスする最良の方法は何ですか?

  1. を渡さないでください。メモリfltを使用してください。__const__異なるスレッドが異なるデータにアクセスするときの const メモリの速度はわかりません。
  2. 上記のように使用します。スレッドが異なるデータにアクセスするため、均一ロードは使用されません。それにもかかわらず、キャッシュのために高速になりますか?
  3. 共有メモリにコピーし、共有メモリ配列を使用します。
  4. テクスチャを使用します。テクスチャを使用したことがない...このアプローチは高速ですか?

共有メモリの場合は、おそらく配列を転置した方がよいでしょうflt。つまり、バンクの競合を避けるために、次の方法で配列にアクセスします。

float fj = flt_shared[j * 8 + ind]; // where j = 0, ..., 7

PS: ターゲット アーキテクチャは Fermi と Kepler です。

4

1 に答える 1

1

「最良の」方法は、作業しているアーキテクチャにも依存します。Fermi と Keplerでのランダム アクセス (マッピングを使用しているため、アクセスは一種のランダムのように見えますinds[id]) に関する私の個人的な経験では、L1 が非常に高速になったため、多くの場合、共有メモリではなくグローバル メモリを使用し続ける方がよいということです。メモリまたはテクスチャ メモリ。

グローバル メモリ ランダム アクセスの高速化: L1 キャッシュ ラインの無効化

Fermi および Kepler アーキテクチャは、グローバル メモリからの 2 種類のロードをサポートします。フル キャッシングがデフォルト モードで、L1、L2、GMEM の順にヒットを試行し、負荷の粒度は 128 バイト ラインです。L2-onlyは L2 でヒットを試み、次に GMEM で、負荷の粒度は 32 バイトです。特定のランダム アクセス パターンでは、L1 を無効にし、L2 のより低い粒度を活用することで、メモリ効率を向上させることができます。–Xptxas –dlcm=cgこれは、にオプションを付けてコンパイルすることで実行できますnvcc。

グローバル メモリ アクセスを高速化するための一般的なガイドライン: ECC サポートの無効化

Fermi および Kepler GPU はエラー訂正コード (ECC) をサポートしており、ECC はデフォルトで有効になっています。ECC は、ピーク メモリ帯域幅を削減し、医療画像処理や大規模なクラスター コンピューティングなどのアプリケーションでデータの整合性を強化するために必要です。不要な場合は、Linux の nvidia-smi ユーティリティ (リンクを参照) を使用するか、Microsoft Windows システムのコントロール パネルを使用して、パフォーマンスを向上させるために無効にすることができます。ECC のオンとオフを切り替えるには、再起動が必要であることに注意してください。

Kepler でグローバル メモリ アクセスを高速化するための一般的なガイドライン: 読み取り専用データ キャッシュの使用

Kepler は、関数の実行中は読み取り専用であることがわかっているデータ用に 48KB のキャッシュを備えています。読み取り専用パスの使用は、共有/L1 キャッシュ パスをオフロードし、フル スピードのアライメントされていないメモリ アクセスをサポートするため、有益です。読み取り専用パスの使用は、コンパイラによって自動的に管理される (const __restrictキーワードを使用する) か__ldg()、プログラマによって明示的に管理される (組み込みを使用する) ことができます。

于 2013-03-02T07:18:51.927 に答える