私の CUDA プログラムは、統合されていないグローバル メモリ アクセスに苦しんでいます。idx 番目のスレッドは配列の [idx] 番目のセルのみを処理しますが、以下に示すように間接的なメモリ アクセスが多数あります。
int idx=blockDim.x*blockIdx.x+threadIdx.x;
.... = FF[m_front[m_fside[idx]]];
m_fisde[idx] についてはアクセスを合体させましたが、実際に必要なのは FF[m_front[m_fside[idx]]] です。2 レベルの間接アクセスがあります。
これを直接シーケンシャル アクセスにするために、m_front や m_fsied のデータのパターンをいくつか見つけようとしましたが、ほとんど「ランダム」であることがわかりました。
これに対処する方法はありますか？
グローバル メモリ ランダム アクセスの高速化: 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()、プログラマによって明示的に管理される (組み込みを使用する) ことができます。