В публикации «Сказка о потерянном времени» пользователь crea7or рассказал, как он опровергал Гипотезу Эйлера на современном CPU.
Мне же было интересно узнать как покажет себя GPU, и я сравнил однопоточный код с многопоточным для CPU и совсем многопоточным для GPU, с помощью архитектуры параллельных вычислений CUDA.
Для начала я распараллелил алгоритм для CPU. Для этого в функцию передается диапазон для перебора значений внешнего цикла a. Затем весь диапазон 1..N разбивается на кусочки и скармливается ядрам процессора.
В моем случае было 4 ядра, и вычисления я выполнял через чуть-чуть доработанный пул потоков, который взял за основу тут. Так как никаких зависимостей между данными нет, скорость возрастает практически пропорционально количеству ядер.
Для GPU получается немного сложнее. Внешние два цикла перебора (a и b) будут развернуты, а в функции останется только перебор двух внутренних циклов (c и d). Можно себе представить что программа будет выполняться параллельно для всех значений а = 1..N и b = 1..N. На самом деле конечно это не так, все исполняться параллельно они все-таки не смогут, но это уже забота CUDA максимально распараллелить выполнение, этому можно помочь, если правильно настроить конфигурацию блоков и потоков.
Код для GPU получается вполне прямолинейный и очень похожий на CPU, только бинарный поиск приходится сделать руками.
Так же я реализовал оптимизации, подсказанные тут.
Замеры для CPU делались по два раза, так как они оказались намного стабильнее GPU, которых я делал уже по семь и выбирал лучшее время. Разброс для GPU мог быть двукратным, это я объясняю тем что для CUDA-расчетов использовался единственный в системе видеоадаптер.
И этим адреналином будет Profile-guided optimization.
Для PGO я привожу только время CPU, потому что GPU мало изменилось, и это ожидаемо.
Видно что PGO смогло ускорить оптимизацию #1 на целых 18%, для остальных прирост оказался скромнее.
Забыл упомянуть забавную особенность. Сначала я искал первое решение, и поставил return сразу после printf. Так вот это снижало производительность на порядок. Когда стал искать все решения, то производительность резко скакнула вверх.
Для CPU попробовать развернуть циклы, считать по два числа за раз, использовать векторные команды процессора.
Для GPU поиграть с порядком вычислений, попробовать по-экономить на регистрах, лучше подобрать параметры блоков и потоков.
1. Писать на CUDA просто.
2. Без ухищрений, прямолинейный код на CUDA оказался быстрее CPU реализации в десятки раз.
3. В числодробильных задачах GPU сильно быстрее CPU за те же деньги.
4. GPU требуется много времени на «раскачку» и для совсем коротких задач ускорения может не быть.
5. Более сложный алгоритм может оказаться быстрее для CPU, но медленнее для GPU.
6. Профильная оптимизация хорошо ускоряет код и уменьшает размер файла одновременно.
Традиционно можно потрогать код на GitHub
Мне же было интересно узнать как покажет себя GPU, и я сравнил однопоточный код с многопоточным для CPU и совсем многопоточным для GPU, с помощью архитектуры параллельных вычислений CUDA.
Железо и компиляторы
CPU Core i5 4440
GPU GeForce GTX 770
MSVS 2015 update 3
tookit CUDA 8.
Конечно, сборка была релизная и 64 битная, с оптимизацией /02 и /LTCG.
GPU GeForce GTX 770
MSVS 2015 update 3
tookit CUDA 8.
Конечно, сборка была релизная и 64 битная, с оптимизацией /02 и /LTCG.
Отключение сброса видеоадаптера
Так как вычисления длятся более двух секунд, видеодрайвер завершает CUDA-программу. Чтобы этого не происходило, нужно указать достаточное время до сброса через ключ реестра
HKEY_LOCAL_MACHINE\SYSTEM\CurrentControlSet\Control\GraphicsDrivers\TdrDelay и перезагрузить компьютер.
HKEY_LOCAL_MACHINE\SYSTEM\CurrentControlSet\Control\GraphicsDrivers\TdrDelay и перезагрузить компьютер.
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 максимально распараллелить выполнение, этому можно помочь, если правильно настроить конфигурацию блоков и потоков.
Блоки и потоки
blocks_x — это диапазон перебора первого цикла a.
А вот перебор второго цикла b пришлось сделать сложнее из-за ограничения видеокарты на количество потоков (1024), и поместить счетчик одновременно в threads и blocks_y:
a и b вычисляется так:
Это точно не оптимальная конфигурация, так как получается много холостых циклов. Но дальше подстраивать значения я не стал.
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);
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%, для остальных прирост оказался скромнее.
Тут водятся драконы
Забыл упомянуть забавную особенность. Сначала я искал первое решение, и поставил return сразу после printf. Так вот это снижало производительность на порядок. Когда стал искать все решения, то производительность резко скакнула вверх.
__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);
return; <- это дало замедление на порядок
}
}
}
}
Что можно сделать еще
Для CPU попробовать развернуть циклы, считать по два числа за раз, использовать векторные команды процессора.
Для GPU поиграть с порядком вычислений, попробовать по-экономить на регистрах, лучше подобрать параметры блоков и потоков.
Выводы
1. Писать на CUDA просто.
2. Без ухищрений, прямолинейный код на CUDA оказался быстрее CPU реализации в десятки раз.
3. В числодробильных задачах GPU сильно быстрее CPU за те же деньги.
4. GPU требуется много времени на «раскачку» и для совсем коротких задач ускорения может не быть.
5. Более сложный алгоритм может оказаться быстрее для CPU, но медленнее для GPU.
6. Профильная оптимизация хорошо ускоряет код и уменьшает размер файла одновременно.
Исходники
Традиционно можно потрогать код на GitHub