2010-08-19 2 views
45

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

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


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

Это правильно?

ответ

49

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

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

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

+12

Только одно дополнение: на новых устройствах имеется 32 (Compute Capability 2.0) или 48 (2.1) сердечников CUDA на SM. Фактическое число на самом деле не имеет большого значения для программирования, размер основы равен 32 и имеет тот же смысл (т. Е. Выполняется в режиме блокировки). – Tom

+8

И на самом деле Compute Capability 3.0 (Kepler) теперь увеличивает сердечники/SM чрезвычайно - до 192! – Edric

+1

Я все еще не понимаю. Таким образом, это всегда 1 warp на ядро, а количество перекосов на SM равно количеству ядер на SM? И как блоки потоков блокируются? Блоки всегда состоят из целых чисел перекосов? Если, например, каждый блок содержит 3 искажения, значит ли это, что я использую 3 ядра на заданном SM? –

15

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

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

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

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

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

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

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

Вы можете скачать расчет размещения лист от Nvidia Occupancy Calculation sheet (by Nvidia).

2

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

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

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

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

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

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

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

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

Смежные вопросы