Сложение векторов произвольной длины на GPU

Опубликовано апреля 29, 2018 в Технология CUDA

Мы не были до конца откровенны, когда первый раз обсуждали запуск параллельных блоков на GPU.

Помимо ограничения на количество нитей, существует также аппаратное ограничение на количество блоков (хотя и гораздо более либеральное). Как мы уже упоминали, количество блоков вдоль любого измерения сетки не должно превышать 65 535.

Но тогда в последней реализации сложения векторов возникает проблема. Если запускать N/128 блоков, то мы столкнемся с ошибкой, когда число элементов в каждом векторе окажется больше 65 535 * 128 = 8 388 480. Вроде бы большое число, но при нынешних объемах памяти от 1 до 4 Гб наиболее мощные графические процессоры легко справляются с хранением данных, совокупный размер которых на несколько порядков больше, чем 8 миллионов элементов.

К счастью, есть очень простое решение. Сначала изменим ядро.

https://vsesam.org

Здесь мы тоже обходили данные в цикле while. Вспомните, мы говорили, что в версии для нескольких CPU или процессорных ядер индекс массива следовало бы увеличивать не на 1, а на количество процессоров. Тот же принцип мы сейчас применим в версии для GPU.

В реализации для GPU мы можем рассматривать параллельно работающие нити как процессоры. Хотя в реальном GPU процессорных блоков может быть меньше (или больше), мы считаем, что каждая нить логически исполняется параллельно, а планирование фактического выполнения поручаем оборудованию.

Отделение распараллеливания от истинного способа аппаратного исполнения — бремя, которое CUDA снимает с плеч программиста. И это действительно манна небесная, если учесть, что в современных процессорах NVIDIA количество АЛУ варьируется от 8 до 480!

Теперь, разобравшись с принципом, лежащим в основе реализации, мы должны только понять, как вычислить начальный индекс каждой параллельно исполняемой нити и на сколько его увеличивать в цикле. Мы хотим, чтобы каждая нить начинала работу с разных индексов, поэтому нужно линеаризовать пару (номер блока, номер нити), как показано в разделе «Сложение более длинных векторов на GPU».

Мы почти у цели! Осталось только подправить код запуска. Если помните, мы начали этот разговор с того, что вызов add<<<(N+127)/128,128>>>( dev_a, dev_b, dev_c ) завершится ошибкой, если (N+127)/128 больше 65 535. Чтобы гарантированно не запускать слишком много блоков, нужно лишь заранее выбрать разумное их количество. Ну и чтобы не вводить лишних чисел, остановимся на 128:128 блоков по 128 нитей в каждом.

Можете взять любые другие значения, лишь бы они не превышали вышеупомянутых ограничений. Позже мы обсудим влияние того или иного выбора на производительность, но пока будем считать, что 128 блоков по 128 нитей достаточно. Теперь мы можем складывать векторы произвольной длины и ограничены лишь объемом памяти, доступной GPU.