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

Чтобы получить индексы цикла из идентификаторов потоков и идентификаторов блоков, вы должны сначала хорошо понять, что такое блок потока, сетка потока, идентификатор потока, идентификатор блока и размер блока и т. Д.

Что такое блокировка потока?

Один блок потока состоит из набора потоков. Эти потоки могут быть 1D, 2D или 3D. Когда мы рассматриваем блок потока, стандартные переменные threadIdx и blockDim в CUDA могут считаться очень важными.

threadIdx = используется для доступа к индексу потока внутри блока потока

threadIdx.x = Индекс потока внутри блока в направлении X

threadIdx.y = Индекс потока внутри блока в направлении Y

threadIdx.z = Индекс потока внутри блока в направлении Z

blockDim = Количество потоков в блоке для определенного направления

blockDim.x = Количество потоков в блоке для направления X

blockDim.y = Количество потоков в блоке для направления Y

blockDim.z = Количество потоков в блоке для направления Z

Давайте посмотрим, как они используются в контексте.

1D

Для потока 0 threadIdx.x = threadIdx.y = threadIdx.z = 0. Но для потока 3 threadIdx.x = 4 и threadIdx.y = threadIdx.z = 0. И blockDim.x = 5 и blockDim.y = blockDim.z = 1

2D

Для потока 1 threadIdx.x = threadIdx.y = threadIdx.z = 0. Для потока 6 threadIdx.x = 2, threadIdx.y = 1 и threadIdx.z = 0. А также blockDim.x = 3 и blockDim.y = 3.

3D

Здесь блок ниток - это кубоид нитей. Надеюсь, вы сможете представить себе ситуацию. Это не что иное, как потоки во всех направлениях x, y и z.

Итак, что такое сетка резьбы?

Как и раньше, сетка резьбы представляет собой набор блоков резьбы. Блоки также могут быть в 1D, 2D или 3D (представьте себе замену резьбы блоками резьбы в предыдущем пояснении для блоков резьбы). Когда дело доходит до сетки резьбы, важны следующие переменные.

blockIdx = Используется для доступа к индексу блока потока внутри сетки потока

blockIdx.x = Индекс блока протектора в направлении X

blockIdx.y = Индекс блока протектора в направлении Y

blockIdx.z = Индекс блока протектора в направлении Z

gridDim = количество блоков резьбы в определенном направлении.

gridDim.x = Количество блоков резьбы в направлении X

gridDim.y = Количество блоков резьбы в направлении Y

gridDim.z = Количество блоков резьбы в направлении Z

Когда ядро ​​запущено, мы можем установить размеры для блоков потоков и сеток потоков. Мы увидим это позже.

Один цикл for

Начнем с одного цикла for. Наша цель - получить значение i внутри ядра CUDA. Вот код для этого.

Вот результат работы программы.

Не гарантируется, что порядок i будет таким же, как для CPU. Размер сетки установлен на (1, 1, 1). Это означает, что создается только один блок потока. Какого размера блок резьбы? Это было указано как (4, 1, 1). Это означает, что в направлении x создается 4 потока. Следовательно,

blockDim.x = 4, blockDim.y = blockDim.z = 1.

blockIdx.x = blockIdx.y = blockIdx.z = 0.

threadIdx.y = threadIdx.z = 0, threadIdx.x варьируется от 0 до 3 (включительно).

Как я принял значение i? Я добавил индекс потока для направления X к умножению индекса блока в направлении X и размерности блока для направления X.

Два вложенных цикла for

Теперь мы можем перейти к вложенным циклам for. Рассмотрим две петли. Внутри самого внутреннего цикла у нас есть комбинации i и j.

Здесь я рассмотрел возможность использования сетки (1, 1, 1) только с одним двумерным блоком (4, 4, 1). Это означает, что в этой конфигурации threadIdx.x и threadIdx.y изменяются от 0 до 3. Когда ядро ​​запускается, будут вызваны 16 потоков (4 * 4) в двухмерном блоке потоков.

Мы можем добиться того же, увеличив количество блоков в 2D-пространстве. Следующее показывает это.

Здесь я создал два блока (это означает, что размер сетки установлен на (2, 2, 1)). Следовательно, в этом подходе

threadIdx.x и threadIdx.y меняются от 0 до 1.

blockIdx.x и blockIdx.y также изменяются от 0 до 1.

Например, как взяты i = 3 и j = 2:

Для i

threadIdx.x = 1, blockIdx.x = 1 и blockDim.x = 2. Тогда i = 3 {1 + 1 * 2}.

Для j

threadIdx.y = 0, blockIdx.y = 1 и blockDim.y = 2. Тогда j = 2 {0 + 1 * 2}.

Три вложенных цикла for

Давайте посмотрим на представление трех вложенных циклов for в трехмерном пространстве.

Как и раньше, мы можем использовать два блока вместо одного.

Здесь я использовал 8 блоков потоков размером (1, 1, 1). Но значения i, j и k берутся такими же, как и раньше. Это очень важно, потому что мы не можем выделять потоки для блока больше определенного предела (1024 потока на блок). Если мы хотим, чтобы количество потоков превышало лимит, мы можем увеличить количество блоков, как указано выше.

Во всех вышеупомянутых подходах, когда есть только один цикл for, мы распределяли потоки только в одном направлении (X). Когда есть два вложенных цикла for, мы выделяем потоки в измерениях X и Y. Когда есть три вложенных цикла for, мы выделяем потоки во всех трех измерениях (X, Y и Z). Но когда дело доходит до 2D или 3D, мы можем рассмотреть возможность выделения потоков только в одном измерении в GPU. Давай посмотрим на это.

Рассмотрим следующие петли.

Как видите, оба дают одинаковые значения i и j. Это означает, что мы можем представлять вложенные циклы, используя только одно измерение. Потому что один цикл for может быть представлен одним измерением в CUDA.

Ниже показано то же самое с тремя вложенными циклами for.

Давайте посмотрим на пример того, как это сделать с ядром CUDA.

В приведенном выше подходе создается только один блок со всем необходимым количеством потоков в направлении x. Но есть ограничение (1024) на количество потоков в блоке. Следовательно, вам может потребоваться создать набор блоков, когда количество требуемых потоков превышает лимит. В следующем коде мы можем установить количество потоков на блок. Я установил 1024.

Это работает для любого значения limitX, limitY и limitZ. Поскольку при этом подходе можно создать больше потоков, чем требуется, мы должны проверить, выйдет ли threadId за пределы нашего предела, потому что в противном случае возможно, что мы можем получить значения для i, j и k больше, чем требуется.

Думаю, я рассмотрел различные подходы к распределению потоков и получению индексов цикла с использованием индексов потоков и индексов блоков. Просто попробуйте сами и лучше поймите, как это делается. Я надеюсь, что мои примеры кода будут в этом очень полезны.

Удачного кодирования 😃