1

CUDAでニューラルネットワークをアクティブ化するためのコードを書いていますが、問題が発生しています。特定のニューロンに入る重みの正しい合計が得られません。

これがカーネルコードです。変数を使ってもう少しわかりやすく説明します。

__global__ void kernelSumWeights(float* sumArray, float* weightArray, int2* sourceTargetArray, int cLength)
{
int nx = threadIdx.x + TILE_WIDTH*threadIdx.y;
int index_in = (blockIdx.x + gridDim.x*blockIdx.y)*TILE_WIDTH*TILE_WIDTH + nx;
 if(index_in < cLength)
 {

  sumArray[sourceTargetArray[index_in].y] += fabs(weightArray[index_in]);
  //__threadfence();
  __threadfence_block();

 }

}

まず、ネットワーク内の接続数はですcLength。すべての接続には、ソースニューロンとターゲットニューロン、およびその接続の重みがあります。SourceTargetArrayその情報が含まれています。したがって、のインデックスiはsourceTargetArray接続のソースニューロンインデックスであり、接続iのターゲットニューロンインデックスですi。にはweightArray重み情報が含まれています（したがって、のインデックスiはweightArray接続に対応しますi）。

ご覧のとおり、SumArrayここに合計を保存しています。したがって、カーネルはsumArray（接続のターゲットニューロンインデックスでi）接続の重みの絶対値だけ増分しますi。直感的には、ニューロンへのすべての着信接続について、すべての重みを合計します。私がこのカーネルでやろうとしているのはこれだけです。最終的には、この合計を使用して重みを正規化します。

問題はそれが間違っているということです。私はこれを連続して行いましたが、答えは異なります。答えは異なり、通常は約12〜15倍です（したがって、正しい答えは700.0になり、私が得ているのは50年代の範囲です）。

私が追加したことがわかります__threadfence()（そして__threadfence_block()、書き込みがすべてのスレッドによって同時に行われていないことを確認するために）。これが私のコードの問題かどうかはわかりません。ウェイト配列がテストしたシリアルバージョンと同一であり、ソース/ターゲット情報も同一であることを確認しました。私は何が間違っているのですか？

編集：参考までに、__threadfence()usagedはCUDAプログラミングガイドv3.1付録B.5メモリフェンス機能で説明されています

4

2 に答える 2

4

+=アトミックではありません=>スレッドセーフではありません。atomicAddを使用します。

また、同じメモリセルへの書き込みは避けてください。問題は、これらの呼び出しがシリアル化され、スレッドが並んで待機することです。この操作を回避できない場合は、アルゴリズムを個別の計算とマージの2つのフェーズに分割してみてください。並列マージは非常に効率的に実装できます。

于 2010-09-14T16:52:32.530 に答える
3

あなたは削減を行う必要があります。

各スレッドに割り当てられた要素を合計し、結果を配列に配置します。cache[threadsPerBlock]、次に__Syncthreads

次に、隣接する小計を連続して追加して、結果の小計を減らします。

int cacheIndex = threadIdx.x;
int i = blockDim.x / 2;
while (i != 0)
{
    if (cacheIndex < i)
        cache[cacheIndex] += cache[cacheIndex] + 1;
        __syncthreads;
        i /= 2;
    }
}

次のデッキでは、これについて詳しく説明しています。

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

このためのサンプルコードはここにあります：

http://www.nvidia.com/object/cuda_sample_data-parallel.html

また、「CUDA BYExample」（コードフラグメントの出所）でもよく説明されています。

このアプローチには1つの大きな注意点があります。追加は、シリアルコードの場合と同じ順序では発生しません。フロートの追加は可換ではないため、丸め誤差によって結果がわずかに異なる場合があります。

于 2010-10-01T06:19:13.483 に答える