CUDAのマルチスレッドカーネルで無駄な書き込み操作が効率的ではありませんか？

Question

CUDAコードにカーネルがあります。共有メモリのいくつかの部分で多くのスレッドを実行したい（グローバルメモリで実行するよりもはるかに速いため）、その結果をグローバルメモリに書き込みます（私はそれ以降のカーネルで使うことができます）。カーネルはこのようなものになります。

__global__ void calc(float * globalmem) 
{ 
    __shared__ float sharemem; //initialize shared memory 
    sharemem = 0; //set it to initial value 
    __syncthreads(); 

    //do various calculations on the shared memory 
    //for example I use atomicAdd() to add each thread's 
    //result to sharedmem... 

    __syncthreads(); 
    *globalmem = sharedmem;//write shared memory to global memory 
} 

私は本当に一度だけそれを書くために必要がある場合、すべて単一のスレッドが、グローバルメモリへの共有から外にデータを書き込んでいるという事実を、私に怪しい感じています。私はまた、すべてのスレッドがコードの開始時に共有メモリをゼロに初期化するという事実から同じ感情を得る。現在の実装よりも速い方法がありますか？

Answer 1

ワープ・レベルでは、冗長な読取りまたは書込み対単一スレッドの実行の間にパフォーマンスの差はほとんどないでしょう。

しかし、スレッドブロック内に複数のワープを置くことで、冗長な読み取りまたは書き込み（1つのスレッドに対して）を実行することで、パフォーマンスの差が大きくなることが予想されます。

むしろ冗長に比べて、単一のスレッドを持つことにより、これらの問題に対処し、読み取りを行うか、書くのに十分であるべきである：共有にあなたがthreadblock内アトミックを使用して、それについて尋ねませんでしたが

__global__ void calc(float * globalmem) 
{ 
    __shared__ float sharemem; //initialize shared memory 
    if (!threadIdx.x) sharemem = 0; //set it to initial value 
    __syncthreads(); 

    //do various calculations on the shared memory 
    //for example I use atomicAdd() to add each thread's 
    //result to sharedmem... 

    __syncthreads(); 
    if (!threadIdx.x) *globalmem = sharemem;//write shared memory to global memory 
} 

メモリは、おそらく共有メモリ削減方法による（おそらくより良い性能のための）交換可能であり得る。