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

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

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

CPU

Для начала я распараллелил алгоритм ^[8] для 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 ядра, и вычисления я выполнял через чуть-чуть доработанный пул потоков, который взял за основу тут ^[9]. Так как никаких зависимостей между данными нет, скорость возрастает практически пропорционально количеству ядер.

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

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

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

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

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

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

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

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

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

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

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

Выводы

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

Исходники

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

Автор: drBasic

Источник ^[13]