CUDA：統合グローバルメモリに類似したテクスチャメモリのアクセス時間ですか？

Question

私のカーネルスレッドは、線形文字配列に合体してアクセスします。配列 をテクスチャにマップするとスピードアップが表示されません。実行時間はほとんど同じで です。私は、計算能力2.0のTesla C2050に取り組んでいて、どこかにグローバルアクセスがキャッシュされているところで を読んでいます。本当？おそらくそれで私は の実行時間の違いを見ていないのです。

メインプログラムの配列は

char *dev_database = NULL; 
cudaMalloc((void**) &dev_database, JOBS * FRAGMENTSIZE * sizeof(char)); 

であり、私はそれが

cudaBindTexture(NULL, texdatabase, dev_database, JOBS * FRAGMENTSIZE * sizeof(char)); 

各スレッドは、次にID がthreadIdx.x + blockIdx.x * blockDim.xとpあるch = tex1Dfetch(texdatabase, p + id)オフセットの文字を読み取るとtexture<char> texdatabaseをテクスチャに結合します。

私は一度だけ結合しており、dev_databaseは大きな配列です。実際には私が見つけた サイズが大きすぎる場合、バインドに失敗します。バインドする配列 のサイズに制限はありますか？どうもありがとう。

Answer 1

パフォーマンスに差異がない理由はいくつかありますが、このメモリへのアクセスははボトルネックではありません。ボトルネックになっていない場合は、パフォーマンスを上げてもパフォーマンスは向上しません。

キャッシュについて：この場合、バイトだけを読み取っているため、各ワープは32バイトを読み込みます。つまり、4つのワープグループがそれぞれのキャッシュラインにマップされます。キャッシュの競合が少ないと仮定すると、キャッシュから最大4倍の再利用が可能になります。したがって、このメモリアクセスがボトルネックである場合、テクスチャキャッシュが汎用キャッシュよりも恩恵を受けていない可能性があります。

まず、帯域幅が制限されているかどうか、このデータアクセスが原因であるかどうかを判断する必要があります。これを済ませたら、メモリアクセスを最適化してください。考慮するもう1つの方法は、一度に飛行中のメモリトランザクションの数を増やすためにスレッドごとに1つではなく、1ロードあたり4〜16文字（バイトパッキング/アンパックでchar4またはint4構造体を使用）にアクセスすることです。大域メモリバスを飽和させる。

あなたが見たいかもしれないa good presentation by Paulius Micikevicius from GTC 2010があります。分析主導の最適化と、飛行中のメモリトランザクションの具体的な概念の両方を対象としています。