Comments 52
спасибо, интересно предлагаю деление в части SIMD превратить
2.0f * ((width)^(-1)) #^ в степени
2.0f * ((height)^(-1))у симда есть операция приведения к степени
может такие <</>> штуки сработают
хотя на первом приближении они по годболту вроде одно и тоже по ассемблеру делают - спорно
поидее передаем данные с ГПУ на ЦПУ и обратно, потомучто нужно синхронизировать вычисления, если вы смотрели Вулкан там такая же ситуация
память привязана к обьектам над которыми ведутся вычисления, видимо надо отслеживать их в памяти на видеокарте тоесть она должна вернуть память поидее(мне так кажется) (смотрите тему GPGPU - Vulkan -draw without overhead)
Всё это очень интересно, но есть в таких статьях фундаментальная проблема.
Все эти быстрые технологии обычно подразумевают 32-битные числа одинарной точности, тогда как на практике, чтобы нарисовать действительно качественного Мандельброта, даже 64-битных категорически недостаточно.
Конечно, если это рассматривать как учебную задачу для изучения разных технологий - вопросов нет. Но рисовать красивые фракталы так не получится.
постойте что вы имеете ввиду? тоесть если показывают демку фрактала бесконечного в 3д то там вы ситаете другая точность? но ведь пространство укладывается по крайней мере у меня на компе в одинарную точность, можно поподробнее пожалуйста (там и на шейдере и на стороне компилятора флоаты, могут быть, точность пространства тогда надо рассматривать, но ведь бесконечность не отрисовать вдаль она будет отсекаться так и так как не крути, ну если только Raytraycing - наверно)
Если вы хотите нарисовать ММ просто в "общем виде" - проблем не будет. Но это неинтересно, намного интереснее его зумить и рисовать увеличенные фрагменты. И вот тогда, по мере увеличения, вы в какой-то момент заметите, что детализация фрактала падает, он как бы замыливается - это потому что точности не хватает.
У меня было несколько попыток подхода к ММ, но все уже довольно давно. Точных цифр уже не помню, но кажется на зумах порядка от 30 (каждый шаг зума равен x2) двойной точности уже не хватает - хотя с точки зрения разноообразия рисунков самое интересное там только начинается.
спасибо, понял
(немного не в тему блин это имба просто skeletal_animation_in_vulkan_after_struggling_for я просто посматриваю паралельно что там в вулкане и попал на сообщество улкана ))
Спасибо, согласен. На приближениях более чем в 1000 раз float'ам уже перестаёт хватать точности.
Приближение в ~15000 раз
А про то, как отрисовывать гифки типо таких, надо писать уже отдельную статью, на совершенно другую тематику ;)
камеру будете двигать?, как раз и масштаб красиво кладётся
если делать просто скейл будет порог, а если еще пользоваться движением порог будет всегда впереди тоесть предел
а предел проекции камеры это (там участвует тангенс/(или векторный расчет)) вперед
общее приближение (зум-т.е. масштаб)15000 раз + с каждым приближением (камера)1.0f можно делать градацию относительно приближения, и пересчитывать обратно предел передний
Там, наверно, длинные рациональные числа используются? Хотя, может и длинные числа с плавающей запятой.
С кудой есть такая проблема - можно перейти на double, но... на игровых картах на double зарезана производительность. Хитрая Nvidia хочет, чтобы мы покупали их Теслы за бешеные деньги.
Так статья не про ММ, а про способы параллелизации алгоритмов на современном железе. ММ тут только как пример
только ММ нету вкатки в мат аппарат(векторный), потомучто мат аппарат даёт ясность что надо ускорить(как при библиотеках, так и при рейкастах), мат аппарат даёт условия относительности обьектов и камеры относительно обьектов на сцене в том числе
(где мат аппарат это (2д <-> 3д))
на плоской поверхности без мат аппарата считать теми большими преобразованиями возможно, на первое время кажется лучше(но векторное пространство еще лучше, тоесть не придётся делать те большие расчеты, а просто делать базовые операции обусловленные пространством)
тоесть базовые принципы плоских в 3д и общие моменты в 3д, это всё тема что ускорять
Ответ, что надо ускорить, даёт не абстрактный матаппарат, а вполне конкретный запуск задачи с профайлером. Что касается данной статьи, то если вы замените здесь ММ на, например, перемножение матриц, её смысл совершенно не поменяется. Просто ММ - это красиво
а мат аппарат не абстрактный а вполне описуемый ведь там симмулируется и 3д и 2д и рисуется на плоской поверхности
нет тут приведён конкретный случай ММ, перемножение матриц выглядит по другому и тут не показано участия перемножения матриц
в мат аппарате всё можно на avx2 перекинуть или sse4.2 или avx512, в библиотеках по ускорению всё для этого есть, и dotproduct и прочее
Ну даже если и так, то эти всплывающие ограничения для меня были неприятным сюрпризом в свое время.
То привык к двойной точности и это было как бы само собой разумеющимся - а теперь вдруг нельзя. И методы обхода этой проблемы довольно нетривиальны.
если ММ можно отобразить в пространстве 3д, я предположу что можно, то там запускаются решения тривиальные для работы в 3д пространстве, 2д картинка в 3д имеет так же масштаб, с позиции 3д это тривиально, просто в 3д видимо зависимость от дистанции придётся лучами генерировать картинку или проекцию ММ, типо Camera как обьект двигается вперед, а перед ней происходит ММ, плюс есть пороги приближения(она не просто так аппарат называется, там просто крутейший интерфейс чтобы понять всё укладывается в логику и работает как аппарат), можно назвать базовый аппарат трансформацийCoordinate-Systems обьектов в 3д(потомучто 2д в нем можно проецировать матрицей ортогональной)
А на какой карточке Nvidia запускался cuda-тест? Стоит добавить модель для полноты картины
"PCI-E 16x — шины, по которой передаются данные с GPU. Это примерно 16 Гбайт/с" - судя по цифрам, это PCI-E 3.0 ? Какая-то древняя карта получается) Но версию PCI-E хорошо бы уточнить
"группа потоков. Может иметь от одного до трех измерений. В нашем случае он (16, 16)" - а почему всё же 16? Я бегло глянул ссылку и как-то понимания не возникло. Например, современные карты Nvidia поддерживают 1024 потока в треде, почему бы не поставить 32x32 ?
В остальном статья отличная, такая, какими в идеале должны быть статьи на Хабре.
Спасибо, очень приятно
Статью поправил. Модель GPU на каком-то этапе редактирования статьи действительно затерялась. Использовалась старушка RTX 2060 super с PCI-E 3.0 и архитектурой Turing (SM 7.5), на ней SM может запускать до 2048 потоков.
Из вариантов были группировки (32x8), (8x32), (32x32), (16x16), но эмпирически победил вариант (16x16). Вероятно, потому что 1024 потока требуют больше ресурсов SM, которых могло не хватать. Или причина могла лежать в неравномерной загрузке варпов.
Как только будет возможность, проведу тесты повторно и оставлю результаты здесь
Ок, это понятный ответ - попробовали разные варианты, 16х16 оказался лучшим. Для GPU сходу нельзя однозначно сказать, какая конфигурация количества тредов в блоке будет оптимальной, поэтому имеет смысл попробовать разные варианты. А можно ради интереса табличку с цифрами для разных вариантов группировок? Насколько там разница большая в перфе?
Наконец то добрался до компьютера с видеокартой
Конфигурация
$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Fri_Feb_21_20:23:50_PST_2025
Cuda compilation tools, release 12.8, V12.8.93
Build cuda_12.8.r12.8/compiler.35583870_0$ nvidia-smi
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 570.133.07 Driver Version: 570.133.07 CUDA Version: 12.8 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA GeForce RTX 2060 ... Off | 00000000:01:00.0 On | N/A |
| 31% 35C P8 9W / 184W | 479MiB / 8192MiB | 3% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+В этот раз подошёл серьезнее к тестированию: зафиксировал частоту, включил persistence mode, отключил desktop environment, следил за температурой:
sudo nvidia-smi -lgc 2115,2115 # Lock to 2115 MHz
sudo nvidia-smi -pm 1 # Enable persistence modeНу и сами результаты, организовал их по убыванию fps:
+---------+----------+
|blocksize| FPS |
+---------+----------+
| 16x8 | 6360±60 |
+---------+----------+
| 16x16 | 6210±50 |
+---------+----------+
| 32x8 | 6160±40 |
+---------+----------+
| 8x8 | 6060±40 |
+---------+----------+
| 16x32 | 5920±40 |
+---------+----------+
| 32x16 | 5880±40 |
+---------+----------+
| 8x32 | 5810±40 |
+---------+----------+
| 32x32 | 5380±40 |
+---------+----------+И победителем оказался 16x8, обогнав 32x32 почти на 20%
UPD: "поддерживают 1024 потока в треде," - 1024 потока в блоке, конечно же
а почему всё же 16?
Варп на 32 потока, 16х16 это 8 варпов на SM, в доках сказано минимум 2 варпа нужно чтоб скрыть задержки для float4 fma, для скаляров в 4 раза больше, то есть 8 варпов.
Это всё фантазии ваши.
Во-первых, латенси нужно скрывать, если нет ILP, а там он есть. Там только одна пара операций друг за другом с зависимостью идут, если я ничего не проглядел.
Во-вторых, увеличение варпов должно как раз лучше скрадывать латенси, и Нвидиа как раз рекомендует максимально увеличивать количество потоков. Поэтому 32х32 должно быть строго не хуже не 16х16.
Это всё фантазии ваши.
Доки https://docs.nvidia.com/cuda/turing-tuning-guide/#instruction-scheduling. Единственное я забыл про 64 fp32 блока на SM, поэтому должно быть 32x16.
Поэтому 32х32 должно быть строго не хуже не 16х16.
Хуже. Если достаточный минимум 32х16, то на 32х32 будет 16 варптов стоять без дела, это не страшно если все SM загружены, а если нет, то теряем производительность.
Хуже. Если достаточный минимум 32х16, то на 32х32 будет 16 варптов стоять без дела, это не страшно если все SM загружены, а если нет, то теряем производительность
Нет, всё несколько не так.
Во-первых, я вам выше уже указал, что написанное в документации Nvidia описывают ситуацию сглаживания RAW зависимости для FMA операции в предельном случае, когда у вас в коде все операции зависят друг от друга. А в примере из статьи это не так, поэтому достаточный минимум будет не 32x16, а меньше. Другими словами, у вас может быть всего 4 варпа (по 32 треда, по одному на каждый WS в SM), и вы утилизируете полностью SM. Такой пример легко написать. Обратите на этот момент внимание, он важен для понимания того, что написано в доке Nvidia, на которую вы ссылаетесь
Во-вторых, давайте вернёмся к нашему конкретному примеру.
16х16 - это 8 варпов
32х16 - это 16 варпов
32х32 - это 32 варпа
Блок планируется на SM. SM состоит из 4 PU(у каждого свой WS и очередь варпов). Варпы распределяются по PU по кругу - т.е. на нулевую PU попадут 0-ой, 3-ий, 7-ой и т.д. варпы. На первую - 1-ый, 4-ый, 8-ой и т.д. Ну и т.д.
На каждой PU количество одновременно работающих варпов у Nvidia либо 8, либо 16. На Тьюринге по-моему 8, буду исходить из этого (если 16, это никак не поменяет суть рассчётов).
Значит, одновременно на SM могут работать 32 варпа (4 PU x 8 варп слотов).
Иными словами, для всех вариантов разбиение по тредам у вас всегда будет полная утилизация, никаких "на 32х32 будет 16 варпов стоять". Просто в одном случае у вас будет работать 4 блока на SM, а в другом 1 блок.
32 варпа на SM нужно чтобы скрыть доступ к памяти.
В варианте где чисто float4 FMA мне хватает 32х2 как и говорится в доках - нужно х2 варпов, потому что FMA выполняется 2 такта. Ниже в доках говорится о максимум 16 тредблоках на SM, так что размер группы не так уж влияет, важно еще количество регистров и общей памяти.
В общем 16х16 достаточно для скаляров как я написал в самом начале, то что для SM рекомендуется минимум 32х16 так это драйвер дальше решает сколько отправить, если другие лимиты не превышены, он накидает еще варпов.
FMA выполняется 2 такта
FMA исполняется за 4 такта на Тьюринге
Ниже в доках говорится о максимум 16 тредблоках на SM, так что размер группы не так уж влияет, важно еще количество регистров и общей памяти
Это всё вообще непричём в контексте обсуждения данного конкретного теста.
В общем 16х16 достаточно для скаляров как я написал в самом начале
Так и 16х8 может быть достаточно. И 16х16 ничем не лучше 32х32 исходя из общей теории. Вопрос как бы в этом, и пока я не услышал ни одного аргумента в пользу того, почему 16х16 может быть лучше 32х32.
FMA исполняется за 4 такта на Тьюринге
Я про эту часть: "Instructions are performed over two cycles"
и пока я не услышал ни одного аргумента в пользу того, почему 16х16 может быть лучше 32х32.
Например на мобилках может быть доступно 512 потоков на группу, но доступно только 32 регистра, а если поставить 256 потоков то доступны все 64. На десктопных картах регистров намного больше и проблема возникает на большом объеме кода.
Это недостатки приводят к тому, что под каждую архитектуру и под каждую ширину векторных инструкций придётся долго и нудно переписывать код.
Вообще-то нет, если использовать SIMD как вектор из скаляров, то все легко портируется, проблемы возникают со старым подходом, когда SSE вектор использовали как float4 тип и мучались с перестановкой компонентов, например для реализации dot product.
Итого, мы получили прирост в 6.7 раз. Конечно, хотелось x8, но не стоит забывать, что время вычисления 8ми пикселей определяется временем самого медленного. Мы хорошо приблизились к теоритическому пределу.
Zen2 поддерживает 2 FMA инструкции за такт, так что теоретический предел 16х, но реально он достигается только с SMT.
попробуйте избегать gather/scatter путем прямой установки в вектор
__mm<>_set_<>{наивысший..наименьший} там в ассемблере будет как нужно без перетасовок прям как укажете просто правый будет индекс 0
мне это сократило 100 строк кода
Вообще-то нет
Вообще-то да, потому что у вас в зависимости от версии CPU может существенно отличаться набор SIMD-команд, да и простое изменение длины вектора оказаться совсем не простым, т.к. логика остального кода должна быть написана соответствующе + детали внутренней микроархитектуры процессора также могут добавить неприятные сюрпризы.
Zen2 поддерживает 2 FMA инструкции за такт, так что теоретический предел 16х, но реально он достигается только с SMT.
В реальности всё сложнее. Теоретический предел определяется количеством скалярных FPU юнитов и шириной векторных FPU юнитов внутри процессора. И это соотношение совсем в теории может быть от 0 до бесконечности. Но на практике в данном случае цифра будет скорее всего 4 или 8. И полученные результаты говорят о том, что это 8.
И полученные результаты говорят о том, что это 8.
Так может код плохо оптимизирован? Как можно по одному тесту судить о пределах железа?
https://ashvardanian.com/posts/cpu-ports/ тут например используют сложение на 4х портах, то есть 2 сложения за такт выполняется на пути для сложения и еще 2 выполняются на пути для FMA. Получаем 4 операции за такт. На GPU тоже можно выполнять 2 сложения за такт, либо параллелить fp32 и i32, но сильно зависит от железа.
у вас в зависимости от версии CPU может существенно отличаться набор SIMD-команд, да и простое изменение длины вектора оказаться совсем не простым
Длина вектора как раз не проблема, я же говорил про скаляры. Какая разница что у нас 4 скаляра за такт, что 8, что 16 ? Код компилируется под определенный набор инструкций и размер вектора. В SVE вообще перешли на переменный размер вектора.
Так может код плохо оптимизирован? Как можно по одному тесту судить о пределах железа?
Это суждение не на основе одного теста, а на основе хорошего понимания того, как внутри устроено железо, понимания алгоритма теста + цифр данного теста. Можно залезть в доки AMD Ryzen 5 5600H и дать точный ответ (если там будет нужная инфа), но честно говоря, не настолько это важный вопрос, чтобы тратить на это столько времени. Просто нет пока ни одного повода к тому, чтобы степень параллелизации была х16. Хотя в теории это возможно, конечно же.
тут например используют сложение на 4х портах, то есть 2 сложения за такт выполняется на пути для сложения и еще 2 выполняются на пути для FMA
Во-первых, сами по себе порты ещё ничего не говорят. Например, важно, сколько ваш vfmadd132ps будет свой порт занимать. Другими словами, нужно смотреть на микроархитектурную картину целиком.
Во-вторых, вы начинаете рассуждать о том, что можно ещё как-то пооптимизировать код, и сделать его ещё быстрее - очевидно это возможно. Но в статье был вполне понятный контекст, когда говорилось про х8, и в рамках написанного - рассуждения вполне корректные. А вы уже начинаете говорить о немного других вещах.
На GPU тоже можно выполнять 2 сложения за такт
GPU может выполнять намного больше чем 2 сложения за такт - от нескольких десятков до нескольких тысяч. В этом его суть - массовой мультипоточности.
Код компилируется под определенный набор инструкций и размер вектора
Код компилируется под определённый набор инструкций, которого у вас может не оказаться на CPU. А вы использовали интринсик и у вас всё сломается.
Можно залезть в доки AMD Ryzen 5 5600H и дать точный ответ (если там будет нужная инфа),
Для FMA 32 FLOPS/cy в доках, это 2 FLOPS так как сложение и вычитание, х8 так как AVX и х2 потому что 2 порта. Ну и 32 FLOPS/cy * 4.2GHz я получал на своем Zen2.
GPU может выполнять намного больше чем 2 сложения за такт - от нескольких десятков до нескольких тысяч. В этом его суть - массовой мультипоточности.
Так и CPU многопоточные, я же про один поток говорил. Даже уточню - про 1 поток в варпе.
Код компилируется под определённый набор инструкций, которого у вас может не оказаться на CPU. А вы использовали интринсик и у вас всё сломается.
У меня не сломается, я проверяю какие есть инструкции и загружаю нужный бинарник.
Для FMA 32 FLOPS/cy в доках, это 2 FLOPS так как сложение и вычитание, х8 так как AVX и х2 потому что 2 порта. Ну и 32 FLOPS/cy * 4.2GHz я получал на своем Zen2.
Вы всё про порты, хотя я вам выше написал, что в контексте данного обсуждения это ничего не говорит. Из написанного можно утверждать, что при максимальной векторизации мы можем исполнять 16 FMA в такт. Но в коде выше нет FMA. И вопрос в том, сколько add/mul single-presicion мы исполняем в такт при максимальной векторизации? И верно ли ваше утвержение, что у нас 4 FP юнита, каждый из которых может исполнять add/mul, и только 2 юнита исполняющих FMA. Исходя из ответов на эти вопросы можно уже понять сколько должна давать AVX 256-bit векторизация.
У меня не сломается, я проверяю какие есть инструкции и загружаю нужный бинарник.
Ну вот видите, уже необходима своя версия кода для каждого случая. О чём и была речь
Вообще-то нет, если использовать SIMD как вектор из скаляров, то все легко портируется
Да, если использовать вектор из скаляров, то всё легко портируется, вы правы. Однако этот комментарий относился к архитектурно зависимой оптимизации на интринсиках. Следующим шагом в статье я делаю именно то, что вы и описываете.
Есть еще вариант ускорить на 40% за счет ILP.
Одна команда выполняется 4 такта, но две независимые команды - 5 тактов, 3 команды - 6 тактов и тд, так получается 1 операция на тект. Но если есть зависимость по данным, то две команды выполняются 8 тактов.
Для этого надо развернуть цикл на 4:
__m256 radius2 = mm256_add_ps(_z_x2, zy2);
// разворачивается в
__m256 radius2[0] = mm256_add_ps(_z_x2[0], zy2[0]);
__m256 radius2[1] = mm256_add_ps(_z_x2[1], zy2[1]);
__m256 radius2[2] = mm256_add_ps(_z_x2[2], zy2[2]);
__m256 radius2[3] = mm256_add_ps(_z_x2[3], zy2[3]);
И так каждую строчку в цикле. Дополнительно: можно накапливать пиксели в uint32x4 векторе и записывать в память одной командой.
стойте допустим окно 800 на 600 , придумываем оптимальную формулу на разбиение по односторонним квадратикам, разбиваем оптимизированно уместно по процессору или еще какому-то признаку, и просто паралельно одновременно пишем эти квадратики 1 проходом, соотв мы имеем поситанный закешированную строку - буффер - картинку, а считаем тоже в ячейках например, получается считаем в квадратике и пишем в квадратик, например 8 квадратиков или 16
тут тогда вырисовывается облик чанков на плоской(тогда нужна ситуация для прыжков по квантам картинки тоесть по 16 отрезкам одной формулы )
вообще конечно в 3д это было бы 1 расчет на процессоре и заброс текстуры в ГПУ, а на ГПУ поидее воркеры сделают своё дело, какой-нить компута-шейдер или меш-шейдерmes-seidery (почему меш, потомучто говорят это будет новый тип пайплайна с ускоренным доступам к старым стадиям видов пайплайна)
Это всё достаточно бессмысленно, т.к. современные процессоры умеют сами извлекать ILP и делать разрыв ложных зависимостей за счёт переименования регистров. Только если вам повезёт и вы нарвётесь на какой-то косяк в аппаратуре. Но сомневаюсь, что AMD Ryzen 5 5600H на таком простом коде слажает
В коде используется только ALU, запись пикселей можно не считать - Zen2 записывает 32байта за такт, тогда откуда ускорение от SMT?
Это уже другой вопрос. Ваша идея понятна - раз SMT даёт улучшение, значит мы недоутилизируем ALU, что может говорить о нехватке ILP. Но это лишь один из вариантов, в реальности причин может быть много. Просто в нашем случае реальный ILP по сути одинаковый во всех вариантах, т.к. задача embarrassingly parallel. И как я написал выше, современные процессоры умеют извлекать ILP в рамках окна исполнения, соответствующему размеру ROB, что составляет сотни инструкций. Поэтому ваш анроллинг здесь ничего не даст, процессор все эти операции сам увидит.
Действительно, в статье есть пробел. Спасибо, что заметили!
Судя по всему, -O3 настолько раздувает размер ассемблера (godbolt), что процессор уже не может увидеть ILP. Попробовал разные уровни анроллинга. Это делается увеличением константы VEC_SIZE = UNROLL * AVX_SIZE:
+------+-------------+
|Unroll| FPS |
+--------------------+
| 1 | 46.8 (x6.7) |
+--------------------+
| 2 | 78.3 (x11.2)|
+--------------------+
| 3 | 93.2 (x13.3)|
+--------------------+
| 4 | 75.3 (x10.8)|
+--------------------+
| 5 | 51.5 (x7.4) |
+--------------------+Итого ускорение в 13.3 раза, то есть дополнительные 67% производительности относительно изначальной реализации.
Также попробовал заанролить многопоточную реализацию, и, как ожидалось, никакого прироста не последовало.
Возможно даже стоит написать продолжение к статье на эту тему
Оптимизация по ILP будет полезна для новых интелов, где нет гипертрединга.
Еще вместо
if (mask == 0) break;
можно попробовать:
if (mask != 0) [[likely]] {} else break;
Это дает +10% за счет более точных предсказаний ветвления. [[unlikely]] не сработает, так как слишком короткий код, поэтому приходится хитрить.
Еще clang чуть лучше оптимизирует, чем MSVC, тоже +10%.
Еще вместо
if (mask == 0) break;
можно попробовать:
if (mask != 0) [[likely]] {} else break;
Попробовал, clang'у это ничего не дало.
Еще clang чуть лучше оптимизирует, чем MSVC, тоже +10%.
В статье он и используется.
Но раз уж на то пошло, попробовал еще PGO с инструментацией. Получил +8%, итого 101 fps и x14.4 относительно наивной имплементации, приятно.
Плоховата производительность для первого варианта (однопоточный CPU без SIMD). Моя реализация в разрешении 800x600 с 128 итераций выдаёт по ощущениям не меньше 20 FPS.
Поставил настройки в точности как у вас: разрешение, количество итераций, зум и смещение. Наивная реализация выдаёт 35fps.
Это неудивительно, учитывая, что нужно обрабатывать в ~4 раза меньше пикселей (1920x1080 vs 800x600), а на чёрные участки тратить в 2 раза меньше итераций (256 vs 128 итераций).
Разгон Мандельброта: SIMD с бубнами, OpenMP и CUDA