Привет, я делал доклад на тему «Эмуляция запуска GPU-ядер на CPU или нестандартные приёмы отладки». Выкладываю презентацию, в которой рассказываю про реализацию симулятора, воспроизводящего модель выполнения GPU (Grid → Block → Thread) с полной интеграцией в GDB.
В отличие от CUDA, где есть много иструментов для отладки и анализа, экосистема Vulkan и OpenCL практически лишена инструментов пошаговой отладки общего назначения. Презентация показывает трюки, позволяющие разрабатывать и отлаживать ядра без физического ускорителя.
Ключевые архитектурные решения:
• Lockstep-исполнение на корутинах (C++23): Каждый поток реализуется через std::generator. Вызов co_yield эмулирует аппаратный барьер, а внешний диспетчер синхронно продвигает все генераторы, обеспечивая детерминированное блокирующее выполнение, идентичное __syncthreads().
• Иерархия контекста и памяти: Контекст инкапсулирует координаты потоков и размеры сетки/блока. Выделяются именованные области разделяемой памяти с семантикой, повторяющей shared, для кооперативной работы потоков внутри блока.
• Трекер доступа к памяти: Прокси-обёртка над буферами логирует каждое чтение/запись с привязкой к координатам потока. Постобработка группирует доступы, проверяет последовательность адресов для соседних потоков и рассчитывает метрику эффективности группировки и распределение шагов (strides).
• GDB Python-инфраструктура: Кастомные pretty-printers и специализированные команды позволяют инспектировать состояния блоков, разделяемую память и визуализировать линейные буферы как 2D-матрицы прямо в сеансе отладки.
Презентация
Код проекта
#VulkanCompute #Cpp #Coroutines #GDB #GPUDebugging #SIMT
В отличие от CUDA, где есть много иструментов для отладки и анализа, экосистема Vulkan и OpenCL практически лишена инструментов пошаговой отладки общего назначения. Презентация показывает трюки, позволяющие разрабатывать и отлаживать ядра без физического ускорителя.
Ключевые архитектурные решения:
• Lockstep-исполнение на корутинах (C++23): Каждый поток реализуется через std::generator. Вызов co_yield эмулирует аппаратный барьер, а внешний диспетчер синхронно продвигает все генераторы, обеспечивая детерминированное блокирующее выполнение, идентичное __syncthreads().
• Иерархия контекста и памяти: Контекст инкапсулирует координаты потоков и размеры сетки/блока. Выделяются именованные области разделяемой памяти с семантикой, повторяющей shared, для кооперативной работы потоков внутри блока.
• Трекер доступа к памяти: Прокси-обёртка над буферами логирует каждое чтение/запись с привязкой к координатам потока. Постобработка группирует доступы, проверяет последовательность адресов для соседних потоков и рассчитывает метрику эффективности группировки и распределение шагов (strides).
• GDB Python-инфраструктура: Кастомные pretty-printers и специализированные команды позволяют инспектировать состояния блоков, разделяемую память и визуализировать линейные буферы как 2D-матрицы прямо в сеансе отладки.
Презентация
Код проекта
#VulkanCompute #Cpp #Coroutines #GDB #GPUDebugging #SIMT
🔥4