Глава 7: Дизайн вычислительного многопроцессора в GPU-дизайне
Вычислительный многопроцессор (SM) является основным строительным блоком архитектур GPU NVIDIA. Каждый SM содержит набор ядер CUDA, которые выполняют инструкции в режиме SIMT (Single Instruction, Multiple Thread). SM отвечает за управление и планирование варпов, обработку расходящихся ветвлений и обеспечение быстрого доступа к разделяемой памяти и кэшам. В этой главе мы исследуем микроархитектуру SM, включая её конвейеры, механизмы планирования варпов, конструкцию регистрового файла, а также организацию разделяемой памяти и L1-кэша.
Микроархитектура и конвейеры SM
SM представляет собой высокопараллельный и многоконвейерный процессор, предназначенный для эффективного выполнения сотен потоков одновременно. На рисунке 7.1 показана упрощённая блок-схема SM в архитектуре NVIDIA Volta.
Кэш инструкций
|
v
Планировщик варпов
|
v
Блок распределения (4 варпа)
| | | |
v v v v
Ядро CUDA (FP64/FP32/INT)
Ядро CUDA (FP64/FP32/INT)
Ядро CUDA (FP64/FP32/INT)
...
Тензорное ядро
Тензорное ядро
...
Блок загрузки/сохранения
Блок загрузки/сохранения
...
Блок специальных функций
^
|
Регистровый файл (64 КБ)
^Общая память / кэш L1 (96 КБ)
Рисунок 7.1: Упрощенная блок-схема SM в архитектуре NVIDIA Volta.
Основные компоненты SM включают:
-
Кэш инструкций: Хранит часто используемые инструкции для уменьшения задержки и повышения пропускной способности.
-
Диспетчер варпов: Выбирает варпы, готовые к выполнению, и отправляет их на доступные исполнительные блоки.
-
Блок распределения: Извлекает и декодирует инструкции для до 4 варпов за цикл и отправляет их на соответствующие исполнительные блоки.
-
CUDA-ядра: Программируемые исполнительные блоки, поддерживающие широкий спектр целочисленных и плавающих операций. Каждый SM в Volta содержит 64 CUDA-ядра.
-
Тензорные ядра: Специализированные исполнительные блоки, предназначенные для ускорения задач глубокого обучения и ИИ. Каждый SM в Volta содержит 8 тензорных ядер.
-
Блоки загрузки/сохранения: Выполняют операции с памятью, включая загрузку и сохранение в глобальную память, общую память и кэши.
-
Блоки специальных функций: Выполняют трансцендентные и другие сложные математические операции.
-
Регистровый файл: Обеспечивает быстрый доступ к частным регистрам потоков. Каждый SM в Volta имеет 64 КБ регистрового файла.
-
Общая память / кэш L1: Настраиваемое пространство памяти, которое может использоваться как управляемый программным обеспечением кэш (общая память) или как управляемый аппаратно кэш L1 данных.
Конвейер SM разработан для максимизации пропускной способности, позволяя выполнять несколько варпов одновременно и скрывать задержку памяти. Рисунок 7.2 иллюстрирует упрощенное представление конвейера SM.
Выборка инструкций
|
v
Декодирование инструкций
|
v
Сбор операндов
|
v
Выполнение (CUDA-ядра, тензорные ядра, блоки загрузки/сохранения, блоки специальных функций)
|
v
Запись результатовРисунок 7.2: Упрощенный конвейер SM.
Этапы конвейера:
-
Выборка инструкций: Диспетчер варпов выбирает варп, готовый к выполнению,Вот перевод на русский язык с сохранением оригинального кода:
-
Выборка инструкции: Процессор SM (Streaming Multiprocessor) выбирает следующую инструкцию для этого варпа (группы потоков) из кэша инструкций.
-
Декодирование инструкции: Выбранная инструкция декодируется для определения типа операции, операндов и регистров назначения.
-
Сбор операндов: Необходимые операнды для инструкции собираются из регистрового файла или общей памяти.
-
Выполнение: Инструкция выполняется на соответствующем блоке исполнения (CUDA-ядре, Tensor-ядре, модуле загрузки/сохранения или специальном функциональном блоке).
-
Запись результата: Результат выполнения записывается обратно в регистровый файл или общую память.
Для достижения высокой производительности SM использует несколько техник для максимального использования ресурсов и маскировки задержек:
- Двойная выдача: SM может выдавать две независимые инструкции за один такт для одного варпа, что позволяет увеличить параллелизм на уровне инструкций.
- Конвейерное выполнение: Блоки исполнения имеют конвейерную архитектуру, что позволяет SM начинать новую операцию, пока предыдущая еще не завершилась.
- Маскировка задержек: SM может переключаться между варпами на каждом такте, что позволяет скрыть задержки от обращений к памяти и длительных операций, выполняя инструкции из других варпов.
Пример 7.1 показывает простое ядро CUDA, выполняющее поэлементное сложение двух векторов.
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
c[tid] = a[tid] + b[tid];
}
}Пример 7.1: Ядро CUDA для сложения векторов.
В этом примере каждый поток в ядре вычисляет сумму соответствующих элементов из входных векторов a и b и записывает результат в выходной вектор c. SM выполняет это ядро, назначая каждый поток на CUDA-ядро и планируя варпы потоков для выполнения на доступных ядрах. Модули загрузки/сохранения используются для получения входных данных из глобальной памяти и записи результатов обратно.
Планирование варпов и обработка расхождений
EfФайл на русском языке:
Эффективное планирование варпов имеет решающее значение для максимизации производительности SM. Планировщик варпов отвечает за выбор готовых к выполнению варпов и их отправку на доступные вычислительные блоки. Основная цель планировщика варпов - занять вычислительные блоки, обеспечивая наличие готовых к выполнению варпов.
SM использует двухуровневый механизм планирования варпов:
-
Планирование варпов: Планировщик варпов выбирает готовые к выполнению варпы на основе политики планирования, такой как round-robin или oldest-first. Выбранные варпы затем отправляются на доступные вычислительные блоки.
-
Планирование инструкций: Внутри каждого варпа SM планирует инструкции на основе их зависимостей и доступности вычислительных блоков. SM может выдавать несколько независимых инструкций из одного и того же варпа за один цикл для максимизации параллелизма на уровне инструкций.
На рисунке 7.3 иллюстрируется двухуровневый механизм планирования варпов.
Пул варпов
Варп 1 (готов)
Варп 2 (ожидает)
Варп 3 (готов)
...
Варп N (готов)
|
v
Планировщик варпов
|
v
Блок распределения
|
v
Вычислительные блокиРисунок 7.3: Двухуровневый механизм планирования варпов.
Одной из ключевых проблем в планировании варпов является обработка расхождения ветвлений. В модели SIMT все потоки в варпе выполняют одну и ту же инструкцию синхронно. Однако, когда варп сталкивается с инструкцией ветвления (например, с оператором if-else), некоторые потоки могут выбрать ветку if, а другие - ветку else. Эта ситуация называется расхождением ветвлений.
Для обработки расхождения ветвлений SM использует технику предикации. Когда варп сталкивается с расходящейся ветвью, SM выполняет обе ветви ветвления последовательно, маскируя потоки, которые не выбирают каждую ветвь. Результаты затем объединяются с использованием предикатных регистров, чтобы убедиться, что каждый поток получает правильный результат.
Пример 7.2 показывает CUDA-ядро с расходящимся ветвлением.Вот перевод на русский язык:
__global__ void divergentKernel(int *data, int *result) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (data[tid] > 0) {
result[tid] = data[tid] * 2;
} else {
result[tid] = data[tid] * 3;
}
}Пример 7.2: CUDA-ядро с дивергентной ветвью.
В этом примере условие ветвления data[tid] > 0 может заставить некоторые потоки в варпе выполнять путь if, а другие - путь else. SM обрабатывает эту дивергенцию, последовательно выполняя оба пути и отключая неактивные потоки в каждом пути.
Рисунок 7.4 иллюстрирует процесс предикации для варпа с дивергентными потоками.
Warp (32 потока)
Поток 1: data[1] = 5, result[1] = 10
Поток 2: data[2] = -3, result[2] = -9
...
Поток 32: data[32] = 7, result[32] = 14
Дивергентная ветвь:
if (data[tid] > 0) {
result[tid] = data[tid] * 2;
} else {
result[tid] = data[tid] * 3;
}
Предикация:
Шаг 1: Выполнить путь `if` с маской
Поток 1: result[1] = 10
Поток 2: (отключен)
...
Поток 32: result[32] = 14
Шаг 2: Выполнить путь `else` с маской
Поток 1: (отключен)
Поток 2: result[2] = -9
...
Поток 32: (отключен)
Окончательный результат:
Поток 1: result[1] = 10
Поток 2: result[2] = -9
...
Поток 32: result[32] = 14Рисунок 7.4: Процесс предикации для варпа с дивергентными потоками.
Используя предикацию, SM может обрабатывать дивергенцию ветвей без необходимости явных инструкций ветвления или дивергенции управляющих потоков. Однако дивергентные ветви все еще могут влиять на производительность, поскольку SM должен последовательно выполнять оба пути, снижая эффективный параллелизм.
Регистровый файл и операторы сбора операндов
Регистровый файл является критически важным компонентом SM, обеспечивающим быстрый доступ к частным регистрам потоков. Каждый SM имеет большой регистровый файл, чтобы поддерживать множество активных потоков и обеспечивать эффективное переключение контекста между варпами.Вот русский перевод файла с комментариями к коду:
В архитектуре NVIDIA Volta каждый SM имеет 64 КБ регистрового файла, организованного в 32 банка по 2 КБ каждый. Регистровый файл разработан для обеспечения высокой пропускной способности и низкой латентности для поддержки большого количества одновременно выполняющихся потоков.
Для минимизации конфликтов банков и улучшения производительности, SM использует технику под названием "сбор операндов". Коллекторы операндов являются специализированными блоками, которые собирают операнды из банков регистрового файла и передают их в блоки исполнения. Используя коллекторы операндов, SM может снизить влияние конфликтов банков и повысить использование блоков исполнения.
На рисунке 7.5 показана упрощенная схема регистрового файла и коллекторов операндов в SM.
Регистровый файл (64 КБ)
Банк 1 (2 КБ)
Банк 2 (2 КБ)
...
Банк 32 (2 КБ)
|
v
Коллекторы операндов
|
v
Блоки исполненияРисунок 7.5: Регистровый файл и коллекторы операндов в SM.
Коллекторы операндов работают, собирая операнды из нескольких инструкций и нескольких пучков, позволяя SM выдавать инструкции из разных пучков в блоки исполнения за один цикл. Это помогает скрыть латентность доступа к регистровому файлу и улучшает общую производительность SM.
Пример 7.3 показывает CUDA-ядро, которое выполняет операцию скалярного произведения двух векторов.
__global__ void dotProduct(float *a, float *b, float *result, int n) {
// Выделение разделяемой памяти для частичных сумм
__shared__ float partialSum[256];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x + threadIdx.x;
partialSum[tid] = 0; // Инициализация частичной суммы
// Вычисление частичных сумм в каждом потоке
while (i < n) {
partialSum[tid] += a[i] * b[i];
i += blockDim.x * gridDim.x;
}
// Синхронизация потоков
__syncthreads();
// Последовательное суммирование частичных сумм
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
partialSum[tid] += partialSum[tid + s];
}
__syncthreads();
}
// Запись финальной суммы в глобальную память
if (tid == 0) {
result[blockIdx.x] = partialSum[0];
}
}В этом примере каждый поток вычисляет частичную сумму скалярного произведения, используя назначенные ему элементы.Ниже приведен перевод на русский язык, где комментарии к коду были переведены, а код остался без изменений:
Элементы из входных векторов. Частичные суммы хранятся в массиве разделяемой памяти partialSum. После того, как все потоки вычислили свои частичные суммы, выполняется параллельное уменьшение для суммирования частичных сумм и получения окончательного результата скалярного произведения.
Коллектор операндов играет важную роль в этом примере, эффективно собирая операнды для обращений к разделяемой памяти и арифметических операций. Это помогает избежать конфликтов банков и повысить использование блоков исполнения.
Вывод
Многопроцессорный стример является основной вычислительной единицей в современных архитектурах GPU. Его конструкция сосредоточена на максимизации пропускной способности и сокрытии латентности памяти через комбинацию мелкозернистого многопоточного выполнения, выполнения SIMT и эффективного сбора операндов.
Ключевые компоненты SM включают планировщик варп, который выбирает варпы для выполнения; стек SIMT, который обрабатывает дивергенцию и конвергенцию ветвлений; регистровый файл и коллекторы операндов, которые обеспечивают быстрый доступ к потоковым частным регистрам; и разделяемая память и кэш L1, которые позволяют обмениваться данными и повторно использовать их с низкой задержкой.
По мере развития архитектур GPU исследования в областях, таких как обработка дивергенции ветвлений, планирование варпов и конструкция регистрового файла, будут иметь решающее значение для улучшения производительности и эффективности будущих GPU. Новые методики, такие как динамическое формирование варпов, компактирование блоков потоков и кэши повторного использования операндов, имеют потенциал значительно расширить возможности SM и обеспечить новые уровни производительности в параллельных вычислительных рабочих нагрузках.