Предисловие
Эта статья, по-сути, краткое описание моего опыта разработки приложений на OpenCL/SYCL. По большому счёту, это просто «рассуждения о жизни» в рамках того, что я успел подметить во время работы с гетерогенными программами.
Вряд-ли этот текст представляет большую ценность, чем просто ознакомление с некоторыми концепциями. Тут будет очень мало конкретики в угоду лирики. Я намеренно не буду давать пояснения многим вещам, ввиду того, что эта статья не призвана объяснить, что такое гетерогенные системы, а лишь заинтересовать потенциального читателя в теме.
Все примеры из данной статьи опубликованы в репозитории на GitHub – вы можете самостоятельно изучить их и воспроизвести на своём железе.
Представленные сравнения являются лишь демонстрацией на уровне разбираемых примеров и не призваны быть бенчмарком для железа.
Видеокарты
Вычисления на видеокартах ассоциируются преимущественно с графикой, научными рассчётами и с недавних пор с нейросетями.
Но чаще всего с графикой. И это вполне закономерно, ведь как правило мы используем видеокарты, чтобы выводить картинку на монитор, покупаем их с расчётом на системные требования игр или графических приложений, да и само название – GPU, то есть Graphics Processing Unit, или как его ещё называют в русскоязычной литературе – Графический процессор, напрямую отсылает к рендерингу чего-то, даже если это просто графический интерфейс ОС.
Тем не менее, графические процессоры обладают свойствами, за счёт которых их очень удобно использовать в задачах, напрямую не связанных с перечисленным выше. И главное из этих свойств – массовый параллелизм.
От чего вычисление матриц, битоническая сортировка, свёртки и старые добрые векторы так удобно считать на видеокартах – они делают это в разы быстрее, чем обычный процессор, хоть и с оговорками по точности результатов.
Разумеется в CPU никто не отменял параллельные вычисления, SIMD и кучу соответствующих расширений архитектур (AVX, NEON, RVV). Однако они почти всегда проигрывают GPU в скорости, если речь о массово-параллельной задаче.
И тут мы плавно переходим к GPGPU или GPU-compute, или «гетерогенному программированию». Строго говоря, все эти понятия не одно и тоже, но я в той или иной степени использую их в относительно одинаковом смысле.
Гетерогенное программирование
Самый краткий ввод в гетерогенное программирование: у нас есть две роли – хост и девайс, задача хоста – формировать задачи и отдавать (offload) их на девайс, в то время как задача девайса обработать их и вернуть результат. При этом хост и девайс могут быть одним вычислительным устройством (это не обязательно CPU + GPU, так как CPU может отдавать задачи сам себе).
Теперь к сути.
OpenCL и его несовершенство
У нас есть программа-драйвер на OpenCL C++, которая ищет первый попавшийся GPU, главное, чтобы у него было хотя бы 1 вычислительное ядро, и отдаёт ему на исполнение kernel с программой перемножения матриц (matrix_localmem.cl). Она разбивает данные на тайлы (по 16) и использует локальную память. Хостовая часть выполняет аналогичный алгоритм на CPU (правда, одноядерный – да-да, сравнение не совсем честное) и замеряет время исполнения программы на GPU и на CPU. При этом на GPU отдельно учитывается общее время (с накладными расходами) и время исполнения исключительно kernel'а.
И того, на GPU мы имеем общее время исполнения в 8 секунд и время исполнения kernel'а в 5 секунд, против 131 секунды на CPU (изображение 1).
Собственно, результат вы видите сами: GPU уделывает CPU в 16 раз, если брать общее время и в 26 раз, если брать только время выполнения на видеокарте. При том, что матрица довольно маленькая (1024). Даже если сделать скидку на однопоточность CPU-программы, то отрыв бы сократился, но не намного.
В качестве заметки, возможно мне стоит написать хороший многопоточный алгоритм для CPU и провести более релевантное сравнение для подобного рода вычислений.
Также есть версия kernel'а с довольно наивным и откровенно тупым перемножением матриц (matrix_simple.cl), которую я написал для сравнения скорости выполнения между наивным и оптимизированным алгоритмом.
Этот kernel выглядит так:
__kernel void matrixmult(__global const float* A, __global const float* B, __global float* C, const unsigned int N) { unsigned int row = get_global_id(0); unsigned int col = get_global_id(1); if (row < N && col < N) { float sum = 0.0f; for (unsigned int k = 0; k < N; ++k) { sum += A[row * N + k] * B[k * N + col]; } C[row * N + col] = sum; } }
Не правда ли мило? И даже он справляется в разы быстрее CPU, что, в общем-то, очевидная истина.
Опять же, это не умаляет возможностей CPU при правильном подходе – я хоть и акцентирую внимание на скорости, но делаю это для того, чтобы показать тот отрыв, который предполагают массово-параллельные устройства в соответствующих задачах.
Однако у OpenCL есть несколько проблем.
Во-первых, это Си'шный API – да, я оборачиваю его в высокоуровневые C++ классы через OpenCL-CLHPP, но де-факто мы каждый раз делаем низкоуровневые C API вызовы.
Во-вторых, kernel на OpenCL – это по-сути просто... текст, который мы отдаём на видеокарту, где он компилируется: да, это можно воспринимать как рабочий си'шный код, но по границам интерфейса он не значит ровным счётом ничего – у нас нет ни контроля типов, ни чего-то ещё.
Это боль моделей с разделённым исходным кодом, с которой люди не смирились.
Поэтому тут мы плавно переходим к SYCL – модели с совмещённым исходным кодом (первой была CUDA, конечно же), которая решает проблему разделения функциональности хостовой и девайсной части программы.
SYCL и память
Возвращаясь к оптимизациям, стоит сказать, что чаще всего под оптимизацией подобных вычислений подразумевается рациональное использование памяти устройства – насколько часто вам придётся перебрасывать данные от хоста к девайсу и обратно.
Память – это одна из ключевых и самых сложных концепций в гетерогенных системах.
Дело в том, что когда мы имеем дело с памятью на GPU, необходимо помнить, что модель памяти на видеокартах несколько отличается от привычной и имеет ��вои подводные камни.
В SYCL модель памяти можно разделить на локальную (доступную одному рабочему элементу (work item) внутри одной рабочей группы (work-group) программы) и глобальную (доступную всем рабочим элементам и группам), где первая – содержит в себе приватную память (private memory) для конкретного рабочего элемента.
Проблема в том, что во время итераций вычисления, каждый раз лезть в глобальную память довольно дорого, так как она, в общем-то, далеко – чаще всего это видеопамять карты.
Поэтому очень часто мы хотим оптимизировать этот процесс. В качестве примера таких оптимизаций, рассмотрим несколько вариаций.
У нас есть ещё одна программа для перемножения матриц и несколько kernel'ов. Мы так же находим первый попавшийся GPU, на котором есть runtime для SYCL. И у нас есть очень простой алгоритм перемножения матриц (помечен как SIMPLE), который работает из глобальной памяти.
На первый взгляд, самой простой оптимизацией будет использование так называемой приватной памяти (private memory) – это небольшое итерационное пространство внутри отдельного compute unit'а, от чего она кажется самой эффективной.
Однако мы компилируем этот пример (помечен как PRIVATE), запускаем и видим, что результат довольно плачевный – он сопоставим со скалярными вычислениями на CPU.
И вот это самое плохое в видеокартах: когда вы думаете, что сделали оптимизацию, но на деле всё стало только хуже.
Давайте всё же попытаемся это оптимизировать: просто разбиваем данные на тайлы, записываем из глобальной памяти и перемножаем в локальной. После чего пишем результат обратно в глобальную память (помечен как LOCAL).
И результат значительно улучшается. Единственное, не стоит забывать втыкать барьеры, дабы избежать рассинхронизации данных.
Послесловие
Пожалуй, главной сложностью при работе с гетерогенными системами и видеокартами в частности, является необходимость начать мыслить в парадигме асинхронного программирования, о чём было упомянуто вскользь.
В заключение хочется сказать, что это довольно обширная сфера, и то, что я рассматриваю в этой статье – даже не верхушка айсберга.
