Какова связь между ядром CUDA, потоковым мультипроцессором и моделью CUDA блоков и потоков?
Что сопоставляется с тем, что и что распараллеливается и как? и что более эффективно, максимизируйте количество блоков или количество потоков?
Мое настоящее понимание заключается в том, что на один процессор приходится 8 ядер cuda. и что каждое ядро cuda сможет выполнить один блок cuda за раз. и все потоки в этом блоке выполняются последовательно в этом конкретном ядре.
Правильно ли это?
Макет потока/блока подробно описан в руководстве по программированию CUDA. В частности, в главе 4 говорится:
Архитектура CUDA построена вокруг масштабируемого массива многопоточных потоковых многопроцессоров (SMs). Когда программа CUDA на главном процессоре вызывает сетку ядра, блоки сетки перечисляются и распределяются по мультипроцессорам с доступной пропускной способностью. Нити блока потока выполняются одновременно на одном мультипроцессоре, и несколько блоков потоков могут выполняться одновременно на одном многопроцессорном компьютере. По мере того, как блоки потоков прекращаются, на освобожденных мультипроцессорах запускаются новые блоки.
Каждый SM содержит 8 ядер CUDA, и в любой момент времени они выполняют одну основную цепочку из 32 потоков - так что требуется 4 такта для выдачи одной инструкции для всего warp. Вы можете предположить, что потоки в любой заданной деформации выполняются в режиме блокировки, но для синхронизации между перекосами вам нужно использовать __syncthreads().
Для GTX 970 есть 13 потоковых мультипроцессоров (SM) с 128 ядрами Cuda каждый. Cuda Cores также называются потоковыми процессорами (SP).
Вы можете определить сетки, которые сопоставляют блоки с графическим процессором.
Вы можете определить блоки, которые отображают потоки для потоковых процессоров (128 Cuda Cores для 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 может выбрать основы для выполнения. Все потоки выполненных деформаций выполняются параллельно.
Здесь вы найдете номера для разных вычислительных возможностей и архитектур графических процессоров: https://en.wikipedia.org/wiki/CUDA#Limitations
Вы можете загрузить лист расчета занятости из Nvidia лист расчета занятости (от Nvidia).
Compute Work Distributor будет планировать блок потоков (CTA) на SM, только если SM имеет достаточно ресурсов для блока потоков (разделяемая память, деформации, регистры, барьеры,...). Ресурсы уровня блока потоков, такие как разделяемая память, выделяются. Выделение создает достаточные перекосы для всех потоков в блоке потоков. Менеджер ресурсов распределяет деформации с использованием циклического перебора для подразделов SM. Каждое подразделение SM содержит планировщик деформации, файл реестра и исполнительные блоки. Как только деформация будет выделена подразделу, она останется в этом подразделении до тех пор, пока не завершится или не будет заменена переключением контекста (архитектура Pascal). При переключении контекста при восстановлении основы будет восстановлен тот же SM с тем же идентификатором деформации.
Когда все потоки в деформации завершены, планировщик деформации ожидает завершения всех невыполненных инструкций, выданных деформацией, и затем менеджер ресурсов освобождает ресурсы уровня деформации, которые включают идентификатор деформации и файл реестра.
Когда все деформации в блоке потока завершены, ресурсы уровня блока высвобождаются, и SM уведомляет распространителя вычислительных работ о завершении блока.
После того, как деформация выделена подразделу и все ресурсы распределены, деформация считается активной, что означает, что планировщик деформации активно отслеживает состояние деформации. В каждом цикле планировщик деформации определяет, какие активные деформации останавливаются, а какие имеют право выдавать инструкцию. Планировщик деформации выбирает деформацию с наивысшим приоритетом и выдает 1-2 последовательных инструкции из деформации. Правила двойного выпуска специфичны для каждой архитектуры. Если деформация вызывает загрузку памяти, она может продолжать выполнять независимые инструкции, пока не достигнет зависимой инструкции. Деформация будет сообщать об остановке до завершения загрузки. То же самое верно для зависимых математических инструкций. Архитектура SM предназначена для сокрытия задержки ALU и памяти путем переключения между циклами между перекосами.
В этом ответе не используется термин "ядро CUDA", поскольку он вводит неверную ментальную модель. Ядра CUDA представляют собой конвейерные блоки с плавающей запятой одинарной точности/целочисленные. Частота проблем и задержка зависимости зависят от конкретной архитектуры. Каждое подразделение SM и SM имеют другие исполнительные блоки, включая блоки загрузки/хранения, блоки с плавающей запятой двойной точности, блоки с плавающей запятой половинной точности, блоки ветвления и т.д.
Чтобы максимизировать производительность, разработчик должен понимать, как соотносятся блоки, деформации и регистры/потоки.
Термин "занятость" - это отношение активных деформаций к максимальным деформациям на SM. Архитектура Kepler - Pascal (кроме GP100) имеет 4 планировщика деформации на SM. Минимальное количество деформаций на SM должно быть по меньшей мере равным количеству планировщиков деформаций. Если архитектура имеет зависимую задержку выполнения, равную 6 циклам (Максвелла и Паскаля), то вам потребуется по меньшей мере 6 перекосов на планировщик, что составляет 24 на SM (24/64 = 37,5% занятости), чтобы покрыть задержку. Если потоки имеют параллелизм на уровне команд, это можно уменьшить. Почти все ядра выдают команды с переменной задержкой, такие как загрузка памяти, которая может занимать 80-1000 циклов. Это требует более активных деформаций на планировщик деформации, чтобы скрыть задержку. Для каждого ядра существует компромисс между количеством перекосов и другими ресурсами, такими как разделяемая память или регистры, поэтому не рекомендуется оптимизировать загрузку на 100%, так как, скорее всего, будут сделаны некоторые другие жертвы. Профилировщик CUDA может помочь определить частоту проблем с инструкциями, занятость и причины задержки, чтобы помочь разработчику определить этот баланс.
Размер блока потока может повлиять на производительность. Если ядро имеет большие блоки и использует барьеры синхронизации, то это может стать причиной сбоя в барьере. Это может быть облегчено путем уменьшения перекосов в каждом блоке потока.
На одном устройстве установлено несколько потоковых мультипроцессоров.
SM может содержать несколько блоков. Каждый блок может содержать несколько потоков.
У SM есть несколько ядер CUDA (как разработчик, вы не должны заботиться об этом, потому что это абстрагируется от деформации), которые будут работать в потоке. СМ всегда работает на перекосе потоков (всегда 32). Деформация будет работать только с потоком из одного блока.
SM и блок имеют ограничения по количеству потоков, количеству регистров и общей памяти.