6

1.3小さな二次元配列が発生する場所について二次元配列をスキャンすることを含む計算機能 GPU を備えた CUDA を使用するアプリケーションを実装しています。これまで、合体のためのメモリ アラインメント要件を満たすために、両方の配列が を使用して割り当てられ、 を使用cudaMallocPitch()して転送されていました。cudaMemcpy2D()

最初の最適化ステップでは、データを共有メモリにまとめて読み取ることにより、メモリ アクセスをグローバル メモリに統合しようとしています。最適化されていないコードでのテスト (たとえば、発散​​分岐があり、グローバル メモリへのメモリ アクセスが結合されていない場合) を使用してより大きな配列を割り当てたcudaMalloc()ところ、パフォーマンスが最大 倍に向上したことがわかりました50%。これはどのように可能ですか？

4

2 に答える 2

3

cudaMallocPitch() は、2 次元配列 (行優先) の各行の開始アドレスが 2^N の倍数であることを確認します (N は、計算能力に応じて 7 ~ 10 です)。

アクセスがより効率的かどうかは、データ アライメントだけでなく、計算能力、グローバル メモリ アクセス方法、場合によってはキャッシュ構成にも依存します。

このブログでは、初期のコンピューティング機能でのデータ アクセスのミスアラインメントによる帯域幅の大幅な削減について説明しています。

https://developer.nvidia.com/content/how-access-global-memory-effectively-cuda-cc-kernels

パフォーマンスは多くの要因に依存するため、さらに調査できるように、デバイス モジュールの種類とカーネル コードも投稿する必要がある場合があります。

于 2013-02-06T02:26:07.997 に答える