Расщепление параллельных блоков

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

В предыдущей статье мы видели, как на GPU исполняется параллельный код. Для этого мы сообщали исполняющей среде CUDA о том, сколько нужно запустить параллельных экземпляров ядра. Эти параллельные экземпляры мы назвали блоками.

Исполняющая среда CUDA позволяет расщепить блоки на нити. Напомним, что при запуске нескольких параллельных блоков мы указывали их количество в первом аргументе внутри угловых скобок. Так, в задаче о сложении векторов мы запускали по одному блоку для каждого элемента вектора размером N:

И снова о сложении векторов

Мы собираемся решить ту же задачу, что и в предыдущей главе, то есть, имея два вектора на входе, получить третий вектор, содержащий их сумму. Но на этот раз вместо блоков мы воспользуемся нитями.

лучшая игровая клавиатура

Возможно, вам не понятно, какие преимущества дает использование нитей вместо блоков. Пока никаких достойных обсуждения преимуществ не видно. Однако параллельные нити внутри блока могут делать такие вещи, которые параллельным блокам не под силу. Так что наберитесь терпения и позвольте нам переделать программу из предыдущей главы, заменив блоки нитями.
Все просто, не правда ли? В следующем разделе мы поговорим об ограничении, присущем решению, которое основано на одних лишь потоках. А затем расскажем, зачем вообще нужно расщеплять блоки на более мелкие параллельные компоненты.

Сложение более длинных векторов на GPU

В предыдущей главе мы упомянули об аппаратном ограничении на количество блоков в одной сетке — 65 535 по каждому измерению. Существует также ограничение и на количество нитей в одном блоке. Точнее, оно не должно превышать величины в поле maxThreadsPerBlock в структуре, описывающей свойства устройства. Для многих современных графических процессоров количество нитей в блоке не должно превышать 512. И как же воспользоваться подходом на основе нитей для сложения векторов длиннее 512? Для этого понадобится комбинация нитей и блоков.

Как и раньше, мы должны будем внести два изменения: способ вычисления индекса внутри ядра и запуск самого ядра. Когда имеются несколько блоков и нитей, индексация напоминает стандартный метод перехода от двумерного индекса к линейному.

Здесь мы пользуемся еще одной встроенной переменной, blockDim. Она одинакова во всех блоках и содержит количество нитей вдоль каждой размерности блока. Поскольку мы работаем с одномерным блоком, то интерес представляет только величина blockDim. х. Напомним, что имеется похожая переменная gridDim, в которой хранится количество блоков вдоль каждого измерения сетки. Правда, переменная g ridDim двумерная, a blockDim — трехмерная. Иначе говоря, исполняющая среда CUDA позволяет запускать двумерную сетку блоков, каждый из которых представляет собой трехмерный массив нитей. Да уж, целая куча измерений, и маловероятно, что все они вам понадобятся, но если что — они к вашим услугам.

Вычисление индекса в линейном массиве по приведенной выше формуле интуитивно очевидно. Если вам так не кажется, то представьте себе, что набор блоков, включающих нити, расположен в пространстве как двумерный массив пикселей.

Если нитям соответствуют столбцы, а блокам — строки, то для получения индекса конкретной нити нужно умножить номер блока на количество нитей в одном блоке и прибавить к произведению номер нити в блоке. Точно так же в задаче о фрактале Джулия мы вычисляли смещение от начала буфера, зная двумерный индекс пикселя.

А теперь предположим, что DIM — размерность блока (измеренная в нитях), у -номер блока, ах — номер нити в блоке.

Второе изменение следует внести в код самого ядра. Нам, как и раньше, нужно запустить N параллельных нитей, только теперь мы хотим распределить их по нескольким блокам, чтобы не выйти за рамки ограничения 512 нитей на блок. Одно из возможных решений — произвольно выбрать фиксированное количество нитей в блоке, например 128. Тогда для получения N нитей нужно будет запустить N/128 блоков.

Беда в том, что результатом целочисленного деления N/128 является целое число. В частности, если N равно 127, то N/128 равно нулю, а запустив 0 нитей, мы ничего не вычислим. Да и вообще, если N не кратно 128, то мы запустим слишком мало нитей. Это плохо. Хотелось бы, чтобы частное от деления округлялось с избытком.

Теперь если индекс оказывается за концом массива — а так будет всегда, когда длина массива не кратна 128, — мы не производим никаких вычислений. И, что еще важнее, не пытаемся читать или записывать данные в область памяти вне массивов.