Просмтор этого файла доступен только зарегистрированным пользователям. Но у нас супер быстрая регистрация: достаточно только электронной почты!

Текст из файла (страница 4)

График 3.3 Сетка 256x256, 1000 шагов.

График 3.4 Сетка 512x512, 1000 шагов.

График 3.5 1024x1024, 1000 шагов.

Из полученных измерений можем сделать вывод, что минимальное время достигается, если количество узлов есть 1/5 от характерного размера сетки. Настоящей причины этого факта установить не удалось. Есть предположение о том, что при таком значении обрабатываемых узлов каждым потоком достигается наилучшее соотношение между временем вычисления и временем копирования в память.

Теперь можно построить график зависимости ускорения от характерного размера сетки.

График 3.6

Хотя из График 3.6 мы видим, что ускорение увеличилось, после подбора наилучшего числа вычисляемых в одном потоке узлов. Но все же нет уверенности в правильном подборе константы 1/5. И все же пока не использована разделяемая память, с помощью которой ожидается добиться еще большего ускорения.

Как результат, можно сказать, что мы уже добились ускорения до 30 раз. Конечно, при увеличении размера сетки увеличивается ускорение. Как говорилось ранее, это связано с некоторой инициализацией работы устройства.

Теперь, сделав первое приближение алгоритма, продолжим его улучшать и используем разделяемую память.

1. Распараллеливание решения 2-х мерной задачи упругости с помощью технологии CUDA и использования разделяемой(shared) памяти.

Напомним, что разделяемая память видна только потокам из одного блока. Объем той памяти зависит от GPU, но в нашем случае он равен 48 Кб. Этого вполне достаточно для блока, ведь основное ограничение, это объем общей памяти на GPU.

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

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

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

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

График 3.7 Ускорение для float по сравнению с CPU.

Из График 3.7 можно сделать вывод, что ускорение увеличивается при увеличении сетки, с этим эффектом мы уже сталкивались.

Зафиксируем результат. Получилось добиться ускорения в 65 раз на самом деле не на большой сетке. Чем больше сетка, тем большее ускорение будет.

Теперь рассмотрим, какое ускорение можно получить для double.

График 3.8 Зависимость ускорения для double по сравнению с CPU.

Из График 3.8 можно сделать вывод, что ускорение для double не очень сильно увеличивается с увеличением размера сетки. Но в целом, ускорение в 30 раз, это довольно хороший результат для double.

Теперь, когда программа работает достаточно быстро, стоит рассмотреть ее применимость к реальным задачам.

Основная проблема в том, что в реальных задачах нужны большие сетки, например, сетка узлов это вполне нормально. Рассмотрим, сколько у нас узлов может быть максимально.

На устройстве nVidia Tesla s2050 реально возможно использовать около 2GB памяти. Хотя в спецификации написана цифра 2,37 GB, но все же реально использовать получается меньше.

Для двухмерного случая:

Таблица 5.1 Расчет максимального размера матрицы.

Конечно, корректность работы программы на таких сетках проверена, хотя вывод программы, для одного шага занимает порядка 100 MB места на жестком диске. Что конечно довольно много, ведь даже сохранение 100 шагов, это уже GB.

Как видно из Таблицы 3.1, размер полученных сеток совсем небольшой, т.е. даже для данных типа float мы получим всего лишь узлов. Что для реальных расчетов является небольшим размером сетки.

Но зададимся вопросом, какой смысл использовать то, что не имеет практического применения? Или имеет, но в весьма ограниченной области. Как можно решать эту проблему?

Есть 2 варианта:

1. Считать матрицу кусками, т.е. подсчет, например, по оси X, что сейчас мы осуществляем как 1 вызов программы ядра (вызов функции, которая выполняется на GPU), можно разбить на несколько. Т.е. организуется итерационный процесс, загрузили первые 100 строк, подсчитали, выгрузили.

2. Воспользоваться несколькими GPU устройствами, т.е. постараться распараллелить наш алгоритм на несколько GPU.

Теперь рассмотрим, какие проблемы существуют у этих вариантов, чтобы мы могли сделать вывод в каком направлении двигаться.

Рис. 5.3

На Рис. 3.3 представлено как изменится работа программы, если матрицы уже не будут полностью помещаться в память GPU.

Давайте оценим время, просто на то, чтобы при каждом шаге выгружать и загружать данные:

Таблица 5.2

Хочется заметить, что сейчас технологии позволяют делать подсчет и загрузку/выгрузку параллельно, но это поддерживается только в последних устройствах. Но даже если учитывать этот факт, то несложным подсчетом получим, что на это будет уходить неоправданно большое время. Такая архитектура нас не устраивает.

Второй вариант более привлекателен. Нам не надо постоянно копировать данные на GPU и обратно. В случае подсчета на нескольких GPU имеется 2 основных проблемы:

1. Синхронизация устройств между собой и хостом.

2. Обмен информации между устройствами.

Из всего вышесказанного можно заключить, что хотя у обоих вариантов есть свои минусы, все же гораздо выгоднее использовать 2-ой. Т.е. попытаться распараллелить программу на несколько GPU.

1. Распараллеливание решения 2-х мерной задачи упругости с помощью технологии CUDA на нескольких GPU

Проблемы, с которыми мы столкнемся в разработке этого алгоритма, описаны в предыдущей главе.

Начнем с проблемы обмена информации между устройствами.

Есть несколько способов, первый из них, это обмениваться данными через хост. Конечно, это относительно долго. Система совершает 1 лишнее копирование. Но с другой стороны, это работает при любой конфигурации GPU устройств, т.е. не важно, как они соединены. Есть 2-ой способ, заключается он в прямом копировании с одного GPU устройства на другое (PeerToPeer Copy). PeerToPeerCopy – функция, появившаяся в CUDA 4.0 [8]. Она позволяет копировать данные с одной карты на другую, если они соединены через шину PCI-Express. Важно, что при этом не надо перегонять данные через host. Это позволяет значительно увеличить размер исходной сетки.

Но это не всегда работает, потому что необходимым условием для такой работы является то, что все GPU подсоединены к одной шине PCI-Express. При этом мы избегаем лишнего копирования.

Разрабатывался 2-ой вариант. Сервер содержит 8 GPU устройств (NVidia Tesla s2050). И 6 из них удовлетворяют необходимым требованиям. Поэтому все результаты, которые будут получены, исследованы на 6 GPU устройствах.

Рассмотрим проблему синхронизации. Все вызовы функций GPU асинхронны с хостом. Однако возможность синхронизации естественно есть. Осуществляется она с помощью функции cudaDeviceSynchronize(). И должна вызываться, когда cudaSetDevice выставлена как раз на устройство, которое мы хотим синхронизировать.

Изобразим ход программы, и в каких местах мы планируем делать синхронизацию.

Рис. 5.4 Работа алгоритма на нескольких GPU.

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

Для подсчета на нескольких GPU матрица разбивается достаточно тривиальным образом. Получаем, сколько рядов матрицы приходится на каждый GPU, и остаток отдаем последнему.

Как уже говорилось, для получения нового значения в узле, нам нужно знать значения, 4 соседних узлов, 2-х справа и 2-х слева. Поэтому, при обмене информации между GPU нам достаточно только пересылать два последних ряда следующему GPU и получить от него 2 ряда, которые для него считаются первыми. Следующий рисунок лучше иллюстрирует происходящее:

Рис. 5.5

Каждая из матриц, изображенных на Рис. 3.5 обрабатывается отдельным GPU, в данном случае верхняя обрабатывается GPU под номером N-1, средняя GPU под номером N, и нижняя GPU под номером N+1. Зеленым цветом обозначены данные, которые не пересчитываются на GPU, а передаются с соседнего GPU. На соседнем GPU, данные, которые были пересчитаны, и теперь должны передаваться обозначены желтым светом.

Таким образом, мы минимизировали количество информации, передаваемой между GPU.

Понятно, что при использовании одного GPU устройства никак не изменилась скорость работы при сравнении с алгоритмом, описанным в главе 3.4.

Теперь рассмотрим, полученные результаты:

График 3.9 Зависимость ускорения от количества устройств, по сравнению со временем работы на 1 устройстве для типа данных float.

График 3.10 Зависимость ускорения от количества устройств, по сравнению со временем работы на 1 устройстве для типа данных double.

На графиках 3.9 и 3.10 в легенде указаны характерные размеры сетки. Тестирование проводилось на квадратных сетках. Итак, как видно из графиков, числа немного отличаются. Но тенденция совпадает в обоих случаях. Она заключается в том, что при больших размерах сетки, мы меньше времени, относительно остальных операций, тратим на обмен данных между графическими устройствами и синхронизацию между ними и хостом, т.е. реальные вычисления на GPU становятся более тяжеловесными. А это означает, что ускорение будет расти.

Мы получили ускорение до 4 раз на 6 карточках, этот результат можно считать приемлемым.

Но, главная цель не в этом. Основная цель была приблизиться к размеру сетки, которая бы могла применяться в реальных расчетах.

Таблица 5.3 Максимальный размер сетки при использовании нескольких GPU.

Как результат, здесь можно зафиксировать, что мы получили максимальное количество узлов на сетки, при использовании float узлов. Что конечно уже можно использовать в практических целях.

1. Оптимизация

Исследование возможностей оптимизации проводилось с помощью NVidia Profiler.

Так как на сервере не было возможности использовать графические средства для профилирования, использовалась программа nvprof. После анализа работы программы, она записывает эти данные в файлы, которые можно открыть с помощью уже графического анализатора NVidia Profiler.

В нашей задаче, анализ проводился так же на карточке GPU GT9800, которая имеет compute compatibility 1.1, что в свою очередь сильно сужает количество анализируемых данных о состоянии системы.

Основная оптимизация, сделанная в работе, описана в главе 5.4.

1. Моделирование влияния головной ударной волны Челябинского метеорита на поверхность Земли

   1. Введение

15 февраля 2013 года на Землю упал Челябинский метеорит. Явление таких масштабов происходит достаточно редко, потому что при взрыве в атмосфере, ударная волна нанесла серьезные повреждения в Челябинске. При анализе этого события было получено огромное количество данных, которые позволяют исследовать это явление и схожие с ним. Челябинский метеорит является одним из крупнейших небесных тел, упавших на Землю, после Тунгусского метеорита. Наиболее полные исследования явления падения Челябинского метеорита были описаны в [10, 11].

1. Начальные данные

Задача решалась следующим образом. Первым шагом было решение газодинамической задачи с помощью уровнений Эйлера с учетом теплорадиационного обмена. Для этой задачи задавались начальные значения, а именно:

1. характерный размер метеорита

2. его плотность

3. высота взрыва падающего небесного тела

4. угол падения

5. скорость падения

Так как скорость метеорита и угол, под которым происходило падение этого небесного тела известны достаточно точно, то для цели уточнения модели их варирование невозможно.

На выходе газодинамической задачи получается функция зависимости давления на поверхности Земли (дневной поверхности) от времени. Эти данные были записаны в файл, который в дальнейшем передавался в виде начальных данных для сейсмической задачи.

Характеристики

Список файлов ВКР