CUDA Global Barrier - работает на Кеплера, а не на Ферми

Question

Следующий глобальный барьер работает на Kepler K10, а не на Fermi GTX580:

__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();
  }
}

Итак, просмотрев выходные данные, отправленные обратно в ЦП, я заметил, что один поток (1-й или последний) выходит за пределы барьера и возобновляет выполнение раньше, чем другие. Я использую CUDA 5.0. количество блоков также всегда меньше, чем количество SM (в моем наборе прогонов).

Есть идеи, почему один и тот же код не работает на двух архитектурах? Что нового в Kepler, который помогает этой глобальной синхронизации?

Answer 1

Поэтому я подозреваю, что сам барьерный код, вероятно, работает так же. Похоже, это то, что происходит с другими структурами данных, не связанными с самой функциональностью барьера.

Ни у Кеплера, ни у Ферми есть кэши L1, которые связаны друг с другом. То, что вы обнаружили (хотя это не связано с самим вашим барьерным кодом), - это то, что поведение кэша L1 отличается у Кеплера и Ферми.

В частности, кеш Kepler L1 не работает при глобальных нагрузках, как описано в приведенной выше ссылке, и поэтому поведение кеширования обрабатывается на уровне L2, который является общесистемным, и, следовательно, согласованным. Когда Kepler SMX считывает свои глобальные данные, он получает согласованные значения от L2.

С другой стороны, в Fermi есть кэши L1, которые также участвуют в глобальных нагрузках (по умолчанию - хотя такое поведение можно отключить), а кэши L1, как описано в приведенной выше ссылке, уникальны для каждого SM Fermi и не согласованы с L1 кеширует в других SM. Когда SM Fermi читает свои глобальные данные, он получает значения от L1, которые могут быть не согласованы с другими кэшами L1 в других SM.

В этом заключается разница в "когерентности", которую вы видите, в данных, которыми вы манипулируете до и после барьера.

Как я уже упоминал, я считаю, что сам барьерный код, вероятно, работает одинаково на обоих устройствах.