Как посчитать вспомогательные потоки.
Растеризация происходит квадратами 2х2, это нужно для работы деривативов (fwidth и прочие), чаще всего для выбора мипуровня текстуры. Если при растеризации полигон не попадает целиком в квадрат, то все равно запускается квадрат, но часть потоков в нем помечается как вспомогательные. Узнать об этом можно через gl_HelperInvocation.
Вспомогательный поток никак не может модифицировать память, поэтому кажется, что нет способа подсчитать потоки с gl_HelperInvocation. Но решение есть - как минимум один поток может писать в память, а сумма
Способ используется для отладки, чтобы найти слишком мелкую геометрию, которую желательно перевести на софтварную растеризацию.
Растеризация происходит квадратами 2х2, это нужно для работы деривативов (fwidth и прочие), чаще всего для выбора мипуровня текстуры. Если при растеризации полигон не попадает целиком в квадрат, то все равно запускается квадрат, но часть потоков в нем помечается как вспомогательные. Узнать об этом можно через gl_HelperInvocation.
Вспомогательный поток никак не может модифицировать память, поэтому кажется, что нет способа подсчитать потоки с gl_HelperInvocation. Но решение есть - как минимум один поток может писать в память, а сумма
subgroupQuadBroadcast( gl_HelperInvocation, 0..3 )
даст количество вспомогательных потоков в группе.Способ используется для отладки, чтобы найти слишком мелкую геометрию, которую желательно перевести на софтварную растеризацию.
Использование дериватив для тонких линий.
Деривативы нужны чтобы узнать изменение пространства между соседними пикселями. Они не требуются только в 2D пространстве без искажений, тогда деривативы всегда будут возвращать одинаковую разницу.
На картинке зеленая линия uv - координаты в пространстве, они изменяются равномерно, без перегибов. Красная линия sdf - дистанция до линии или любой другой формы, заданной SDF функцией, дистанция измеряется в том же пространсве, что и uv.
Часто sdf идет с перегибами, здесь пик оказался между пикселями и потерялся, поэтому минимальное значение sdf смещается на расстояние между пикселями md.
В 3D на расстоянии или под большим углом шаг становится слишком большим и теряются детали в sdf, тогда сглаживание перестает работать и приходится делать затухание.
Деривативы нужны чтобы узнать изменение пространства между соседними пикселями. Они не требуются только в 2D пространстве без искажений, тогда деривативы всегда будут возвращать одинаковую разницу.
На картинке зеленая линия uv - координаты в пространстве, они изменяются равномерно, без перегибов. Красная линия sdf - дистанция до линии или любой другой формы, заданной SDF функцией, дистанция измеряется в том же пространсве, что и uv.
Часто sdf идет с перегибами, здесь пик оказался между пикселями и потерялся, поэтому минимальное значение sdf смещается на расстояние между пикселями md.
В 3D на расстоянии или под большим углом шаг становится слишком большим и теряются детали в sdf, тогда сглаживание перестает работать и приходится делать затухание.
Неправильная микрооптимизация.
Наткнулся на оптимизированный вариант
Почему так произошло: внутри ядра ГП есть отдельные пайплайны, специализированные под разные операции, например FMA (fused multiply add) pipe, SFU (special function unit) pipe - для div,sqrt,log,sin и тд, CVT pipe для конвертации типов.
У многих ГП соотношение 1 SFU на 4 потока, поэтому его использование в 4 раза медленнее FMA, но вызов sqrt только один раз задействует SFU, а "оптимизированный" - 2 раза на деление плюс несколько вызовов FMA, отсюда и потеря производительности.
Возможно ЦП менее заточены под графические задачи и на них деление работает быстрее, чем sqrt, тогда оптимизация имеет смысл.
#gpu_opt
Наткнулся на оптимизированный вариант
sqrt
и cbrt
на shadertoy и конечно же тесты показали падение производительности в 2 раза на sqrt и до 5 раз на cbrt, в том числе на смартфонах 8-ми летней давности.Почему так произошло: внутри ядра ГП есть отдельные пайплайны, специализированные под разные операции, например FMA (fused multiply add) pipe, SFU (special function unit) pipe - для div,sqrt,log,sin и тд, CVT pipe для конвертации типов.
У многих ГП соотношение 1 SFU на 4 потока, поэтому его использование в 4 раза медленнее FMA, но вызов sqrt только один раз задействует SFU, а "оптимизированный" - 2 раза на деление плюс несколько вызовов FMA, отсюда и потеря производительности.
Возможно ЦП менее заточены под графические задачи и на них деление работает быстрее, чем sqrt, тогда оптимизация имеет смысл.
#gpu_opt
sRGB текстуры
Все что мы видим на экране это цвета в пространстве sRGB.
Все что хранится в текстурах и без изменений отображается на экран это тоже sRGB.
Но все манипуляции с цветом (освещение, фильтрация, блендинг, тонемапинг, предумноженная альфа и тд) должны происходить в линейном пространстве, то есть:
В чем преимущество sRGB формата текстур - запись и чтение идет в линейном пространстве, конвертация в/из sRGB происходит автоматически, для 8 бит на канал sRGB формат дает большую точность.
Но есть и недостатки вроде сложных манипуляций, чтоб менять текстуру в компьют шейдере, а так же случаи когда конвертация не применяется.
Конвертация при чтении происходит при сэмплинге в шейдере и при блите (BlitSrc).
Конвертация при записи происходит при рендере в текстуру, при блите (BlitDst) и при очистке (Clear).
Копирование (vkCmdCopy...) работает как memcpy, поэтому читает sRGB значение.
Вывод на экран (Present) тоже читает sRGB.
Не-sRGB формат (RGBA8_UNorm, RGBA16F и тд), хранящий sRGB цвет, работает некорректно - фильтрация и блендинг происходят в sRGB пространстве вместо линейного.
RGBA8_UNorm формат с линейным цветом дает меньшую точность чем sRGBA8, но работает корректно, главное перед выводом на экран перевести в sRGB пространство. А RGBA16F формат с линейным цветом компенсирует потерю точности.
Все что мы видим на экране это цвета в пространстве sRGB.
Все что хранится в текстурах и без изменений отображается на экран это тоже sRGB.
Но все манипуляции с цветом (освещение, фильтрация, блендинг, тонемапинг, предумноженная альфа и тд) должны происходить в линейном пространстве, то есть:
sRGB-to-Linear -> op -> Linear-to-sRGB -> Present
В чем преимущество sRGB формата текстур - запись и чтение идет в линейном пространстве, конвертация в/из sRGB происходит автоматически, для 8 бит на канал sRGB формат дает большую точность.
Но есть и недостатки вроде сложных манипуляций, чтоб менять текстуру в компьют шейдере, а так же случаи когда конвертация не применяется.
Конвертация при чтении происходит при сэмплинге в шейдере и при блите (BlitSrc).
Конвертация при записи происходит при рендере в текстуру, при блите (BlitDst) и при очистке (Clear).
Копирование (vkCmdCopy...) работает как memcpy, поэтому читает sRGB значение.
Вывод на экран (Present) тоже читает sRGB.
Не-sRGB формат (RGBA8_UNorm, RGBA16F и тд), хранящий sRGB цвет, работает некорректно - фильтрация и блендинг происходят в sRGB пространстве вместо линейного.
RGBA8_UNorm формат с линейным цветом дает меньшую точность чем sRGBA8, но работает корректно, главное перед выводом на экран перевести в sRGB пространство. А RGBA16F формат с линейным цветом компенсирует потерю точности.
Быстрое чтение и запись больших данных в шейдере
Есть массив данных, которые надо загрузить в шейдер, обработать и записать обратно. Эта задача для GPGPU и на старых рендерилках может решаться нетривиально.
Зато на современной RTX 2080 вызовы imageLoad/Store упирается только в пропускную способность памяти. Но можно ли сделать быстрее?
В Vulkan появился input attachment, который позволяет читать и писать в один и тот же пиксель. Таким способом можно приспособить фрагментный шейдер под задачи GPGPU и это включает сжатие данных, что увеличивает пропускную способность памяти минимум на 10%, а максимум - в 3 раза.
Больше не нужно думать какой там z-curve у текстуры, чтобы перераспределить потоки в компьют шейдере, драйвер сам все сделает.
Способ работает на NV RTX, PowerVR B-Series, Mali Valhall.
На AMD RDNA архитектуре сжатие работает и в компьют шейдере, поэтому такая оптимизация не требуется.
Главный недостаток - графические задачи не параллелятся, поэтому нужно нагружать все ядра.
#gpu_opt
Есть массив данных, которые надо загрузить в шейдер, обработать и записать обратно. Эта задача для GPGPU и на старых рендерилках может решаться нетривиально.
Зато на современной RTX 2080 вызовы imageLoad/Store упирается только в пропускную способность памяти. Но можно ли сделать быстрее?
В Vulkan появился input attachment, который позволяет читать и писать в один и тот же пиксель. Таким способом можно приспособить фрагментный шейдер под задачи GPGPU и это включает сжатие данных, что увеличивает пропускную способность памяти минимум на 10%, а максимум - в 3 раза.
Больше не нужно думать какой там z-curve у текстуры, чтобы перераспределить потоки в компьют шейдере, драйвер сам все сделает.
Способ работает на NV RTX, PowerVR B-Series, Mali Valhall.
На AMD RDNA архитектуре сжатие работает и в компьют шейдере, поэтому такая оптимизация не требуется.
Главный недостаток - графические задачи не параллелятся, поэтому нужно нагружать все ядра.
#gpu_opt
В Vulkan добавили расширения для более явного поведения сабгрупп и квадгрупп.
По стандарту вспомогательный поток (helper invocation) может не создаватся, даже если используются операции с сабгруппами.
Например на Mali Valhall код
Это можно проверить вызвав
Для корректной работы дериватив нужно чтобы все потоки квадгруппы выполняли одинаковый код, для этого добавили функции
#vk
По стандарту вспомогательный поток (helper invocation) может не создаватся, даже если используются операции с сабгруппами.
Например на Mali Valhall код
subgroupQuadBroadcast( gl_HelperInvocation, 0 ) + ...
всегда возвращает 0, то есть вспомогательные потоки не вызываются, если явно или неявно не используются деривативы.Это можно проверить вызвав
subgroupQuadBroadcast( 1, 0 ) + ...
, результат будет менее 4, на углах треугольника. Что даст некорректный результат при использовании квадгрупп, поэтому и появилось расширение GLSL_EXT_shader_quad, где layout (full_quads) in;
не дает драйверу оптимизировать вспомогательные потоки.Для корректной работы дериватив нужно чтобы все потоки квадгруппы выполняли одинаковый код, для этого добавили функции
quadAny()
и quadAll()
.#vk
Обзор C++26
9.39 - Pattern matching
27.45 - Рефлексия. Появятся новые конструкции, зато будет кодогенерация.
#cpp #rus
9.39 - Pattern matching
27.45 - Рефлексия. Появятся новые конструкции, зато будет кодогенерация.
#cpp #rus
Основное отличие ЦП от ГП
На ЦП изначально пишется однопоточный код и для распараллеливания нужно приложить усилия - создать поток, подключить либу с планировщиком и тд.
На ГП иначе - все задачи выполняются параллельно насколько это возможно, а задача программиста - раставить синхронизации, чтобы данные обрабатывались последовательно.
У ЦП есть определенная иерархия кэшей и они примерно одинаковы у всех производителей, ведь должны одинаково справляться с тем же кодом. Даже размер кэшлинии одинаковый у большинства ЦП.
У ГП во времена обычных рендерилок архитектура могла быть любой, задача была только как можно быстрее все отрисовать. Затем появились GPGPU и прямой доступ к памяти (image/buffer storage), потом еще и управление страницами памяти (sparse memory). Тогда для оптимизации сложных алгоритмов понадобилось знать детали каждой архитектуры, но они уже во многом отличались. Например распределение потоков внутри потокового процессора (SM) отличается у всех производителей, хотя NV/AMD/Intel используются для одних и тех же задач, также как Mali/Adreno/PVR. А порядок потоков влияет на последовательность доступа к памяти и на работу кэшей.
С увеличением разрешения экранов увеличилась и нагрузка на память, поэтому в ГП добавили сжатие данных при рендере в текстуру, на ЦП такого сжатия нет.
ГП использует SIMT - одна скалярная операция выполняется для группы потоков, а на ЦП есть SIMD - операции над векторами, но если развернуть вектора, то получится тот же SIMT, например сначала сложить 4
Скорость выполнения инструкций подгоняется под задачи, на ГП это add, mul, abs, min/max, clamp и операции над векторами dot, cross, length для fp16, fp32 типов. Остальные типы данных будут медленее в разы. На ЦП SIMD оптимизируется под вектора с fp32, fp64 типами и под int типы разной битности.
В более свежик ЦП/ГП есть инструкции для NPU - перемножение больших матриц для i8, fp16 типов.
На ЦП изначально пишется однопоточный код и для распараллеливания нужно приложить усилия - создать поток, подключить либу с планировщиком и тд.
На ГП иначе - все задачи выполняются параллельно насколько это возможно, а задача программиста - раставить синхронизации, чтобы данные обрабатывались последовательно.
У ЦП есть определенная иерархия кэшей и они примерно одинаковы у всех производителей, ведь должны одинаково справляться с тем же кодом. Даже размер кэшлинии одинаковый у большинства ЦП.
У ГП во времена обычных рендерилок архитектура могла быть любой, задача была только как можно быстрее все отрисовать. Затем появились GPGPU и прямой доступ к памяти (image/buffer storage), потом еще и управление страницами памяти (sparse memory). Тогда для оптимизации сложных алгоритмов понадобилось знать детали каждой архитектуры, но они уже во многом отличались. Например распределение потоков внутри потокового процессора (SM) отличается у всех производителей, хотя NV/AMD/Intel используются для одних и тех же задач, также как Mali/Adreno/PVR. А порядок потоков влияет на последовательность доступа к памяти и на работу кэшей.
С увеличением разрешения экранов увеличилась и нагрузка на память, поэтому в ГП добавили сжатие данных при рендере в текстуру, на ЦП такого сжатия нет.
ГП использует SIMT - одна скалярная операция выполняется для группы потоков, а на ЦП есть SIMD - операции над векторами, но если развернуть вектора, то получится тот же SIMT, например сначала сложить 4
x
компонент векторов, затем 4 y
компонент и тд.Скорость выполнения инструкций подгоняется под задачи, на ГП это add, mul, abs, min/max, clamp и операции над векторами dot, cross, length для fp16, fp32 типов. Остальные типы данных будут медленее в разы. На ЦП SIMD оптимизируется под вектора с fp32, fp64 типами и под int типы разной битности.
В более свежик ЦП/ГП есть инструкции для NPU - перемножение больших матриц для i8, fp16 типов.
Ветвление в шейдере
Есть 3 способа сделать ветвление в шейдере:
1. Ветви через
2. Выполнить все ветви, умножить их на 0 или 1 и сложить.
3. Заменить множество ветвей на умножение на матрицу.
ГП раскидывает задачи по варпам, где 32 потока выполняют одну инструкцию для 32 значений. Ветвление сделано битовой маской, по одному биту на поток. По очереди выполняются все ветви, но записать значение может только активный поток. Похожий механизм используется и для вспомогательного потока (helper invocation). Но есть и оптимизация - если все потоки варпа идут по одному пути, то выполняется только одна ветвь и производительность улучшается, это называется uniform control flow.
Результаты теста:
На Apple M1, Mali Midgard, Intel 9.5 gen разница между uniform и non-uniform менее 30%, значит плохо работает оптимизация ветвления под uniform control flow.
На AMD GCN4 разница 60%, остальные: Adreno 5xx, Adreno 6xx, Mali Valhall, NV Turing, PowerVR B-series дают более 100%, значит хорошо оптимизируют ветвление.
Более старые ГП лучше справляются с умножением матриц, тот же Mali Midgard использует векторную архитектуру. Поэтому вариант "один раз построить матрицу и умножать на нее" работает быстрее ветвления, но построение матрицы не оптимизировано под вектора и сильно проигрывает при неоднородном выполнении.
В среднем вариант с ветвлением работает быстрее, особенно если заранее не известно будет ли выполняться по одному пути или по разным.
#gpu_opt
Есть 3 способа сделать ветвление в шейдере:
1. Ветви через
if
, у некоторых есть предубеждение, что это медленно.2. Выполнить все ветви, умножить их на 0 или 1 и сложить.
3. Заменить множество ветвей на умножение на матрицу.
ГП раскидывает задачи по варпам, где 32 потока выполняют одну инструкцию для 32 значений. Ветвление сделано битовой маской, по одному биту на поток. По очереди выполняются все ветви, но записать значение может только активный поток. Похожий механизм используется и для вспомогательного потока (helper invocation). Но есть и оптимизация - если все потоки варпа идут по одному пути, то выполняется только одна ветвь и производительность улучшается, это называется uniform control flow.
Результаты теста:
На Apple M1, Mali Midgard, Intel 9.5 gen разница между uniform и non-uniform менее 30%, значит плохо работает оптимизация ветвления под uniform control flow.
На AMD GCN4 разница 60%, остальные: Adreno 5xx, Adreno 6xx, Mali Valhall, NV Turing, PowerVR B-series дают более 100%, значит хорошо оптимизируют ветвление.
Более старые ГП лучше справляются с умножением матриц, тот же Mali Midgard использует векторную архитектуру. Поэтому вариант "один раз построить матрицу и умножать на нее" работает быстрее ветвления, но построение матрицы не оптимизировано под вектора и сильно проигрывает при неоднородном выполнении.
В среднем вариант с ветвлением работает быстрее, особенно если заранее не известно будет ли выполняться по одному пути или по разным.
#gpu_opt
Что не так с интерполяцией?
Линейная интерполяция текстур происходит с 8-битной точностью. Обычно это не заметно, так как градиент занимает несколько пикселей на экране. Но проявляется на картах высот и при использовании текстур с шумом, даже в таком случае на глаз разница незаметна, пока не нужно расчитать нормали, вот они выделят все ступеньки.
Хорошо что есть
В Vulkan есть свойство subTexelPrecisionBits, которое на большинстве ГП равно 8 битам, а в DX прописано в документации: 8 бит и округление в большую сторону. Исправляем
Линейная интерполяция текстур происходит с 8-битной точностью. Обычно это не заметно, так как градиент занимает несколько пикселей на экране. Но проявляется на картах высот и при использовании текстур с шумом, даже в таком случае на глаз разница незаметна, пока не нужно расчитать нормали, вот они выделят все ступеньки.
Хорошо что есть
textureGather()
, где сразу возвращается 4 текселя, применяешь интерполяцию и готово. Но и тут есть 8-битная субпиксельная точность из-за которой fract(uv * dim + 0.5)
округлит не в ту сторону и получится неверный результат интерполяции.В Vulkan есть свойство subTexelPrecisionBits, которое на большинстве ГП равно 8 битам, а в DX прописано в документации: 8 бит и округление в большую сторону. Исправляем
fract(... + 0.5/256)
и теперь интерполяция работает правильно.Как приостановить поток. Тесты на Mac M1 и Android.
Реализация
Кроме STD доступны и Unix функции
Функция
Функция
Для коротких пауз есть
На Windows компилятор инлайнит вызов
На ARM64 появилась инструкция WFE - Wait For Event (
Получаются предсказуемые интервалы:
30нс-3мкс для
1-30мкс для
500мкс для
15мс для
#cpp #threading
Реализация
std::this_thread::sleep_for()
и sleep_until()
под MacOS работает точнее, минимальное время 4-10мкс, далее результаты близки к требуемым, но погрешность может доходить до 20%. Например для 15мс результат 15.9мс, для 100мкс - 120мкс.Кроме STD доступны и Unix функции
nanosleep
и usleep
.Функция
nanosleep()
со стандартными настройками минимальное время пазуы: 6мкс, для 10мкс - 17мкс, для 100мкс - 130мкс. Чем больше время ожидания, тем лучше точность, но всегда больше 10%.Функция
usleep()
помечена как устаревшая и по точности не лучше nanosleep()
.Для коротких пауз есть
__builtin_arm_yield
(аналоги __yield
, asm volatile("yield")
) что занимает около 30нс. В отличие от x86 с гипертредингом, эта инструкция не совершает никаких действий, только информирует ЦП, что не нужно снижать частоту. На Android девайсах с более низкой частотой инструкция выполняется дольше, так на Cortex A53 занимает 1-3мкс, а на A55 - 90нс. Вызов инструкции внутри цикла ничего не меняет.На Windows компилятор инлайнит вызов
_mm_pause()
даже из cpp файла, а на Mac и Android - нет, поэтому при оборачивании инструкции в функцию следует разместить ее в хэдере и пометить как forceinline.На ARM64 появилась инструкция WFE - Wait For Event (
__builtin_arm_wfe
, __wfe
), которая переводит ЦП в экономичный режим в ожидании события. Если не использовать механизм событий, то функция приостановит поток на фиксированное время - ARM_BOARD_WFE_TIMEOUT_NS
, что около 1мкс.Получаются предсказуемые интервалы:
30нс-3мкс для
__builtin_arm_yield
1-30мкс для
__builtin_arm_wfe
500мкс для
nanosleep(420us)
15мс для
sleep_for(15ms)
Тесты на Windows#cpp #threading
Обертка над Vulkan, управление ресурсами
Удаление ресурсов должно происходить, когда ГП закончил все операции над ресурсом, то есть командный буфер завершил выполнение. Для этого есть несколько подходов, в FG в командном буфере увеличивается счетчик ссылок для каждого используемого ресурса и добавляется в список, по завершению выполнения идет проход по списку и уменьшается счетчик ссылок. Это увеличивает нагрузку на ЦП, но не так существенно - трейс Doom 2016 выполнялся на FG на 60 fps в одном потоке. Похожий подход используется внутри Metal API.
Но можно сделать лучше - если обертка используется для рилтайм графики, то работает с фиксированной частотой кадров, значит при начале нового кадра 'кадр - 2' уже завершился и можно освобождать ресурсы. Теперь счетчик ссылок управляется только пользователем, когда он дойдет до нуля, ресурс добавится в список на удаление и по завершении кадра удалится. Так в AE нагрузка на ЦП стала меньше.
В DE сделано аналогично, но ресурс удаляется при завершении кадра на каждой из очередей (graphics, compute).
Кроме удаления ресурсов есть еще staging buffer, который скрывает от пользователя работу с памятью под промежуточный буфер. В FG командный буфер захватывал отдельный буфер, например на 16МБ, когда память заканчивалась захватывался дополнительный буфер. При долгом выполнении командного буфера, промежуточные буферы продолжают висеть как используемые.
В AE благодаря привязке к кадрам механизм улучшился - выделяется один статичный буфер на кадр (обычно до 4Мб) и атомарными операциями смещается текущая позиция в нем. Когда требуется загрузить больше данных, то выделяются динамические буферы, они могут переиспользоваться в течение нескольких секунд, иначе удаляются. Пользователь может задать лимит памяти на каждый кадр, это нужно для стабильного fps. При превышении лимита память не выделяется, а пользователю возвращается размер данных, которые не получилось скопировать.
Каждый кадр данные перегоняются по PCI-E шине, пропускная способность у 3й версии - 16Гб/с, на 60fps это 273Мб/с при условии, что копирование произойдет в начале кадра. Если превысить этот лимит или копировать в середине/конце кадра, то время кадра увеличится и fps станет нестабильным. Этот механизм позволяет стримить сцены без async transfer очереди без влияния на рендер, драйвер сам парллелит копирование пока нет синхронизаций.
В DE память выделяется по мере необходимости и переиспользуется, но при превышении лимита каждый кадр происходит выделение и освобождение памяти, что очень медленно - до 10fps в однопотоке. Отдельно выделяется статичный буфер под юниформы (обычно 8Мб), при превышении лимита вызывается vkDeviceWaitIdle и память безопасно переиспользуется.
#vk
Удаление ресурсов должно происходить, когда ГП закончил все операции над ресурсом, то есть командный буфер завершил выполнение. Для этого есть несколько подходов, в FG в командном буфере увеличивается счетчик ссылок для каждого используемого ресурса и добавляется в список, по завершению выполнения идет проход по списку и уменьшается счетчик ссылок. Это увеличивает нагрузку на ЦП, но не так существенно - трейс Doom 2016 выполнялся на FG на 60 fps в одном потоке. Похожий подход используется внутри Metal API.
Но можно сделать лучше - если обертка используется для рилтайм графики, то работает с фиксированной частотой кадров, значит при начале нового кадра 'кадр - 2' уже завершился и можно освобождать ресурсы. Теперь счетчик ссылок управляется только пользователем, когда он дойдет до нуля, ресурс добавится в список на удаление и по завершении кадра удалится. Так в AE нагрузка на ЦП стала меньше.
В DE сделано аналогично, но ресурс удаляется при завершении кадра на каждой из очередей (graphics, compute).
Кроме удаления ресурсов есть еще staging buffer, который скрывает от пользователя работу с памятью под промежуточный буфер. В FG командный буфер захватывал отдельный буфер, например на 16МБ, когда память заканчивалась захватывался дополнительный буфер. При долгом выполнении командного буфера, промежуточные буферы продолжают висеть как используемые.
В AE благодаря привязке к кадрам механизм улучшился - выделяется один статичный буфер на кадр (обычно до 4Мб) и атомарными операциями смещается текущая позиция в нем. Когда требуется загрузить больше данных, то выделяются динамические буферы, они могут переиспользоваться в течение нескольких секунд, иначе удаляются. Пользователь может задать лимит памяти на каждый кадр, это нужно для стабильного fps. При превышении лимита память не выделяется, а пользователю возвращается размер данных, которые не получилось скопировать.
Каждый кадр данные перегоняются по PCI-E шине, пропускная способность у 3й версии - 16Гб/с, на 60fps это 273Мб/с при условии, что копирование произойдет в начале кадра. Если превысить этот лимит или копировать в середине/конце кадра, то время кадра увеличится и fps станет нестабильным. Этот механизм позволяет стримить сцены без async transfer очереди без влияния на рендер, драйвер сам парллелит копирование пока нет синхронизаций.
В DE память выделяется по мере необходимости и переиспользуется, но при превышении лимита каждый кадр происходит выделение и освобождение памяти, что очень медленно - до 10fps в однопотоке. Отдельно выделяется статичный буфер под юниформы (обычно 8Мб), при превышении лимита вызывается vkDeviceWaitIdle и память безопасно переиспользуется.
#vk
О безопасности в Windows
Часто встречается новость, что в каком-то приложении найдена уязвимость - запись за пределы памяти. Но насколько это опасно?
Самое главное - атакующий должен передать данные в программу и спровоцировать запись за пределы памяти, да так, чтобы подменить указатель на инструкцию, чтобы дальше начал выполняться его код.
То есть опасны данные, получаемые снаружи. Также у атакующего должен быть доступ к бинарникам программы, чтобы найти саму уязвимость и способ ей воспользоваться. И даже воспользовавшись уязвимостью атакующий получит только доступ к памяти приложения и все разрешения, которое оно запросило. Насколько это критично зависит от конкретного приложения, атакующим интересны приложения с правами администратора или хранящие персональные данне, платежную информацию или хотя бы с доступом к камере.
Data Execution Prevention (DEP) запрещает исполнение кода не помеченного как исполняемый, включен всегда, так что атакующему придется записать свой код в память, выделенную под выполнение, а это есть не в каждом приложении.
Arbitrary code guard (ACG) запрещает помечать память для выполнения, но это не совместимо с JIT компиляцией, поэтому включается опционально.
Control flow guard (CFG) - код компилируется с опцией /guard:cf и каждый непрямой вызов функции, когда адрес берется из регистра, проверяется на валидность, для этого при запуске приложения в ядро ОС загружаются валидные адреса. Так атакующие не смогут вызвать свой код, но для JIT компиляции вся динамическая память, помеченная как исполняемая, является валидным адресом для вызова.
У атакующих остается еще одна возможность - вызвать уже существующий код в бинарнике, это может быть код для выделения памяти под JIT компиляцию, а затем скопировать туда свой бинарник и запустить его - не простая задача. Для предотвращения этого используется Address Space Layout Randomization (ASLR), код линкуется с опцией /DYNAMICBASE (и опционально /HIGHENTROPYVA) и при каждом запуске будут случайные адреса функций и выделяемой динамической памяти.
Подробнее в Exploit protection reference
Часто встречается новость, что в каком-то приложении найдена уязвимость - запись за пределы памяти. Но насколько это опасно?
Самое главное - атакующий должен передать данные в программу и спровоцировать запись за пределы памяти, да так, чтобы подменить указатель на инструкцию, чтобы дальше начал выполняться его код.
То есть опасны данные, получаемые снаружи. Также у атакующего должен быть доступ к бинарникам программы, чтобы найти саму уязвимость и способ ей воспользоваться. И даже воспользовавшись уязвимостью атакующий получит только доступ к памяти приложения и все разрешения, которое оно запросило. Насколько это критично зависит от конкретного приложения, атакующим интересны приложения с правами администратора или хранящие персональные данне, платежную информацию или хотя бы с доступом к камере.
Data Execution Prevention (DEP) запрещает исполнение кода не помеченного как исполняемый, включен всегда, так что атакующему придется записать свой код в память, выделенную под выполнение, а это есть не в каждом приложении.
Arbitrary code guard (ACG) запрещает помечать память для выполнения, но это не совместимо с JIT компиляцией, поэтому включается опционально.
Control flow guard (CFG) - код компилируется с опцией /guard:cf и каждый непрямой вызов функции, когда адрес берется из регистра, проверяется на валидность, для этого при запуске приложения в ядро ОС загружаются валидные адреса. Так атакующие не смогут вызвать свой код, но для JIT компиляции вся динамическая память, помеченная как исполняемая, является валидным адресом для вызова.
У атакующих остается еще одна возможность - вызвать уже существующий код в бинарнике, это может быть код для выделения памяти под JIT компиляцию, а затем скопировать туда свой бинарник и запустить его - не простая задача. Для предотвращения этого используется Address Space Layout Randomization (ASLR), код линкуется с опцией /DYNAMICBASE (и опционально /HIGHENTROPYVA) и при каждом запуске будут случайные адреса функций и выделяемой динамической памяти.
Подробнее в Exploit protection reference
Пример уязвимого кода:
Буфер выделен на стэке, по смещению
Вот так просто, но Data Execution Prevention уже не даст выполнить код из стэка. DEP включен всегда, начиная с Win Vista, а флаг
А вот реальный пример уязвимости в 7z 2018 года: 7-Zip: Multiple Memory Corruptions via RAR and ZIP
В 7z подсовывается RAR архив, который инициализирует память декодера, а потом вызывает ошибку из-за чего бросается исключение и память не очищается. Далее второй RAR архив использует уже инициализированную атакующим память, за счет этого получается сместить один указатель за пределы его памяти и попасть на подготовленную атакующим память, эта память затем используется для копирования в буфер на стэке, при выходе за пределы буфера получается перенаправить выполнение кода.
В итоге: помять не очищается, границы памяти не проверяются для указателя, границы буфера не проверяются, не используются флаги
Защита улучшается при компиляции с /GS, тогда значение регистра EIP копируется до локальных переменных, а при выходе из функции проверяются на повреждение. Обещают незначительное влияние на производительность. Но эта опция не дает полную защиту - выключается при оптимизации или если нет буфера на стэке, а также не защищает от изменения vtable.
void function(char *str) {
char buffer[16];
strcpy(buffer,str);
}
Буфер выделен на стэке, по смещению
buffer[16+x]
расположен адрес возврата из функции (EIP), если его заменить на buffer[16+y]
, то вызовется произвольный код. Подробный разбор: smashstack, Win32 Buffer Overflows.Вот так просто, но Data Execution Prevention уже не даст выполнить код из стэка. DEP включен всегда, начиная с Win Vista, а флаг
/NXCOMPAT
, который требуется для DEP, появился в VS2005.А вот реальный пример уязвимости в 7z 2018 года: 7-Zip: Multiple Memory Corruptions via RAR and ZIP
В 7z подсовывается RAR архив, который инициализирует память декодера, а потом вызывает ошибку из-за чего бросается исключение и память не очищается. Далее второй RAR архив использует уже инициализированную атакующим память, за счет этого получается сместить один указатель за пределы его памяти и попасть на подготовленную атакующим память, эта память затем используется для копирования в буфер на стэке, при выходе за пределы буфера получается перенаправить выполнение кода.
В итоге: помять не очищается, границы памяти не проверяются для указателя, границы буфера не проверяются, не используются флаги
/NXCOMPAT
и /DYNAMICBASE
.Защита улучшается при компиляции с /GS, тогда значение регистра EIP копируется до локальных переменных, а при выходе из функции проверяются на повреждение. Обещают незначительное влияние на производительность. Но эта опция не дает полную защиту - выключается при оптимизации или если нет буфера на стэке, а также не защищает от изменения vtable.
Instruction-level parallelism (ILP) - параллелизм на уровне команд.
Одна инструкция выполняется за несколько этапов, процессор может параллельно выполнять разные этапы для нескольких инструкций. Но если следующая инструкция зависит от результата предыдущей, то параллельное выполнение невозможно. Иногда компилятор и процессор переставляют инструкции для лучшего распараллеливания, но часто вручную получается добиться лучших результатов.
Вместо последовательного выполнения операций над одной переменной оптимальнее выполнять одну инструкцию для 4х переменных, тогда не будет зависимостей по памяти и выполнится максимально быстро.
ILP актуально для SSE/AVX/NEON, а также на ГП начиная с AMD RDNA архитектуры, на NV появилось раньше.
Подробнее на algorithmica, RDNA Architecture, Better Performance at Lower Occupancy.
#gpu_opt #cpu_opt
Одна инструкция выполняется за несколько этапов, процессор может параллельно выполнять разные этапы для нескольких инструкций. Но если следующая инструкция зависит от результата предыдущей, то параллельное выполнение невозможно. Иногда компилятор и процессор переставляют инструкции для лучшего распараллеливания, но часто вручную получается добиться лучших результатов.
Вместо последовательного выполнения операций над одной переменной оптимальнее выполнять одну инструкцию для 4х переменных, тогда не будет зависимостей по памяти и выполнится максимально быстро.
ILP актуально для SSE/AVX/NEON, а также на ГП начиная с AMD RDNA архитектуры, на NV появилось раньше.
Подробнее на algorithmica, RDNA Architecture, Better Performance at Lower Occupancy.
#gpu_opt #cpu_opt
Возможно ли ускорить memcpy?
Внутри memcpy проверяется выравнивание памяти, поддержка AVX/SSE и в зависимости от этого выбирается наилучший вариант копирования. Но есть нюанс.
В SSE/AVX есть два варианта копирования: store и stream. Store записывает значение в кэш, а stream - нет, результат сразу пишется в RAM и скорость ограничена пропускной способностью памяти. В memcpy используют оба варианта, но для разных случаев, так store используется для диапазона 512B - 2MB, после 2MB идет stream, а до 512B - не SIMD версия. 2MB выбрано как примерный размер L3 кэша, но реальный размер L3 сильно зависит от модели ЦП.
Получается на диапазоне от 2MB до 8MB (на Ryzen 3900X) можно обогнать memcpy используя
Многопоточный тест.
На 4х потоках store после 2MB сильно замедляется, а memcpy показывает максимальную производительность. Похоже все ЦП на x64 имеют кэш, настоенный под лимит в 2MB. Но нашелся старый AMD Phenom II на котором memcpy после 2MB выдает 4ГБ/с, а
#cpp #cpu_opt
Внутри memcpy проверяется выравнивание памяти, поддержка AVX/SSE и в зависимости от этого выбирается наилучший вариант копирования. Но есть нюанс.
В SSE/AVX есть два варианта копирования: store и stream. Store записывает значение в кэш, а stream - нет, результат сразу пишется в RAM и скорость ограничена пропускной способностью памяти. В memcpy используют оба варианта, но для разных случаев, так store используется для диапазона 512B - 2MB, после 2MB идет stream, а до 512B - не SIMD версия. 2MB выбрано как примерный размер L3 кэша, но реальный размер L3 сильно зависит от модели ЦП.
Получается на диапазоне от 2MB до 8MB (на Ryzen 3900X) можно обогнать memcpy используя
_mm256_storeu_si256
.Многопоточный тест.
На 4х потоках store после 2MB сильно замедляется, а memcpy показывает максимальную производительность. Похоже все ЦП на x64 имеют кэш, настоенный под лимит в 2MB. Но нашелся старый AMD Phenom II на котором memcpy после 2MB выдает 4ГБ/с, а
_mm_stream_si128
- 6ГБ/с, по какой-то причине в SSE версии memcpy используется только _mm_storeu_si128
и здесь можно ускорить.#cpp #cpu_opt
Instruction-level parallelism на ГП
Для примера возьму NV Turing. Один SM выполняет до 32 варпов, при условии что они умещаются в 64К регистров и используют до 64КБ общей памяти. SM выполняет 64 fp32 FMA операций за цикл, это всего 2 варпа. Инструкция выполняется 2 цикла, поэтому нужно 4 варпа для полной нагрузки SM.
Если есть зависимость между инструкциями, то следующая запускается через 4 цикла, это называется instruction issue latency. Теперь нужно 16 варпов на SM, чтобы не было простоев. При ILP=4 минимальное количество варпов снижается до 4, но увеличивается количество регистров на поток, а значит SM сможет вместить меньше варпов.
Намного больше времени тратится на доступ к памяти, пока один варп ждет данные из памяти, другие выполняют вычисления, потом меняются. Чем больше обращений к памяти, тем больше нужно варпов на SM и тем меньше регистров они должны занимать. Но это важнее для мобильных ГП, где намного меньше регистров.
Также ГП давно перешли на скалярную архитектуру, где операция с float4 это 4 инструкции - по одной на каждый элемент, тут уже ILP=4 и оптимизировать не требуется. Часто код шейдеров достаточно простой, чтобы компилятор смог его оптимизировать и улучшить ILP.
В итоге ILP слабо влияет на производительность ГП. Вот SIMD на ЦП намного более чувствительны к ILP, но об этом позже.
Подробнее: instruction-scheduling, occupancy.
#gpu_opt
Для примера возьму NV Turing. Один SM выполняет до 32 варпов, при условии что они умещаются в 64К регистров и используют до 64КБ общей памяти. SM выполняет 64 fp32 FMA операций за цикл, это всего 2 варпа. Инструкция выполняется 2 цикла, поэтому нужно 4 варпа для полной нагрузки SM.
Если есть зависимость между инструкциями, то следующая запускается через 4 цикла, это называется instruction issue latency. Теперь нужно 16 варпов на SM, чтобы не было простоев. При ILP=4 минимальное количество варпов снижается до 4, но увеличивается количество регистров на поток, а значит SM сможет вместить меньше варпов.
Намного больше времени тратится на доступ к памяти, пока один варп ждет данные из памяти, другие выполняют вычисления, потом меняются. Чем больше обращений к памяти, тем больше нужно варпов на SM и тем меньше регистров они должны занимать. Но это важнее для мобильных ГП, где намного меньше регистров.
Также ГП давно перешли на скалярную архитектуру, где операция с float4 это 4 инструкции - по одной на каждый элемент, тут уже ILP=4 и оптимизировать не требуется. Часто код шейдеров достаточно простой, чтобы компилятор смог его оптимизировать и улучшить ILP.
В итоге ILP слабо влияет на производительность ГП. Вот SIMD на ЦП намного более чувствительны к ILP, но об этом позже.
Подробнее: instruction-scheduling, occupancy.
#gpu_opt
Vertical SIMD
Для float3 вектора достаточно функционала SSE (128бит), но давно появился AVX (256бит) с 2х производительностью, а за ним и AVX512 и SVE на ARM. В большинстве открытых движков AVX почти не используются, только в сторонних либах для физики и хэшей. В UE5 AVX используется для транспонирования матриц и для double4 вектора. В O3DE AVX используется только в masked occlusion culling от Intel.
К тому же писать код на SSE/AVX сложно, часто требуется применять shuffle инструкцию, чтобы переставить компоненты для нужного порядка сложения, как это сделано для реализации dot, cross. А последовательное преобразование снижает ILP, задержка инструкций для SSE/AVX составляет 4 цикла для новых ЦП и 5 для более старых.
Намного проще становится при переворачивании SIMD операций, так для AVX получается аналог варпа на 8 потоков, которые одновременно выполняют одну операцию со скаляром. То же самое сделали ГП при переходе на скалярную архитектуру (SIMT).
Было:
стало:
Плюсы такого подхода: легче переносить скалярные операции на вектора, лучше ILP, экономится 25% памяти для float3 типов, легко переключаться на разную длину вектора и тип команд SSE/NEON.
Недостатки тоже есть: ограничение в 16 регистров, потребуется переход на структуру с массивами (SoA) вместо массива структур (AoS), сложно перемещать отдельные элементы.
В тестах получается ускорение в 3 раза для Dot и в 1.7 раза для Cross.
Похожий способ используется в JoltPhysics: RayAABox4, RayTriangle4.
#cpu_opt
Для float3 вектора достаточно функционала SSE (128бит), но давно появился AVX (256бит) с 2х производительностью, а за ним и AVX512 и SVE на ARM. В большинстве открытых движков AVX почти не используются, только в сторонних либах для физики и хэшей. В UE5 AVX используется для транспонирования матриц и для double4 вектора. В O3DE AVX используется только в masked occlusion culling от Intel.
К тому же писать код на SSE/AVX сложно, часто требуется применять shuffle инструкцию, чтобы переставить компоненты для нужного порядка сложения, как это сделано для реализации dot, cross. А последовательное преобразование снижает ILP, задержка инструкций для SSE/AVX составляет 4 цикла для новых ЦП и 5 для более старых.
Намного проще становится при переворачивании SIMD операций, так для AVX получается аналог варпа на 8 потоков, которые одновременно выполняют одну операцию со скаляром. То же самое сделали ГП при переходе на скалярную архитектуру (SIMT).
Было:
A{x,y,z,w} + B{x,y,z,w}
стало:
X{A,..} Y{A,..}
+X{B,..} +Y{B,..}
Плюсы такого подхода: легче переносить скалярные операции на вектора, лучше ILP, экономится 25% памяти для float3 типов, легко переключаться на разную длину вектора и тип команд SSE/NEON.
Недостатки тоже есть: ограничение в 16 регистров, потребуется переход на структуру с массивами (SoA) вместо массива структур (AoS), сложно перемещать отдельные элементы.
В тестах получается ускорение в 3 раза для Dot и в 1.7 раза для Cross.
Похожий способ используется в JoltPhysics: RayAABox4, RayTriangle4.
#cpu_opt
GDC2015: SIMD at Insomniac Games
pdf, video
15. Авто-векторизация и когда она не работает.
25. ISPC - компилятор от Intel для C-подобного языка, сделан специально для автовекторизации.
38. Vec4 тип на SSE это не лучшее решение.
47. Переходим на структуру из массивов (SoA).
60. Иногда AoS лучше. При поиске одного элемента получаем один кэш промах, вместо промахов на каждый массив.
63. Реальный пример. Замена ООП на SoA + SIMD.
104. Как выбрать определенные элементы из вектора. Проблемы с перемещением элементов в лево.
187. Branchless вместо ветвления.
#cpu_opt
pdf, video
15. Авто-векторизация и когда она не работает.
25. ISPC - компилятор от Intel для C-подобного языка, сделан специально для автовекторизации.
38. Vec4 тип на SSE это не лучшее решение.
47. Переходим на структуру из массивов (SoA).
60. Иногда AoS лучше. При поиске одного элемента получаем один кэш промах, вместо промахов на каждый массив.
63. Реальный пример. Замена ООП на SoA + SIMD.
104. Как выбрать определенные элементы из вектора. Проблемы с перемещением элементов в лево.
187. Branchless вместо ветвления.
#cpu_opt
Harnessing Wave Intrinsics For Good (And Evil)
Подробно расмотрено несколько алгоритмов с использованием операций над сабгруппой.
Разбирается случай waterfall loop, когда одна операция компилируется в цикл с перебором.
#gpu_opt
Подробно расмотрено несколько алгоритмов с использованием операций над сабгруппой.
Разбирается случай waterfall loop, когда одна операция компилируется в цикл с перебором.
#gpu_opt