Из Википедии, бесплатной энциклопедии
  (Перенаправлено из Warp (CUDA) )
Перейти к навигации Перейти к поиску

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

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

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

Размеры [ править ]

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

Поскольку многие параллельные приложения включают многомерные данные, блоки потоков удобно организовывать в одномерные, двухмерные или трехмерные массивы потоков. Блоки в сетке должны иметь возможность выполняться независимо, поскольку связь или взаимодействие между блоками в сетке невозможны. 'Когда ядро ​​запускается, количество потоков на блок потока и количество блоков потока указывается, это, в свою очередь, определяет общее количество запущенных потоков CUDA. [4] 'Максимальные размеры блока по x, y и z составляют 1024, 1024 и 64, и он должен быть распределен таким образом, чтобы x × y × z ≤ 1024, что является максимальным количеством потоков на блок. [4] Блоки могут быть организованы в одно- или двухмерные сетки до 65 535 блоков в каждом измерении. [4]Ограничение на количество потоков в блоке фактически наложено, потому что количество регистров, которые могут быть выделены для всех потоков, ограничено. [4]

Индексирование [ править ]

1D-индексация [ править ]

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

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

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

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

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. Соответствующий код в потоке показан ниже: [6]

__global__ void  vecAddKernel  ( float  * A  ,  float  * B  ,  float  *  C  ,  int  n ) {  int  index  =  blockIdx . x  *  blockDim . x  +  threadIdx . х ;  if  ( index  <  n )  {  C [ index ]  =  A [ index ]  +  B [ index]  ;  } }

2D-индексация [ править ]

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

[7]

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

Хотя мы заявили иерархию потоков, мы должны отметить, что потоки, блоки потоков и сетка, по сути, представляют собой точку зрения программиста. Чтобы получить полную суть блока потока, важно знать его с точки зрения оборудования. Оборудование группирует потоки, которые выполняют одну и ту же инструкцию, в деформации. Несколько основ составляют блок резьбы. Несколько потоковых блоков назначены потоковому мультипроцессору (SM). Несколько SM составляют весь блок GPU (который выполняет всю Kernel Grid).

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

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

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

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

Для достижения этой цели SM содержит следующее: [9]

  • Ядра исполнения. (блоки с плавающей запятой одинарной точности, блоки с плавающей запятой двойной точности, блоки специальных функций (SFU)).
  • Кеши:
  1. Кэш L1 . (для уменьшения задержки доступа к памяти).
  2. Общая память . (для общих данных между потоками).
  3. Постоянный кеш (для трансляции чтений из постоянной памяти).
  4. Кеш текстур . (для агрегирования пропускной способности из текстурной памяти).
  • Планировщики перекосов. (они предназначены для выдачи инструкций по деформации на основе определенных политик планирования).
  • Значительное количество регистров. (SM может одновременно запускать большое количество активных потоков, поэтому необходимо иметь тысячи регистров.)

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

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

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

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

Деформации [ править ]

На аппаратной стороне блок резьбы состоит из «основ». Деформация - это набор из 32 потоков в блоке потока, так что все потоки в деформации выполняют одну и ту же инструкцию. Эти потоки последовательно выбираются SM. [13]

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

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

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

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

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

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

См. Также [ править ]

  • Параллельные вычисления
  • CUDA
  • Поток (вычисления)
  • Блок обработки графики

Ссылки [ править ]

  1. ^ «Обзор CUDA» . cuda.ce.rit.edu . Проверено 21 сентября 2016 .
  2. ^ https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
  3. ^ {url = https://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf%7Ctitle=Глава 4. Аппаратная реализация. Потоки блока потока выполняются одновременно на одном многопроцессоре, и несколько блоков потоков могут выполняться одновременно на один мультипроцессор. }
  4. ^ a b c d "Модель потока CUDA" . www.olcf.ornl.gov . Проверено 21 сентября 2016 .
  5. ^ «Иерархия потоков в программировании CUDA» . Проверено 21 сентября 2016 .
  6. ^ Кирк, Дэвид; Хву, Вэнь-мэй (28 января 2010 г.). Программирование параллельных процессоров: практический подход .
  7. ^ «Шпаргалка по индексированию потоков» (PDF) . Проверено 21 сентября 2016 .
  8. ^ «Оптимизация потоков (Университет Мейленда)» (PDF) .
  9. ^ a b Уилт, Николас (2013). Справочник CUDA: Всеобъемлющее руководство по программированию на GPU .
  10. ^ «Оптимизация потоков (Университет Мейленда)» (PDF) .
  11. ^ «Оптимизация потоков (Университет Мейленда)» (PDF) .
  12. ^ "Вычисления на GPU с помощью CUDA Лекция 2 - Воспоминания CUDA" (PDF) .
  13. ^ «Использование примитивов уровня деформации CUDA» . Nvidia . 2018-01-15 . Проверено 8 апреля 2020 . Графические процессоры NVIDIA выполняют группы потоков, известные как деформации в режиме SIMT (одиночная инструкция, многопоточность).
  14. ^ «Проблемы с памятью в CUDA и планирование выполнения в CUDA» (PDF) .
  15. ^ a b «Влияние выборки инструкций и планирования памяти на производительность графического процессора» (PDF) .
  16. ^ «CAWS: критично-ориентированное планирование деформации для рабочих нагрузок GPGPU» (PDF) .