CUDAコードでの定数メモリの使用

Question

私自身は理解できません。私のカーネルで使用されているメモリを確実にする最良の方法は一定です。 http://stackoverflow...r-pleasant-wayに同様の質問があります。 私はGTX580で作業しており、2.0の機能のみをコンパイルしています。別の方法からSRCポインタを削除し、ファイルを.cuする

__constant__ src[size]; 

を追加することです

cudaMalloc(src, size); 
cudaMemcpy(src, hostSrc, size, cudaMemcpyHostToDevice); 
Foo<<<...>>>(src, result); 

：私のカーネルは、私がホスト上で次のコードを実行

__global__ Foo(const int *src, float *result) {...} 

のように見えますカーネルと実行する

cudaMemcpyToSymbol("src", hostSrc, size, 0, cudaMemcpyHostToDevice); 
Foo<<<...>>>(result); 

これらの2つの方法は同じですか、または最初のものがグローバルメモリではなく定数メモリの使用を保証していませんか？ サイズは動的に変更されるため、私の場合は2番目の方法は便利ではありません。

Answer 1

2番目の方法は、配列がCUDA定数メモリにコンパイルされ、定数メモリキャッシュを介して正しくアクセスされるようにする唯一の方法です。しかし、スレッドのブロック内でその配列の内容がどのようにアクセスされるのか、自分自身に尋ねるべきです。すべてのスレッドがアレイに一様にアクセスする場合、定数メモリキャッシュからのブロードキャストメカニズムが存在するため、定数メモリを使用すると性能が向上します（定数メモリはオフチップDRAMに格納され、キャッシュDRAMのトランザクション数を削減します）。しかし、アクセスがランダムであれば、ローカルメモリへのアクセスがシリアライズされ、パフォーマンスに悪影響を及ぼします。

__constant__メモリに適した典型的なものは、実行時に設定する必要があるモデルの係数、重み、およびその他の定数値です。フェルミGPUでは、カーネル引数リストは、例えば、定数メモリに格納される。しかし、内容が不均一にアクセスされ、メンバーのタイプまたはサイズがコールからコールまで一定でない場合、通常のグローバル・メモリーが推奨されます。

また、GPUのコンテキストごとに64 KBという制限がありますので、非常に大量のデータを定数メモリに保存することは現実的ではありません。キャッシュ付きの読み取り専用ストレージが大量に必要な場合は、データをテクスチャにバインドして、パフォーマンスがどのようなものかを確認する価値があります。プリフェルミカードでは、通常、便利なパフォーマンスが得られます。フェルミでは、そのアーキテクチャのキャッシュレイアウトが向上しているため、結果はグローバルメモリに比べて予測しにくい可能性があります。

Answer 2

最初の方法は、メモリが関数Fooの内部で一定であることを保証します。 2つは等価ではなく、2つ目はそれが初期化された後でそれを保持することを保証します。あなたが最初のやり方に似た何かを使うのは元気でなければならない。