1

これは、ブロック内で最大50個の値の配列を見つけるためにリダクションを実行しようとしている私のコードです。配列を64にパディングしました。

スレッド1〜31の場合、正しいmaxValの出力がありますが、スレッド32〜49の場合、これは完全に乱数です。何が間違っているのかわかりません。

ところで。展開時にすべての行を_syncする必要はないと思いましたが、どうやらそうする必要があります。それについて何か提案はありますか？

助けてくれてありがとう。 

//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);
//}
4

3 に答える 3

3

これは、より効率的で（少なくともFermi 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]一時を保持することによりt、コード行ごとに1つの共有メモリ負荷を節約できます。また、Fermiは、レジスタを命令オペランドとしてのみ使用できるロード/ストアアーキテクチャであるため、1行あたり1命令、つまり合計6命令を節約できます（全体で約25％になると思います）。

古いT10/GT200アーキテクチャ以前では、コード（volatileを使用し、__ 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];
}
于 2011-06-30T04:06:58.917 に答える
2

__syncthreads()発散コードでは使用しないでください！特定のブロックからのすべてのスレッドまたはスレッドなしのいずれかが__syncthreads()、同じ場所のすべてに到達する必要があります。

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

の重要性を詳しく説明し__threadfence_block()ます。次の2行を検討してください。

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の場合は明らかに失敗します。

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

__threadfence()グローバルメモリの可視性を確保するための同様の関数が存在します。

于 2011-06-28T20:10:41.460 に答える
1

私がしたように、将来このスレッドに遭遇するすべての人のために、ここにハリズムの答えに加えてアドバイスがあります-シャッフル操作を検討することはパフォーマンスの観点から価値があるかもしれません、それで更新されたコードは64から最大を取得しますシングルワープを使用する要素は次のようになります。

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回の書き込みだけなので、かなりきれいです。

于 2014-02-13T00:54:33.270 に答える