10

ワープのすべてのスレッドは定義上同期しているので、ワープベースの並列削減について考えました。

したがって、同期を必要とせずに、入力データを64分の1に減らすことができる（各スレッドは2つの要素を減らす）という考え方でした。

Mark Harrisによる元の実装と同じように、削減はブロックレベルで適用され、データは共有メモリにあります。 http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf

彼のバージョンとワープベースのバージョンをテストするためにカーネルを作成しました。
カーネル自体は、BLOCK_SIZE要素を共有メモリに完全に同じように格納し、その結果を出力配列の一意のブロックインデックスに出力します。

アルゴリズム自体は正常に機能します。「カウント」をテストするために、自分の完全な配列でテストされました。

実装の関数本体：

/**
 * Performs a parallel reduction with operator add 
 * on the given array and writes the result with the thread 0
 * to the given target value
 *
 * @param inValues T* Input float array, length must be a multiple of 2 and equal to blockDim.x
 * @param targetValue float 
 */
__device__ void reductionAddBlockThread_f(float* inValues,
    float &outTargetVar)
{
    // code of the below functions
}

1.彼のバージョンの実装：

if (blockDim.x >= 1024 && threadIdx.x < 512)
    inValues[threadIdx.x] += inValues[threadIdx.x + 512];
__syncthreads();
if (blockDim.x >= 512 && threadIdx.x < 256)
    inValues[threadIdx.x] += inValues[threadIdx.x + 256];
__syncthreads();
if (blockDim.x >= 256 && threadIdx.x < 128)
    inValues[threadIdx.x] += inValues[threadIdx.x + 128];
__syncthreads();
if (blockDim.x >= 128 && threadIdx.x < 64)
    inValues[threadIdx.x] += inValues[threadIdx.x + 64];
__syncthreads();

//unroll last warp no sync needed
if (threadIdx.x < 32)
{
    if (blockDim.x >= 64) inValues[threadIdx.x] += inValues[threadIdx.x + 32];
    if (blockDim.x >= 32) inValues[threadIdx.x] += inValues[threadIdx.x + 16];
    if (blockDim.x >= 16) inValues[threadIdx.x] += inValues[threadIdx.x + 8];
    if (blockDim.x >= 8) inValues[threadIdx.x] += inValues[threadIdx.x + 4];
    if (blockDim.x >= 4) inValues[threadIdx.x] += inValues[threadIdx.x + 2];
    if (blockDim.x >= 2) inValues[threadIdx.x] += inValues[threadIdx.x + 1];

    //set final value
    if (threadIdx.x == 0)
        outTargetVar = inValues[0];
}

リソース：

4つの同期スレッドが使用されました12if
ステートメントが使用されまし
た11読み取り+追加+書き込み操作
1最終書き込み操作
5レジスタの使用

パフォーマンス：

5回のテスト実行の平均：〜19.54ミリ秒

2.ワープベースのアプローチ:(上記と同じ機能本体）

/*
 * Perform first warp based reduction by factor of 64
 *
 * 32 Threads per Warp -> LOG2(32) = 5
 *
 * 1024 Threads / 32 Threads per Warp = 32 warps
 * 2 elements compared per thread -> 32 * 2 = 64 elements per warp
 *
 * 1024 Threads/elements divided by 64 = 16
 * 
 * Only half the warps/threads are active
 */
if (threadIdx.x < blockDim.x >> 1)
{
    const unsigned int warpId = threadIdx.x >> 5;
    // alternative threadIdx.x & 31
    const unsigned int threadWarpId = threadIdx.x - (warpId << 5);
    const unsigned int threadWarpOffset = (warpId << 6) + threadWarpId;

    inValues[threadWarpOffset] += inValues[threadWarpOffset + 32];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 16];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 8];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 4];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 2];
    inValues[threadWarpOffset] += inValues[threadWarpOffset + 1];
}

// synchronize all warps - the local warp result is stored
// at the index of the warp equals the first thread of the warp
__syncthreads();

// use first warp to reduce the 16 warp results to the final one
if (threadIdx.x < 8)
{
    // get first element of a warp
    const unsigned int warpIdx = threadIdx.x << 6;

    if (blockDim.x >= 1024) inValues[warpIdx] += inValues[warpIdx + 512];
    if (blockDim.x >= 512) inValues[warpIdx] += inValues[warpIdx + 256];
    if (blockDim.x >= 256) inValues[warpIdx] += inValues[warpIdx + 128];
    if (blockDim.x >= 128) inValues[warpIdx] += inValues[warpIdx + 64];

    //set final value
    if (threadIdx.x == 0)
        outTargetVar = inValues[0];
}

リソース：

1つのsyncthreadが使用されました7if
ステートメント
10読み取り追加書き込み操作
1最終書き込み操作
5レジスタの使用

5ビットシフト
1追加
1サブ

パフォーマンス：

5回のテスト実行の平均：〜20.82ミリ秒

256mbのfloat値を持つGeforce8800GT512mbで両方のカーネルを複数回テストします。そして、ブロックあたり256スレッドでカーネルを実行します（100％の占有率）。

ワープベースのバージョンは、約1.28ミリ秒遅くなります。

将来のカードでより大きなブロックサイズが許可される場合、最大値が4096で64に減少し、最終的なワープによって1に減少するため、ワープベースのアプローチではそれ以上の同期ステートメントは必要ありません。

なぜそれは速くないのですか？またはアイデアの欠陥はどこにありますか、カーネル？

リソースの使用法から、ワープアプローチを先に進める必要がありますか？

編集1：スレッドの半分だけがアクティブであり、範囲外の読み取りが発生しないようにカーネルを修正し、新しいパフォーマンスデータを追加しました

4

2 に答える 2

12

あなたのコードが私のコードよりも遅い理由は、私のコードでは、最初のフェーズで各 ADD に対して半分のワープがアクティブになっているためだと思います。コードでは、最初のフェーズのすべてですべてのワープがアクティブになっています。したがって、全体として、コードはより多くのワープ命令を実行します。CUDA では、1 回のワープで実行される命令の数だけでなく、実行される「ワープ命令」の総数を考慮することが重要です。

また、ワープの半分だけを使用しても意味がありません。ワープを 2 つの分岐を評価して終了させるためだけにワープを起動すると、オーバーヘッドが生じます。

もう 1 つの考えは、unsigned charand を使用するとshort実際にパフォーマンスが低下する可能性があるということです。よくわかりませんが、レジスタは単一の 32 ビット変数にパックされていないため、確実に保存されません。

また、元のコードでは、blockDim.x をテンプレート パラメーター BLOCKDIM に置き換えました。これは、5 つの実行時 if ステートメントしか使用しないことを意味します (第 2 段階の if はコンパイラーによって削除されます)。

ところで、あなたを計算する安価な方法threadWarpIdは

const int threadWarpId = threadIdx.x & 31;

他のアイデアについては、この記事をチェックしてください。

編集: これは代替のワープベースのブロック削減です。

template <typename T, int level>
__device__
void sumReduceWarp(volatile T *sdata, const unsigned int tid)
{
  T t = sdata[tid];
  if (level > 5) sdata[tid] = t = t + sdata[tid + 32];
  if (level > 4) sdata[tid] = t = t + sdata[tid + 16];
  if (level > 3) sdata[tid] = t = t + sdata[tid +  8];
  if (level > 2) sdata[tid] = t = t + sdata[tid +  4];
  if (level > 1) sdata[tid] = t = t + sdata[tid +  2];
  if (level > 0) sdata[tid] = t = t + sdata[tid +  1];
}

template <typename T>
__device__
void sumReduceBlock(T *output, volatile T *sdata)
{
  // sdata is a shared array of length 2 * blockDim.x

  const unsigned int warp = threadIdx.x >> 5;
  const unsigned int lane = threadIdx.x & 31;
  const unsigned int tid  = (warp << 6) + lane;

  sumReduceWarp<T, 5>(sdata, tid);
  __syncthreads();

  // lane 0 of each warp now contains the sum of two warp's values
  if (lane == 0) sdata[warp] = sdata[tid];

  __syncthreads();

  if (warp == 0) {
    sumReduceWarp<T, 4>(sdata, threadIdx.x);
    if (lane == 0) *output = sdata[0];
  }
}

これは、最初のステージで起動されたすべてのワープを使用し、追加のブランチ、共有ロード/ストア、および__syncthreads()新しい中間ステージを犠牲にして、最後のステージ内に分岐がないため、少し高速になるはずです。このコードはテストしていません。実行する場合は、そのパフォーマンスを教えてください。元のコードで blockDim のテンプレートを使用すると、さらに高速になる可能性がありますが、このコードの方が簡潔だと思います。

tFermi およびそれ以降のアーキテクチャは純粋なロード/ストア アーキテクチャを使用するため、一時変数が使用されることに注意してください。そのため+=、共有メモリから共有メモリに追加のロードが発生します (sdataポインタは揮発性でなければならないため)。テンポラリに明示的に一度ロードすると、これを回避できます。G80 では、パフォーマンスに違いはありません。

于 2012-10-04T23:37:34.560 に答える
0

SDK の例も確認する必要があります。いくつかの削減方法を実装した非常に良い例を 1 つ覚えています。それらの少なくとも 1 つは、ワープ ベースのリダクションも使用します。

（他のマシンにしかインストールされていないため、現在名前を調べることができません）

于 2012-10-05T00:06:17.537 に答える