Глава 2: Модели программирования для GPU
Графические процессорные устройства (GPU) эволюционировали от фиксированных функций графических ускорителей до высокопараллельных, программируемых вычислительных двигателей, способных ускорять широкий спектр приложений. Для того, чтобы программисты могли эффективно использовать массивный параллелизм в GPU, были разработаны несколько моделей параллельного программирования и API, таких как NVIDIA CUDA, OpenCL и DirectCompute. Эти модели программирования предоставляют абстракции, которые позволяют программистам выражать параллелизм в своих приложениях, при этом скрывая низкоуровневые детали аппаратного обеспечения GPU.
В этой главе мы исследуем ключевые концепции и принципы за моделями параллельного программирования для GPU, сосредоточившись на модели выполнения, архитектурах набора инструкций (ISA) GPU, ISA NVIDIA GPU и ISA Graphics Core Next (GCN) AMD. Мы также предоставим примеры, чтобы проиллюстрировать, как эти концепции применяются на практике.
Модель выполнения
Модель выполнения современных моделей программирования GPU основана на концепции ядер, которые являются функциями, выполняемыми параллельно большим количеством потоков на GPU. При запуске ядра программист указывает количество потоков, которые будут созданы, и как они организованы в иерархию сеток, блоков (или массивов кооперативных потоков - CTAs) и отдельных потоков.
- Сетка представляет всё пространство проблемы и состоит из одного или нескольких блоков.
- Блок - это группа потоков, которые могут сотрудничать и синхронизироваться друг с другом через разделяемую память и барьеры. Потоки внутри блока выполняются на одном ядре GPU (называемом мультипроцессором потоков или вычислительным блоком).
- Каждый поток имеет уникальный идентификатор внутри своего блока и сетки, который может использоваться для вычисления адресов памяти и принятия решений управления потоком.
Эта иерархическая организация позволяет программистам выражать как параллелизм данных (когда одна и та же операция применяется к нескольким элементам данных), так и параллелизм задач (когда различные задачи выполняются параллельно).
РисунокВот перевод на русский язык с сохранением неизмененным кода:
e 2.1 иллюстрирует иерархию потоков в модели исполнения GPU.
Сетка
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Блок |
| | | |
Поток Поток ...Рисунок 2.1: Иерархия потоков в модели исполнения GPU.
Выполнение SIMT
Модели программирования GPU, такие как CUDA и OpenCL, следуют модели выполнения Single-Instruction, Multiple-Thread (SIMT). В модели SIMT потоки выполняются группами, называемыми warps (терминология NVIDIA) или wavefronts (терминология AMD). Все потоки внутри warp'а выполняют одну и ту же инструкцию одновременно, но каждый поток работает с различными данными.
Однако, в отличие от традиционной модели Single-Instruction, Multiple-Data (SIMD), где все вычислительные элементы выполняются синхронно, SIMT позволяет потокам иметь независимые пути выполнения и расходиться при ветвлениях. Когда warp встречает инструкцию ветвления, аппаратное обеспечение GPU оценивает условие ветвления для каждого потока в warp'е. Если все потоки следуют по одному и тому же пути (сошлись), warp продолжает выполнение как обычно. Если некоторые потоки следуют по разным путям (расошлись), warp разделяется на два или более подгруппы, каждая из которых следует по своему пути. Аппаратное обеспечение GPU последовательно выполняет расходящиеся пути, отключая неактивные потоки в каждой подгруппе. Когда все пути завершены, подгруппы снова сходятся и продолжают выполнение синхронно.
Рисунок 2.2 иллюстрирует выполнение SIMT с расходящимся управляющим потоком.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Ветвление |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
СходимостьРисунок 2.2: Выполнение SIMT с расходящимся управляющим потоком.
Этот механизм обработки расхождения позволяет SIMT поддерживать более гибкий контрольный поток tПредоставляю перевод на русский язык для данного markdown-файла. Для кода переведены только комментарии, а сам код остался без изменений.
Ускорение SIMD в Han, но это происходит за счет снижения эффективности SIMD при расхождении. Программистам следует стремиться к минимизации расхождений внутри warpu для достижения оптимальной производительности.
Иерархия памяти
GPU имеют сложную иерархию памяти для поддержки высокой пропускной способности и низкой задержки параллельных рабочих нагрузок. Иерархия памяти, как правило, состоит из:
- Глобальная память: Самая большая, но самая медленная область памяти, доступная всем потокам в ядре. Глобальная память, как правило, реализуется с использованием высокоскоростной памяти GDDR или HBM.
- Разделяемая память: Быстрая, встроенная в чип память, доступная всем потокам в блоке. Разделяемая память используется для межпотоковой связи и обмена данными внутри блока.
- Константная память: Только для чтения пространство памяти, используемое для широковещания данных только для чтения для всех потоков.
- Текстурная память: Только для чтения пространство памяти, оптимизированное для пространственной локальности и доступное через кэш текстур. Текстурная память чаще используется в графических рабочих нагрузках.
- Локальная память: Частное пространство памяти для каждого потока, используемое для перелива регистров и больших структур данных. Локальная память, как правило, отображается на глобальную память.
Эффективное использование иерархии памяти имеет решающее значение для достижения высокой производительности на GPU. Программисты должны стремиться к максимальному использованию разделяемой памяти и минимизации обращений к глобальной памяти, чтобы снизить задержку и пропускную способность памяти.
На рисунке 2.3 иллюстрируется иерархия памяти GPU.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________Рис.Вот перевод на русский язык:
Архитектуры набора инструкций GPU
Архитектуры набора инструкций GPU (GPU ISA) определяют низкоуровневый интерфейс между программным и аппаратным обеспечением. Они указывают на инструкции, регистры и режимы адресации памяти, поддерживаемые GPU. Понимание архитектур GPU ISA является essential для разработки эффективного кода GPU и оптимизации производительности.
В этом разделе мы исследуем ISA двух основных поставщиков GPU: NVIDIA и AMD. Мы сосредоточимся на PTX (Parallel Thread Execution) и SASS ISA NVIDIA, а также на ISA Graphics Core Next (GCN) AMD.
Архитектуры GPU ISA NVIDIA
Графические процессоры NVIDIA поддерживают два уровня ISA: PTX (Parallel Thread Execution) и SASS (Streaming ASSembler). PTX является виртуальной ISA, которая обеспечивает стабильную цель для компиляторов CUDA, в то время как SASS является родным ISA графических процессоров NVIDIA.
PTX (Parallel Thread Execution)
PTX - это низкоуровневая виртуальная ISA, разработанная для графических процессоров NVIDIA. Она похожа на LLVM IR или Java байт-код тем, что обеспечивает стабильную, независимую от архитектуры цель для компиляторов. Программы CUDA обычно компилируются в код PTX, который затем переводится в собственные инструкции SASS драйвером графического процессора NVIDIA.
PTX поддерживает широкий спектр арифметических, памятных и управляющих инструкций. Она имеет неограниченное количество виртуальных регистров и поддерживает предикацию, что позволяет эффективно реализовывать управление потоком. PTX также предоставляет специальные инструкции для синхронизации потоков, атомарных операций и выборки текстур.
Вот пример кода PTX для простого ядра векторного сложения:
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
ld.param.u64 %rd1, [vecAdd_param_0]; // Загрузить первый параметр
ld.param.u64 %rd2, [vecAdd_param_1]; // Загрузить второй параметр
ld.param.u64 %rd3, [vecAdd_param_2]; // Загрузить третий параметр
ld.param.u32 %r1, [vecAdd_param_3]; // Загрузить четвертый параметр
cvta.to.global.u64 %rd4, %rd1; // Преобразовать указатель в глобальный
cvta
```Вот русский перевод этого markdown-файла:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;
ret; }
Этот код PTX определяет функцию-ядро `vecAdd`, которая принимает четыре параметра: указатели на входные и выходные векторы, а также размер векторов. Ядро вычисляет глобальный идентификатор потока, загружает соответствующие элементы из входных векторов, выполняет сложение и сохраняет результат в выходной вектор.
#### SASS (Streaming ASSembler)
SASS - это родной ISA (набор инструкций) графических процессоров NVIDIA. Это низкоуровневый, аппаратно-специфичный ISA, который непосредственно сопоставляется с оборудованием GPU. Инструкции SASS генерируются драйвером GPU NVIDIA из кода PTX и, как правило, не видны программистам.
Инструкции SASS закодированы в компактном формате для уменьшения пропускной способности памяти и размера кэша инструкций. Они поддерживают широкий спектр типов операндов, включая регистры, немедленные значения и различные режимы адресации для доступа к памяти.
Ниже приведен пример кода SASS для ядра сложения векторов:
```sass
code_version_number 90
@P0 LDG.E R2, [R8];
@P0 LDG.E R0, [R10];
@P0 FADD R0, R2, R0;
@P0 STG.E [R12], R0;
EXIT;Этот код SASS соответствует коду PTX, показанному ранее. Он загружает элементы входных векторов из глобальной памяти (LDG.E), выполняет сложение (FADD), сохраняет результат обратно в глобальную память (STG.E) и завершает работу ядра.
Архитектура и ISA AMD Graphics Core Next
Графические процессоры AMD используют архитектуру и ISA Graphics Core Next (GCN). GCN - это RISC-based ISA, которая поддерживает как графические, так и вычислительные рабочие нагрузки. Она разработана для высокой производительности, масштабируемости и энергоэффективности.
GCN вводит несколько ключевых функций, таких как:
- Масштабируемая архитектура, поддерживающая различные количества вычислительных ядер
- Векторные инструкции SIMD для эффективного выполнения параллельных вычислений
- Улучшенная память и иерархия кэша для повышения пропускной способности памяти
- Поддержка различных форматов данных для графических и вычислительных приложений
Общая стратегия GCN - обеспечить высокую производительность и энергоэффективность для широкого спектра приложений.Русский перевод:
alar ALU для эффективного выполнения скалярных операций и управления потоком.
- Векторный ALU для параллельного выполнения операций параллельной обработки данных.
- Высокоскоростная система памяти с поддержкой атомарных операций и доступа с малой задержкой к общей памяти.
- Гибкий режим адресации для операций с памятью, поддерживающий адресацию по base+offset и скалярный+векторный адрес.
Вот пример кода ядра GCN ISA для векторного сложения:
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
s_load_dwordx4 s[0:3], s[4:5], 0x0 # Загрузка аргументов ядра из памяти
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
v_lshlrev_b32 v0, 2, v0 # Сдвиг индекса v0 влево на 2 бита
v_add_u32 v1, vcc, s1, v0 # Вычисление адреса источника 1
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc # Вычисление адреса источника 2
flat_load_dword v1, v[1:2] # Загрузка элементов источника 1 из памяти
v_add_u32 v3, vcc, s0, v0 # Вычисление адреса результата
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc
flat_load_dword v0, v[3:4] # Загрузка элементов источника 2 из памяти
v_add_f32 v0, v0, v1 # Сложение элементов источников
flat_store_dword v[3:4], v0 # Сохранение результата в память
s_endpgmЭтот код GCN загружает элементы входных векторов с помощью flat_load_dword, выполняет сложение с помощью v_add_f32 и сохраняет результат в память с помощью flat_store_dword. Инструкции s_load_dwordx4 и s_load_dword используются для загрузки аргументов ядра из памяти.
Сопоставление алгоритмов с архитектурами GPU
Эффективное сопоставление алгоритмов с архитектурой GPU имеет ключевое значение для достижения высокой производительности. Основные моменты:
Обеспечение достаточного параллелизма
Алгоритм должен быть разделен на множество мелкозернистых потоков, которые могут выполняться параллельно, чтобы полностью использовать параллельные вычислительные возможности GPU. Это часто включает в себя выявление частей алгоритма с параллельной обработкой данных, которые могут выполняться независимо на различных элементах данных.
Минимизация расхождения в ветвлении
Расходящееся управление потоком внутри warpa/wavefronts может привести к сериализации и снижению эффективности SIMD. Алгоритмы должны быть структурированы так, чтобы минимизировать расхождение в ветвлении там, где это возможно. Этого можно добиться, уменьшая использование зависящего от данных условного перехода.Вот русский перевод этого файла, при этом комментарии к коду переведены, а сам код не переводился:
Эксплуатация иерархии памяти
Обращение к глобальной памяти является дорогостоящей операцией. Алгоритмы должны максимально использовать общую память и регистры, чтобы уменьшить обращения к глобальной памяти. Данные также должны быть организованы в памяти таким образом, чтобы обеспечить коалесцированные обращения к памяти, когда потоки в варпе обращаются к смежным ячейкам памяти. Эффективное использование иерархии памяти может значительно снизить задержку памяти и проблемы с пропускной способностью.
Балансировка вычислений и обращений к памяти
Алгоритмы должны иметь высокое соотношение арифметических операций к операциям с памятью, чтобы эффективно скрывать задержку памяти и достигать высокой вычислительной производительности. Это можно достичь за счет максимизации повторного использования данных, предварительной загрузки данных и перекрытия вычислений с обращениями к памяти.
Минимизация передач данных между хостом и устройством
Передача данных между памятью хоста (CPU) и памятью устройства (GPU) является медленной операцией. Алгоритмы должны минимизировать такие передачи, выполняя как можно больше вычислений на GPU. Данные должны передаваться на GPU большими пакетами и оставаться на устройстве как можно дольше, чтобы амортизировать накладные расходы на передачу.
При разработке ядер GPU часто используются следующие параллельные шаблоны алгоритмов:
-
Отображение (Map): Каждый поток выполняет одну и ту же операцию над различными элементами данных, позволяя простую параллельную обработку больших наборов данных.
-
Уменьшение (Reduce): Параллельное уменьшение используется для эффективного вычисления одного значения (например, суммы, максимума) из большого входного набора данных. Потоки выполняют локальные уменьшения, которые затем объединяются для получения окончательного результата.
-
Сканирование (Scan): Также известное как префиксная сумма, сканирование используется для вычисления бегущей суммы элементов в массиве. Эффективные параллельные алгоритмы сканирования являются ключевыми строительными блоками для многих приложений, ускоренных на GPU.
-
Шаблон (Stencil): Каждый поток вычисляет значение на основе соседних элементов данных. Вычисления со шаблоном широко применяются в научных симуляциях и обработке изображений.Вот перевод на русский язык этого файла Markdown. Для кода не переводятся сами фрагменты кода, а только комментарии к ним.
-
Сбор/Рассеивание: Потоки читают из (сбор) или записывают в (рассеивание) произвольные места в глобальной памяти. Для эффективности требуется тщательный подход к компоновке данных и схеме доступа к ним.
На рисунке 3.20 показан пример шаблона "map", где каждый поток применяет функцию (например, квадратный корень) к отдельному элементу входного массива.
Входной массив:
| | | | | | | |
v v v v v v v v
______________________________
Потоки: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Выходной массив: Рисунок 3.20: Пример шаблона "map" в программировании GPU.
Заключение
Модели программирования GPU, такие как CUDA и OpenCL, открывают разработчикам параллельные вычислительные возможности современных GPU, позволяя ускорять широкий спектр приложений. Эти модели программирования предоставляют абстракции, позволяющие эффективно отображать мелкозернистые параллельные вычислительные нагрузки на аппаратные средства GPU.
Понимание модели выполнения, иерархии памяти и примитивов синхронизации, предоставляемых этими моделями программирования, является ключевым для написания высокопроизводительного кода GPU. Разработчики должны тщательно учитывать такие факторы, как организация потоков, расхождение ветвлений, схемы доступа к памяти и дизайн алгоритмов, чтобы в полной мере использовать вычислительную мощность GPU.
По мере развития архитектур GPU модели и инструменты программирования также должны совершенствоваться, чтобы позволить разработчикам эффективно использовать новые аппаратные возможности. Продолжающиеся исследования в областях, таких как дизайн языков программирования, оптимизация компиляторов и автоматическая настройка, будут иметь решающее значение для повышения продуктивности программистов и переносимости производительности в эпоху гетерогенных вычислений.