0

次のグローバル バリアは、Fermi GTX580 ではなく Kepler K10 で機能します。

__global__ void cudaKernel (float* ref1, float* ref2, int* lock, int time, int dim) {
  int gid  = blockIdx.x * blockDim.x + threadIdx.x;
  int lid  = threadIdx.x;                          
  int numT = blockDim.x * gridDim.x;               
  int numP = int (dim / numT);                     
  int numB = gridDim.x;

  for (int t = 0; t < time; ++t) {
    // compute @ time t
    for (int i = 0; i < numP; ++i) {
      int idx  = gid + i * numT;
      if (idx > 0 && idx < dim - 1)
        ref2 [idx]  = 0.333f * ((ref1 [idx - 1] + ref1 [idx]) + ref1 [idx + 1]);
    }

    // global sync
    if (lid == 0){
      atomicSub (lock, 1);
      while (atomicCAS(lock, 0, 0) != 0);
    }
    __syncthreads();

    // copy-back @ time t
    for (int i = 0; i < numP; ++i) {
      int idx  = gid + i * numT;
      if (idx > 0 && idx < dim - 1)
        ref1 [idx]  = ref2 [idx];
    }

    // global sync
    if (lid == 0){
      atomicAdd (lock, 1);
      while (atomicCAS(lock, numB, numB) != numB);
    }
    __syncthreads();
  }
}

そのため、CPU に送り返された出力を見ると、1 つのスレッド (最初または最後のスレッド) がバリアを回避し、他のスレッドよりも早く実行を再開していることに気付きました。CUDA5.0を使用しています。ブロックの数も常にSMの数よりも小さくなります（私の実行セットでは）。

同じコードが 2 つのアーキテクチャで機能しない理由がわかりませんか? このグローバル同期を支援する Kepler の新機能は何ですか?

4

1 に答える 1

1

したがって、バリアコード自体はおそらく同じように機能しているのではないかと思います。問題になっているのは、バリア機能自体に関連付けられていない他のデータ構造で起こっていることのようです。

ケプラーもフェルミも、互いにコヒーレントな L1 キャッシュを持っていません。あなたが発見したこと (バリア コー​​ド自体には関連付けられていませんが) は、L1 キャッシュの動作がKeplerとFermiの間で異なるということです。

特に、上記のリンクで説明されているように、Kepler の L1 キャッシュはグローバル ロードでは使用されないため、キャッシュ動作はデバイス全体の L2 レベルで処理されるため、一貫性があります。Kepler SMX がグローバル データを読み取るとき、L2 から一貫した値を取得しています。

一方、Fermi にはグローバル ロードにも参加する L1 キャッシュがあり (デフォルトでは -- この動作はオフにすることができます)、上記のリンクで説明されている L1 キャッシュは各 Fermi SM に固有であり、非一貫性があります。他の SM の L1 キャッシュ。Fermi SM がそのグローバル データを読み取るとき、L1 から値を取得します。これは、他の SM の他の L1 キャッシュと一貫性がない可能性があります。

これは、バリアの前後で操作しているデータの「一貫性」の違いです。

前述したように、バリア コー​​ド自体はおそらく両方のデバイスで同じように機能していると思います。

于 2013-01-09T22:55:40.367 に答える