
Говорят, настоящий программист должен уметь писать под GPU. Но каждый раз, когда дело доходит до изучения всех grid/block-семантик, kernel launch и ad-hoc-синхронизации...
Похоже, в AMD это поняли и решили сделать шаг к соединению двух миров. hipThreads – библиотека, которая позволяет переносить многопоточный C++-код с CPU на GPU, меняя std::thread на hip::thread и добавляя пару строк обвязки. Остальное – магия ROCm и persistent-планировщика.
Исходный код уже доступен на GitHub, а подробный разбор с примерами опубликован в официальном блоге AMD.
hipThreads – это C++‑библиотека параллелизма для AMD GPU, которая приносит знакомую потоковую абстракцию в мир видеокарт. Вместо того чтобы осваивать конфигурацию ядер и модель блоков, разработчик просто пишет hip::thread, hip::mutex, hip::lock_guard, hip::condition_variable – совсем как в выражении std::thread. Библиотека заполняет нишу между мощными, но высокоуровневыми инструментами вроде rocThrust и низкоуровневыми примитивами libhipcxx.
С точки зрения архитектуры hipThreads запускает на GPU один долгоживущий поток-менеджер, который ставит в очередь виртуальные потоки и выполняет их совместно. Каждый такой hip::thread может содержать несколько фибр – параллельных исполнительных линий, использующих SIMD‑ширину GPU. Если вы знакомы с AVX‑512, представьте, что теперь у вас не 16, а 32 слота для данных и всё это упаковано в привычный потоковый интерфейс.
Чтобы оценить простоту, достаточно взглянуть на пример SAXPY (одинарной точности α·X + Y), который авторы блога разобрали по шагам. Сначала идёт обычный CPU‑код с std::thread, разбивающий массивы на чанки. Затем – минимальные правки: включаем <hip/thread> вместо <thread>, добавляем управление памятью через thrust::unique_ptr и заменяем std::thread на hip::thread. Уже на этом этапе код улетает на GPU, хоть и без использования фибр.
На третьем шаге добавляется указание ширины потока (hip::thread::max_width()), а цикл переписывается с учётом идентификатора фибры:
for(uint32_t i = hip::this_thread::get_fiber_id(); i < n; i += hip::this_thread::get_width()) { // обработка элементов }
Итоговые изменения – всего 16 строк, а 96% кода остаются прежними.
В AMD привели результаты замеров на тестовой системе с Ryzen 9 9900X и Radeon AI PRO R9700. Даже без глубокой оптимизации hipThreads показывает впечатляющий прирост:
Платформа | Реализация | Время, мс | SIMD-ширина | Относительно CPU |
AMD Ryzen 9 9900X | std::thread + AVX-512 | 271,88 | 16 | 1,00× |
AMD Radeon AI PRO R9700 | hip::thread (32 фибры) | 42,60 | 32 | 6,4× быстрее |
Проверка на других задачах подтверждает тенденцию:
Разреженное умножение матриц: 3,6× быстрее.
Трассировщик лучей InOneWeekend: 2,9× быстрее.
Причём всё это достигнуто без кардинальной смены парадигмы – просто знакомые потоки, теперь живущие на GPU.
Для работы hipThreads понадобится Linux (рекомендуется Ubuntu 22.04+), CMake 3.21+, ROCm 7.0.2 (именно эта версия, никакая другая) и libhipcxx v2.7. Библиотека ставится стандартно через CMake и прописывается в /opt/rocm. В своих проектах достаточно добавить find_package(hipthreads REQUIRED) и линковаться с hipthreads::hipthreads.
Разработчики обещают развивать интеграцию с асинхронным API, добавлять новые примитивы синхронизации и упрощать типовые паттерны портирования.
Делегируйте часть рутинных задач вместе с BotHub! Для доступа к сервису не требуется VPN и можно использовать российскую карту. По ссылке вы можете получить 300 000 бесплатных токенов для первых задач и приступить к работе с нейросетями прямо сейчас!
