Русский
Русский
English
Статистика
Реклама

Gpu вычисления

Я вас графония принес! Как нейросеть может улучшить разрешение в старых играх до HD

25.06.2020 18:18:39 | Автор: admin

Данная статья является вольным переводом моей статьи на Medium

Я детстве я любил играть на компьютере. Совсем маленьким я застал несколько игра на кассетном ZS Spectrum, однако настоящим открытием стали красочные DOS игры 90x годов. Тогда же и зародилось большинство существующих жанров. Немного поностальгировав, я решил вспомнить молодость и запустить одну из старых игр на эмуляторе Dosbox и был неприятно поражен гигантскими пикселями и низким разрешением. Хотя в крупнопиксельной старой графике может быть свое очарование, многих сейчас не устраивает такое качество.
Для повышения разрешения и избавления от угловатости в играх в настоящее время используются различные алгоритмы постпроцессинга и сглаживания (подробно можно почитать, например тут zen.yandex.ru/media/id/5c993c6021b68f00b3fe919c/kak-rabotaet-sglajivanie-v-kompiuternyh-igrah-5c9b3e76d82a083cc9a0f1a7 ), но алгоритмы сглаживания приводят к всем ненавистной мыльной картинке, которая часто еще менее предпочтительна, чем угловатость больших пикселей.

При этом улучшение качества графики часто критично для геймеров. Перерисовка текстур для HD игре Heroes 3 заняла около полугода в 2014 году у компании Ubisoft и вызвало всплеск интереса к данной игре. Из недавних новостей переиздание первой CNC в hd графике. Подробно можно увидеть рос интереса к переизданным в hd графике на google trends:
trends.google.com/trends/explore?date=2012-06-01%202020-06-25&q=%2Fm%2F056_97,%2Fm%2F065pfn

Не для каждой игры имеет смысл заморачиваться и переиздавать в высоком качестве перерисовка текстур -занятие затратное. Но можно попробовать улучшить качество графики использую технологию суперразрешения (superresolution). Идея superresolution лежит в улучшении разрешения изображения путем дорисовки недостающих пикселей нейросетью на основании имеющихся данных. Сейчас достигнуты впечатляющие результаты, вызывающие ассоциацию со разобранной на мемы сценой улучшения изображения из фильм bladerunner https://www.youtube.com/watch?v=hHwjceFcF2Q .

Технология superresolution улучшает визуальное восприятие картинки, например вот github.com/tg-bomze/Face-Depixelizer, однако привносит новую информацию в изображение. И может быть использована для улучшения качества фильмов
,
. Однако большинство алгоритмов ресурсоемки, а мне хотелось создать скрипт для улучшения игры в реальном времени.
Немного теории
Все нижесказанное будет относиться к сверточным нейросетям подтипу нейросетей, использумому для работы с изображениями. Для начала рассмотрим, как работает нейросеть для решения задачи superresolution. Задача очень похожа на решение задачи автоэнкодера . На вход сети необходимо подать изображение, на выходе получить такое же изображение. Однако автоэнкодеры обычно используют для решения задачи эффективного сжатия данных, поэтом особенностью их архитектуры является Bottleneck бутылочное горлышко, то есть слой сети с небольшим количеством нейронов. Наличие такого слоя затставляет оставшиеся части обучаться для эффективного кодирования и раскодирования информации. Для обучения сетей superresolution разрешение высококачественного изображения сначала намеренно уменьшается и подается на вход нейросети. На выходе ожидается исходное изображение в высоком качестве. Задача superresolution определяет архитектуру используемых сетей:
обычно в них присутствует связь между входными и выходными данными (skip-connection), сильно ускоряющая обучение. Размер пикселя входных данных увеличивается и прибавляется к выходу сверточной сети. Таким образом, фактически не нужно обучать сеть превращать изображение в почти такое же. Нужно лишь обучить ее дорисовывать разницу между увеличенным при помощи увеличения размера пикселя и реальным изображением. Идея наличия skip connections различного уровня и через разное количество слоев чрезвычайно эффективна и привела к появлению класса сетей Residual Networks. Сейчас подобные связи используются почти во всех популярных арзитектурах. Неплохой обзор state-of-art архитектур для решения задач superresolution можно посмотреть тут (http://personeltest.ru/aways/github.com/krasserm/super-resolution ). Моей же задачей было создать нейросеть для решения задачи superresolution в реальном времени.

Сначала мной была выбрана архитектура edsr с 4 блоками (обычно используют более 16 слоев из блоков) ResNet c увеличением разрешения в 4 раза. Тут я частично воспользовался наработками из github.com/krasserm/super-resolution и генератором данных из этого же проекта.
Общая архитектура сети показана на схеме. Каждый блок изображение X*Y*N, где ширина соответствует числу каналов. Переходы соответствуют сверткам 3x3 (в случае res блоков с последующей активацией ReLU для нелинейности). Шаг Upscaling увеличение размерности за счет уплощения каналов.

Если приглядеться к такой небольшой сети, то можно заметить, что последние блоки линейные и наличие нескольких шагов upscaling неоправданно. Можно без потери качества заменить их на один слой с гораздо меньшим количеством фильтров без потери качества работы сети:

Оказалось, что такая сеть практически так же справляется с задачей superresolution, но гораздо быстрее и может быть запущена для обработки небольшого по разрешению видео в реальном времени.
Высококачественные фото были скачаны из data.vision.ee.ethz.ch/cvl/DIV2K. В принципе можно использовать любые фото для обучения.
import tensorflow as tffrom model import resolve_singlefrom model. rtsr import rtsrimport numpy as npfrom utils import load_image, plot_samplefrom data import DIV2Kfrom tensorflow.keras.optimizers.schedules import PiecewiseConstantDecay# Model creationrtsr_model= rtsr(scale=4, num_res_blocks=4)#Data loading and generator creationtrain_loader = DIV2K(scale=4,                                  downgrade='bicubic',                     subset='train')                           train_ds = train_loader.dataset(batch_size=16,                                random_transform=True,                                repeat_count=None)  valid_loader = DIV2K(scale=4,                       downgrade='bicubic',                     subset='valid')                       valid_ds = valid_loader.dataset(batch_size=16,           # use batch size of 1 as DIV2K images have different size                                random_transform=True, # use DIV2K images in original size                                repeat_count=40)         # 1 epoch#Saving checkpoints during trainingcbcks=tf.keras.callbacks.ModelCheckpoint('callback_x4{epoch:02d}-{loss:.2f}', monitor='loss', verbose=0, save_best_only=False)learning_rate=PiecewiseConstantDecay(boundaries=[20000], values=[1e-4, 5e-5])#model compilation before trainingrtsr_model.compile(loss='mae', optimizer=tf.keras.optimizers.Adam(learning_rate=learning_rate))#Trainingrtsr_model.fit(train_ds,validation_data=valid_ds, steps_per_epoch=3000, epochs=100, validation_steps=10,callbacks=[cbcks])#saving weightsrtsr_model.save_weights('weights/ myedsr-lastdiv2_4block-x4_weights.h5')


На выходе получаем относительно небольшую (2.6 МБ) нейросеть с простой архитектурой. При этом проверка дает малозаметное отличие от предобученной 16 блочной сети:

Изображение слева исходное, справа улучшенное при помощи 16 слойной edsr, в середине с помощью rtsr. Разница между 2 и 3 с моей точки зрения несущественна.

Инференс

github.com/Alexankharin/RTSR
Полученную сеть запустим на видеокарте (у меня GTX 1060) с поддержкой cudnn (http://personeltest.ru/aways/developer.nvidia.com/cudnn ) для высокой производительности.
Pipeline для инференса выглядит следующим образом:

1. Захват изображения из области или из окна
2. Улучшение изображения
3. Отрисовка улучшенного изображения в новом окне

При тестировании я обнаружил, хотя при запуске большинства игр в эмуляторе dosbox разрешение составляет 640x480 пикселей, однако чаще всего это просто увеличенные в размере пиксели разрешения 320x240(и позже нашел подробности www.dosgamers.com/dos/dosbox-dos-emulator/screen-resolution ), что приводит ук необходимости в некоторых случаях делать downscale в 2 раза перед обработкой.
Захват скриншотов производится с использованием библиотеки mss (для OS ubuntu) или d3dshot (быстрее для windows).
Отображение с использованием opencv-python. Для закрытия окна необходимо сделать его активным и нажать на клавишу q. Управлением захватом улучшаемой области при помощи клавиш WSAD и IKJL.
Скрипт написан в файле superres_win.py
Результаты:
на ноутбуке с GTX 1060 (3Gb) и OS Windows10 скрипт выдает 14-15 FPS, что достаточно для квестов или стратегий, но немного маловато для платформеров. Кроме того, при запуске RTSR на стационарном ПК с OS Ubuntu и такой же видеокартой падал до 10-12 (почему пока не разбирался). Судя по бенчмаркам ai-benchmark.com/ranking_deeplearning_detailed.html 1080 на схожих задачах должна дать FPS около 25, что близко к оригинальным значениям и достаточно для комфортной игры. Пример улучшения графики можно увидеть на видео:


Примеры улучшения:
MegaManX:


Legend of Kyrandia



Wolf3d


Heroes of might and magic


Больше примеров можете попробовать сами

Как запустить?
Для быстрой работы необходима видеокарта с поддержкой cuda и cudnn (http://personeltest.ru/aways/developer.nvidia.com/cuda-gpus ) и установленными библиотеками cuda/cudnn. Нужен установленный ptyhon 3.7 и tensorflow (версия выше 2.0 с поддержкой gpu). Это может быть сложной задачей, и могут возникнуть проблемы совместимости (http://personeltest.ru/aways/www.tensorflow.org/install/source_windows ).
Простейшим способом может быть установка дистрибутива Anaconda (http://personeltest.ru/aways/www.anaconda.com/products/individual ), а затем в установка tensorflow-gpu При помощи conda:
conda install tensorflow-gpu

Если не получится из-за конфликтов, то можно попробовать
conda install cudnn
pip install tensorflow-gpu


должно сработать.
Остальные библиотеки можно установить при помощи pip:
pip install opencv-python
pip install pywin32
pip install mss


Дальее необходимо запустить скрипт командой
python superres.py

Управление окном захвата проводится при помощи клавиш wsad и ijkl (изменение размера). q- закрытие окна. 0 -включение и выключение режима superresolution. Цифры 1 и 2 режим изображения (2 по умолчанию означает, что в игре используется увеличение за счет увеличения пикселя в 2 раза)
Подробнее..

Game of Life с битовой магией, многопоточностью и на GPU

03.07.2020 20:16:17 | Автор: admin

Всем привет!


Недавняя статья на Хабре в очередной раз показала неостывающий интерес к игре Жизнь в частности и всевозможным оптимизациям в общем. Статья и комментарии к ней, особенно любопытство к вычислениям на GPU, вдохновили меня на то, чтобы поделиться своими изысканиями на данном поприще и, забегая вперёд, скажу, что повествование пойдёт о расчётах на GPU, битовой магии, многопоточности и огромных полях для игры Жизнь, порядка миллиарда клеток.



О себе


Пару слов о себе и проекте. Вот уже несколько лет моим хобби и pet-проектом является написание игры Жизнь, в ходе которого я разбираюсь в интересующих меня темах. Так, сперва я сделал её с помощью библиотеки glut и прикрутил многопользовательский режим, будучи вдохновлённым архитектурой peer-to-peer. А около двух лет назад решил начать всё с нуля, используя Qt и вычисления на GPU.

Масштабы


Идея нового проекта заключалась в том, чтобы сделать кроссплатформенное приложение, в первую очередь для мобильных устройств, в котором пользователи смогут окунуться в игру Жизнь, познакомиться с разнообразием всевозможных паттернов, наблюдать за причудливыми формами и комбинациями этой игры, выбрав скорость от 1 до 100 мс на шаг и задав размер игрового поля вплоть до 32'768 x 32'768, то есть 1'073'741'824 клетки.

На всякий случай напомню об игре Жизнь. Игра происходит на плоскости, поделённой на множество клеток. Клетки квадратные, соответственно, для каждой клетки есть 8 соседей. Каждая клетка может быть либо живой, либо мёртвой, то есть пустой. В рамках игры пользователь заполняет клетки жизнью, выставляя на поле заинтересовавшие его комбинации клеток паттерны.
Далее процесс шаг за шагом развивается по заданным правилам:
  • в пустой клетке, рядом с которой ровно 3 живые клетки, зарождается жизнь
  • если у живой клетки есть 2 или 3 живых соседки, то эта клетка продолжает жить
  • если у живой клетки меньше 2 или больше 3 живых соседки, клетка умирает

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

В моей реализации поле зациклено, то есть после 32'767-й клетки вновь следует 0-я. Таким образом, перемещающиеся паттерны способны обогнуть поле и оказаться в точке, где они были размещены изначально. Забавный эффект случается с паттерном планёрное ружьё, когда выпущенные им планёры в конечном итоге врезаются в само ружьё и разрушают его, этакое самоубийство в мире клеточных автоматов:

Gosper glider gun


Архитектура


Что касается реализации, то в первую очередь хочется отметить, что каждая клетка в игровом поле представляет собой всего лишь 1 бит в памяти. Так, при выборе размера игрового поля в памяти выделяется буфер размером n^2 / 8 байт. Это улучшает локальность данных и снижает объём потребляемой памяти, а главное позволяет применить ещё более хитроумные оптимизации, о которых поговорим ниже. Игровое поле всегда квадратное и его сторона всегда степени 2, в частности затем, чтобы без остатка осуществлять деление на 8.

Архитектурно выделяются два слоя, ответственные за вычисления:
  • низкий уровень платформозависимый; на текущий момент существует реализация на Metal API графическое API от Apple позволяет задействовать GPU на устройствах от Apple, а также fallback-реализация на CPU. Одно время существовала версия на OpenCL. Планируется реализация на Vulkan API для запуска на Android
  • высокий уровень кроссплатформенный; по сути класс-шаблонный метод, делегирующий реализацию нужных методов на низкий уровень

Низкий уровень


Задача низкого уровня заключается непосредственно в расчёте состояния игрового поля и устроена следующим образом. В памяти хранятся два буфера n^2 / 8 байт. Первый буфер служит как input текущее состояние игрового поля. Второй буфер как output, в который записывается новое состояние игрового поля в процессе расчётов. По завершении вычислений достаточно сделать swap буферов и следующий шаг игры готов. Такая архитектура позволяет распараллелить расчёты в силу константности input. Дело в том, что каждый поток может независимо обрабатывать клетки игрового поля.

За счёт многопоточности GPU-подход добивается максимальной эффективности. На CPU также происходит распараллеливание, но, соответственно, с гораздо меньшим числом потоков и меньшей эффективностью, что отражено в секции бенчмарков ниже. Сам алгоритм идентичен для всех реализаций и заключается в следующем.

Назовём атомарной операцией расчёт одного байта буфера, хранящего игровое поле, то есть расчёт восьми клеток. Этот расчёт совершается одним потоком и требует информации о соседних восьми байтах. Попробуем схематично изобразить один байт следующим образом, где точка это бит:

[........]


А участок игрового поля с интересующим нас байтом и его восемью соседними байтами изобразим так:

[........][........][........]
[........][********][........]
[........][........][........]


Исходя из рисунка, можно догадаться, что все соседние клетки можно уложить в один uint64_t (назовём его neighbours), а ещё в одном uint8_t (назовём его self), будет храниться информация о соседях в самом обрабатываемом байте.

Рассмотрим на примере расчёт 0-го бита целевого байта. Звёздочкой отметим интересующий бит, а нулём соседние для него биты:

[.......0][00......][........]
[.......0][*0......][........]
[.......0][00......][........]


Пример посложнее. Здесь звёздочкой отмечены 0-й, 3-й и 7-й биты целевого байта и соответствующими числами соседние биты:

[.......0][00333.77][7.......]
[.......0][*03*3.7*][7.......]
[.......0][00333.77][7.......]


Думаю, кто-то из читателей уже догадался, в чём заключается магия. Имея указанную информацию, остаётся лишь составить битовые маски для каждого бита целевого байта и применить их к neighbours и self соответственно. Результатом станут 2 значения, сумма единичных бит которых будет характеризовать число соседей, что можно интерпретировать как правила игры Жизнь: 2 или 3 бита для поддержания жизни в живой клетке и 3 бита для зарождения новой жизни в пустой клетке. В противном случае клетка остаётся/становится пустой.

Описанные вычисления на нижнем уровне происходят каждый шаг игры, и поле всегда обрабатывается целиком вне зависимости от его заполненности и без трюков с мемоизацией а-ля Hashlife.

После заполнения всего выходного буфера игровое поле считается рассчитанным.
Код Metal-шейдера по обработке 1 байта
// Бросается в глаза C++-подобный синтаксис#include <metal_stdlib>#include <metal_integer>using namespace metal;// Вспомогательные функции для вычисления позиции клеткиushort2 pos(uint id) {   return ushort2(id % WIDTH, id / HEIGHT); }uint idx(ushort2 pos) {   return pos.x + pos.y * HEIGHT; }ushort2 loopPos(short x, short y) {   return ushort2((x + WIDTH) % WIDTH, (y + HEIGHT) % HEIGHT); }// Битовые маски для вычисления соседей интересующего битаtemplate<uint Bit> struct Mask {   constant constexpr static uint c_n_e_s_w = 0x70007 << (Bit - 1);   constant constexpr static uint c_nw_ne_se_sw = 0x0;   constant constexpr static uint c_self = 0x5 << (Bit - 1); };template<> struct Mask<0> {   constant constexpr static uint c_n_e_s_w = 0x80030003;   constant constexpr static uint c_nw_ne_se_sw = 0x80000080;   constant constexpr static uint c_self = 0x2; };template<> struct Mask<7> {   constant constexpr static uint c_n_e_s_w = 0xC001C0;   constant constexpr static uint c_nw_ne_se_sw = 0x10100;   constant constexpr static uint c_self = 0x40; };// Для указанного бита функция вычисляет состояние клетки в зависимости от её соседей, применяя соответствующие биту маскиtemplate<uint Bit>uint isAlive(uint self, uint n_e_s_w, uint nw_ne_se_sw) {   /*  [.......0][00333.77][7.......]  [.......0][*03*3.7*][7.......]  [.......0][00333.77][7.......]  */  // До определённой версии в Metal не было 64-битного целого, поэтому составляются две маски  uint neighbours = popcount(Mask<Bit>::c_n_e_s_w & n_e_s_w)     + popcount(Mask<Bit>::c_nw_ne_se_sw & nw_ne_se_sw)     + popcount(Mask<Bit>::c_self & self);   return static_cast<uint>((self >> Bit & 1) == 0     ? neighbours == 3     : neighbours == 2 || neighbours == 3) << Bit;}// Язык Metal даже умеет в шаблонную магиюtemplate<uint Bit>uint calculateLife(uint self, uint n_e_s_w, uint nw_ne_se_sw) {   return isAlive<Bit>(self, n_e_s_w, nw_ne_se_sw)     | calculateLife<Bit - 1>(self, n_e_s_w, nw_ne_se_sw); }template<>uint calculateLife<0>(uint self, uint n_e_s_w, uint nw_ne_se_sw){  return isAlive<0>(self, n_e_s_w, nw_ne_se_sw); }// Главная функция compute-шейдера. На вход подаются два буфера, о которых речь шла выше - константный input и output, а также id - координата целевого байтаkernel void lifeStep(constant uchar* input [[buffer(0)]],                             device uchar* output [[buffer(1)]],                             uint id [[thread_position_in_grid]]) {   ushort2 gid = pos(id * 8);   // Вычисляем соседние байты  uint nw = idx(loopPos(gid.x - 8, gid.y + 1));   uint n  = idx(loopPos(gid.x,     gid.y + 1));   uint ne = idx(loopPos(gid.x + 8, gid.y + 1));   uint e  = idx(loopPos(gid.x + 8, gid.y    ));   uint se = idx(loopPos(gid.x + 8, gid.y - 1));   uint s  = idx(loopPos(gid.x    , gid.y - 1));   uint sw = idx(loopPos(gid.x - 8, gid.y - 1));   uint w  = idx(loopPos(gid.x - 8, gid.y    ));   // Вычисляем байт с целевым битом  uint self = static_cast<uint>(input[id]);   // Подготавливаем битовые маски с соседями  // north_east_south_west  uint n_e_s_w = static_cast<uint>(input[n >> 3]) << 0 * 8     | static_cast<uint>(input[e >> 3]) << 1 * 8     | static_cast<uint>(input[s >> 3]) << 2 * 8     | static_cast<uint>(input[w >> 3]) << 3 * 8;   // north-west_north-east_south-east_south-west  uint nw_ne_se_sw = static_cast<uint>(input[nw >> 3]) << 0 * 8     | static_cast<uint>(input[ne >> 3]) << 1 * 8     | static_cast<uint>(input[se >> 3]) << 2 * 8     | static_cast<uint>(input[sw >> 3]) << 3 * 8;     // В этой строчке рассчитываются все 8 клеток обрабатываемого байта  output[id] = static_cast<uchar>(calculateLife<7>(self, n_e_s_w, nw_ne_se_sw)); };


Высокий уровень


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

Отрисовка клеток выполняется средствами Qt, а именно посредством заполнения пикселей в QImage. Интерфейс выполнен в QML. Пиксели заполняются лишь для небольшой области видимого игроку игрового поля. Таким образом удаётся избежать лишних затрат ресурсов на отрисовку.

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

Бенчмарки


Результаты замеров скорости работы игры для разных машин в зависимости от размеров игрового поля приведены в таблицах. Полученные значения в таблицах указаны в миллисекундах. При проведении тестов игровое поле заполнялось жизнью на 50% с помощью генератора случайных чисел.

В таблицах есть замеры скорости работы нижнего и верхнего уровня, причём берутся минимальные значения. А для более общей картины приведены замеры скорости вычисления 100 шагов игры и среднего времени выполнения шага. Последнее более показательно, так как включает в себя время на отрисовку и прочие сопутствующие операции.

MacBook Pro 2014
Processor 2,6 GHz Dual-Core Intel Core i5
Memory 8 GB 1600 MHz DDR3
Graphics Intel Iris 1536 MB

GPU реализация
1024 2048 4096 8192 16384 32768
Низкий уровень (min) 0 0 2 9 43 170
Высокий уровень (min) 0 0 0 1 12 55
100 шагов 293 446 1271 2700 8603 54287
Время на шаг (avg) 3 4 13 27 86 542

CPU реализация
1024 2048 4096 8192 16384 32768
Низкий уровень (min) 3 25 117 552 2077 21388
Высокий уровень (min) 0 0 0 1 4 51
100 шагов 944 3894 15217 65149 231260 -
Время на шаг (avg) 9 39 152 651 2312 -

MacBook Pro 2017
Processor 2.8 GHz Intel Core i7
Memory 16 GB 2133 MHz LPDDR3
Graphics Intel HD Graphics 630 1536 MB

GPU реализация
1024 2048 4096 8192 16384 32768
Низкий уровень (min) 0 0 0 2 8 38
Высокий уровень (min) 0 0 0 0 3 13
100 шагов 35 67 163 450 1451 5886
Время на шаг (avg) 0 1 2 5 15 59

CPU реализация
1024 2048 4096 8192 16384 32768
Низкий уровень (min) 1 7 33 136 699 2475
Высокий уровень (min) 0 0 0 0 3 18
100 шагов 434 1283 4262 18377 79656 264711
Время на шаг (avg) 4 13 43 184 797 2647

Видно, что на стареньком макбуке процессор совсем не справляется с огромным игровым полем и приложение становится неюзабельным. Даже процессор нового макбука едва вытягивает такой объём вычислений. Зато видеокарта значительно превосходит по производительности процессор как на старом, так и на новом железе. С использованием GPU новый макбук предлагает вполне комфортный пользовательский опыт даже на огромных игровых пространствах.

Итог


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

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

В будущем хочется провести бенчмарки на мобильных устройствах. Проект находится на ранней стадии разработки, а потому Apple-Developer аккаунта для проведения таких тестов сейчас у меня нет.

Спасибо за внимание! Буду рад любым комментариям к статье и к проекту:
Код на GitHub.

Код Metal реализации нижнего уровня

Код CPU реализации нижнего уровня

Код верхнего уровня
Подробнее..

Boost.Compute или параллельные вычисления на GPUCPU. Часть 2

16.08.2020 16:17:46 | Автор: admin

Вступление


Привет, Хабр!

Предыдущая часть понравилась многим, поэтому я снова перелопатил половину документации boost и нашёл о чем написать. Очень странно что вокруг boost.compute нету такого же ажиотажа как и вокруг boost.asio. Ведь достаточно, того эта библиотека кроссплатформенная, так ещё и предоставляет удобный (в рамках c++) интерфейс взаимодействия с параллельными вычислениями на GPU и CPU.

Все части




Содержание


  • Платформы
  • Асинхронные операции
  • Пользовательские функции
  • Сравнение скорости работы разных устройств в разных режимах
  • Заключение


Платформы


Не смотря на то что boost.compute кроссплатформенный инструмент, работает он не со всеми вычислительными устройствами. Возможно всё зависит от версии OpenCL, точно сказать не могу, но проверить у себя, с чем вы сможете работать, можно следующим кодом:
auto platforms = compute::system::platforms();for (size_t i = 0; i < platforms.size(); i++){cout << platforms[i].name() << endl;}


Я получил такой результат:


Экспериментальным путём я выяснил что собранные на моём ноутбуке программы будут запускаться лишь на устройствах таких же производителей.

Асинхронные операции


Казалось бы, куда ещё быстрее? один из способов ускорить работу с контейнерами пространства имён compute это использование асинхронных функций. Boost.compute предоставляет нам несколько инструментов. Из них класс compute::future для контроля использования функций и функции copy_async(), fill_async() для копирования или заполнения массива. Конечно, существуют ещё и инструменты для работы с событиями, но их рассматривать нет необходимости. Дальше будет пример использования всего выше перечисленного:
auto device = compute::system::default_device();auto context = compute::context::context(device);auto queue = compute::command_queue(context, device);std::vector<int> vec_std = {1, 2, 3};compute::vector<int> vec_compute(vec_std.size(), context);compute::vector<int> for_filling(10, context);int num_for_fill = 255;compute::future<void> filling = compute::fill_async(for_filling.begin(), for_filling.end(), num_for_fill, queue); // асинхронно заполняет заданный векторcompute::future<void> copying = compute::copy_async(vec_std.begin(), vec_std.end(), vec_compute.begin(), queue); // асинхронно копирует следующий векторfilling.wait();copying.wait();


Пояснять тут особо нечего. Первые три строчки стандартная инициализация необходимых классов, потом два векторы для копирования, вектор для заполнения, переменная которой будем заполнять предыдущий вектор и непосредственно функции для заполнения и копирования соответственно. Потом дожидаемся их выполнения.
Для тех, кто работал с std::future из STL, тут абсолютно всё тоже самое, только в другом пространстве имён и нет аналога std::async().

Пользовательские функции для вычислений



В предыдущей части я сказал, что поясню как использовать свои собственные методы для обработки массива данных. Я насчитал 3 способа как это можно сделать: использовать макрос, использовать make_function_from_source<>() и использовать специальный фреймворк для лямбда выражений.

Начну с самого первого варианта макроса. Сначала приложу пример кода а потом поясню как работает.
BOOST_COMPUTE_FUNCTION(float, add,(float x, float y),{ return x + y; });

Первым аргументом указываем тип возвращаемого значения, потом название функции, её аргументы и тело функции. Дальше под именем add, данную функцию можно использовать например в функции compute::transform(). Использование этого макроса очень похоже на обычное лямбда выражение, но я проверял, они работать не будут.

Второй и, наверное, самый сложный способ очень похож на первый. Я смотрел код предыдущего макроса и оказалось, что он использует именно второй способ.
compute::function<float(float)> add = compute::make_function_from_source<float(float)>("add", "float add(float x, float y) { return x + y; }");

Здесь всё очевидней чем может показаться на первый взгляд, функция make_function_from_source(), использует всего два аргумента, один из которых название функции, а второй её реализация. После объявления функции её можно использовать так же как и после реализации макросом.

Ну и последний вариант это фреймворк для лямбда выражений. Пример использования:
compute::transform(com_vec.begin(),      com_vec.end(),      com_vec.begin(),      compute::_1 * 2,      queue);

Четвёртым аргументом мы указываем что хотим умножить каждый элемент из первого вектора на 2, всё достаточно просто и делается на месте.

Этим же способом можно указывать логические выражения. Например в методе compute::count_if():
std::vector<int> source_std = { 1, 2, 3 };compute::vector<int> source_compute(source_std.begin() ,source_std.end(), queue);auto counter = compute::count_if(source_compute.begin(), source_compute.end(), compute::lambda::_1 % 2 == 0, queue);

Таким образом мы посчитали все чётные числа в массиве, counter будет равен единице.

Сравнение скорости работы разных устройств в разных режимах



Ну и последнее, про что я хотел бы написать в этой статье, это сравнение скорости обработки данных на разных устройствах и в разных режимах(только для CPU). это сравнение докажет, когда есть смысл использовать GPU для вычислений и параллельные вычисления в целом.

Тестировать я буду так: с помощью compute для всех устройств вызову функцию compute::sort() для того чтоб отсортировать массив из 100 млн. значений типа float. Для теста однопоточного режима вызову std::sort для массива такого же размера. Для каждого устройства засеку время в миллисекундах с помощью стандартной библиотеки chrono и выведу всё в консоль.

Получился такой результат:


Теперь сделаю всё тоже самое только для тысячи значений. На этот раз время будет в микросекундах.


На этот раз процессор в однопоточном режиме опередил всех. Из этого делаем вывод что такого рода операции стоит делать только когда речь идёт о действительно больших данных.

Заключение


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

Буду рад позитивным отзывам, спасибо за уделённое время!

Всем удачи!
Подробнее..
Категории: C++ , Gpu вычисления , Boost::compute

Сравнение времени выполнения алгоритма на CPU и GPU

31.10.2020 16:04:39 | Автор: admin

Использование CUDA Runtime API для вычислений. Сравнение CPU и GPU вычислений.


В данной статье я решил провести сравнение выполнения алгоритма написанного на C++ на центральном и графическом процессоре(выполнение вычислений с помощью Nvidia CUDA Runtime API на поддерживаемом GPU Nvidia). CUDA API позволяет выполнение некоторых вычислений на графическом процессоре. Файл c++ использующий cuda, будет иметь расширение .cu.
Схема работы алгоритма приведена ниже.

Задача алгоритма состоит в том, что найти возможные числа X, при возведении которых в степень degree_of, будет получатся исходное число max_number. Сразу отмечу, что все числа которые будут передаваться GPU, будут хранится в массивах. Алгоритм, выполняемый каждым потоком, имеет приблизительно следующий вид:

intdegree_of=2;intdegree_of_max=Number_degree_of_max[0];//Массивхранящийзначениемаксимальнойстепеничислаintx=thread;//номервыполняемогопотокаintmax_number=INPUT[0];//Массивхранящийчисло,котороенеобходимополучитьintNumber=1;intDegree;boolBREAK=false;//Переменнаядлязавершенияwhilewhile(degree_of<=degree_of_max&&!BREAK){Number=1;for(inti=0;i<degree_of;i++){Number*=x;Degree=degree_of;}if(Number==max_number){OUT_NUMBER[thread]=X;//OUT_NUMBERМассивхранящийчислакоторыенеобходимовозвестивстепеньDegreeдляполученияисходногочислаOUT_DEGREE[thread]=Degree;//OUT_DEGREEМассивхранящийстепеньвкоторуюнужновозвестичислоXдляполученияисходногочисла}degree_of++;//Вслучаевыходазапредел:if(degree_of>degree_of_max||Number>max_number){BREAK=true;}}

Код для выполнения на CPU
#include <iostream>#include<vector>#include<string>//необходимо для getline#include<thread>#include<fstream>using namespace std;int Running_thread_counter = 0;void Upload_to_CPU(unsigned long long  *Number, unsigned long long  *Stepn, bool *Stop,unsigned long long  *INPUT, unsigned long long  *max, int THREAD);void Upload_to_CPU(unsigned long long  *Number, unsigned long long  *Stepn, bool *Stop,unsigned long long  *INPUT, unsigned long long  *max, int THREAD) {int thread = THREAD;Running_thread_counter++;unsigned long long  MAX_DEGREE_OF = max[0];int X = thread;unsigned long long  Calculated_number = 1;unsigned long long  DEGREE_OF = 2;unsigned long long   INP = INPUT[0];Stop[thread] = false;bool BREAK = false;if (X != 0 && X != 1) {while (!BREAK) {if (DEGREE_OF <= MAX_DEGREE_OF) {Calculated_number = 1;for (int counter = 0; counter < DEGREE_OF; counter++) {Calculated_number *= X;}if (Calculated_number == INP) {Stepn[thread] = DEGREE_OF;Number[thread] = X;Stop[thread] = true;BREAK = true;}DEGREE_OF++;}else { BREAK = true; }}}}void Parallelize_to_threads(unsigned long long  *Number, unsigned long long  *Stepn, bool *Stop,unsigned long long  *INPUT, unsigned long long  *max, int size);int main(){int size = 1000;unsigned long long  *Number = new unsigned long long[size], *Degree_of = new unsigned long long[size];unsigned long long  *Max_Degree_of = new unsigned long long[1];unsigned long long  *INPUT_NUMBER = new unsigned long long[1];Max_Degree_of[0] = 7900;INPUT_NUMBER[0] = 216 * 216 * 216;ifstream inp("input.txt");if (inp.is_open()) {string t;vector<unsigned long long>IN;while (getline(inp, t)) {IN.push_back(stol(t));}INPUT_NUMBER[0] = IN[0];//исходное числоMax_Degree_of[0] = IN[1];//значение максимальной степени}else {ofstream error("error.txt");if (error.is_open()) {error << "No file " << '"' << "input.txt" << '"' << endl;error << "Please , create a file" << '"' << "input.txt" << '"' << endl;error << "One read:input number" << endl;error << "Two read:input max stepen" << endl;error << "." << endl;error.close();INPUT_NUMBER[0] = 1;Max_Degree_of[0] = 1;}}//расскометрируйте следующий код , если хотите видеть исходные значения в окне консоли //cout << INPUT[0] << endl;bool *Elements_that_need_to_stop = new bool[size];Parallelize_to_threads(Number, Degree_of, Elements_that_need_to_stop, INPUT_NUMBER, Max_Degree_of, size);vector<unsigned long long>NUMBER, DEGREEOF;for (int i = 0; i < size; i++) {if (Elements_that_need_to_stop[i]) {if (Degree_of[i] < INPUT_NUMBER[0] && Number[i] < INPUT_NUMBER[0]) {//проверка на ошибки NUMBER.push_back(Number[i]);DEGREEOF.push_back(Degree_of[i]);}}}//расскометрируйте следующий код , если хотите вывести результаты в консоль//это может замедлить программу /*for (int f = 0; f < NUMBER.size(); f++) {cout << NUMBER[f] << "^" << DEGREEOF[f] << "=" << INPUT_NUMBER[0] << endl;}*/ofstream out("out.txt");if (out.is_open()) {for (int f = 0; f < NUMBER.size(); f++) {out << NUMBER[f] << "^" << DEGREEOF[f] << "=" << INPUT_NUMBER[0] << endl;}out.close();}}void Parallelize_to_threads(unsigned long long  *Number, unsigned long long  *Stepn, bool *Stop,unsigned long long  *INPUT, unsigned long long  *max, int size) {thread *T = new thread[size];Running_thread_counter = 0;for (int i = 0; i < size; i++) {T[i] = thread(Upload_to_CPU, Number, Stepn, Stop, INPUT, max, i);T[i].detach();}while (Running_thread_counter < size - 1);//дождаться завершения выполнения всех потоков }


Для работы алгоритма необходим текстовый файл с исходным числом и максимальной степенью.
Код для выполнения вычислений на GPU
//библиотеки cuda_runtime.h и device_launch_parameters.h//для работы с cyda#include "cuda_runtime.h"#include "device_launch_parameters.h"#include<vector>#include<string>//для getline#include <stdio.h>#include<fstream>using namespace std;__global__ void Upload_to_GPU(unsigned long long  *Number,unsigned long long  *Stepn, bool *Stop,unsigned long long  *INPUT,unsigned long long  *max) {int thread = threadIdx.x;unsigned long long  MAX_DEGREE_OF = max[0];    int X = thread;unsigned long long  Calculated_number = 1;unsigned long long  Current_degree_of_number = 2;    unsigned long long   Original_numberP = INPUT[0];Stop[thread] = false;bool BREAK = false;if (X!=0&&X!=1) {while (!BREAK) {if (Current_degree_of_number <= MAX_DEGREE_OF) {Calculated_number = 1;for (int counter = 0; counter < Current_degree_of_number; counter++) { Calculated_number*=X;}if (Calculated_number == Original_numberP) {Stepn[thread] = Current_degree_of_number;Number[thread] = X;Stop[thread] = true;BREAK = true;}Current_degree_of_number++;}else { BREAK = true; }}}}cudaError_t Configure_cuda(unsigned long long *Number, unsigned long long  *Stepn, bool *Stop,unsigned long long  *INPUT, unsigned long long  *max,unsigned int size);int main(){int size = 1000;    unsigned long long  *Number=new unsigned long long [size], *Degree_of=new unsigned long long [size];unsigned long long  *Max_degree_of = new unsigned long long [1];unsigned long long  *INPUT_NUMBER = new unsigned long long [1];   Max_degree_of[0] = 7900;ifstream inp("input.txt");if (inp.is_open()) {string text;vector<unsigned long long>IN;while (getline(inp, text)) {IN.push_back( stol(text));}INPUT_NUMBER[0] = IN[0];Max_degree_of[0] = IN[1];}else {ofstream error("error.txt");if (error.is_open()) {error<<"No file "<<'"'<<"input.txt"<<'"'<<endl;error<<"Please , create a file" << '"' << "input.txt" << '"' << endl;error << "One read:input number" << endl;error << "Two read:input max stepen" << endl;error << "." << endl;error.close();INPUT_NUMBER[0] = 1;Max_degree_of[0] = 1;}}bool *Elements_that_need_to_stop = new bool[size];    // Загрузка массивов в cudacudaError_t cudaStatus =  Configure_cuda(Number, Degree_of, Elements_that_need_to_stop, INPUT_NUMBER, Max_degree_of, size);    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "addWithCuda failed!");        return 1;    }vector<unsigned long long>NUMBER, DEGREEOF;for (int i = 0; i < size; i++) {if (Elements_that_need_to_stop[i]) {NUMBER.push_back(Number[i]);//занести в вектор числоDEGREEOF.push_back(Degree_of[i]);//занести в вектор степень числа}}//раскоментируйте следующий код , чтобы вывести результаты в консоль/*for (int f = 0; f < NUMBER.size(); f++) {cout << NUMBER[f] << "^" << DEGREEOF[f] << "=" << INPUT_NUMBER[0] << endl;}*/ofstream out("out.txt");if (out.is_open()) {for (int f = 0; f < NUMBER.size(); f++) {out << NUMBER[f] << "^" << DEGREEOF[f] << "=" << INPUT_NUMBER[0] << endl;}out.close();}    //Очистить ресурсы связанные с устройством    cudaStatus = cudaDeviceReset();    if (cudaStatus != cudaSuccess) {        fprintf(stderr, "cudaDeviceReset failed!");        return 1;    }    return 0;}cudaError_t  Configure_cuda(unsigned long long  *Number, unsigned long long *Degree_of, bool *Stop,unsigned long long *INPUT, unsigned long long *max,unsigned int size) {unsigned long long *dev_Number = 0;unsigned long long *dev_Degree_of = 0;unsigned long long *dev_INPUT = 0;unsigned long long *dev_Max = 0;bool *dev_Elements_that_need_to_stop;cudaError_t cudaStatus;// УСТАНОВКА ИСПОЛЬЗУЕМОГО GPU cudaStatus = cudaSetDevice(0);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");goto Error;}// РЕЗЕРВИРОВАНИЕ МЕСТА В ПАМЯТИ ПОД ДАННЕcudaStatus = cudaMalloc((void**)&dev_Number, size * sizeof(unsigned long long));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!dev_Number");goto Error;}cudaStatus = cudaMalloc((void**)&dev_Degree_of, size * sizeof(unsigned long long));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!dev_Degree_of");goto Error;}cudaStatus = cudaMalloc((void**)&dev_Max, size * sizeof(unsigned long long int));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!dev_Max");goto Error;}cudaStatus = cudaMalloc((void**)&dev_INPUT, size * sizeof(unsigned long long));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!dev_INPUT");goto Error;}cudaStatus = cudaMalloc((void**)&dev_Elements_that_need_to_stop, size * sizeof(bool));if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMalloc failed!dev_Stop");goto Error;}// ПЕРЕМЕЩЕНИЕ ДАННХ В ПАМЯТЬ GPUcudaStatus = cudaMemcpy(dev_Max, max, size * sizeof(unsigned long long), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}cudaStatus = cudaMemcpy(dev_INPUT, INPUT, size * sizeof(unsigned long long), cudaMemcpyHostToDevice);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}Upload_to_GPU<<<1, size>>>(dev_Number, dev_Degree_of, dev_Elements_that_need_to_stop, dev_INPUT, dev_Max);// Проверка сбоев ядраcudaStatus = cudaGetLastError();if (cudaStatus != cudaSuccess) {fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));goto Error;}// Ожидание завершения операций , выполняемых ядромcudaStatus = cudaDeviceSynchronize();if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);goto Error;}// Перемещение данных из памяти GPU в системную памятьcudaStatus = cudaMemcpy(Number, dev_Number, size * sizeof(unsigned long long), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}cudaStatus = cudaMemcpy(Degree_of, dev_Degree_of, size * sizeof(unsigned long long), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}cudaStatus = cudaMemcpy(Stop, dev_Elements_that_need_to_stop, size * sizeof(bool), cudaMemcpyDeviceToHost);if (cudaStatus != cudaSuccess) {fprintf(stderr, "cudaMemcpy failed!");goto Error;}Error://Освобождение памяти GPU от данныхcudaFree(dev_INPUT);cudaFree(dev_Degree_of);cudaFree(dev_Max);cudaFree(dev_Elements_that_need_to_stop);cudaFree(dev_Number);return cudaStatus;}



Идентификатор
__global__  
в .cu файле указывает, что функция выполняется на уровне ядра GPU.
Для работы с cyda, перед вызовом функции, нужно зарезервировать память под массив и перенести элементы в память GPU. Это увеличивает объем кода, но позволяет разгрузить CPU, так как вычисления производятся на GPU.Поэтому ,cuda, дает как минимум возможность разгрузить процессор для выполнения других нагрузок, не использующих cuda.
В случае примера на cuda, задача процессора заключается лишь в загрузке инструкций на GPU и обработке результатов пришедших с GPU; В то время как в коде для CPU, процессор обрабатывает каждый поток. Стоит отметить, что cyda имеет ограничения по количеству запускаемых потоков, поэтому в обоих алгоритмах я взял одинаковое количество потоков, равное 1000. Также, в случае с CPU я использовал переменную
intRunning_thread_counter=0;

чтобы считать количество уже выполненных потоков и дожидаться, пока все потоки не выполнятся.
Тестируемая конфигурация
  • CPU :amd ryzen 5 1400(4core,8thread)
  • ОЗУ:8гбDDR4 2666
  • GPU:Nvidia rtx 2060

  • OS:windows 10 version 2004
  • Cuda:
    • Compute Capability 7.5
    • Threads per Multiprocessor 1024
    • CUDA 11.1.70

  • GPU-Z:version 2.35.0
  • Visual Studio 2017

Сведения о cyda были взяты из GPU-Z

Для тестирования алгоритма я использовал
следующий код на C#
usingSystem;usingSystem.Collections.Generic;usingSystem.Linq;usingSystem.Text;usingSystem.Threading.Tasks;usingSystem.Diagnostics;usingSystem.IO;namespaceConsoleAppTESTSTEPEN_CPU_AND_GPU_{classProgram{staticstringUpload(Int64number,Int64degree_of){stringOUT="";string[]Chord_values=newstring[2];Int64Degree_of=degree_of;Int64Number=number;Chord_values[0]=Number.ToString();Chord_values[1]=Degree_of.ToString();File.WriteAllLines("input.txt",Chord_values);//файлвходныхданныхOUT+="inputnumber:"+Number.ToString()+"\n";OUT+="inputdegreeofnumber:"+Degree_of.ToString()+"\n";DateTimerunning_CPU_application=DateTime.Now;//записатьвремязапускапрограммыProcessproc=Process.Start("ConsoleApplication29.exe");//exeреализацияалгоритманаc++x64использующаяCPUдлявычисленийwhile(!proc.HasExited);//дождатсязавершенияпрограммыDateTimestop_CPU_application=DateTime.Now;//записатьвремяостановкипрограммыstring[]outs=File.ReadAllLines("out.txt");//получитьрезультатыFile.Delete("out.txt");OUT+="CPU:"+"\n";if(outs.Length>0){for(intj=0;j<outs.Length;j++){OUT+=outs[j]+"\n";}}else{OUT+="novalues"+"\n";}OUT+="running_CPU_application:"+running_CPU_application.ToString()+"\n";OUT+="stop_CPU_application:"+stop_CPU_application.ToString()+"\n";OUT+="GPU:"+"\n";//альтернативныедействиядляреализацииалгоритмаkorenXN.exex64использующегодлявычисленийGPUDateTimerunning_GPU_application=DateTime.Now;ProcessprocGPU=Process.Start("korenXN.exe");while(!procGPU.HasExited);DateTimestop_GPU_application=DateTime.Now;string[]outs2=File.ReadAllLines("out.txt");File.Delete("out.txt");if(outs2.Length>0){for(intj=0;j<outs2.Length;j++){OUT+=outs2[j]+"\n";}}else{OUT+="novalues"+"\n";}OUT+="running_GPU_application:"+running_GPU_application.ToString()+"\n";OUT+="stop_GPU_application:"+stop_GPU_application.ToString()+"\n";returnOUT;//возвратитьрезультат}staticvoidMain(){Int64start=36*36;//начальноезначениевходногочислаInt64degree_of_strat=500;//начальноезначениемаксимальнойстепениintsize=20-5;//количествоэлементоввмассивеInt64[]Number=newInt64[size];//массиввходныхчиселInt64[]Degree_of=newInt64[size];//массивмаксимальныхстепенейstring[]outs=newstring[size];//масссиврезультатовfor(intn=0;n<size;n++){if(n%2==0){Number[n]=start*start;}else{Number[n]=start*degree_of_strat;Number[n]-=n+n;}start+=36*36;Degree_of[n]=degree_of_strat;degree_of_strat+=1000;}for(intn=0;n<size;n++){outs[n]=Upload(Number[n],Degree_of[n]);Console.WriteLine(outs[n]);}System.IO.File.WriteAllLines("result.txt",outs);//записатьрезультатывфайлresult.txt}}}


, который создавал файл с исходными данными, затем последовательно запускал exe файлы алгоритмов использующих CPU или GPU и замерял время их работы, затем заносил это время и результаты работы алгоритмов в файл result.txt. Для замера загруженности процессора использовался диспетчер задач windows.
Результаты теста превидены в таблице:

Как видно из таблицы, время выполнения алгоритма на GPU немного больше, чем на CPU.
Однако, отмечу, что вовремя работы алгоритма использующего для вычислений GPU загрузка им CPU, в Диспетчере задач, не превышала 30%, в то время как алгоритм использующий для вычислений CPU, загружал его на 68-85%, что в свою очередь иногда приводило к замедлению других приложений. Также, ниже приведен график, показывающий различие во
времени выполнения (по оси Y)CPU и GPU в зависимости от входного числа(по оси X).
график


Далее я решил провести тестирование при процессоре нагруженном другими приложениями. Процессор был нагружен так, что запущенный в приложение тест, не занимал больше 55% ресурсов процессора. Результаты теста приведены ниже:

График


Как видно из таблицы, в случае с нагруженным CPU, выполнение вычислений на GPU, дает прирост производительности, так как загруженность процессора в 30% укладывается в лимит 55%, а в случае использования CPU для вычислений, его загрузка составляет 68-85% , что тормозит работу алгоритма, если CPU нагружен другими приложениями.

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


Подробнее..

Категории

Последние комментарии

  • Имя: Макс
    24.08.2022 | 11:28
    Я разраб в IT компании, работаю на арбитражную команду. Мы работаем с приламы и сайтами, при работе замечаются постоянные баны и лаги. Пацаны посоветовали сервис по анализу исходного кода,https://app Подробнее..
  • Имя: 9055410337
    20.08.2022 | 17:41
    поможем пишите в телеграм Подробнее..
  • Имя: sabbat
    17.08.2022 | 20:42
    Охренеть.. это просто шикарная статья, феноменально круто. Большое спасибо за разбор! Надеюсь как-нибудь с тобой связаться для обсуждений чего-либо) Подробнее..
  • Имя: Мария
    09.08.2022 | 14:44
    Добрый день. Если обладаете такой информацией, то подскажите, пожалуйста, где можно найти много-много материала по Yggdrasil и его уязвимостях для написания диплома? Благодарю. Подробнее..
© 2006-2024, personeltest.ru