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

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

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

Железо и компиляторы
CPU Core i5 4440
GPU GeForce GTX 770
MSVS 2015 update 3
tookit CUDA 8.
Конечно, сборка была релизная и 64 битная, с оптимизацией /02 и /LTCG.
Отключение сброса видеоадаптера
Так как вычисления длятся более двух секунд, видеодрайвер завершает CUDA-программу. Чтобы этого не происходило, нужно указать достаточное время до сброса через ключ реестра
HKEY_LOCAL_MACHINE\SYSTEM\CurrentControlSet\Control\GraphicsDrivers\TdrDelay и перезагрузить компьютер.


CPU


Для начала я распараллелил алгоритм для CPU. Для этого в функцию передается диапазон для перебора значений внешнего цикла a. Затем весь диапазон 1…N разбивается на кусочки и скармливается ядрам процессора. 
void Algorithm_1(const uint32_t N, const std::vector& 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(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(N));
  int blocks_y = (N + threads - 1) / threads;
  dim3 grid(blocks_x, blocks_y);
  NaiveGPUKernel<<>>(N);

blocks_x — это диапазон перебора первого цикла a.
А вот перебор второго цикла b пришлось сделать сложнее из-за ограничения видеокарты на количество потоков (1024), и поместить счетчик одновременно в threads и blocks_y:
a и b вычисляется так: 
  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 %d\n", a, b, c, d, s);
      }
    }
  }
}

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

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


Замеры для CPU делались по два раза, так как они оказались намного стабильнее GPU, которых я делал уже по семь и выбирал лучшее время. Разброс для GPU мог быть двукратным, это я объясняю тем что для CUDA-расчетов использовался единственный в системе видеоадаптер.
N CPU время, мс CPU (4 потока) время, мс GPU время, мс
@antonkrechetov 100 58.6 16.7 14.8
Базовый 100 45.3 13.0 10.7
Базовый оптимизация #1 100 6.3 2.1 12.7
Базовый оптимизация #2 100 1.4 0.7 0.8
@antonkrechetov 250 2999.7 782.9 119.0
Базовый 250 2055.6 550.1 90.9
Базовый оптимизация #1 250 227.2 60.3 109.2
Базовый оптимизация #2 250 42.9 11.9 6.0
@antonkrechetov 500 72034.2 19344.1 1725.83
Базовый 500 38219.7 10200.8 976.7
Базовый оптимизация #1 500 3725.1 926.5 1140.36
Базовый оптимизация #2 500 630.7 170.2 48.7
@antonkrechetov 750 396566.0 105003.0 11521.2
Базовый 750 218615.0 57639.2 5742.5
Базовый оптимизация #1 750 19082.7 4736.8 6402.1
Базовый оптимизация #2 750 3272.0 846.9 222.9
Базовый оптимизация #2 1000 10204.4 2730.3 1041.9
Базовый оптимизация #2 1250 25133.1 6515.3 2445.5
Базовый оптимизация #2 1500 51940.1 14005.0 4895.2

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


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

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

N CPU
время, мс		 CPU (4 потока)
время, мс		 CPU+PGO
время, мс		 CPU+PGO
(4 потока)
время, мс
@antonkrechetov 100 58.6 16.7 55.3 15.0
Базовый 100 45.3 13.0 42.2 12.1
Базовый оптимизация #1 100 6.3 2.1 5.2 1.9
Базовый оптимизация #2 100 1.4 0.7 1.3 0.8
@antonkrechetov 250 2999.7 782.9 2966.1 774.1
Базовый 250 2055.6 550.1 2050.2 544.6
Базовый оптимизация #1 250 227.2 60.3 200.0 53.2
Базовый оптимизация #2 250 42.9 11.9 40.4 11.4
@antonkrechetov 500 72034.2 19344.1 68662.8 17959.0
Базовый 500 38219.7 10200.8 38077.7 10034.0
Базовый оптимизация #1 500 3725.1 926.5 3132.9 822.2
Базовый оптимизация #2 500 630.7 170.2 618.1 160.6
@antonkrechetov 750 396566.0 105003.0 404692.0 103602.0
Базовый 750 218615.0 57639.2 207975.0 54868.2
Базовый оптимизация #1 750 19082.7 4736.8 15496.4 4082.3
Базовый оптимизация #2 750 3272.0 846.9 3093.8 812.7
Базовый оптимизация #2 1000 10204.4 2730.3 9781.6 2565.9
Базовый оптимизация #2 1250 25133.1 6515.3 23704.3 6244.1
Базовый оптимизация #2 1500 51940.1 14005.0 48717.5 12793.5

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

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


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

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

Выводы


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

Исходники


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

Комментарии (0)

© Habrahabr.ru