9

私はCudaを学んでいますが、まだ並列処理を理解しています。私が現在抱えている問題は、値の配列に最大削減を実装することです。これは私のカーネルです

__global__ void max_reduce(const float* const d_array,
                     float* d_max,
                     const size_t elements)
{
    extern __shared__ float shared[];

    int tid = threadIdx.x;
    int gid = (blockDim.x * blockIdx.x) + tid;

    if (gid < elements)
        shared[tid] = d_array[gid];
    __syncthreads();

    for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
    {
        if (tid < s && gid < elements)
            shared[tid] = max(shared[tid], shared[tid + s]);
        __syncthreads();
    }

    if (gid == 0)
        *d_max = shared[tid];
}

同じ方法 (max 関数を min に置き換える) を使用して min reduce を実装しましたが、これは正常に機能します。

カーネルをテストするために、シリアル for ループを使用して最小値と最大値を見つけました。最小値と最大値は常にカーネル内で同じ結果になりますが、一致するのは最小リデュースのみです。

私が行方不明/間違っていることは明らかですか？

4

2 に答える 2

19

削除された回答の主な結論は正しかったです。投稿したカーネルは、そのカーネル実行の最後に全体的な削減をかなり行ったという事実を理解していませんが、結果は完全ではありません。各ブロックの結果を (何らかの方法で) 組み合わせる必要があります。コメントで指摘されているように、コードには他にもいくつかの問題があります。変更されたバージョンを見てみましょう。

__device__ float atomicMaxf(float* address, float val)
{
    int *address_as_int =(int*)address;
    int old = *address_as_int, assumed;
    while (val > __int_as_float(old)) {
        assumed = old;
        old = atomicCAS(address_as_int, assumed,
                        __float_as_int(val));
        }
    return __int_as_float(old);
}


__global__ void max_reduce(const float* const d_array, float* d_max, 
                                              const size_t elements)
{
    extern __shared__ float shared[];

    int tid = threadIdx.x;
    int gid = (blockDim.x * blockIdx.x) + tid;
    shared[tid] = -FLOAT_MAX;  // 1

    if (gid < elements)
        shared[tid] = d_array[gid];
    __syncthreads();

    for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
    {
        if (tid < s && gid < elements)
            shared[tid] = max(shared[tid], shared[tid + s]);  // 2
        __syncthreads();
    }
    // what to do now?
    // option 1: save block result and launch another kernel
    if (tid == 0)        
        d_max[blockIdx.x] = shared[tid]; // 3
    // option 2: use atomics
    if (tid == 0)
      atomicMaxf(d_max, shared[0]);
}
  1. Pavan が示したように、共有メモリ配列を初期化する必要があります。gridDim.x*blockDim.xが より大きい場合、起動された最後のブロックは「完全な」ブロックではない可能性がありますelements。
  2. この行では、 ( gid) を操作するスレッドが 未満でelementsあることを確認していますが、共有メモリにインデックスを付けるために を追加sするとgid、最後のブロックで、共有メモリにコピーされた正当な値の外側でもインデックスを作成できることに注意してください。したがって、注 1 に示す共有メモリの初期化が必要です。
  3. すでに発見したように、最後の行は正しくありませんでした。各ブロックは独自の結果を生成するため、何らかの方法でそれらを組み合わせる必要があります。起動されたブロックの数が少ない場合に検討できる 1 つの方法 (これについては後で詳しく説明します) は、atomicsを使用することです。通常、アトミックは実行時間の点で「コストがかかる」ため、アトミックを使用しないように人々を導きます。しかし、私たちが直面しているもう 1 つのオプションは、ブロックの結果をグローバル メモリに保存し、カーネルを終了してから、場合によっては別のカーネルを起動して個々のブロックの結果を結合することです。最初に多数のブロック (たとえば 1024 以上) を起動した場合、この方法論に従うと、2 つの追加のカーネルを起動することになる可能性があります。したがって、アトミックの考察。おっしゃる通り、ネイティブはいません。atomicMaxフロートの関数ですが、ドキュメントに示されているように、 を使用atomicCASして任意のアトミック関数を生成できatomicMaxfますfloat。

しかし、1024 以上のアトミック関数 (ブロックごとに 1 つ) を実行するのが最善の方法でしょうか? おそらくそうではありません。

スレッドブロックのカーネルを起動するときは、マシンをビジー状態に保つのに十分な数のスレッドブロックを起動するだけで済みます。経験則として、SM ごとに少なくとも 4 ～ 8 のワープが動作する必要があります。しかし、最初に何千ものスレッドブロックを起動しても、マシン使用率の観点から特にメリットはありません。SM あたり 8 スレッドブロックなどの数を選択し、GPU に最大で 14 ～ 16 の SM がある場合、これは比較的少数の 8*14 = 112 スレッドブロックになります。良い丸め数として 128 (8*16) を選びましょう。これは魔法のようなものではなく、GPU をビジー状態に保つのに十分です。これらの 128 個のスレッドブロックのそれぞれに追加の作業を行わせて、全体を解決するとします。問題が発生した場合、(おそらく) ペナルティをあまり支払うことなくアトミックの使用を活用し、複数のカーネルの起動を回避できます。では、これはどのように見えるでしょうか?:

__device__ float atomicMaxf(float* address, float val)
{
    int *address_as_int =(int*)address;
    int old = *address_as_int, assumed;
    while (val > __int_as_float(old)) {
        assumed = old;
        old = atomicCAS(address_as_int, assumed,
                        __float_as_int(val));
        }
    return __int_as_float(old);
}


__global__ void max_reduce(const float* const d_array, float* d_max, 
                                              const size_t elements)
{
    extern __shared__ float shared[];

    int tid = threadIdx.x;
    int gid = (blockDim.x * blockIdx.x) + tid;
    shared[tid] = -FLOAT_MAX; 

    while (gid < elements) {
        shared[tid] = max(shared[tid], d_array[gid]);
        gid += gridDim.x*blockDim.x;
        }
    __syncthreads();
    gid = (blockDim.x * blockIdx.x) + tid;  // 1
    for (unsigned int s=blockDim.x/2; s>0; s>>=1) 
    {
        if (tid < s && gid < elements)
            shared[tid] = max(shared[tid], shared[tid + s]);
        __syncthreads();
    }

    if (tid == 0)
      atomicMaxf(d_max, shared[0]);
}

この変更されたカーネルでは、カーネル起動を作成するときに、全体のデータ サイズに基づいて起動するスレッドブロックの数を決定していません ( elements)。代わりに、固定数のブロック (たとえば、128、この数を変更して最も速く実行されるものを見つけることができます) を起動し、各スレッドブロック (したがってグリッド全体) をメモリを介してループさせ、各要素の部分的な最大操作を計算します共有メモリ。gid次に、コメント 1 でマークされた行で、変数を初期値に再設定する必要があります。gridDim.x*blockDim.xこれは実際には不要であり、グリッド ( ) のサイズが 未満であることを保証すれば、ブロック削減ループ コードをさらに単純化できますelements。これは、カーネルの起動時に行うのは難しくありません。

このアトミック メソッドを使用する場合は、結果 (*d_maxこの場合) を のような適切な値に初期化する必要があることに注意してください-FLOAT_MAX。

繰り返しになりますが、私たちは通常、アトミックな使用から人々を誘導しますが、この場合、慎重に管理するかどうかを検討する価値があり、追加のカーネル起動のオーバーヘッドを節約できます.

高速な並列リダクションを行う方法の忍者レベルの分析については、関連するCUDA サンプルで利用できる Mark Harris の優れたホワイトペーパーをご覧ください。

于 2013-06-29T15:40:36.107 に答える