CUDA推力ファクタGMEMアクセス：ctorデータコピー対ctor dev ptr arg

Question

スラストファンクタがグローバルに非ベクトル化不均一アクセスの読取り専用状態にアクセスできるようにする2つの方法があります。残念ながら、カーネルの実行時間には100倍の違いがあります。なぜ私の2つの戦略に違いはありますか？

さらに一般的には、スラストファンクターにこれらのグローバル変数へのアクセスを提供する正式な方法はありますか？

私の最初の方法は、私のグローバルデータのコピーをファンクタに入れることです。

// assuming barData is a float[] 
foo<N>(barData); 

私の第二の方法は、自分自身が推力を使用してデバイスにアップロードを実行することである

推力:: for_eachを使用して呼び出され

// functor containing a copy of array dependency 
template<size_t BARSIZE> 
struct foo1_func 
{ 
    __align__(16) float bar[BARSIZE]; 
    foo1_func(float _bar[BARSIZE]) { memcpy(bar,_bar,BARSIZE*sizeof(float)); } 
    __host__ __device__ operator()(float &t) { t = do_something(t, bar); } 
} 

...：スラスト機械は、デバイス上にアップロードしてキャッシングを実行するように見えます::アップロードしたデータのデバイスメモリポインタを私のファンクタにコピーして渡すだけです。この方法はかなり遅いように見える：感謝して受け入れられた正規のか、独自のファンクタのパターンを示すソースへ

推力:: for_eachを使用して呼び出され

// functor containing device pointers to array in GMEM 
struct foo2_func 
{ 
    float *bar; 
    foo2_func(float* _bar) { bar = bar; } 
    __host__ __device__ operator()(float &t) { t = do_something(t, bar); } 
} 

...

// assuming d_bar is a thrust::device_vector 
foo(thrust::raw_pointer_cast(d_bar.data())); 

リンク。

Answer 1

最初の方法では、foo1_func構造体をカーネル関数パラメータとして渡して、実際には配列bar全体をGPUレジスタに配置しようとしています。

​​

barの大きさは、抵抗者に置かれるのに十分小さい場合、barへのランダムアクセスが実際に登録するためのランダムアクセスです。

しかし、2番目の方法では、構造体を通してグローバルメモリポインタしか渡されませんでした。したがって、barへのランダムアクセスは、グローバルメモリへのランダムアクセスです。

これは、2番目の方法がはるかに遅い理由です。

どちらの方法もユースケースを持っています。 barのサイズと、barのキャッシングにどれくらいのレジスタを費やしたいか、どちらを選択してもかまいません。