0

私はCUDAが提供するプログラミングガイドで自分自身にCUDAを教えています。練習のために、配列の最大値を決定してCPUに返す単純なカーネルを作成しました。

  __global__ void getTheMaximum(float* myArrayFromCPU, float* returnedMaximum) {
    // Store my current value in shared memory.
    extern __shared__ float sharedData[];
    sharedData[threadIdx.x] = myArrayFromCPU[threadIdx.x];

    // Iteratively calculate the maximum.
    int halfScan = blockDim.x / 2;
    while (halfScan > 0 && threadIdx.x < halfScan) {
      if (sharedData[threadIdx.x] < sharedData[threadIdx.x + halfScan]) {
        sharedData[threadIdx.x] = sharedData[threadIdx.x + halfScan];
      }
      halfScan = halfScan / 2;
    }

    // Put maximum value in global memory for later return to CPU.
    returnedMaximum[0] = sharedData[0];
  }

myArrayFromCPUはサイズ1024のfloat値の配列です。returnedMaximumは、単一の項目を含む単純な配列です。計算された最大値。

このアルゴリズムの私の考えは、ブロックサイズの半分から現在の値を超えて値を削り落とすときに、最大値を繰り返し決定するというものです。

ただし、このコードを実行すると、信頼性の低い出力が得られます。返される最大値は異なります。何故ですか？単一のアルゴリズムで毎回異なる値を生成するにはどうすればよいですか？

アップデート：

また、1つのブロックで実行しています。X=1024の1次元ブロックサイズを設定することでこれを保証します。

4

1 に答える 1

2

ブロック全体のすべてのスレッドがまったく同時に実行されるとは限りません。これにより、1 つのワープ (32 スレッドのグループ) 内にのみ存在することが保証されます。

ブロック内での同時実行の危険を回避する__syncthreads()には、組み込み関数を使用して、すべてのスレッドがポイントに到達するまでブロックに到達するのを停止できます。__syncthreads()すべてのスレッドが均一にスポットに到達することを保証できない分岐コードを挿入しないでください。

次のループを試してください。

__syncthreads();
while (halfScan > 0) {
  if (threadIdx.x < halfScan) {
    if (sharedData[threadIdx.x] < sharedData[threadIdx.x + halfScan]) {
      sharedData[threadIdx.x] = sharedData[threadIdx.x + halfScan];
    }
  }
  __syncthreads();
  halfScan = halfScan / 2;
}

すべてのスレッドを同じ場所で同じ回数実行しthreadIdx.x < halfScanたいので、while ループから条件を削除したことに注意してください。__syncthreads()

また、ループの前に、ループの開始前に (すべてのスレッドに対して)__syncthreads()からのロードが完了していることを確認するのに役立つ場合があります。myArrayFromCPU

于 2012-11-03T20:35:11.763 に答える