Блок потоков (программирование CUDA)

Абстракция программирования

Блок потоков — это абстракция программирования, представляющая группу потоков , которые могут выполняться последовательно или параллельно. Для лучшего сопоставления процессов и данных потоки группируются в блоки потоков. Количество потоков в блоке потоков раньше было ограничено архитектурой до 512 потоков на блок, но по состоянию на март 2010 года с вычислительными возможностями 2.x и выше блоки могут содержать до 1024 потоков. Потоки в одном блоке потоков выполняются на одном и том же потоковом процессоре. ^[1] Потоки в одном блоке могут взаимодействовать друг с другом через общую память , барьерную синхронизацию или другие примитивы синхронизации, такие как атомарные операции.

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

CUDA — это платформа параллельных вычислений и модель программирования , которую языки более высокого уровня могут использовать для использования параллелизма. В CUDA ядро ​​выполняется с помощью потоков. Поток — это абстрактная сущность , которая представляет выполнение ядра. Ядро — это функция, которая компилируется для запуска на специальном устройстве. Многопоточные приложения используют множество таких потоков, работающих одновременно, для организации параллельных вычислений. Каждый поток имеет индекс, который используется для расчета адресов памяти, а также для принятия управляющих решений.

Размеры

CUDA работает на основе модели гетерогенного программирования , которая используется для запуска прикладных программ хост-устройства. Он имеет модель выполнения, аналогичную OpenCL . В этой модели мы начинаем выполнение приложения на главном устройстве, которым обычно является ядро ​​ЦП . Устройство представляет собой устройство, ориентированное на пропускную способность, т. е. ядро ​​графического процессора , выполняющее параллельные вычисления. Функции ядра используются для выполнения этих параллельных выполнения. Как только эти функции ядра выполняются, управление передается обратно хост-устройству, которое возобновляет последовательное выполнение.

Поскольку многие параллельные приложения используют многомерные данные, блоки потоков удобно организовывать в одномерные, двумерные или трехмерные массивы потоков. Блоки в сетке должны иметь возможность выполняться независимо, поскольку связь или сотрудничество между блоками в сетке невозможны. 'При запуске ядра указывается количество потоков на блок потоков, а также количество блоков потоков, что, в свою очередь, определяет общее количество запущенных потоков CUDA. ^[2] ' Максимальные размеры блока по x, y и z составляют 1024, 1024 и 64, и он должен быть распределен так, чтобы x × y × z ≤ 1024, что является максимальным количеством потоков на блок. ^[3] Блоки могут быть организованы в одно-, двух- или трехмерные сетки, содержащие до 2 ³¹ -1, 65 535 и 65 535 блоков в измерениях x, y и z соответственно. ^[3] В отличие от максимального количества потоков на блок, не существует ограничения количества блоков на сетку, отличного от максимальных размеров сетки.

Индексирование

1D-индексация

Каждый поток в CUDA связан с определенным индексом, поэтому он может вычислять и получать доступ к ячейкам памяти в массиве.

Рассмотрим пример, в котором имеется массив из 512 элементов. Одна из организационных структур представляет собой сетку с одним блоком, имеющим 512 потоков. Предположим, что существует массив C из 512 элементов, который состоит из поэлементного умножения двух массивов A и B, каждый из которых состоит из 512 элементов. Каждый поток имеет индекс i, и он выполняет умножение i ^-го элемента A и B, а затем сохраняет результат в i ^-м элементе C. i вычисляется с использованием blockIdx (который в данном случае равен 0, поскольку существует только один поток). блок), blockDim (в данном случае 512, поскольку блок содержит 512 элементов) и threadIdx, который варьируется от 0 до 511 для каждого блока.

Иерархия потоков в программировании CUDA ^[4]

Индекс резьбы i рассчитывается по следующей формуле:

${\displaystyle i=blockIdx.x*blockDim.x+threadIdx.x}$

blockIdx.x — идентификатор блока измерения x.

blockDim.x — это размер x размера блока.

threadIdx.x — это размер x идентификатора потока.

Таким образом, «i» будет иметь значения от 0 до 511, что охватывает весь массив.

Если мы хотим рассмотреть вычисления для массива размером более 1024, мы можем иметь несколько блоков по 1024 потока каждый. Рассмотрим пример с 2048 элементами массива. В данном случае у нас есть 2 блока потоков по 1024 потока каждый. Таким образом, значения идентификаторов потоков будут варьироваться от 0 до 1023, идентификатор блока будет варьироваться от 0 до 1, а размерность блока будет равна 1024. Таким образом, первый блок получит значения индекса от 0 до 1023, а последний будет иметь значения индекса. с 1024 по 2047 год.

Таким образом, каждый поток сначала вычисляет индекс памяти, к которой он имеет доступ, а затем продолжает расчет. Рассмотрим пример, в котором элементы из массивов A и B добавляются параллельно с использованием потоков, а результаты сохраняются в массиве C. Соответствующий код в потоке показан ниже: ^[5]

__global__ void vecAddKernel ( float * A , float * B , float * C , int n ) { int index = blockIdx . х * блокДим . х + идентификатор потока . Икс ; если ( индекс < n ) { C [ индекс ] = A [ индекс ] + B [ индекс ] ; } }                                 

2D-индексация

Таким же образом в особенно сложных сетках идентификатор блока, а также идентификатор потока необходимо рассчитывать для каждого потока в зависимости от геометрии сетки. Рассмотрим двумерную сетку с двумерными блоками. ThreadId и BlockId будут рассчитываться по следующим формулам:

${\displaystyle blockId=blockIdx.x+blockIdx.y*gridDim.x;}$ ${\displaystyle threadId=blockId*(blockDim.x*blockDim.y)+(threadIdx.y*blockDim.x)+threadIdx.x;}$^[6]

Аппаратная перспектива

Хотя мы установили иерархию потоков, следует отметить, что потоки, блоки потоков и сетка — это, по сути, точка зрения программиста. Чтобы получить полное представление о блоке потоков, важно знать его с аппаратной точки зрения. Аппаратное обеспечение группирует потоки, выполняющие одну и ту же инструкцию, в варпы. Несколько основ составляют блок нити. Несколько блоков потоков назначены потоковому мультипроцессору (SM). Несколько SM составляют целый блок графического процессора (который выполняет всю сетку ядра). ^{[ нужна цитата ]}

Наглядная корреляция точки зрения программиста и аппаратной точки зрения блока потоков в графическом процессоре ^[7]

Потоковые мультипроцессоры

Каждая архитектура графического процессора (скажем, Kepler или Fermi ) состоит из нескольких SM или потоковых мультипроцессоров. Это процессоры общего назначения с низкой тактовой частотой и небольшим кэшем. SM может выполнять несколько блоков потоков параллельно. Как только один из его блоков потоков завершает выполнение, он занимает следующий по порядку блок потока. В общем, SM поддерживают параллелизм на уровне инструкций, но не прогнозирование ветвей . ^[8]

Иллюстрация потокового мультипроцессора и его ресурсов ^[9]

Для достижения этой цели СМ содержит следующее: ^[8]

• Исполняющие ядра. (блоки с плавающей запятой одинарной точности, блоки с плавающей запятой двойной точности, специальные функциональные блоки (SFU)).
• Кэши:

1. Кэш L1 . (для уменьшения задержки доступа к памяти).
2. Общая память . (для общих данных между потоками).
3. Кэш констант (для трансляции операций чтения из постоянной памяти).
4. Кэш текстур . (для агрегирования пропускной способности из текстурной памяти).

• Планировщики для варпов. (они предназначены для выдачи инструкций варпам на основе определенной политики планирования).
• Значительное количество регистров. (SM может одновременно запускать большое количество активных потоков, поэтому необходимо иметь регистры в тысячах.)

Аппаратное обеспечение планирует блокировку потоков для SM. Обычно SM может обрабатывать несколько блоков потоков одновременно. Всего SM может содержать до 8 блоков потоков. Идентификатор потока назначается потоку соответствующим SM.

Всякий раз, когда SM выполняет блок потоков, все потоки внутри блока потоков выполняются одновременно. Следовательно, чтобы освободить память блока потоков внутри SM, крайне важно, чтобы весь набор потоков в блоке завершил выполнение. Каждый блок потока разделен на запланированные блоки, известные как варп. Они подробно обсуждаются в следующем разделе.

Иллюстрация планировщика двойной деформации, реализованного в микроархитектуре Fermi от Nvidia ^[10]

Планировщик варпов SM решает, какой из варпов будет приоритетным во время выдачи инструкций. ^[11] Некоторые политики определения приоритетов варпа также обсуждались в следующих разделах.

Искажения

С аппаратной стороны блок потоков состоит из «перекосов». (Этот термин происходит от плетения . ^[12] ). Деформация — это набор из 32 потоков внутри блока потоков, при котором все потоки в деформации выполняют одну и ту же инструкцию. Эти нити последовательно выбираются SM. ^[13]

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

Рассмотрим деформацию из 32 потоков, выполняющих инструкцию. Если один или оба его операнда не готовы (например, еще не были выбраны из глобальной памяти), происходит процесс , называемый « переключением контекста », который передает управление другому варпу. ^[14] При переключении с определенного варпа все данные этого варпа остаются в файле регистров, так что его можно быстро возобновить, когда его операнды станут готовы. Когда инструкция не имеет невыполненных зависимостей данных, то есть оба ее операнда готовы, соответствующая деформация считается готовой к выполнению. Если для выполнения допускается более одного варпа, родительский SM использует политику планирования варпа , чтобы решить, какой варп получит следующую полученную инструкцию.

Ниже обсуждаются различные политики планирования варпов, которые могут быть выполнены: ^[15]

1. Круговой перебор (RR) — инструкции выбираются по круговому принципу. RR следит за тем, чтобы SM были заняты и не тратились тактовые циклы на задержки памяти.
2. Наименее недавно выбранное (LRF). В этой политике варп, для которого инструкция не извлекалась в течение самого длительного времени, получает приоритет при выборке инструкции.
3. Честный (FAIR) ^[15] — в этой политике планировщик гарантирует, что всем варпам будет предоставлена ​​«справедливая» возможность по количеству выбранных для них инструкций. Он извлекает инструкции в варп, для которого было получено минимальное количество инструкций.
4. CAWS на основе блоков потоков ^[16] (планирование деформации с учетом критичности). Акцент этой политики планирования делается на улучшении времени выполнения блоков потоков. Он выделил больше временных ресурсов для варпа, выполнение которого займет больше всего времени. Отдавая приоритет наиболее критической деформации, эта политика позволяет блокам потоков завершаться быстрее, поэтому ресурсы становятся доступными быстрее.

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

Чтобы воспользоваться преимуществами варп-архитектуры, языки программирования и разработчики должны понимать, как объединить доступ к памяти и как управлять расхождением потоков управления. Если каждый поток в варпе выбирает другой путь выполнения или если каждый поток обращается к значительно отличающейся памяти, то преимущества архитектуры варпа теряются и производительность значительно снижается.

Рекомендации

1. «Глава 4. Аппаратная реализация. Потоки блока потоков выполняются одновременно на одном мультипроцессоре, и несколько блоков потоков могут выполняться одновременно на одном мультипроцессоре».
2. «Модель потока CUDA» . www.olcf.ornl.gov . Архивировано из оригинала 23 сентября 2016 г. Проверено 21 сентября 2016 г.
3. «Документация по набору инструментов CUDA: функции и технические характеристики» . docs.nvidia.com . Проверено 24 мая 2022 г.
4. «Иерархия потоков в программировании CUDA» . Проверено 21 сентября 2016 г.
5. Кирк, Дэвид; Ху, Вэнь-мэй В (28 января 2010 г.). Программирование массово-параллельных процессоров: практический подход .
6. «Шпаргалка по индексированию потоков» (PDF) . Проверено 21 сентября 2016 г.
7. «Оптимизация потоков (Университет Мэйленда)» (PDF) .
8. Уилт, Николас (2013). Справочник по CUDA: Комплексное руководство по программированию на графическом процессоре .
9. «Оптимизация потоков (Университет Мэйленда)» (PDF) .
10. «Оптимизация потоков (Университет Мэйленда)» (PDF) .
11. «Вычисления на графическом процессоре с использованием CUDA, лекция 2 — Воспоминания CUDA» (PDF) .
12. «Выполнение параллельного потока ISA версии 6.0» . Зона разработчика: Документация по набору инструментов CUDA . Корпорация NVIDIA. 22 сентября 2017 года. Архивировано из оригинала 28 октября 2017 года . Проверено 27 октября 2017 г.
13. «Использование примитивов уровня деформации CUDA» . Нвидия . 15 января 2018 г. Проверено 8 апреля 2020 г. Графические процессоры NVIDIA выполняют группы потоков, известные как деформации, в режиме SIMT (одна инструкция, несколько потоков).
14. «Проблемы с памятью в CUDA и планирование выполнения в CUDA» (PDF) .
15. «Влияние выборки инструкций и планирования памяти на производительность графического процессора» (PDF) .
16. «CAWS: планирование деформации с учетом критичности для рабочих нагрузок GPGPU» (PDF) .