Программирование массово параллельных процессоров: Практический подход

📚 Краткое содержание

Этот курс предоставляет всестороннее введение в вычисления на графических процессорах (GPU) и параллельное программирование с использованием среды CUDA C. Охватываются архитектуры GPU, данные параллелизм, управление потоками, оптимизация памяти и продвинутые вопросы производительности, проиллюстрированные реальными кейсами, такими как реконструкция МРТ и визуализация молекул.

Освойте искусство высокопроизводительных параллельных вычислений с практическим пошаговым руководством по CUDA и архитектуре GPU.

Авторы: Дэвид Б. Кирк, Вэнь-Мэй В. Ху

Благодарности: Иэн Бак, Джон Николлс, команда NVIDIA DevTech, Дженсен Хуанг, Дэвид Лубке, Билл Бин, Саймон Грин, Марк Харрис, Манжу Хедж, Надим Мухаммад, Брент Остер, Питер Ширли, Эрик Йоунг и Сильвий Зеллер.

🎯 Цели обучения

1. Отличать философии проектирования и траектории производительности многопроцессорных центральных процессоров (CPU) и многоядерных графических процессоров (GPU).
2. Определять ключевые компоненты современной архитектуры GPU, включая потоковые мультипроцессоры (SM) и структуры памяти.
3. Применять закон Амдала для расчета теоретического ускорения и выявлять влияние последовательных узких мест.
4. Сравнить архитектурные различия между фиксированными функциональными конвейерами и программируемыми унифицированными массивами процессоров.
5. Объяснить роль «GPGPU» как промежуточного этапа и ограничения ранних моделей программирования шейдеров.
6. Проанализировать, как аппаратные возможности, такие как атомарные операции, синхронизация барьеров и поддержка двойной точности, способствовали переходу к масштабируемым вычислениям общего назначения.
7. Выявлять и использовать параллелизм данных в алгоритмах умножения матриц.
8. Реализовывать управление памятью устройств, включая выделение, передачу данных между хостом и устройством, и освобождение.
9. Создавать и запускать ядра CUDA с соответствующей индексацией потоков и конфигурациями сетки/блока.
10. Спроектировать многомерные иерархии потоков (сетки и блоки), чтобы отобразить сложные структуры данных на аппаратную платформу GPU.

🔹 Урок 1: Введение в параллельные вычисления и архитектуру GPU

Обзор: Этот урок исследует фундаментальный сдвиг от последовательных к параллельным вычислениям, обусловленный расхождением философий проектирования CPU и GPU. Учащиеся рассмотрят траектории «многопроцессорных» и «многоядерных» систем, поймут аппаратную архитектуру, позволяющую достигать огромной пропускной способности, и изучат математические ограничения ускорения по закону Амдала.

Результаты обучения:

• Отличать философии проектирования и траектории производительности многопроцессорных центральных процессоров (CPU) и многоядерных графических процессоров (GPU).
• Определять ключевые компоненты современной архитектуры GPU, включая потоковые мультипроцессоры (SM) и структуры памяти.
• Применять закон Амдала для расчета теоретического ускорения и выявлять влияние последовательных узких мест.

🔹 Урок 2: Эволюция и будущее вычислений на GPU

Обзор: Этот урок прослеживает архитектурную эволюцию графического процессора (GPU) от специализированного фиксированного аппаратного обеспечения для рендеринга треугольников до современного мощного, унифицированного параллельного процессора общего назначения. Учащиеся исследуют переход от жестких графических конвейеров к программируемым шейдерам, появление движения GPGPU и современные масштабируемые архитектуры, лежащие в основе современных научных и инженерных симуляций.

Результаты обучения:

• Сравнить архитектурные различия между фиксированными функциональными конвейерами и программируемыми унифицированными массивами процессоров.
• Объяснить роль «GPGPU» как промежуточного этапа и ограничения ранних моделей программирования шейдеров.
• Проанализировать, как аппаратные возможности, такие как атомарные операции, синхронизация барьеров и поддержка двойной точности, способствовали переходу к масштабируемым вычислениям общего назначения.

🔹 Урок 3: Архитектура программы CUDA и управление памятью

Обзор: Этот урок охватывает основную архитектуру программы CUDA, делая акцент на различии между выполнением на хосте (ЦПУ) и устройстве (GPU). Учащиеся узнают, как выявлять параллелизм данных в матричных операциях, управлять отдельными пространствами памяти с помощью API CUDA, и организовывать параллельное выполнение через иерархию сеток, блоков и потоков в стиле однопрограммного множественного данных (SPMD).

Результаты обучения:

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

🔹 Урок 4: Расширенное многопоточное программирование и планирование в CUDA

Обзор: Этот урок исследует иерархическую организацию потоков в CUDA, сосредоточившись на том, как многомерная индексация отображается на физические данные и аппаратные ресурсы. Он детально описывает механизмы синхронизации барьеров и прозрачную масштабируемость, завершаясь принципами архитектуры назначения потоков и планирования по группам (warps), используемыми для достижения терпимости к задержкам в высокопроизводительных вычислениях.

Результаты обучения:

• Спроектировать многомерные иерархии потоков (сетки и блоки), чтобы отобразить сложные структуры данных на аппаратную платформу GPU.
• Реализовать точную индексацию данных с использованием встроенных переменных CUDA (blockIdx, threadIdx, blockDim).
• Применить синхронизацию барьеров для обеспечения целостности данных при сохранении прозрачной масштабируемости на разных архитектурах GPU.

🔹 Урок 5: Оптимизация памяти и тайлинг в разделяемой памяти

Обзор: Этот урок исследует, как пропускная способность памяти и ограничения ресурсов действуют как основные узкие места в параллельных вычислениях. Он подробно описывает использование «тайлинга» для снижения трафика глобальной памяти и объясняет критическую роль барьеров синхронизации (__syncthreads()) и стратегического выбора между регистрами и разделяемой памятью для оптимизации производительности.

Результаты обучения:

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

🔹 Урок 6: Анализ производительности и выполнение по модели SIMT

Обзор: Этот урок исследует архитектурные и алгоритмические аспекты, необходимые для оптимизации ядер CUDA. Переходя от базовых моделей выполнения — в частности, единицы одиночной команды, множественных потоков (SIMT) и разделения варпов — к продвинутым методам настройки производительности, таким как коалиция памяти, тайловое умножение матриц и динамическое разделение ресурсов потокового мультипроцессора (SM).

Результаты обучения:

• Проанализировать отображение многомерных блоков потоков на линейный порядок выполнения варпов аппаратной части.
• Оценить и минимизировать расхождение управления потоками в параллельных алгоритмах сокращения.
• Оптимизировать пропускную способность глобальной памяти, реализуя коалицию памяти и шаблоны доступа к тайловым данным.

🔹 Урок 7: Арифметика с плавающей точкой и числовая точность

Обзор: Этот урок охватывает фундаментальную архитектуру чисел с плавающей точкой, с акцентом на компоненты стандарта IEEE 754: знак, экспонента с избыточным кодированием и нормализованная мантисса. Учащиеся исследуют, как эти битовые шаблоны отображаются на дискретной числовой оси, и как ограничения этой представления влияют на точность сложных алгоритмов, таких как суммирование больших массивов.

Результаты обучения:

• Разобрать формат числа с плавающей точкой для вычисления числовых значений из битовых шаблонов с использованием нормализованного представления и избыточного кодирования.
• Визуализировать распределение представляемых чисел на числовой оси и объяснить влияние распределения битов между экспонентой и мантиссой.
• Количественно оценить числовую неточность с использованием единиц последнего разряда (ULP) и определить, как различные режимы округления способствуют ошибкам.

🔹 Урок 8: Кейс-стади: Параллелизация реконструкции МРТ

Обзор: Этот урок исследует параллелизацию продвинутой реконструкции магнитно-резонансной томографии (МРТ) на GPU. Основное внимание уделяется итерационному процессу реконструкции для непрямоугольных траекторий, особенно оптимизации вычислительно затратного ядра F^H d с помощью преобразований циклов, управления постоянной памятью, перестройки структуры данных и использования аппаратно ускоренных тригонометрических функций.

Результаты обучения:

• Понять переход от реконструкции на основе быстрого преобразования Фурье (FFT) к итерационным алгоритмам решателей линейных уравнений для непрямоугольных данных k-пространства.
• Применить разбиение циклов и перестановку циклов для преобразования последовательного кода на языке С в структуру, подходящую для массового отображения на потоки CUDA.
• Оптимизировать пропускную способность памяти с помощью чанкирования постоянной памяти и структур типа «массив структур» (AoS).

🔹 Урок 9: Кейс-стади: Визуализация молекул и многопроцессорное выполнение

Обзор: Этот урок исследует практическое применение вычислений на GPU для визуализации молекул, конкретно с использованием метода прямого суммирования Кулона (DCS) для расчета карт электростатического потенциала. Учащиеся переходят от базовой реализации ядра к высокооптимизированным версиям, использующим развертывание инструкций, коалицию памяти и дополнение (padding).

Результаты обучения:

• Реализовать ядро прямого суммирования Кулона (DCS) с использованием постоянной памяти CUDA и техник скрытия задержек глобальной памяти.
• Оптимизировать производительность ядра с помощью развертывания инструкций и повторного использования вычислений общих координат.
• Применить стратегии коалиции памяти и дополнения для выравнивания обращений к глобальной памяти GPU с целью максимальной пропускной способности.

🔹 Урок 10: Компьютерное мышление и выбор параллельных алгоритмов

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

Результаты обучения:

• Определить основные цели параллельного программирования и рассчитать теоретическое ускорение с использованием закона Амдала.
• Различать декомпозицию на уровне задач и на уровне данных, применяя стратегии, ориентированные на атомы (рассеивание) и на сетку (собирание).
• Оценивать и выбирать параллельные алгоритмы на основе таких критериев, как пропускная способность памяти, вычислительная сложность и архитектурные ограничения.

🔹 Урок 11: Введение в модель программирования OpenCL

Обзор: Этот урок представляет OpenCL как фреймворк для гетерогенного параллельного программирования, с акцентом на его модель данных параллелизма и иерархическую абстракцию аппаратных средств. Учащиеся научатся отображать структуры NDRange и памяти OpenCL на эквиваленты в CUDA, а также освоить управление устройствами со стороны хоста через динамическую модель компиляции.

Результаты обучения:

• Отображать параллелизм и иерархии памяти OpenCL на специфические архитектуры CUDA (например, отображать рабочие группы на блоки и локальную память на разделяемую память).
• Реализовывать функции ядер OpenCL и управлять средой выполнения со стороны хоста с использованием контекстов и очередей команд.
• Выполнять динамический процесс компиляции для создания ядер из исходного кода во время выполнения.

🔹 Урок 12: Современные функции GPU и перспективы будущего

Обзор: Этот урок исследует эволюцию архитектуры и функциональных возможностей GPU, сосредоточившись на переходе к сложной системе управления памятью, расширенным возможностям выполнения ядер и увеличению производительности ядер. Учащиеся рассмотрят, как такие функции, как единое пространство памяти устройства и вызовы функций на уровне ядра, трансформируют GPU в процессор общего назначения.

Результаты обучения:

• Объяснить значение эволюции архитектуры памяти и перехода к 64-битному единому пространству памяти устройства.
• Проанализировать, как расширенные атомарные операции и вызовы функций на уровне ядра позволяют реализовывать сложные структуры данных и алгоритмы.
• Оценить влияние производительности при одновременном выполнении нескольких ядер, улучшении скорости двойной точности и эффективности управления потоками в современных средах GPU.