Многопоточная сказка о потерянном времени

В публикации «Сказка о потерянном времени» пользователь crea7or рассказал, как он опровергал Гипотезу Эйлера на современном CPU.

Мне же было интересно узнать как покажет себя GPU, и я сравнил однопоточный код с многопоточным для CPU и совсем многопоточным для GPU, с помощью архитектуры параллельных вычислений CUDA.

CPU

Для начала я распараллелил алгоритм для CPU. Для этого в функцию передается диапазон для перебора значений внешнего цикла a. Затем весь диапазон 1..N разбивается на кусочки и скармливается ядрам процессора.

void Algorithm_1(const uint32_t N, const std::vector<uint64_t>& powers, const uint32_t from, const uint32_t to) {
  for (uint32_t a = from; a < to; ++a) {
    for (uint32_t b = 1; b < a; ++b) {
      for (uint32_t c = 1; c < b; ++c) {
        for (uint32_t d = 1; d < c; ++d) {
          const uint64_t sum = powers[a] + powers[b] + powers[c] + powers[d];
          if (std::binary_search(std::begin(powers), std::end(powers), sum)) {
            const auto it = std::lower_bound(std::begin(powers), std::end(powers), sum);
            uint32_t e = static_cast<uint32_t>(std::distance(std::begin(powers), it));
            std::cout << a << " " << b << " " << c << " " << d << " " << e << "n";
          }
        }
      }
    }
  }
}

В моем случае было 4 ядра, и вычисления я выполнял через чуть-чуть доработанный пул потоков, который взял за основу тут. Так как никаких зависимостей между данными нет, скорость возрастает практически пропорционально количеству ядер.

GPU

Для GPU получается немного сложнее. Внешние два цикла перебора (a и b) будут развернуты, а в функции останется только перебор двух внутренних циклов (c и d). Можно себе представить что программа будет выполняться параллельно для всех значений а = 1..N и b = 1..N. На самом деле конечно это не так, все исполняться параллельно они все-таки не смогут, но это уже забота CUDA максимально распараллелить выполнение, этому можно помочь, если правильно настроить конфигурацию блоков и потоков.

  int blocks_x = N;
  int threads = std::min(1024, static_cast<int>(N));
  int blocks_y = (N + threads - 1) / threads;
  dim3 grid(blocks_x, blocks_y);
  NaiveGPUKernel<<<grid, threads>>>(N);

  const int a = blockIdx.x + 1;
  const int b = threadIdx.x + blockIdx.y * blockDim.x + 1;

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

__constant__ uint64_t gpu_powers[8192];

inline __device__ int gpu_binary_search(const uint32_t N, const uint64_t search) {
  uint32_t l = 0;
  uint32_t r = elements_count - 1;
  uint32_t m;
  while (l <= r) {
    m = (l + r) / 2;
    if (search < gpu_powers[m])
      r = m - 1;
    else if (search > gpu_powers[m])
      l = m + 1;
    else
     return l;
  }
  return -1;
}

__global__ void NaiveGPUKernel(const uint32_t N) {
  const int a = blockIdx.x + 1;
  const int b = threadIdx.x + blockIdx.y * blockDim.x + 1;
  if (b >= a)
    return;
  for (uint32_t c = 1; c < b; ++c) {
    for (uint32_t d = 1; d < c; ++d) {
      const uint64_t sum = gpu_powers[a] + gpu_powers[b] + gpu_powers[c] + gpu_powers[d];
      const auto s = gpu_binary_search(N, sum);
      if (s > 0) {
        printf("%d %d %d %d %dn", a, b, c, d, s);
      }
    }
  }
}

Так же я реализовал оптимизации, подсказанные тут.

Замеры скорости

Замеры для CPU делались по два раза, так как они оказались намного стабильнее GPU, которых я делал уже по семь и выбирал лучшее время. Разброс для GPU мог быть двукратным, это я объясняю тем что для CUDA-расчетов использовался единственный в системе видеоадаптер.

А теперь вколем немного адреналина в CPU!

И этим адреналином будет Profile-guided optimization.

Для PGO я привожу только время CPU, потому что GPU мало изменилось, и это ожидаемо.

Видно что PGO смогло ускорить оптимизацию #1 на целых 18%, для остальных прирост оказался скромнее.

Что можно сделать еще

Для CPU попробовать развернуть циклы, считать по два числа за раз, использовать векторные команды процессора.

Для GPU поиграть с порядком вычислений, попробовать по-экономить на регистрах, лучше подобрать параметры блоков и потоков.

Выводы

1. Писать на CUDA просто.
2. Без ухищрений, прямолинейный код на CUDA оказался быстрее CPU реализации в десятки раз.
3. В числодробильных задачах GPU сильно быстрее CPU за те же деньги.
4. GPU требуется много времени на «раскачку» и для совсем коротких задач ускорения может не быть.
5. Более сложный алгоритм может оказаться быстрее для CPU, но медленнее для GPU.
6. Профильная оптимизация хорошо ускоряет код и уменьшает размер файла одновременно.

Исходники

Традиционно можно потрогать код на GitHub

Автор: drBasic

Источник