Пример использования иерархии потоков в CUDA
Я изучаю CUDA и читаю, что смысл иерархии потоков заключается в обеспечении масштабируемости и совместимости графических процессоров с разными возможностями. (Я не знаю, является ли это единственным преимуществом).
В документе NVIDIA (Раздел 1.3) есть небольшая картинка, показывающая, как 8-блочную «многопоточную» программу CUDA можно запускать как на графическом процессоре с 2 SM, так и на графическом процессоре с 4 SM. Однако я не понимаю, почему этого нельзя достичь, используя только потоки, а не абстракции сетки и блоков.
Пожалуйста, может ли кто-нибудь привести пример использования иерархии потоков, где наличия только потоков было бы недостаточно?
Это станет ясно, если продолжить чтение после введения. 2.2. Иерархия потоков: «Потоки внутри блока могут взаимодействовать, обмениваясь данными через некоторую общую память и синхронизируя их выполнение для координации доступа к памяти».
@paleonix Спасибо. Я действительно прочитал первые 4 раздела, но не был уверен в том, как использовалась общая память/другие функции, и были ли они связаны с иерархией потоков или просто независимой функцией аппаратного обеспечения. Поэтому я попросил пример. Я изучаю примеры, которыми вы поделились, и думаю, что они очень помогают, так что спасибо. Кроме того, мой вопрос подразумевает подвопрос; невозможно ли просто иметь общую память для всех потоков, а не для блока потоков? Я считаю, что ответ на этот вопрос связан с управлением памятью в варпах...
@paleonix Однако я не верю, что ваш комментарий касается того, как иерархия упрощает перекрестную совместимость?
Когда кто-то, знакомый с CUDA и оборудованием, на котором он должен работать, думает о худшей альтернативе этой масштабируемой парадигме, он, вероятно, думает о прямом программировании оборудования на уровне деформации и SM, а не на уровне потоков и блоков. Наличие только потоков не имеет смысла в этом контексте (даже во введении упоминается взаимодействие между потоками блока). В такой гипотетической модели программирования для написания базовых ядер масштабируемым способом потребуется гораздо больше шаблонов (например, запрос количества SM, чтобы выяснить, какой объем работы должен выполнить каждый SM).
Сетка может (и обычно имеет) больше потоков, чем может одновременно работать на графическом процессоре. Графический процессор может гарантировать только то, что все потоки в одном блоке потоков выполняются одновременно (и на одном потоковом мультипроцессоре; SM). Поэтому они могут легко общаться. Между тем, графические процессоры с разными уровнями производительности (разное количество SM) могут запускать одну и ту же сетку в разное время, выполняя меньше или больше блоков потоков параллельно или друг за другом. Вам, как программисту, не нужно беспокоиться о точном количестве, которое выполняется одновременно.
«Разве невозможно просто иметь общую память для всех потоков, а не для блока потоков?» Глобальная память — это память, совместно используемая всеми блоками потоков. Но доступ (к некэшируемым областям) стоит дорого, поскольку DRAM физически находится вне кристалла. Общая память не только встроена в кристалл (например, кэш L2), но и локальна для SM. Доступ к общей памяти из другого блока на другом SM обычно (при условии, что это возможно) намного дороже. Блок-кластеры в Hopper идут немного в этом направлении, но в целом нет, «глобальная общая память» не имеет смысла, учитывая аппаратную архитектуру.
TL;DR: Кооперация локально для блока (и, следовательно, SM) дешева и поэтому желательна для всего, кроме самых простых ядер. Вам необходимо определить, что является локальным, а что нет, чтобы иметь возможность использовать местоположение. Поэтому иерархия неизбежна. Масштабируемой модель программирования делает не наличие иерархии, а тот факт, что эта иерархия представляет собой абстракцию фактической иерархии аппаратного обеспечения, что позволяет писать и эффективно запускать ядра практически без знания точного оборудования, присутствующего в системе.
Нужно время, чтобы переварить, но все имеет смысл. Спасибо палеониксу и @Homer512.
Ответы 1
Иерархия потоков, как и иерархия памяти, полностью зависит от локальности ссылок.
Варп
Нити сгруппированы в пучки по 32: основа. Потоки в деформации могут обмениваться данными с помощью shfl_[up_,down_,xor_,_]sync , Warp Reduction , Warp Match и (очень полезных) Warp voice инструкций.
Потоки в варпе имеют дешевую синхронизацию: __syncwarp
Блок тем
Следующим шагом является блок потоков, до 32 варпов (1024 потоков), которые взаимодействуют в одном мультипроцессоре (SM). Специальных инструкций для внутриблочного взаимодействия не существует, но потоки в блоке могут использовать __shared__ память для обмена данными. Это почти так же эффективно, как __shfl_sync.
Сетчатый блок
Далее у нас есть мультипроцессоры (SM), в одном графическом процессоре их может быть от 20 до 100. Они (до Блэквелла) могут общаться только через глобальную память. Поток в SM a хранит данные в глобальной памяти, а поток в SM b считывает эти данные.
Два блока сетки обычно работают на разных SM (и если они находятся на одном SM, они не могут совместно использовать __shared__ память).
(В Блэквелле поток SM a может перемещать данные из своей общей памяти в общую память другого SM).
Различные графические процессоры
Следующий уровень иерархии — запуск вашей программы на нескольких графических процессорах на одной машине.
Они могут взаимодействовать только через связь с графическим процессором (имя которого ускользает от меня).
Очевидно, что это еще медленнее и сложнее.
Разные машины последний уровень работает на разных машинах, которые обмениваются данными через сеть (локальную или через Интернет).
Когда что использовать
Очевидно, что один варп может выполнять только 32 одновременные задачи.
Блокировать только 1024 одновременных задачи и так далее.
Если у вас ограничено время и вам нужно использовать все потоки 20x32x32 = 20480 на графическом процессоре с 20 SM, тогда вам нужен блок сетки, и вам придется соответствующим образом спроектировать свою программу.
Каждый шаг вверх по пирамиде потоков выполняется быстрее, чем тот, что ниже нее, поскольку расстояние между потоками на кристалле меньше, что позволяет производителю создавать быстрые межсоединения. Иерархия также упрощает модель программирования, так что разные версии графических процессоров можно программировать одинаково.
Но почему?
Такая конструкция позволяет графическому процессору обеспечивать быстрые возможности на уровне деформации/блока. Если бы у нас были только блоки сетки, то каждый поток был бы вынужден взаимодействовать только через глобальную память.
Следующий пример получения максимального значения (из 32) из битонической последовательности выглядит следующим образом:
/\
/ \ e.g.: 1, 2, 10, 20, 30, 500, 400, 3, 0.
/
__device__ int GetMaxValue(const int value) { //every thread has different value
//the input must be a bitonic sequence
constexpr auto All = -1u; //all threads in the warp take part
const auto Neighbor = __shfl_down_sync(All, value, 1); //thread 31 gets its own value back
const auto mask = __ballot_sync(All, Neighbor > value); //e.g. 0b0000000011111111
const auto MaxBit = __popc(mask); //The thread holding the max value
const auto result = __shfl_sync(All, value, MaxBit); //share max value with all
return result; //all threads return the same result
}
будет намного быстрее, чем эквивалентный код, использующий только глобальную память, если мы сделаем вид, что искажений и блоков не существует:
__device__ int GetMaxValue_NoWarp(int value, int* tempstorage) {
tempstorage[32] = 0;
tempstorage[threadIdx.x] = value;
__syncthreads();
const auto Neighbor = tempstorage[threadIdx.x + 1];
if (threadIdx.x == 31) { Neighbor = value; }
const auto diff = int(Neighbor > value);
atomicOr(&tempstorage[32], diff << threadIdx.x);
__syncthreads();
const auto mask = tempstorage[32];
const auto MaxBit = __popc(mask);
result = tempstorage[MaxBit];
return result;
}
Даже одно чтение из tempstorage занимает больше времени, чем все GetMaxValue, а GetMaxValue_NoWarp требует множества операций чтения и записи, поскольку у него нет быстрого механизма обмена данными между потоками.
Спасибо :). Пример был полезен, но, поскольку я новичок в CUDA, мне было довольно трудно следовать коду. Читателям, находящимся в аналогичном положении, примеры, приведенные в комментариях, оказались в этом плане более простыми.
Однако меня смущает то, как вы объяснили иерархию. Я видел ресурсы, относящиеся к «точке зрения программиста» и «точке зрения аппаратного обеспечения», где они заявляют, что деформации — это аппаратная конструкция, тогда как потоки, блоки и сетки — это программные конструкции. В своем ответе вы используете оба вместе в одной иерархии. Это говорит о том, что все они являются программными конструкциями?
@ThePhysicsOverthinker, Деформации — это программная конструкция. На самом деле 32 потока в варпе поддерживаются только 16 фактическими исполнительными ядрами, которые по очереди выполняют данные. Хотя я не вижу разницы. Блок отображается в мультипроцессоре, так это аппаратная или программная конструкция? Технически вы можете игнорировать деформации и просто использовать threadIdx для разделения работы, но если вы это сделаете, вы потеряете большую производительность. Основная причина, по которой я включаю деформации, заключается в том, что если вы заставите потоки в одной и той же деформации выполнять отличающийся код, вы получите ужасную производительность (до 32 раз медленнее максимальной).
Любое ядро, использующее синхронизацию общей памяти/блоков? Например. сокращения, сканирования... В основном все, что не является тривиально параллельным и, следовательно, извлекает выгоду из локальной связи или даже ядра, которые тривиально параллельны, но не идеальны с точки зрения доступа к памяти/повторного использования данных, например оптимизированные ядра матрично-матричных продуктов.