CUDA、減算を使用して最大値を見つける、エラー

Question

ここで私のコードはブロック内で最大50の値の配列を見つけるために減らそうとしています。私は配列を64にパディングしました。

私は正しいmaxValの印刷がありますが、スレッドの32-49では完全に乱数です。私が間違っていることは分かりません。

btw。私はアンロールですべての行を同期させる必要はないと思ったが、明らかに私はしなければならない。それについての示唆？

ご協力いただきありがとうございます。

//block size = 50 


__syncthreads(); 

if (tid<32){ 

    cptmp[tid]=(cptmp[tid]< cptmp[tid+32]) ? cptmp[tid+32] : cptmp[tid];__syncthreads();  
    cptmp[tid]=(cptmp[tid]< cptmp[tid+16]) ? cptmp[tid+16] : cptmp[tid];__syncthreads(); 
    cptmp[tid]=(cptmp[tid]< cptmp[tid+8]) ? cptmp[tid+8] : cptmp[tid]; __syncthreads();  
    cptmp[tid]=(cptmp[tid]< cptmp[tid+4]) ? cptmp[tid+4] : cptmp[tid]; __syncthreads(); 
    cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid]; __syncthreads();  
    cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid]; __syncthreads(); 

} 

__syncthreads(); 

//if (tid==0) { 
    maxVal=cptmp[0]; 
    if(bix==0 && biy==0) cuPrintf(" max:%f x:%d y:%d\n", maxVal, blockIdx.x, blockIdx.y); 
//} 

Answer 1

これは、（少なくともフェルミGPUでは）より効率的で、volatileを使用している正しいコードです。 Tをあなたのタイプに置き換えてください（またはテンプレートを使用してください）：

if (tid<32) { 
    volatile T *c = cptmp; 
    T t = c[tid]; 
    c[tid] = t = (t < c[tid+32]) ? c[tid+32] : t; 
    c[tid] = t = (t < c[tid+16]) ? c[tid+16] : t; 
    c[tid] = t = (t < c[tid+ 8]) ? c[tid+ 8] : t; 
    c[tid] = t = (t < c[tid+ 4]) ? c[tid+ 4] : t; 
    c[tid] = t = (t < c[tid+ 2]) ? c[tid+ 2] : t; 
    c[tid] = t = (t < c[tid+ 1]) ? c[tid+ 1] : t; 
} 

なぜこれがより効率的ですか？ __syncthreads()がない場合には、共有メモリへの揮発性ポインタを使用する必要があります。しかし、これはコンパイラに共有メモリへのすべての読み書きを「尊重」させることになります。つまり、レジスタ内で何かを最適化して保持することはできません。したがって、常にに常にc[tid]を保存することによって、1行のコードごとに1つの共有メモリの負荷を保存します。フェルミはレジスタを命令オペランドとしてしか使用できないロード/ストアアーキテクチャであるため、1行に6命令（合計で約25％）の命令を保存します。

以前のT10/GT200アーキテクチャー以前では、揮発性で__syncthreads（）を持たないコードも同様に効率的でした。なぜなら、そのアーキテクチャーは命令ごとに1つのオペランドを共有メモリーから直接得ることができたからです。

あなたは?:上ifを好む場合、このコードは同等でなければなりません：私が行ったように、将来的にこのスレッドつまずくだろう、みんなのために

if (tid<32) { 
    volatile T *c = cptmp; 
    T t = c[tid]; 
    if (t < c[tid+32]) c[tid] = t = c[tid+32]; 
    if (t < c[tid+16]) c[tid] = t = c[tid+16]; 
    if (t < c[tid+ 8]) c[tid] = t = c[tid+ 8]; 
    if (t < c[tid+ 4]) c[tid] = t = c[tid+ 4]; 
    if (t < c[tid+ 2]) c[tid] = t = c[tid+ 2]; 
    if (t < c[tid+ 1]) c[tid] = t = c[tid+ 1]; 
} 

Answer 2

ダイバージェントコードで__syncthreads()を使用しないでください！ ブロック内のすべてのスレッドまたはスレッドが、同じ場所にあるすべての__syncthreads()に到達する必要があります。

単一のワープからのすべてのスレッド（32スレッド）は暗黙的に同期されているため、すべてをまとめて入れるのに__syncthreads()は必要ありません。ただし、同じワープの別のスレッドが1つのスレッドの共有メモリ書き込みを見えない場合は、__threadfence_block()を使用してください。

__threadfence_block()の重要性を詳述する。それはこのようなものにコンパイル可能

cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid]; 
cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid]; 

：それは、シングルスレッド・コードの正しいであろうが

int tmp; //assuming that cptmp is an array of int-s 
tmp=cptmp[tid]; 
tmp=(tmp<cptmp[tid+2])?cptmp[tid+2]:tmp; 
tmp=(tmp<cptmp[tid+1])?cptmp[tid+1]:tmp; 
cptmp[tid]=tmp; 

、それは明らかにCUDAのために失敗し、次の2行を考えてみましょう。

cptmpの配列をvolatileと宣言するか、またはその行の間に__threadfence_block()を追加すると、そのような最適化を防ぐことができます。この関数は、関数が存在する前に、同じブロックのすべてのスレッドが現在のスレッドの共有メモリを参照するようにします。

類似性__threadfence()グローバルメモリの可視性を保証する機能があります。

Answer 3

、ここでのアドバイスはharrismの答えに加えている - それは

auto localMax = max(c[tid], c[tid + 32]);  
for (auto i = 16; i >= 1; i /= 2) 
{ 
    localMax = max(localMax, __shfl_xor(localMax, i)); 
} 
c[tid] = localMax; 

は2つだけを読み込み、グローバルメモリから1個の書き込み：更新されたコードは次のようになり、単一のワープを使用して64個の要素のうちの最大値を取得するので、シャッフル操作を考慮することが、性能の観点から価値があるかもしれません必要なので、それはかなりきれいです。