VK_EXT_shader_object proposal
Интересны первые 2 главы, объясняющие зачем понадобилось переделывать API.
Большинство разработчиков вместо того чтобы переписать всю архитектуру заново под PSO, сделали мапу и поиск по ней для каждого рендер объекта, чтоб как и раньше менять состояния. Это приводит к множественным компиляциям пайплайнов и тормозам на стороне ЦП.
При этом драйвер намного лучше может справиться с аналогичной задачей, поэтому и решено было отказаться от больших PSO в пользу раздельных шейдеров и динамических состояний.
Спустя 7 лет с выхода Vulkan многие так сделали быстрый порт на Vulkan без необходимых оптимизаций и кардинальных пересмотров архитектуры.
#vk
Интересны первые 2 главы, объясняющие зачем понадобилось переделывать API.
Большинство разработчиков вместо того чтобы переписать всю архитектуру заново под PSO, сделали мапу и поиск по ней для каждого рендер объекта, чтоб как и раньше менять состояния. Это приводит к множественным компиляциям пайплайнов и тормозам на стороне ЦП.
При этом драйвер намного лучше может справиться с аналогичной задачей, поэтому и решено было отказаться от больших PSO в пользу раздельных шейдеров и динамических состояний.
Спустя 7 лет с выхода Vulkan многие так сделали быстрый порт на Vulkan без необходимых оптимизаций и кардинальных пересмотров архитектуры.
#vk
Far Cry Dunia Engine Shader Pipeline
слайд 18:
- Shader and Pipeline state is fully separate. Shaders are in Data (hlsl) and Pipeline state is in Code (C++)
- Every model could be used with any shader and it was kinda expected to work
- Nearly every object in the world could switch to a fully different and unrelated shaders at any point in game
Типичный пример рендера до Vulkan/DX12, когда можно было менять все состояния без компиляции PSO.
слайд 29:
- Cache all render state; calls from the runtime
- On each Draw Call, hash the whole state + look up PSO in a Dictionary. If not ready, skip the draw call until the PSO is ready at some point in the future.
- Never stall rendering, but certain objects might be missing for few frames
Типичное решение - сделать мапу для поиска нужного PSO под шейдер + рендер стейт.
Чтобы избавиться от простоя ЦП они просто не рисуют объекты, пока не будет готов нужный PSO.
Пока нужное PSO не готово, у них сделан поиск менее подходящей PSO.
Новые расширения в Vulkan позволят избавиться от таких сложных систем.
слайд 18:
- Shader and Pipeline state is fully separate. Shaders are in Data (hlsl) and Pipeline state is in Code (C++)
- Every model could be used with any shader and it was kinda expected to work
- Nearly every object in the world could switch to a fully different and unrelated shaders at any point in game
Типичный пример рендера до Vulkan/DX12, когда можно было менять все состояния без компиляции PSO.
слайд 29:
- Cache all render state; calls from the runtime
- On each Draw Call, hash the whole state + look up PSO in a Dictionary. If not ready, skip the draw call until the PSO is ready at some point in the future.
- Never stall rendering, but certain objects might be missing for few frames
Типичное решение - сделать мапу для поиска нужного PSO под шейдер + рендер стейт.
Чтобы избавиться от простоя ЦП они просто не рисуют объекты, пока не будет готов нужный PSO.
Пока нужное PSO не готово, у них сделан поиск менее подходящей PSO.
Новые расширения в Vulkan позволят избавиться от таких сложных систем.
How mesh shaders are implemented in an AMD driver
* У AMD внутри "страые" шейдеры (вершинный/тесселяции/геометрический) преобразуются в примитив шейдер, аналогичный меш шейдерам.
* Каждый поток в меш шейдере может писать в любую вершину и примитив, но для этого выделяется общая память воркгуппы. Этого можно избежать, если каждый поток пишет только в одну вершину и в один примитив.
*
Task shader driver implementation on AMD HW
* Таск шейдеры выполняются на компьют очереди, а потом графическая очередь выполняет меш шейдеры.
* От использования таск шейдеров есть потеря производительности, поэтому лучше их использовать для больших задач, где потери в драйвере значительно меньше по сравнению с проделанной работой.
#gpu_opt #amd_gpu
* У AMD внутри "страые" шейдеры (вершинный/тесселяции/геометрический) преобразуются в примитив шейдер, аналогичный меш шейдерам.
* Каждый поток в меш шейдере может писать в любую вершину и примитив, но для этого выделяется общая память воркгуппы. Этого можно избежать, если каждый поток пишет только в одну вершину и в один примитив.
*
hw workgroup size = max(api workgroup size, max vertex count, max primitive count) по этой формуле расчитывается сколько потоков будет использовано.Task shader driver implementation on AMD HW
* Таск шейдеры выполняются на компьют очереди, а потом графическая очередь выполняет меш шейдеры.
* От использования таск шейдеров есть потеря производительности, поэтому лучше их использовать для больших задач, где потери в драйвере значительно меньше по сравнению с проделанной работой.
#gpu_opt #amd_gpu
Как расчитать TBN в шейдере
Tangent и Bitangent вектора должны быть направлены в ту же сторону что и текстурные координаты UV.
Это легко сделать в фрагментном шейдере на основе деривативов.
Деривативы
Зная как меняется
Работает отлично на плоских поверхностях и детализированных изогнутых.
Если не работает, то достаточно поменять направление
#blog #cg
Tangent и Bitangent вектора должны быть направлены в ту же сторону что и текстурные координаты UV.
Это легко сделать в фрагментном шейдере на основе деривативов.
Деривативы
dFdx/dFdy вернут разницу переданного значения между соседними пикселями.Зная как меняется
worldPos и uv можно рассчитать нормаль как векторное произведение и касательные (TB) как 3D вектор для uv.float3x3 ComputeTBNinFS (float2 uv, float3 worldPos)
{
float3 wp_dx = dFdx( worldPos );
float3 wp_dy = dFdy( worldPos );
float2 uv_dx = dFdx( uv );
float2 uv_dy = dFdy( uv );
float3 t = normalize( wp_dx * uv_dy.t - wp_dy * uv_dx.t );
float3 b = normalize( -wp_dx * uv_dy.s + wp_dy * uv_dx.s );
float3 n = normalize( cross( wp_dy, wp_dx ));
return float3x3( t, b, n );
}
Работает отлично на плоских поверхностях и детализированных изогнутых.
Если не работает, то достаточно поменять направление
uv.#blog #cg
What is Low Latency C++? CppNow 2023
Part 1 (pdf) (video)
С 44:29 начинаются примеры микрооптимизаций.
56:26 - примеры использования
1:06:15 - немного про новые атрибуты
Part 2 (pdf) (video)
Разбирается что не надо делать в low-latency коде, тут больше про блокировки и нестабильное время выполнения.
Есть ссылка на интересную реализацию спинлока, где ожидание сделано с постепенным увеличением паузы: progressive_backoff_wait
#cpp #threading #lockfree
Part 1 (pdf) (video)
С 44:29 начинаются примеры микрооптимизаций.
56:26 - примеры использования
assume.1:06:15 - немного про новые атрибуты
noalias, unsequenced и тд. Похоже компиляторы ни на что не способны и надо вручную выставлять атрибуты.Part 2 (pdf) (video)
Разбирается что не надо делать в low-latency коде, тут больше про блокировки и нестабильное время выполнения.
Есть ссылка на интересную реализацию спинлока, где ожидание сделано с постепенным увеличением паузы: progressive_backoff_wait
#cpp #threading #lockfree
Optimizing Compute Shaders for L2 Locality using Thread-Group ID Swizzling
Суть в том, что память текстуры расположенна нелинейна, а в виде Z-curve, morton order и тд.
Поэтому при последовательном доступе начинаются кэшпромахи.
Еще в далеком 2018 я написал прогу, которая за счет хака получает данные текстуры с
Сейчас обновил код: detect z-curve.
Шаблон перестановки пикселей зависит от производителя и от поколения ГП, но чаще всего минимальный размер тайла 4х4 для 32-битного формата.
#vk #gpu_opt
Суть в том, что память текстуры расположенна нелинейна, а в виде Z-curve, morton order и тд.
Поэтому при последовательном доступе начинаются кэшпромахи.
Еще в далеком 2018 я написал прогу, которая за счет хака получает данные текстуры с
optimal_layout, в результате чего можно увидеть как идет перестановка данных.Сейчас обновил код: detect z-curve.
Шаблон перестановки пикселей зависит от производителя и от поколения ГП, но чаще всего минимальный размер тайла 4х4 для 32-битного формата.
#vk #gpu_opt
Doing dynamic resolution scaling? Watch out for texture memory size!
Разбираются случаи, когда для большего разрешения текстуры нужно меньше памяти.
Проверил на AMD RX570, Vulkan возвращает разное выравнивание, в одном случае это 256 байт, а в другом 128Кб, поэтому такая разница в размерах. Похоже что у AMD есть оптимизация для тайлов 256x128 для 32-битных форматов.
Воспроизводится только с
Это не единственная проблема, в редких случаях драйвер может возвращать меньшее выравнивание, чем требует спецификация, что приводит к проблемам. Например AMD возвращал выравнивание 4 байта для scratch buffer, где требуется 256 байт.
#vk
Разбираются случаи, когда для большего разрешения текстуры нужно меньше памяти.
Проверил на AMD RX570, Vulkan возвращает разное выравнивание, в одном случае это 256 байт, а в другом 128Кб, поэтому такая разница в размерах. Похоже что у AMD есть оптимизация для тайлов 256x128 для 32-битных форматов.
Воспроизводится только с
usage = transfer_dst | sampled, добавление color_attachment устраняет эту проблему.Это не единственная проблема, в редких случаях драйвер может возвращать меньшее выравнивание, чем требует спецификация, что приводит к проблемам. Например AMD возвращал выравнивание 4 байта для scratch buffer, где требуется 256 байт.
#vk
Explore GPU advancements in M3 and A17 Pro
4:39 - Архитектура ГП.
7:09 - Что-то вроде гипертрединга, пока одна группа потоков (SIMDgroup) ждет чтение из памяти, другая группа потоков делает рассчеты на том же ядре. Динамическая память регистров позволяет разместить 2 потока в одном.
11:14 - 3 разных кэша объединили в один большой. Если какой-то кэш не используется, то память не будет простаивать.
14:18 - ГП оптимизированы под вычисления на fp16. Fp32, fp16, int выполняются параллельно, пока между ними нет зависимости (конвертация типов, общая память и тд).
16:54 - Как улучшилась производительность при переносе трассировки лучей в железо. Добавили перераспределение потоков, аналогично SER от NVidia.
23:12 - Перераспределение потоков работает только при использовании
24:50 - Про меш шейдеры. Теперь меньше данных выгружается в ситемную память и больше остается в памяти чипа. Но говорят только про промежуточные данные мешлета, то есть данные между таск и меш шейдером. Тогда как в тайловой архитектуре примитивы еще нарезаются и выгружаются в системную память, похоже эту часть не оптимизировали.
#apple_gpu #metal
4:39 - Архитектура ГП.
7:09 - Что-то вроде гипертрединга, пока одна группа потоков (SIMDgroup) ждет чтение из памяти, другая группа потоков делает рассчеты на том же ядре. Динамическая память регистров позволяет разместить 2 потока в одном.
11:14 - 3 разных кэша объединили в один большой. Если какой-то кэш не используется, то память не будет простаивать.
14:18 - ГП оптимизированы под вычисления на fp16. Fp32, fp16, int выполняются параллельно, пока между ними нет зависимости (конвертация типов, общая память и тд).
16:54 - Как улучшилась производительность при переносе трассировки лучей в железо. Добавили перераспределение потоков, аналогично SER от NVidia.
23:12 - Перераспределение потоков работает только при использовании
intersector<>::intersect(), и не работает для ray query.24:50 - Про меш шейдеры. Теперь меньше данных выгружается в ситемную память и больше остается в памяти чипа. Но говорят только про промежуточные данные мешлета, то есть данные между таск и меш шейдером. Тогда как в тайловой архитектуре примитивы еще нарезаются и выгружаются в системную память, похоже эту часть не оптимизировали.
#apple_gpu #metal
Как приостановить поток.
Вариантов достаточно много, начнем с самого простого: _mm_pause() на x64, аналогично YieldProcessor() в WinAPI.
Для ЦП с гипертредингом, где одновременно выполняются 2 потока на одном ядре, эта инструкция позволяет второму потоку начать выполнение.
Имеет смысл для циклов со спинлоком, когда второй поток использует тот же спинлок.
Инструкция выполняется за константное время (20-30нс), что позволяет приостанавливать ЦП на очень короткое время.
std::this_thread::yield() зависит от реализации, на WinAPI это соответствует SwitchToThread(), который предлагает ОС заменить текущий поток на другой, который ожидает выполнения.
В отличие от std::this_thread::yield(), SwitchToThread() возвращает bool - произошло ли переключение потоков или нет.
Если переключение потоков не произошло, а приостановить его все равно хочется, то используется Sleep() из WinAPI. Sleep(0) аналогично SwitchToThread(), а Sleep(1) занимает 15мс.
Если не трогать timeBeginPeriod() то ОС может переключать потоки 64 раза в секунду, отсюда и шаг в ~15мс на пробуждение потока.
Более частое переключение приведет к частым сменам контекста (context switch), что очень дорого.
Вызов функций ОС сам по себе занимает время, поэтому SwitchToThread() и Sleep(0) потратит 150нс, а если произойдет переключение потоков, то время увеличивается до 15мс+.
Вариант std::this_thread::sleep_for() и sleep_until() устроены немного сложнее. Для значений более 1мкс там вызывается Sleep(max(t,1)) и получается тот же шаг в ~15мс.
Для значений менее 1мкс поведение меняется и время ожидания снижается до 1мс - 6мс, причем все зависит от настроек компиляции, в дебаге - 1мс, в релизе - 4..6мс.
Более короткие интервалы дают таймеры: CreateWaitableTimerEx() с флагом CREATE_WAITABLE_TIMER_HIGH_RESOLUTION и SetWaitableTimerEx(). Шаг задается в 100нс, но тесты показывают ~10мкс для 100-1000нс и ~0.5мс для 10-100мкс, для >0.5мс погрешность составляет 0..0.5мс.
Флаг CREATE_WAITABLE_TIMER_HIGH_RESOLUTION появился в Win10 1803, без флага шаг таймера увеличивается до ~15мс. На более старых версиях использование флага приведет к ошибке создания таймера.
В итоге под Windows можно получить стабильные паузы:
100нс - 5-6 вызовов YieldProcessor().
0.5мс - SetWaitableTimerEx() с временем до 0.5мс.
15мс - Sleep(1..15), это мало применимо для нагруженных приложений (игр), но сгодится для тестов.
Ссылки:
progressive_backoff_wait - прогрессивное ожидание под x64 и Arm64.
Windows Timer Resolution: The Great Rule Change - про timeBeginPeriod.
#cpp
Вариантов достаточно много, начнем с самого простого: _mm_pause() на x64, аналогично YieldProcessor() в WinAPI.
Для ЦП с гипертредингом, где одновременно выполняются 2 потока на одном ядре, эта инструкция позволяет второму потоку начать выполнение.
Имеет смысл для циклов со спинлоком, когда второй поток использует тот же спинлок.
Инструкция выполняется за константное время (20-30нс), что позволяет приостанавливать ЦП на очень короткое время.
std::this_thread::yield() зависит от реализации, на WinAPI это соответствует SwitchToThread(), который предлагает ОС заменить текущий поток на другой, который ожидает выполнения.
В отличие от std::this_thread::yield(), SwitchToThread() возвращает bool - произошло ли переключение потоков или нет.
Если переключение потоков не произошло, а приостановить его все равно хочется, то используется Sleep() из WinAPI. Sleep(0) аналогично SwitchToThread(), а Sleep(1) занимает 15мс.
Если не трогать timeBeginPeriod() то ОС может переключать потоки 64 раза в секунду, отсюда и шаг в ~15мс на пробуждение потока.
Более частое переключение приведет к частым сменам контекста (context switch), что очень дорого.
Вызов функций ОС сам по себе занимает время, поэтому SwitchToThread() и Sleep(0) потратит 150нс, а если произойдет переключение потоков, то время увеличивается до 15мс+.
Вариант std::this_thread::sleep_for() и sleep_until() устроены немного сложнее. Для значений более 1мкс там вызывается Sleep(max(t,1)) и получается тот же шаг в ~15мс.
Для значений менее 1мкс поведение меняется и время ожидания снижается до 1мс - 6мс, причем все зависит от настроек компиляции, в дебаге - 1мс, в релизе - 4..6мс.
Более короткие интервалы дают таймеры: CreateWaitableTimerEx() с флагом CREATE_WAITABLE_TIMER_HIGH_RESOLUTION и SetWaitableTimerEx(). Шаг задается в 100нс, но тесты показывают ~10мкс для 100-1000нс и ~0.5мс для 10-100мкс, для >0.5мс погрешность составляет 0..0.5мс.
Флаг CREATE_WAITABLE_TIMER_HIGH_RESOLUTION появился в Win10 1803, без флага шаг таймера увеличивается до ~15мс. На более старых версиях использование флага приведет к ошибке создания таймера.
В итоге под Windows можно получить стабильные паузы:
100нс - 5-6 вызовов YieldProcessor().
0.5мс - SetWaitableTimerEx() с временем до 0.5мс.
15мс - Sleep(1..15), это мало применимо для нагруженных приложений (игр), но сгодится для тестов.
Ссылки:
progressive_backoff_wait - прогрессивное ожидание под x64 и Arm64.
Windows Timer Resolution: The Great Rule Change - про timeBeginPeriod.
#cpp
Как работает буферизация файлов в Windows.
Документация: File Buffering
При первом доступе к файлу ОС начинает кэшировать прочитанные/записанные данные в оперативную память. При следующем чтении данные будут загружены не с диска, а из памяти, в большинстве случаев это полезная функция, кроме случаев, когда нужно профилировать работу с файлами.
Иногда кэширование в ОС работает непредсказуемо - одна программа читает файл на 10Гб и он весь кэшируется, другая программа использует FILE_FLAG_OVERLAPPED и кэширование не включается, а в каком-то редком случае может включиться и закэшировать.
Для профилирования есть флаг FILE_FLAG_NO_BUFFERING, но он дает дополнительные ограничения - позиция чтения и размер блока должны быть выровнены по 512-4096 в зависимости от параметров жесткого диска. Флаг также отключает префетчинг FILE_FLAG_SEQUENTIAL_SCAN, что снизит скорость чтения мелких блоков.
В итоге для правильного профилирования работы с файлами придется использовать флаг FILE_FLAG_NO_BUFFERING и писать свой префетчинг.
С асинхронными файлами чаще всего кэширование всего файла не включается, но чтение невыровненых блоков потребует буферизацию в ОС.
Документация: File Buffering
При первом доступе к файлу ОС начинает кэшировать прочитанные/записанные данные в оперативную память. При следующем чтении данные будут загружены не с диска, а из памяти, в большинстве случаев это полезная функция, кроме случаев, когда нужно профилировать работу с файлами.
Иногда кэширование в ОС работает непредсказуемо - одна программа читает файл на 10Гб и он весь кэшируется, другая программа использует FILE_FLAG_OVERLAPPED и кэширование не включается, а в каком-то редком случае может включиться и закэшировать.
Для профилирования есть флаг FILE_FLAG_NO_BUFFERING, но он дает дополнительные ограничения - позиция чтения и размер блока должны быть выровнены по 512-4096 в зависимости от параметров жесткого диска. Флаг также отключает префетчинг FILE_FLAG_SEQUENTIAL_SCAN, что снизит скорость чтения мелких блоков.
В итоге для правильного профилирования работы с файлами придется использовать флаг FILE_FLAG_NO_BUFFERING и писать свой префетчинг.
С асинхронными файлами чаще всего кэширование всего файла не включается, но чтение невыровненых блоков потребует буферизацию в ОС.
std::find() and memchr() Optimizations
Разбирается как оптимизировать memchr с помощью SIMD, но в статье SIMD не дает значительного ускорения.
По моим тестам стандартная реализация find() и memchr по производительности аналогична версии с uint64_t, а более новые ЦП в 4 раза быстрее выполняют AVX2 версию.
Для большей оптимизации можно добавить принудительный инлайнинг, а также аттрибуты likely/unlikely.
Разбирается как оптимизировать memchr с помощью SIMD, но в статье SIMD не дает значительного ускорения.
По моим тестам стандартная реализация find() и memchr по производительности аналогична версии с uint64_t, а более новые ЦП в 4 раза быстрее выполняют AVX2 версию.
Для большей оптимизации можно добавить принудительный инлайнинг, а также аттрибуты likely/unlikely.
for_likely (; i < e; i += 32)
{
__m256i x = _mm256_lddqu_si256(
reinterpret_cast<const __m256i *>( i ));
__m256i r = _mm256_cmpeq_epi8( x, q );
int z = _mm256_movemask_epi8( r );
if_unlikely( z ) {
auto* r2 = i + BitScanForward( z );
return Min( r2, e );
}
}
#cpp #cpu_optПро асинхронные файлы
WinAPI
Открывается файл с FILE_FLAG_OVERLAPPED, создается общий IO-port и IO-port для каждого файла. В общий IO-port приходят события о завершении чтения/записи.
Из минусов - GetQueuedCompletionStatus() это системный вызов, что медленно, поэтому нет быстрого способа проверить, что есть новые события.
POSIX AIO
Внутри использует thread pool, доступен на Linux и MacOS. Смысла использовать эту либу нет, проще написать свой аналог.
Linux AIO
В <linux/aio_abi.h> определены прототипы системных вызовов, отдельно есть обертка libaio.
Рекомендуют использовать один контекст на поток, иначе это повлияет на производительность.
Запись в режиме append может быть синхронной. В некоторых случаях блокировка все равно происходит, флаг O_DIRECT снижает эту вероятность, но не исключает полностью.
Из плюсов - один системный вызов на чтение/запись (io_submit), а чтение завершеных операций можно производить без системного вызова (io_getevents) через доступ к aio_ring.
С помощью счетчика ru_nvcsw в функции getrusage() можно проверить происходит ли блокировка и смена контекста.
io_uring
Замена LinuxAIO в новых версиях, но из-за множества уязвимостей не используется в Android и некоторых Linux.
Пользоваться оберткой liburing намного удобнее, чем низкоуровневым <linux/io_uring.h>.
Рекомендуют использовать один контекст на поток.
В обычных тестах по производительности разницы с AIO нет, может в каких-то специфических случаях не будет блокировок как в AIO.
Linux AIO и io_uring при инициализации принимают максимальное количество событий. На разных системах разный максимальный размер очереди, если его превысить, то вернет ошибку EAGAIN.
Dispatch IO
Замена POSIX AIO на MacOS, также использует потоки. На новых маках стоят быстрые SSD, поэтому синхронные операции работают намного быстрее.
В итоге на MacOS я не нашел настоящих асинхронных файлов.
Выводы.
В одном потоке производительность асинхронного чтения/записи намного выше, конечно же при комбинации с системой тасков или корутинами.
Даже случайный доступ к файлам (SSD, eMMC) работает также быстро как последовательный.
Под MacOS/iOS придется менять логику, лучше комбинировать запросы на чтение/запись или использовать кэш ОС, иначе производительность меньше чем на дешевых Android.
Во многих игровых движках и boost asio используется максимум WinAPI асинхронные файлы, только в движках базданных нашел вариант под Linux.
WinAPI
Открывается файл с FILE_FLAG_OVERLAPPED, создается общий IO-port и IO-port для каждого файла. В общий IO-port приходят события о завершении чтения/записи.
Из минусов - GetQueuedCompletionStatus() это системный вызов, что медленно, поэтому нет быстрого способа проверить, что есть новые события.
POSIX AIO
Внутри использует thread pool, доступен на Linux и MacOS. Смысла использовать эту либу нет, проще написать свой аналог.
Linux AIO
В <linux/aio_abi.h> определены прототипы системных вызовов, отдельно есть обертка libaio.
Рекомендуют использовать один контекст на поток, иначе это повлияет на производительность.
Запись в режиме append может быть синхронной. В некоторых случаях блокировка все равно происходит, флаг O_DIRECT снижает эту вероятность, но не исключает полностью.
Из плюсов - один системный вызов на чтение/запись (io_submit), а чтение завершеных операций можно производить без системного вызова (io_getevents) через доступ к aio_ring.
С помощью счетчика ru_nvcsw в функции getrusage() можно проверить происходит ли блокировка и смена контекста.
io_uring
Замена LinuxAIO в новых версиях, но из-за множества уязвимостей не используется в Android и некоторых Linux.
Пользоваться оберткой liburing намного удобнее, чем низкоуровневым <linux/io_uring.h>.
Рекомендуют использовать один контекст на поток.
В обычных тестах по производительности разницы с AIO нет, может в каких-то специфических случаях не будет блокировок как в AIO.
Linux AIO и io_uring при инициализации принимают максимальное количество событий. На разных системах разный максимальный размер очереди, если его превысить, то вернет ошибку EAGAIN.
Dispatch IO
Замена POSIX AIO на MacOS, также использует потоки. На новых маках стоят быстрые SSD, поэтому синхронные операции работают намного быстрее.
В итоге на MacOS я не нашел настоящих асинхронных файлов.
Выводы.
В одном потоке производительность асинхронного чтения/записи намного выше, конечно же при комбинации с системой тасков или корутинами.
Даже случайный доступ к файлам (SSD, eMMC) работает также быстро как последовательный.
Под MacOS/iOS придется менять логику, лучше комбинировать запросы на чтение/запись или использовать кэш ОС, иначе производительность меньше чем на дешевых Android.
Во многих игровых движках и boost asio используется максимум WinAPI асинхронные файлы, только в движках базданных нашел вариант под Linux.
Подробнее про POSIX AIO.
У каждой ОС есть свои нюансы реализации этого API:
В Android не поддерживается.
В Linux есть дополнительная функция aio_init() в которой описывается реализация API через thread pool, но встречается информация, что есть реализация через Linux AIO с настоящей асинхронщиной.
Например Samba vfs делает асинхронные вызовы, если размер буфера больше лимита, но поддерживает максимум 10 запросов суммарно на чтение и запись.
В BSD системах AIO поддерживается и есть удобная реализация через kqueue (пример).
В macOS, в отличие от других BSD, kqueue не поддерживает AIO, зато это единственное настоящее AIO под macOS который работает быстрее синхронных вызовов. Из минусов - очень маленькая очередь, после 16 вызовов aio_read начинает возвращать ошибку и нужно дождаться завершения предыдущих вызовов, в итоге производительность проседает.
Получить лимиты можно вызовом
Можно поменять параметры, для лучшей производительности:
У каждой ОС есть свои нюансы реализации этого API:
В Android не поддерживается.
В Linux есть дополнительная функция aio_init() в которой описывается реализация API через thread pool, но встречается информация, что есть реализация через Linux AIO с настоящей асинхронщиной.
Например Samba vfs делает асинхронные вызовы, если размер буфера больше лимита, но поддерживает максимум 10 запросов суммарно на чтение и запись.
В BSD системах AIO поддерживается и есть удобная реализация через kqueue (пример).
В macOS, в отличие от других BSD, kqueue не поддерживает AIO, зато это единственное настоящее AIO под macOS который работает быстрее синхронных вызовов. Из минусов - очень маленькая очередь, после 16 вызовов aio_read начинает возвращать ошибку и нужно дождаться завершения предыдущих вызовов, в итоге производительность проседает.
Получить лимиты можно вызовом
sysctl -a | grep aio в терминале. kern.aioprocmax - максимум IO запросов на весь процесс.Можно поменять параметры, для лучшей производительности:
sudo sysctl kern.aiomax=2048 sudo sysctl kern.aioprocmax=1024 < aiomaxsudo sysctl kern.aiothreads=8 = количество ядер ЦПАсинхронные файлы, тесты.
В одном потоке читается файл на 128Мб блоками по 4Кб, отключены все кэши ОС, после чтения данные проверяются на корректность.
Наилучший результат показал Android с LinuxAIO:
Последовательное и случайное асинхронное чтение 128Мб занимает 0.5с.
Синхронное чтение занимает 5с.
Рассчитанная скорость eMMC - 300Мб/с, асинхронное чтение выдает 85% от максимума.
Наихудший результат на macOS с PosixAIO.
Последовательное асинхронное чтение - 0.35с, случайное - 1.17с.
Последовательное синхронное чтение - 0.74с, случайное - 2.06с.
При заявленной скорости SSD в 3.5Гб/с даже асинхронное чтение выдает всего 10% от максимума, случайное чтение в 3 раза медленнее, хотя на других ОС разницы нет.
С параметрами
Последовательное асинхронное чтение - 0.21с, случайное - 0.32с.
macOS с dispath_io оказался в 5 раз медленнее синхронного чтения, так как использует thread pool.
Windows 10.
Последовательное и случайное асинхронное чтение 128Мб занимает 0.5с, при скорости SSD в 550Мб/с, это 50% от максимума.
Последовательное синхронное чтение - 1.6с, случайное - 4.2с.
Ubuntu 23 с io_uring.
Последовательное и случайное асинхронное чтение - 0.12с.
Последовательное синхронное чтение - 0.66с, случайное - 2.66с.
Скорость SSD 3.5Гб/с, асинхронное чтение дает 30% от максимума.
Сам запрос на чтение занимает в 2-3 раза больше, чем на Windows 10, что зависит от ОС и файловой системы, а не от скорости диска.
В одном потоке читается файл на 128Мб блоками по 4Кб, отключены все кэши ОС, после чтения данные проверяются на корректность.
Наилучший результат показал Android с LinuxAIO:
Последовательное и случайное асинхронное чтение 128Мб занимает 0.5с.
Синхронное чтение занимает 5с.
Рассчитанная скорость eMMC - 300Мб/с, асинхронное чтение выдает 85% от максимума.
Наихудший результат на macOS с PosixAIO.
Последовательное асинхронное чтение - 0.35с, случайное - 1.17с.
Последовательное синхронное чтение - 0.74с, случайное - 2.06с.
При заявленной скорости SSD в 3.5Гб/с даже асинхронное чтение выдает всего 10% от максимума, случайное чтение в 3 раза медленнее, хотя на других ОС разницы нет.
С параметрами
kern.aioprocmax=1024 производительность улучшилась:Последовательное асинхронное чтение - 0.21с, случайное - 0.32с.
macOS с dispath_io оказался в 5 раз медленнее синхронного чтения, так как использует thread pool.
Windows 10.
Последовательное и случайное асинхронное чтение 128Мб занимает 0.5с, при скорости SSD в 550Мб/с, это 50% от максимума.
Последовательное синхронное чтение - 1.6с, случайное - 4.2с.
Ubuntu 23 с io_uring.
Последовательное и случайное асинхронное чтение - 0.12с.
Последовательное синхронное чтение - 0.66с, случайное - 2.66с.
Скорость SSD 3.5Гб/с, асинхронное чтение дает 30% от максимума.
Сам запрос на чтение занимает в 2-3 раза больше, чем на Windows 10, что зависит от ОС и файловой системы, а не от скорости диска.
ARM Valhalla 5gen
Самое интересное это Deferred Vertex Shading.
В тайловой архитектуре сначала все треугольники проецируются на экран, нарезаются на тайлы и выгружаются в системную память, а затем тайлы достаются и для них уже выполняется растеризация. Все это плохо работает с большим количеством треугольников, так как увеличивается нагрузка на память, а это дорого.
С DVS тайлер может выбирать какой треугольник сразу отправить на рисование, а какой сохранить в список (в глобальной памяти) для отложеной растеризации.
Маленькие треугольники рисуются сразу, а большие откладываются, так как требуется задействовать много тайлов, что также нагружает память.
Это еще один шаг в сторону меш шейдеров, которые из-за тайловой архитектуры не дают преимущества в производительности.
#mali
Самое интересное это Deferred Vertex Shading.
В тайловой архитектуре сначала все треугольники проецируются на экран, нарезаются на тайлы и выгружаются в системную память, а затем тайлы достаются и для них уже выполняется растеризация. Все это плохо работает с большим количеством треугольников, так как увеличивается нагрузка на память, а это дорого.
С DVS тайлер может выбирать какой треугольник сразу отправить на рисование, а какой сохранить в список (в глобальной памяти) для отложеной растеризации.
Маленькие треугольники рисуются сразу, а большие откладываются, так как требуется задействовать много тайлов, что также нагружает память.
Это еще один шаг в сторону меш шейдеров, которые из-за тайловой архитектуры не дают преимущества в производительности.
#mali
Material Depth Buffer
Техника применяется для отложенного текстурирования в Dawn Engine, где ID материала записывают в 16 битный буфер глубины, а далее рисуется полноэкранные квадраты с разными пайплайнами и тестом глубины equal. ГП умеют это оптимизировать - после теста глубины, прошедшие тест пиксели группируются, чтобы полностью заполнить варп, таким образом 90% ядер ГП загружены.
Сравнение производительности полноэкранного прохода с самым медленным материалом со множеством материалов, если ГП умеет оптимизировать, то падение производительности будет около x2.
На Adreno для полноэкранного прохода не используется TBDR, возможно из-за этого получилось падение производительности.
#gpu_opt
Техника применяется для отложенного текстурирования в Dawn Engine, где ID материала записывают в 16 битный буфер глубины, а далее рисуется полноэкранные квадраты с разными пайплайнами и тестом глубины equal. ГП умеют это оптимизировать - после теста глубины, прошедшие тест пиксели группируются, чтобы полностью заполнить варп, таким образом 90% ядер ГП загружены.
Сравнение производительности полноэкранного прохода с самым медленным материалом со множеством материалов, если ГП умеет оптимизировать, то падение производительности будет около x2.
NV RTX 2080: 2мс / 5мс.
Mac M1: 2.5мс / 6мс.
AMD RX 570: 5мс / 10мс.
Intel UHD 620: 3мс / 12мс.
Mali G57: 20мс / 26мс.
Adreno 660: 3мс / 30мс, 3 материала - 10мс.На Adreno для полноэкранного прохода не используется TBDR, возможно из-за этого получилось падение производительности.
#gpu_opt
REAC 2024 Resource Management Architecture in 4A Engine
pdf, video
Как переделали движок 4A.
Одна из идей - валидация ресурсов как можно раньше, больше проверок кода на этапе компиляции, а также тулзы для ускорения поиска ошибок и быстрых итераций.
Также улучшили работу со стримингом ресурсов, управлением памятью и многопоточностью.
11-13. В коде используют отдельные типы для RT, UAV и тд, в зависимости от типа выводятся флаги для барьеров. Это позволяет делать проверки на этапе компиляции.
37. С bindless моделью нельзя узнать какой ресурс реально используется при рендере, поэтому собирают статистику на ГП (видимо счетчики при доступе к текстуре). Статистика передается на ЦП, где принимается решение о выгрузке данных, загрузке более высокой детализации и тд.
46. Написали тулзы для отладки шейдеров. Они в том числе позволяют делать визуализацию прямо из шейдера, пример кода на слайде.
pdf, video
Как переделали движок 4A.
Одна из идей - валидация ресурсов как можно раньше, больше проверок кода на этапе компиляции, а также тулзы для ускорения поиска ошибок и быстрых итераций.
Также улучшили работу со стримингом ресурсов, управлением памятью и многопоточностью.
11-13. В коде используют отдельные типы для RT, UAV и тд, в зависимости от типа выводятся флаги для барьеров. Это позволяет делать проверки на этапе компиляции.
37. С bindless моделью нельзя узнать какой ресурс реально используется при рендере, поэтому собирают статистику на ГП (видимо счетчики при доступе к текстуре). Статистика передается на ЦП, где принимается решение о выгрузке данных, загрузке более высокой детализации и тд.
46. Написали тулзы для отладки шейдеров. Они в том числе позволяют делать визуализацию прямо из шейдера, пример кода на слайде.
Путеводитель C++ программиста по неопределенному поведению
Очень подробно разбирается множество случаев неопределенного поведения в C++.
loop-counters-signed-vs-unsigned
В дополнение - NVidia рекомендует использовать неопределенное поведение при переполении знакового int чтобы компилятор CUDA лучше оптимизировал циклы.
Если переполнение не происходит (иначе это UB), то компилятор может использовать strength reduction - заменяет
#cpp #rus
Очень подробно разбирается множество случаев неопределенного поведения в C++.
loop-counters-signed-vs-unsigned
В дополнение - NVidia рекомендует использовать неопределенное поведение при переполении знакового int чтобы компилятор CUDA лучше оптимизировал циклы.
Если переполнение не происходит (иначе это UB), то компилятор может использовать strength reduction - заменяет
stride*i на сложение:было:
for (i = 0; i < n; i++)
out[i] = in[offset + stride*i];
стало:
for (i = 0, k = 0; i < n; i++) {
out[i] = in[offset + k];
k += stride;
}#cpp #rus