Потоковые мультипроцессоры, блоки и потоки (CUDA)

Question

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

Что с чем сопоставляется, что распараллеливается и как?и что эффективнее, максимизировать количество блоков или количество потоков?

---

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

Правильно ли это?

Answer 1

Схема потока / блока подробно описана в Руководство по программированию на CUDA.В частности, в главе 4 говорится:

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

Каждый SM содержит 8 ядер CUDA, и в любой момент времени они выполняют один warp из 32 потоков - так что для выдачи одной инструкции для всего warp требуется 4 такта.Вы можете предположить, что потоки в любой заданной деформации выполняются с шагом блокировки, но для синхронизации между деформациями вам необходимо использовать __syncthreads().

Answer 2

Для GTX 970 предусмотрено 13 потоковых мультипроцессоров (SM) с 128 ядрами Cuda каждый.Ядра Cuda также называются потоковыми процессорами (SP).

Вы можете определить сетки, которые отображают блоки на графический процессор.

Вы можете определить блоки, которые сопоставляют потоки с потоковыми процессорами (128 ядер Cuda на SM).

Одна основа всегда формируется 32 нитями, и все нити основы выполняются одновременно.

Чтобы использовать всю возможную мощность графического процессора, вам нужно гораздо больше потоков на SM, чем у SM есть SPS.Для каждой вычислительной возможности существует определенное количество потоков, которые могут находиться в одном SM одновременно.Все определенные вами блоки помещаются в очередь и ожидают, пока у SM появятся ресурсы (количество свободных SP), после чего он загружается.SM начинает выполнять Перекосы.Поскольку одна деформация имеет только 32 нити, а SM имеет, например, 128 SPS, SM может выполнять 4 деформации в данный момент времени.Дело в том, что если потоки обращаются к памяти, поток будет блокироваться до тех пор, пока не будет удовлетворен его запрос на память.В цифрах:Арифметические вычисления в SP имеют задержку 18-22 цикла, в то время как доступ к глобальной памяти без кэширования может занимать до 300-400 циклов.Это означает, что если потоки одной основы ожидают получения данных, то будет работать только подмножество из 128 SPS.Для этого планировщик переключается на выполнение другой деформации, если таковая имеется.И если эта деформация блокируется, она выполняет следующую и так далее.Эта концепция называется скрытием задержки.Количество деформаций и размер блока определяют заполняемость (исходя из того, сколько деформаций SM может выбрать для выполнения).Если заполняемость высокая, то маловероятно, что для SPs не найдется работы.

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

Здесь вы найдете цифры для различных вычислительных возможностей и архитектур графических процессоров:https://en.wikipedia.org/wiki/CUDA#Limitations

Вы можете загрузить таблицу расчета заполняемости с сайта Nvidia Таблица расчета заполняемости (от Nvidia).

Answer 3

Распределитель вычислительных работ запланирует блок потоков (CTA) для SM только в том случае, если SM располагает достаточными ресурсами для блока потоков (общая память, перекосы, регистры, барьеры, ...).Выделяются ресурсы уровня блока потока, такие как общая память.Выделение создает достаточные деформации для всех потоков в блоке потоков.Диспетчер ресурсов распределяет деформации с помощью циклический перебор к подразделам SM.Каждый подраздел SM содержит планировщик warp, файл регистра и исполнительные модули.Как только деформация выделена подразделу, она будет оставаться в подразделе до тех пор, пока не завершится или не будет вытеснена переключателем контекста (архитектура Pascal).При восстановлении контекстного переключения деформация будет восстановлена с тем же SM-идентификатором деформации.

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

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

Как только основа выделена подразделу и все ресурсы распределены, основа считается активной, что означает, что планировщик основы активно отслеживает состояние основы.На каждом цикле планировщик деформации определяет, какие активные деформации остановлены и которые имеют право выдать команду.Планировщик warp выбирает warp с наивысшим приоритетом и выдает 1-2 последовательных инструкции из warp.Правила для двойного выпуска специфичны для каждой архитектуры.Если деформация вызывает загрузку памяти, она может продолжать выполнять независимые инструкции до тех пор, пока не достигнет зависимой инструкции.Затем warp сообщит о остановке до завершения загрузки.То же самое верно и для зависимых математических инструкций.Архитектура SM предназначена для скрытия задержек как в ALU, так и в памяти путем переключения за цикл между перекосами.

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

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

Временная заполняемость - это отношение активных деформаций к максимальным деформациям на SM.Архитектура Kepler - Pascal (кроме GP100) имеет по 4 планировщика деформации на SM.Минимальное количество деформаций на SM должно быть, по крайней мере, равно количеству планировщиков деформации.Если архитектура имеет зависимую задержку выполнения в 6 циклов (Maxwell и Pascal), то вам потребуется не менее 6 деформаций на планировщик, что составляет 24 на SM (24/64 = 37,5% занятости), чтобы покрыть задержку.Если потоки имеют параллелизм на уровне команд, то это может быть уменьшено.Почти все ядра выдают инструкции с переменной задержкой, такие как загрузка памяти, которая может занимать 80-1000 циклов.Это требует более активных деформаций для каждого планировщика деформации, чтобы скрыть задержку.Для каждого ядра существует точка компромисса между количеством деформаций и другими ресурсами, такими как общая память или регистры, поэтому оптимизация для 100%-ной занятости не рекомендуется, поскольку, вероятно, будут принесены какие-то другие жертвы.Профилировщик CUDA может помочь определить частоту проблем с инструкциями, заполняемость и причины задержек, чтобы помочь разработчику определить этот баланс.

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

Answer 4

На одном устройстве имеется несколько потоковых мультипроцессоров.
SM может содержать несколько блоков.Каждый блок может содержать несколько потоков.
SM имеет несколько ядер CUDA (как разработчик, вы не должны заботиться об этом, потому что это абстрагируется warp), которые будут работать с потоком.SM всегда работает с основой нитей (всегда 32).Деформация будет работать только с потоком из того же блока.
SM и block имеют ограничения на количество потоков, количество регистров и разделяемую память.