Наверное, я очень опоздал с изучением CUDA. До недавнего времени даже не знал, что CUDA — это просто C++ с небольшими добавками. Если бы я знал, что изучение её пойдёт как по маслу, я бы столько не медлил. Но если у вас есть багаж привычек C++ , то код на CUDA у вас будет получаться низкокачественным. Поэтому расскажу вам о некоторых уроках, изученных на практике — возможно, мой опыт поможет вам ускорить код.
Слияние блоков памяти
Если у вас множество потоков, работающих над одним массивом в C++, то, вероятно, вы попробуете перебрать его таким образом:
std::vector<T> vec = ...;
size_t per_thread = vec.size() / num_threads;
T * my_slice = vec.data() + per_thread * my_thread_i;
for (size_t i = 0; i < per_thread; ++i) {
do_something(my_slice[i]);
}Таким образом, каждый поток перебирает непрерывный фрагмент памяти. В CUDA эта операция пойдёт медленно, так как предполагается, что все потоки будут загружать память вместе. Таким образом, если поток 0 загружает байты от 0 до 15, то нам нужно, чтобы поток 1 загружал байты от 16 до 31, поток 2 загружал байты от 32 до 47 и т.д. Поэтому цикл будет выглядеть не так, как показано выше, а вот так:
T * data = ...;
size_t num_elements = ...;
for (int i = my_thread_i; i < num_elements; i += num_threads) {
do_something(data[i]);
}Эта техника называется «слиянием блоков памяти» (memory coalescing), где смежные потоки работают со смежными областями памяти. В цикле с малым телом (скалярное произведение) работа идёт втрое быстрее.
Сегодня для максимальной производительности требуется специализированное оборудование
Много лет назад Шон Пэрент показал схему, на которой детально изложено, из чего складывается производительность современного ПК. Я дам её здесь с актуальными числами:
В данном случае сравнивается производительность, теоретически достижимая на ПК с процессорами Ryzen 9950X и RTX 4090. Общая теоретически достижимая производительность составляет ~95 ТФЛОПС. Ещё раз: это теоретический максимум, поэтому, например, на однопоточном ЦП производительность составит всего 5,7 ГГц 4 инструкции на цикл = 22,8 ГФЛОПС. Это синяя линия, которую просто не видно — о настолько малой доле речь. Если задействовать все 32 потока и AVX 512, то можно умножить эту производительность на 3216 = 512, заполнив таким образом красную и жёлтую области графика. Но если вам действительно требуется повысить производительность, то лучше применить графический процессор (GPU), который выведет нас в зелёную часть графика.
Притом, что это всё-таки актуальные цифры, здесь почти полностью упускается производительность, достижимая на GPU. В настоящее время на GPU ставятся специализированные аппаратные компоненты для машинного обучения и трассировки лучей. Добавив всё это на график, получим действительную производительность.
Вот та же самая схема при работе на специализированном оборудовании. Производительность тензорного ядра у меня измеряется в ТФЛОПС, когда я выполняю перемножение матриц BF16. То есть сравнение не совсем честное, так как работа идёт с меньшей точностью (вывод всё равно 32-разрядный). Но все пользуются этим методом при перемножении матриц, думаю, это нормально.
Суть в том, что производительность вашего ПК обеспечивается преимущественно за счёт специализированных чипов. Если вы пишете простой код на CUDA, то большей частью мощностей для производительности просто не успеваете воспользоваться. График выглядит ещё более скособоченным, если построить его для специального графического процессора, рассчитанного на машинное обучение, например H100:
Обратите внимание: теперь по оси X мы заходим за 2000 ТФЛОПС. Если мы не работаем с тензорными ядрами, GPU простаивает в течение более 90 процентов времени. Из-за этого меняются алгоритмы, используемые при глубоком обучении. Если алгоритм A просто может перемножать более крупные матрицы и получать более высококачественные результаты, а алгоритм B достигает лучших результатов, умно выполняя работу фрагментик за фрагментиком, то программист предпочтёт алгоритм A.
Память бывает разная
Память в CUDA устроена сравнительно сложно, но, насколько я понимаю, в CUDA существует условно три различных типа памяти:
Обычная память
Разделяемая память (быстрее)
Регистры (быстрее всего)
Регистры устроены особенно странно. В одном блоке потоков содержится 65536 регистров. Таким образом, в регистрах можно сохранить 256 килобайт данных. То есть больше, чем в разделяемой памяти. Я пытался понять, каким образом некоторые ядра cuDNN могут работать с такой скоростью, какую показывают, и пришёл к выводу, что они держат определённую матрицу полностью в регистрах, где на каждый поток приходится малая часть этой матрицы.
В некоторой степени можно контролировать, сколько у вас регистров. В каждом блоке может быть до 1024 потоков, то есть по умолчанию каждый поток располагает 64 регистрами. Но можно запустить меньше потоков и получить пропорционально больше регистров на поток. Допустим, если вам требуется 150 регистров, чтобы кэшировать некоторые данные, то разделите 65536 на 150 – и получится, что вы можете воспользоваться 436 потоками.
При этом вы по-прежнему пишете на C++, а на этом языке не так просто выразить идею «держи эти данные в регистрах». Как мне видится, лучше всего для этой цели держать в стеке массив фиксированного размера, а затем разматывать командой “#pragma unroll” каждый отдельный цикл, работающий с этим массивом. Этот цикл необходимо разматывать, поскольку на каждой следующей итерации такой размотки цикл должен будет ссылаться уже на другие регистры.
По сравнению с ними разделяемая память устроена просто. В ней вы можете выделить некоторое пространство как кэш под конкретные цели, и данные из этого пространства поступают в совместное распоряжение потоков. Таким образом, его можно использовать в двух целях:
Обеспечивать коммуникацию между потоками
Быстрее загружать данные: если вы хотите загрузить 512 чисел с плавающей точкой, а у вас всего 512 потоков, то каждый из них может загрузить в память одно такое число. Поэтому можно даже обойтись без цикла.
Разделяемость в пределах варпа ~бесплатна
Когда я это увидел впервые, мне это очень понравилось. В CUDA существует понятие «варп» — это группа из 32 потоков, у которых один указатель инструкций на всех. То есть они одновременно работают над общей задачей. Поэтому, например, если вы захотите распараллелить скалярное произведение, то после этого можно будет просуммировать результаты работы всех потоков и получить общий результат за пять шагов при помощи алгоритма параллельного суммирования:
На ЦП от такого алгоритма мало толку, поскольку слишком велики будут издержки на поддержание синхронизации потоков. Но на графическом процессоре потоки синхронизированы, поэтому разделяемость реализуется буквально в пять шагов:
__device__ float add_warp(float x) {
static constexpr const unsigned all = 0xffffffff;
x += __shfl_xor_sync(all, x, 1);
x += __shfl_xor_sync(all, x, 2);
x += __shfl_xor_sync(all, x, 4);
x += __shfl_xor_sync(all, x, 8);
x += __shfl_xor_sync(all, x, 16);
return x;
}Я убедился, что после компиляции каждый из них превращается в две инструкции. Таким образом, имеем 5 инструкций SHFL.BFLY плюс 5 инструкций FADD для сложения. Здесь нет никаких секретных инструкций или барьеров.
Это работает только в пределах варпа (группа из 32 потоков). При работе с блоком потоков, то есть с группой размером вплоть до 1024 потоков, можно пользоваться разделяемой памятью, а в таком случае не обойтись без барьеров, так как автоматически потоки синхронизироваться не будут. Если для работы вам требуется больше потоков, и при этом вы хотите разделять между ними данные, не делайте так. (Часто хочется задействовать гораздо больше потоков, но просто так разделять данные нельзя. Нужно сначала записать результат в память, а затем запустить новый поток, чтобы отдать ему на обработку новые данные).
Первым делом — параллелизм
Оказывается, пробуя интуитивно подобрать нужное количество потоков, я сильно ошибался. Если вы перебираете некоторый набор данных и должны выполнять над ним какие-то нетривиальные операции, то, пожалуй, лучше выделить по одному потоку на каждую операцию, которую вы хотите делать. Так и хочется сказать: «этот поток уже загрузил все интересующие нас данные, осталось выполнить немножко дополнительной работы». Но, ещё раз, в CUDA лучше запускать для выполнения этой дополнительной работы отдельный поток, даже если оба потока будут загружать одни и те же данные. Просто в CUDA гораздо дешевле будет синхронизировать потоки и позволять им совместно использовать свои данные на ЦП.
Когда я прогнал первые пару версий моего кода через Nsight Compute, всю поступавшую обратную связь можно было всегда резюмировать как «графический процессор почти не используется, распараллеливайте задачи».
Это также означает, что вам может понадобиться разделить ваш алгоритм. Если какую-то часть работы можно выполнять в чрезвычайно параллельном режиме (то есть разделить на десятки тысяч потоков), а другая распараллеливается лишь в ограниченном объёме (например, на несколько сотен потоков), то, пожалуй, стоило бы запустить их на отдельных ядрах. Так часть вашей задачи действительно выиграет от сильного распараллеливания, даже если это очень небольшая часть.
Итак, решая какую-либо задачу, первым делом нужно задаться не вопросом «как эту задачу ускорить?», а «как эту задачу распараллелить»? Когда справитесь с этой задачей, ускорить распараллеленный код не составит труда.
Заключение
Конечно, писать на CUDA — это особое ощущение. Такое программирование кажется более замысловатым, поскольку оказывается, что при обычных подходах код использует всего 1% графического процессора. Ситуация напоминает мне работу с TIS-100, особенно фокус с распределением данных по регистрам множества потоков. Но вам придётся не управлять небольшим количеством чипов, а придумать, как подобрать работу для десятков тысяч потоков. Мне представляется, что мы как будто делаем целый флот контейнеровозов, способных двигаться с 10% от скорости света. На них мы доставляем грузы по всему миру. Они настолько быстро ходят, что большая часть работы осуществляется в порту. Поэтому, если бы вы могли загружать и разгружать эти корабли за доли секунды, то он мог бы почти не простаивать между рейсами. Укротить такие системы непросто, но, если у вас получится, то вы сможете выполнять огромные куски работы практически мгновенно.