2016 Ответы на экзаменационные вопросы, страница 4

Параллельная программа на BlueGene/P.

Задача сводится к интегрированию 200 уравнений для ядер и решению 400 трехмерных задач для электронов.

Время одного расчета на 512 узлах BlueGene/P занимает ≈ 20 часов

Использование гибридной схемы распараллеливания MPI/OpenMP

Параллельные стратегии. Несколько уровней распараллеливания

• Крупноблочное распараллеливание на распределенной памяти MPI – распределение коэффициентов волновых функций для всех электронных состояний на все процессоры.

3dFFT Реальное пространство K-пространство

Распределение данных минимизирует число передач с поддержанием загрузки в обоих пространствах.

• OpenMP распараллеливание с общей памятью на узле. Длинные циклы.

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

• Методы копирования. Интегралы по траекториям в молекулярной динамике. Квантовые ядра. Многомерная метадинамика.

Отражение задач на архитектуру суперкомпьютера

• Blue Gene/P. Архитектура трехмерного тора дополненная сетью дерева

• Возникают условия на дизайн FPMD кодов

• Для P< 1024 (BGP) проблема неважна

• Для P=65536 (BGL) хороший выбор отображения дает 60% ускорения

• Оптимизация отображения - N! отображений.

• Проблема использования параллельных библиотек (ScaLAPACK). Узлы не под контролем кода моделирования

• Оптимизация - критическая процедура для больших разбиений

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

• Оптимизированный код Qbox показал производительность 207 TFlop/s (наивысшая производительность на научных приложениях). Blue Gene/L

Производительность CPMD на Blue Gene/P

• CPMD масштабируется до 128000 процессоров

• 100 атомов масштабируются до 2000 процессоров с 70% эффективности и длительностью производственного цикла ~600ps/неделя (МГУ)

• Для систем ~1000 атомов масштабируется на 8000-16000 процессоров с эффективностью 80% и длительностью производственного цикла ~ 20 ps/неделя

• Наибольшая система в настоящий момент 20000 атомов на 16 стойках.

1. Архитектурные особенности графических процессоров, направленные на массивно-параллельные вычисления. Особенности работы с памятью графического процессора.

GPGPU & CUDA

• GPU – Graphics Processing Unit

• GPGPU – General-Purpose computing on GPU. Первые GPU от NVIDIA с поддержкой GPGPU – GeForce восьмого поколения G80 (2006 г.)

• CUDA – Compute Unified Device Architecture, программно-аппаратная архитектура от NVidia, позволяющая производить вычисления с использованием графических процессоров.

Основные преимущества по сравнению с CPU

• Высокое соотношение цена/производительность

• Высокое соотношение производительность/энергопотребление

Программно-аппаратная модель: архитектура GPU NVidia

CPU Intel Core i7

• Небольшое число мощных независимых ядер;

• До 22 ядер, ~3.0 ГГц каждое;

• Поддержка виртуальных потоков (Hyper-Threading)

• 3х уровневый кеш, большой кеш L3 до 50 МБ;

• На каждое ядро L1=32KB (data) + 32KB ( Instructions), L2=256KB;

• Обращения в память обрабатываются отдельно для каждого процесса\нити

GPU Streaming Multiprocessor (SM)

• Потоковый мультипроцессор

• «Единица» построения устройства (как ядро в CPU):

  • 32 скалярных ядра CUDA Core, ~1.5ГГц

  • 2 Warp Scheduler

  • Файл регистров, 128KB

  • 2х уровневый кэш

  • Текстурные юниты

  • 16 x Special Function Unit (SFU) – интерполяция и трансцендентная математика одинарной точности

  • 16 x Load/Store

GPC - Graphics Processing Cluster – Объединение потоковых мультипроцессоров в блоки

Чип поколения Fermi в максимальной конфигурации

• 16 SM

• 512 ядер CUDA Core

• Кеш L2 758KB

• Контроллеры памяти GDDR5

• Интерфейс PCI 2.0

Сравнение GPU и CPU

• Сотни упрощённых вычислительных ядер, работающих на небольшой тактовой частоте ~1.8 ГГц;

• Небольшие кеши на GPU

  • 32 CUDA-ядра разделяют 64 КБ L1

  • L2 общий для всех CUDA ядер 2 MB, L3 отсутствует

• Оперативная память с высокой пропускной способностью и высокой латентностью, оптимизированная для коллективного доступа;

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

Утилизация латентности памяти

• Цель: эффективно загружать CUDA-ядра

• Проблема: латентность памяти

• Решение:

  • CPU: Сложная иерархия кешей;

  • GPU: Много нитей, покрывать обращения одних нитей в память вычислениями в других за счёт быстрого переключения контекста;

• За счёт наличия сотен ядер и поддержки миллионов нитей (потребителей) на GPU легче утилизировать всю полосу пропускания

Вычисления с использованием GPU

Программа, использующая GPU, состоит из:

• Кода для GPU, описывающего необходимые вычисления и работу с памятью устройства;

• Кода для CPU, в котором осуществляется:

  • Управление памятью GPU – выделение / освобождение

  • Обмен данными между GPU/CPU

  • Запуск кода для GPU

  • Обработка результатов и прочий последовательный код

• GPU рассматривается как периферийное устройство, управляемое центральным процессором

  • GPU «пассивно», т.е. не может само загрузить себя работой, но существует исключение!

• Код для GPU можно запускать из любого места программы как обычную функцию

  • «Точечная», «инкрементная» оптимизация программ

Терминология

• CPU Будем далее называть «хостом» (от англ. host )

  • код для CPU - код для хоста, «хост-код» (host-code )

• GPU будем далее называть «устройством» или «девайсом» (от англ. device)

  • код для GPU – «код для устройства», «девайс-код» (device-code )

• Хост выполняет последовательный код, в котором содержатся вызовы функций, побочный эффект которых – манипуляции с устройством.

Код для GPU (device-code)

Код для GPU пишется на C++ с некоторыми расширениями:

• Атрибуты функций, переменных и структур

• Встроенные функции

• Математика, реализованная на GPU

• Синхронизации, коллективные операции

• Векторные типы данных

• Встроенные переменные: threadIdx, blockIdx, gridDim, blockDim

• Шаблоны для работы с текстурами

• …

Компилируется специальным компилятором nvcc

Код для CPU (host-code)

• Код для CPU дополняется вызовами специальных функций для работы с устройством;

• Код для CPU компилируется обычным компилятором gcc/icc/cl;

• Кроме конструкции запуска ядра <<<...>>> !

• Функции для GPU линкуются из динамических библиотек

CUDA Grid

• Хост может запускать на GPU множества виртуальных нитей;

• Каждая нить приписана некоторому виртуальному блоку;

• Грид (от англ. Grid-сетка ) – множество блоков одинакового размера;

• Положение нити в блоке и блока в гриде индексируются по трём измерениям (x, y, z).

• Грид задаётся количеством блоков по [X, Y, Z] (размер грида в блоках) и размерами каждого блока по [X, Y, Z];

• Например, еcли по Z размер грида и блоков равен единице, то получаем плоскую прямоугольную сетку нитей.

CUDA Kernel («Ядро»)

• Каждая нить выполняет копию специально оформленных функций «ядер», компилируемых для GPU.

  • Нет возвращаемого значения (void);

  • Обязательный атрибут __global__.

• Терминология:

  • Хост запускает вычисление ядра на гриде нитей (либо просто хост запускает ядро на GPU).

  • Одно и то же ядро может быть запущено на разных гридах

  • «Ядро» – что делать

  • «Грид» – сколько делать

Запуск ядра

• kernel<<< execution configuration >>>(params);

  • “kernel” – имя ядра,

• “params” – параметры ядра, копию которых получит каждая нить

• execution configuration - Dg, Db, Ns, S

  • dim3 Dg - размеры грида в блоках, Dg.x * Dg.y * Dg.z число блоков

  • dim3 Db - размер каждого блока, Db.x * Db.y * Db.z - число нитей в блоке

  • size_t Ns – размер динамически выделяемой общей памяти (опционально)

  • cudaStream_t S - поток, в котором следует запустить ядро (опционально)

• struct dim3 – стуктура, определённая в CUDA Toolkit,

  • Три поля: unsigned x,y,z

  • Конструктор dim3(unsigned x=1, unsigned y=1, unsigned z=1)

Ориентация нити в гриде

• Осуществляется за счёт встроенных переменных:

  • threaIdx.x threaIdx.y threaIdx.z - индексы нити в блоке

  • blockIdx.x blockIdx.y blockIdx.z – индексты блока в гриде

  • blockDim.x blockDim.y blockDim.z – размеры блоков в нитях

  • gridDim.x gridDim.y gridDim.z – размеры грида в блоках

Host Code

• Выделить память на устройстве

• Переслать на устройство входные данные

• Рассчитать грид

  • Размер грида зависит от размера задачи

• Запустить вычисления на гриде

  • В конфигурации запуска указываем грид

• Переслать с устройства на хост результат

Выделение памяти на устройстве

• cudaError_t cudaMalloc ( void** devPtr, size_t size )

  • Выделяет size байтов линейной памяти на устройстве и возвращает указатель на выделенную память в *devPtr. Память не обнуляется. Адрес памяти выровнен по 512 байт

• cudaError_t cudaFree ( void* devPtr )

  • Освобождает память устройства на которую указывает devPtr.

• Вызов cudaMalloc(&p, N*sizeof(float)) соответствует вызову p = malloc(N*sizeof(float));

Копирование памяти

• cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

  • Копирует count байтов из памяти, на которую указывает src в память, на которую указывает dst, kind указывает направление передачи

    • cudaMemcpyHostToHost– копирование между двумя областями памяти на хосте

    • cudaMemcpyHostToDevice – копирование с хоста на устройство

    • cudaMemcpyDeviceToHost – копирование с устройства на хост

    • cudaMemcpyDeviceToDevice – между двумя областями памяти на устройстве

• Вызов cudaMemcpy() с kind, не соответствующим dst и src , приводит к непредсказуемому поведению

Глобальная память

• Расположена в DRAM GPU

• Объём до 4Gb

  • Параметр устройства totalGlobalMem

• Кешируется в кешах L1 и L2:

  • L1 – на каждом мультипроцессоре

    • максимальный размер 48KB

    • минимальный размер 16KB

  • L2 – на устройстве

    • максимальный размер 768 KB

    • Параметр устройства l2CacheSize

Выделение:

• Динамически с хоста через cudaMalloc()

• Динамически из ядер через malloc()

• Статически - глобальная переменная с атрибутом __device__

• Переменные с атрибутами __device__ и __constant__ находятся в глобальной области видимости и хранятся объектном модуле как отдельные символы

• Память под них выделяется статически при старте приложения, как и под обычные глобальные переменные

• Работать с ними на хосте можно через функции cudaMemcpyToSymbol , cudaMemcpyToSymbolAsync, cudaGetSymbolAddress, cudaMemcpyFromSymbol, cudaMemcpyFromSymbolAsynс, cudaGetSymbolSize

Динамически из ядер

• malloc() из ядра выделяет память в куче

• Не освобождается между запусками ядер

• Освобождение по free() только с устройства

• Компилировать с –arch=sm_20

• Доступны memcpy(), memset()

• Память под кучу выделяется на устройстве при инициализации CUDA runtime и освобождается при завершении программы

  • После создания размер кучи не может быть изменен

  • По-умолчанию 8мб

  • Можно задать до первого вызова ядра c malloc через cudaDeviceSetLimit(cudaLimitMallocHeapSize, N)

Режимы работы кеша L1

• Кеш может работать в двух режимах: 48KB и 16KB

• Переключение режимов:

  • cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig)

Устанавливает режим работы кеша cacheConfig для всего устройства

Возможные режимы:

• cudaFuncCachePreferNone – без предпочтений(по умолчанию). Выбирается последняя использованная конфигурация. Начальная конфигурация – 16KB L1

• cudaFuncCachePreferShared: 16КB L1

• cudaFuncCachePreferL1: 48KB L1

• cudaError_t cudaFuncSetCacheConfig ( const void* func, cudaFuncCache cacheConfig ) Устанавливает режим работы кеша cacheConfig для заданной функции func

  • По умолчанию - cudaFuncCachePreferNone - запускать с режимом устройства

Транзакции

• Глобальная память оптимизирована с целью увеличения полосы пропускания

  • Отдать максимум данных за одно обращение

• Транзакция – выполнение загрузки из глобальной памяти сплошного отрезка в 128 байт, с началом кратным 128 (naturally aligned)

• Инструкция обращения в память выполняется одновременно для всех нитей варпа (SIMT)

  • Выполняется столько транзакций, сколько нужно для покрытия обращений всех нитей варпа

  • Если нужен один байт – все равно загрузится 128

Кеш-линии

• Ядра взаимодействуют не с памятью напрямую, а с кешами

• Транзакция – выполнение загрузки кеш-линии

  • У кеша L1 кеш-линии 128 байт, у L2 - 32 байта, naturally aligned

  • Кеш грузит из памяти всю кеш-линию, даже если нужен один байт

• Можно обращаться в память, минуя кеш L1

  • Транзакции будут по 32 байта

Матрицы и глобальная память

• Матрицы хранятся в линейном виде, по строкам

• Пусть длина строки матрицы – 480 байт (120 float)

  • обращение – matrix[idy*120 + idx]

• Дополним каждую строку до размера, кратного 128 байтам – в нашем случае, 480 + 32 = 512, это наш pitch – фактическая ширина в байтах

• Эти байты никак не будут использоваться, т.е. 32/512=6% лишней памяти будет выделено (Но для больших матриц эта доля будет существенно меньше)

• Зато каждая строка будет выровнена по 128 байт

  • Обращение matrix[idy*128+ idx]

Выделение памяти с «паддингом»

• cudaError_t cudaMallocPitch (void ** devPtr, size_t * pitch, size_t width, size_t height)

  • width – логическая ширина матрицы в байтах

  • Выделяет не менее width * height байтов, может добавить в конец строк набивку, с целью соблюдения выравнивания начала строк

  • сохраняет указатель на память в (*devPtr)

  • сохраняет фактическую ширину строк в байтах в (*pitch)

• Адрес элемента (Row, Column) матрицы, выделенной при помощи cudaMallocPitch:

  • T* pElement = (T*)((char*) devPtr + Row * pitch) + Column;

Копирование в матрицу с padding-ом

• cudaError_t cudaMemcpy2D ( void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, cudaMemcpyKind kind )

  • dst - указатель на матрицу, в которую нужно копировать ,

    • dpitch – фактическая ширина её строк в байтах

  • src - указатель на матрицу из которой нужно копировать,

    • spitch – фактическая ширина её строк в байтах

  • width – сколько байтов каждой строки нужно копировать

  • height – число строк

  • kind – направление копирования (как в обычном cudaMemcpy)

• Данная функция из начала каждой строки исходной матрицы копируется по width байтов. Всего копируется width*height байтов, при этом

  • Адрес строки с индексом Row определяется по фактической ширине:

    • (char*)src + Row* spitch – в матрице-источнике

    • (char*)dst + Row* dpitch – в матрице-получателе

Для обращения к матрице по столбцам её транспонируют