CC BY
22
Астахова И.Ф.,
Воронежский государственный университете, профессор
astachova@list.ru
Коробкин Е.А.
Воронежский государственный университет, аспирант
Evgeniy-voronez@mail.ru
Технология CUDA для метода дискретных элементов в
параллельной среде
Гранулированные материалы, такие как песок, грунт, порошки или зерно, являются неотъемлемой частью нашего мира. Существует потребность в создании симуляции динамики таких материалов. В горнодобывающей промышленности - это изучение самообрушения пород; в сельскохозяйственной отрасли изучение поведения зерна в каналах имеет большое значение.
Метод Дискретного Элемента или МДЭ описывает способ симуляции сыпучих сред. В нем каждая частица материала представляется как тело из механики Ньютона, которое имеет массу и координаты центра, и на которое действуют силы. Таким образом, частицы взаимодействуют друг с другом по определенным законам и возможно наблюдать движение системы в целом. В настоящее время существует несколько комплексов приложений для моделирования частиц с помощью МДЭ. Примером могут служить Open Source проекты LIGGGHTS и Yade, которые решают целый спектр научных задач, связанных с Методом Дискретных Элементов. Устройства находятся в широком доступе, а их мощностей хватает для небольших научных исследований с участием сотен тысяч частиц, тогда как использовать ЦП (центральный процессор) можно для расчета очень небольших систем за приемлемое время.
Технология CUDA, представленная компанией NVIDIA в 2006 г., предоставляет интерфейс низкого уровня для решения задач программирования в параллельной среде с помощью видеокарт NVIDIA и идеально подходит для целей данной работы. В данной работе, используя архитектуру параллельных вычислений CUDA, необходимо реализовать программное и аппаратное обеспечение, позволяющее производить симуляцию гранулированных частиц. Изучить производительность такого программного комплекса относительно классических способов реализации, использующих центральный процессор в качестве основной вычислительной единицы.
Технология CUDA
Технология CUDA использует мощности графических процессоров.
Данные устройства оптимизированы для работы с числами с плавающей точкой одинарной точности; их архитектура сильно отличается от рхитектуры ЦП.
Графические процессоры (ГП) сконцентрированы на вычислениях больше, чем на контроле выполнения (условные переходы) и кэшировании данных. Также ГП ориентируются на параллельные вычисления. Поэтому данные устройства являются периферийными для решения конкретных вычислительных задач, которые возможно выполнять в параллельной среде, при этом основания логика приложения выполняется на ЦП.
Для того чтобы технология CUDA была масштабируема, вводится три ключевых абстракции: иерархия групп потоков, разделенная память и барьерная синхронизация. Программист должен разделить задачу на несколько подзадач, которые выполняются независимо разными блоками потоков, которые в свою очередь выполняют маленькую часть этой подзадачи. Разделенная память и синхронизация позволяют этим потокам работать вместе. Масштабируемость достигается за счет того, что блок может быть выполнен на любом ядре ГП, в любой последовательности, параллельно или последовательно, так что программа может быть выполнена на любом количестве доступных ядер.
CUDA расширяет язык C вводя понятия ядер (kernel) - функций на языке С, которые, когда вызываются, выполняются N раз на N различных потоках CUDA.
Ядро определяется с помощью приставки _global_, а количество потоков определяется с помощью нового синтаксиса конфигурации запуска <<<...>>>. Каждый поток, который выполняет ядро, получает уникальный индекс, доступ к которому можно получить через встроенную переменную threadldx.
Для удобства, переменная threadldx - трехкомпонентный вектор, так что потоки могут быть идентифицированы по одномерному, двумерному и трехмерному индексу, образуя при этом одномерную, двумерную или трехмерную сетку потоков. Таким образом обеспечивается прозрачность в работе с такими объектами как массив, матрица или объемная сетка.
Количество потоков в блоке ограничено из-за того, что все потоки должны быть расположены в пределах одного процессорного ядра. В настоящее время, максимальное количество потоков на блок равняется 1024. Однако ядро может быть выполнено на нескольких блоках, так что общее количество потоков равняется количеству потоков в блоке умноженному на количество блоков.
Блоки также организованны в одномерную, двумерную или трехмерную сетку. Каждый блок в сетке может быть идентифицирован одномерным, двумерным или трехмерным индексом, доступ к которому в ядре можно получить через переменную blockldx. Размерность блока можно узнать из переменной blockDim.
Блоки должны быть независимы: они могут запускаться в любом
порядке, параллельно или последовательно. Это необходимо для того, чтобы каждый блок мог быть запланирован на любом процессорном ядре, не зависимо от их количества; это позволяет писать код, который масштабируется по мере увеличения количества ядер.
Иерархия памяти
Потоки CUDA могут иметь доступ к различным пространствам памяти во время выполнения. Каждый поток имеет свое локальное пространство памяти. Каждый блок потоков имеет разделенную память, которая видима всем потокам этого блока. Все потоки имеют доступ к единой глобальной памяти.
Существует еще два источника данных только для чтения -текстурная память и константная память. Глобальная, текстурная и константная памяти оптимизированы для различных действий. Текстурная память также позволяет производить фильтрацию данных для некоторых форматов.
Центральный процессор называется host (хост), а графический процессор называется device (девайс) в терминах CUDA. Без хоста невозможно функционирование параллельного приложения, так как сама программа по-прежнему выполняется, используя ЦП, и только лишь отдельные фрагменты выполняются на девайсе. Хост и девайс поддерживают свое собственное пространство в памяти - память хоста (host memory) и память девайся (device memory). Так что программа управляет глобальной, константной и текстурной памятью с помощью обращения к API.
Использование CUDA как средство повышения производительности МДЭ
Рассмотрим способ ускорения вычислений, связанных с Методом Дискретных Элементов, с помощью CUDA. Дана система частиц с определенными физическими параметрами, и помещена в некоторое начальное состояние. Используя вычислительные мощности ГП (графического процессора), необходимо переводить систему в новое состояние. Переход системы из одного состояние в новое и далее будем называть циклом системы.
Большинство частиц находятся на относительно большом расстоянии друг от друга. Для данной частицы лишь небольшое количество других частиц находится рядом с ней, а производить расчет сил с другими частицами нет необходимости. Для того чтобы отсекать эти частицы, пространство делится на ячейки одинакового размера, образуя при этом сетку. Каждая такая ячейка содержит объекты, которые в ней находятся. Так как для каждой частицы можно узнать, в какой ячейке она находится, то также легко можно узнать, какие частицы находятся вместе с ней.
Для простоты, количество ячеек в сетке одинаково по всем трем измерениям. Ячейка задается координатами и представляет собой куб со стороной /, которая определяется диаметром. Когда частица находится в
данной ячейке, то необходимо проверять только те частицы, которые находятся в 26 окружающих ее ячейках. Предлагаются следующие способы оптимизации доступа к глобальной памяти:
• Кэширование данных в разделенной памяти;
• Линейный доступ к данным.
Помимо глобальной памяти - основного хранилища данных во время вычислений - существует еще разделенная память, доступная всем потокам в пределах блока. Это память небольшого размера, но доступ к ней осуществляется гораздо быстрее. Встроенные средства синхронизации позволяют использовать разделенную память как кэш данных: каждый поток загружает немного данных из глобальной памяти в разделенную, затем все потоки используют только разделенную память. Это сокращает количество обращений к глобальной памяти и ведет к ускорению работы приложения в целом.
Глобальная память оптимизирована таким образом, чтобы максимизировать производительность линейного доступа к данным. Несколько запросов к данным объединяются в одну транзакцию, если они вмещаются в один сегмент, размер которого зависит от запрашиваемых данных.
Таким образом, если несколько потоков выбирают данные, которые расположены друг за другом, то эти данные выбираются за один раз. Если все данные расположены в случайном порядке, то количество транзакций увеличивается и это ведет к снижению производительности.
В свете вышесказанного, предлагается модифицированный метод хранения частиц в ячейках.
Для работы с технологией CUDA необходим графический процессор, который ее поддерживает. Все ГП компании NVIDIA, начиная с серии 8800, поддерживают данную технологию. Однако более новые модели поддерживают более новую версию CUDA 2.1, которая имеет расширенные возможности. В качестве графического процессора была выбрана видеокарта NVIDIA GeForce GTS 450. Данный ГП имеет относительно невысокую цену, приемлемую производительность и поддерживает последнюю версию технологии.
Также необходимо установить на компьютер CUDA Toolkit - средства разработчиков. Это пакет из программных продуктов, позволяющий создавать параллельные приложения. Он включает в себя компилятор, библиотеки времени выполнения, профайлер, а также исчерпывающую справочную информацию. CUDA Toolkit интегрируется в Microsoft Visual Studio, поэтому это приложение также необходимо. Использовалась версия 2010.
Само приложение является графическим. Трехмерная графика выводится на экран с помощью технологии DirectX 10. Для этого был использован DirectX SDK. Данный выбор был сделан в виду того, что автор
данной работы уже ознакомлен с технологией. В качестве альтернативы можно использовать OpenGL.
Приложение состоит из двух окон: окно консоли и главное графическое окно. В окне консоли выводится информация об ошибках и событиях.
В результате выполненной работы было создано приложение, позволяющее выполнять симуляцию частиц с помощью метода дискретных элементов, используя архитектуру параллельных вычислений CUDA. Было получено, что с помощью данной технологии производительность МДЭ можно увеличить на порядок, по сравнению с классическими решениями, использующими центральный процессор.
В архитектуру CUDA изначально включена возможность масштабирования, а благодаря доступности, данная технология претендует на широкое использование в самых различных сферах научной и повседневной деятельности.
Литература
1. Wikipedia.[Электронный ресурс], http://ru.wikipedia.org/wiki/ Метод_Дискретного_Элемента
2. CUDA. [Электронный ресурс], http://www.nvidia.ru/object/cuda_home_new_ru.html.
