|
Архитектура Аудит Военная наука Иностранные языки Медицина Металлургия Метрология Образование Политология Производство Психология Стандартизация Технологии |
|
|
Архитектура Аудит Военная наука Иностранные языки Медицина Металлургия Метрология Образование Политология Производство Психология Стандартизация Технологии |
Фильтрация данных дефектограмм на основе технологии « CUDA »Стр 1 из 7Следующая ⇒
МИНОБРНАУКИ РОССИИ
Федеральное государственное бюджетное образовательное учреждение высшего образования «Ярославский государственный университет им. П.Г. Демидова»
Кафедра теоретической информатики
Сдано на кафедру «____» ________ 2018 г. Заведующий кафедрой, д. ф.- м. н., профессор ________ В.А. Соколов
Выпускная квалификационная работа
Фильтрация данных дефектограмм на основе технологии « CUDA »
по направлению подготовки 02.03.02 Фундаментальная информатика и информационные технологии
Научный руководитель д. ф.- м. н., профессор ________Е.В. Кузьмин «___» ________ 2018 г.
Студент группы ИТ-41БО ________А.О. Шуманов «___» _________ 2018 г.
Ярославль 2018 г. Реферат Объем 51 с., 3 гл., 11 рис., 2 табл., 15 источников, 2 прил. Дефектограммы, фильтрация данных, технология Cuda, параллельные вычисления, эксперименты, скорость выполнения. Объектом исследования является применимость технологии Cuda к задаче фильтрации дефектограмм. Цель работы – проверка на практике преимущества использования параллельных вычислительных мощностей и нахождение оптимального способа задачи фильтрации больших объемов данных. В работе были рассмотрены варианты ускорения фильтрации данных, реализованы предложенные варианты и проведен сравнительный анализ. В процессе многочисленных тестов было доказано, что скорость работы алгоритма фильтрации дефектограмм с использованием технологии Cuda возростает от 50 до 150 раз.
СОДЕРЖАНИЕ
Введение. 4 1. Задача фильтрации. 8 1.1. Входные данные. 8 1.1. Алгоритм фильтрации. 10 1.2. Варианты реализации. 11 2. Технология CUDA.. 13 2.1. Введение. 13 2.2 Принципиальная разница между CPU и GPU.. 14 2.3 Вычислительная модель GPU.. 15 2.4 Модель памяти CUDA.. 18 2.5 Программирование CUDA.. 21 2.6 Разработка на расширенном Си. 25 3. Результаты экспериментов с технологией CUDA.. 29 3.1. Описание реализации. 29 3.2 Сравнение скорости выполнения. 34 Заключение. 36 Список используемой литературы.. 37 Приложение А. Листинг программы по технологии CUDA.. 39 Приложение. Б Листинг программы без технологии CUDA.. 44
Введение В Российской Федерации железнодорожный транспорт общего и необщего пользования является основополагающим в транспортной системе, оказывая влияющую роль в развитии экономики страны. Железнодорожный путь является важной составляющей транспортной системы РФ. Все элементы железнодорожного пути (верхнее строение - рельсы, стрелочные переводы, подрельсовое основание со скреплениями и балластная призма и нижнее строение - земляное полотно, водоотводные и искусственные сооружения) по прочности, устойчивости и состоянию должны обеспечивать безопасное и плавное движение поездов со скоростями, установленными на данном участке, при безусловном соблюдении требований безопасности движения. Одной из важных составляющих является своевременное обслуживание и удовлетворительное состояние железнодорожных рельсов. При непрерывной эксплуатации железнодорожных путей верхнее строение пути приходит в негодность. Рельсы в пути постоянно работают на сопротивление силам, возникающим от движущихся колес поездов и на температурные напряжения. В результате рельсы подвергаются излому. В весенний период резкого колебания температур повышается хрупкость металла рельсов (т.е. снижается ударная вязкость металла). Как правило, именно весной развиваются такие дефекты как расслоение, трещина и выкрашивание шейки рельса, разлом и коррозия подошвы рельса, которые под давлением подвижного состава приводят к излому рельса и падению вагонов. Особого наблюдения также требуют кривые участки пути, участки пути в случае прохода вагонов с дефектными колесами (с большими ползунами и выбоинами), если произошло буксование или движение вагонов юзом. Также, немаловажным фактором является инспекторский контроль. В 2012 году вступили в силу новые Правила технической эксплуатации железных дорог РФ (ПТЭ) (утв. Приказом Минтранса России от 21.12.2010 г. №286). Согласно новым требованиям, все владельцы ж/д путей необщего пользования должны обеспечивать регулярный осмотр путей и регулярно в соответствии с утвержденной периодичностью проводить дефектоскопию рельсов и стрелочных переводов (см.[14]). Несоблюдение периодичности проведения дефектоскопии рельсов нарушает также требования основного закона «О железнодорожном транспорте» (Федер. закон от 10.01.2003 г. №17-ФЗ, п.1 ст.16) о содержании жд пути в исправном состоянии с обеспечением требований безопасности движения (см. [15]). В данном случае специалисты полагаются на дефектоскопию рельсов с помощью специальной техники, так как визуально определить в каком месте рельса появился дефект, невозможно. Последствием излома рельса в большинстве случаев является сход подвижного состава. И хуже всего, если им окажется, например, цистерна с опасным грузом. Главными преимуществами метода дефектоскопии являются: · отсутствие повреждений и нарушений целостности на исследуемом образце; · высокая достоверность при низкой цене; · возможность исследования в любое время года.
При визуальном методе контроля при всем желании можно выявить лишь наружные дефекты (сколы, выбоины, трещины, вмятины, коррозийные повреждения и пр.). Но известно, что в рельсах постоянно идёт развитие дефектов 1-й, 2-й и 5-й групп (так называемые, контактно-усталостные), обнаружить большую часть которых визуально невозможно, например, расположенные в местах болтовых соединений под стыковыми накладками (см [6]). Дефектоскопия рельса любого вида дает возможность выявлять на ранней стадии подавляющую часть опасных дефектов в районе головки, шейки и подошвы рельса, которые могут развиваться в нем со временем. В процессе контроля устанавливается код дефекта и его местоположение в рельсе. На сегодняшний день основными средствами дефектоскопии на ЖД транспорте являются электромагнитные и акустические дефектоскопы, устанавливаемые на съёмных тележках и на выгонах-дефектоскопах. В области дефектоскопии применительно к объектам ЖД транспорта широко известны работы российских ученых Гурвича А.К., Клюева В.В., Маркова А.А., Ермолова И.Н. К преимуществам электромагнитных методов можно отнести возможность ведения бесконтактного контроля в движении, однако малая глубина проникновения электромагнитного поля в металл не позволяет обнаруживать дефекты на глубине более 6-8мм. К достоинствам акустических методов (AM) контроля можно отнести высокую проникающую способность, что определяет их повсеместное использование для дефектоскопии рельсового пути и узлов ПС. Недостатком существующих AM является необходимость наличия физического контакта между пьезоэлектрическим преобразователем (ПЭП) и объектом контроля (ОК), что не позволяет создавать дистанционные средства дефектоскопии рельсов и узлов ПС в движении. Таким образом, проблема повышения эффективности средств дефектоскопии на ЖД транспорте продолжает оставаться актуальной и требует для своего решения не только усовершенствования существующих средств дефектоскопии, но и разработки принципиально новых методов неразрушающего контроля. В настоящее время решение данной проблемы многими исследователями в области дефектоскопии видится в создании бесконтактных AM НК, основанных на различных физических принципах возбуждения и регистрации акустических колебаний. Изучение технологии Cuda. Задача фильтрации Существуют различные подходы к фильтрации данных. Полученные наработки можно рассматривать для любого из них. Для начала мы определим входные данные нашей задачи, затем алгоритм фильтрации, на основе которого будем рассматривать возможность его ускорения. Входные данные В ходе дефектоскопии получается матрица значений от 0 до 255, значение соответствует амплитуде сигнала от отражающей поверхности, в местах трещин, либо на краях металла значение будет выше, а в однородном металле меньше, причем циклы зондирования происходят через каждые 5 мм пути. Пример матрицы представлен на (рис. 1). Рис. 1 Матрица входных данных Соответственно, массив значений получается очень большим, так как речь идет о сотнях километров путей. Поэтому актуальна задача быстрой обработки больших объемов данных. Обработка в данном случае необходима, так как большая часть дефектограммы заполнена сигналами с малой амплитудой, это могут быть фоновые шумы, незначительные электрические помехи, акустические шумы, загрязненность верхнего слоя железнодорожного полотна или некорректная настройка дефектоскопа. В зависимости от настройки оборудования дефект может быть зарегистрирован на любом из положительных уровней дефектограммы. Задача фильтрации состоит в выявлении сильных, то есть выделяющихся на общем фоне, наклонных связных пачек сигналов. При фильтрации кроме фоновых шумов и шумов типа «снег» должны быть учтены и удаление приконтактные шумовые дорожки, сигналы от выкружки рельса и электрические шумы (чаще всего возникающие при включении рации и имеющие вид частых длинных диагональных полос). На (рис. 2) показана дефектограмма, это так называемые “сырые” данные, полученные с прибора, на них можно четко заметить 2 горизонтальные линии.
Рис. 2 Дефектограмма с изображением болтового соединения
Верхняя линия – это отражение сигнала от головки рельса. Нижняя линия – это сигнал от подошвы рельса. Также можно рассмотреть различные шумы, а вот связная пачка наклонных сигналов – это пример болтового соединения. Оси координат имеют стандартное расположение, начальная точка находится в левом верхнем углу, ось X направлена вправо, ось Y вниз. На дефектограмме по оси Х отмечаются дискреты, каждое значение соответствует 5мм железнодорожного пути. По оси У отмечается время ответа после генерации сигнала в микросекундах. Чем темнее цвет, тем меньше амплитуда сигнала от отражающей поверхности. Таким образом, темный цвет соответствует однородному металлу, отражение от которого дает слабый сигнал (нулевой). Отражение от болтовых стыков, головки и подошвы рельса дает явно выделяющийся светлый цвет. Обычно входные данные разбиваются на файлы, в которых матрица имеет размер 180 строк и 10000 столбцов. Соответственно в каждом таком файле содержится информация о пятидесяти метрах железнодорожного полотна. Причем номер строки означает время отклика в микросекундах, что примерно соответствует расстоянию в миллиметрах. Быстрая обработка полученных данных поможет специалисту в нахождении дефектов, особенно когда нужно проверить многие километры железнодорожного полотна.
Алгоритм фильтрации Алгоритм фильтрации заключается в поэтапном проходе по матрице значений и выделении самых сильных сигналов. Формально можно записать следующим образом: 1) Вдоль каждой строки происходит проход малым “скользящим окном”, длина которого превышает в два-три раза типичную ширину наклонной пачки сигналов. При этом каждый раз в “скользящем окне” остаются наиболее сильные сигналы, остальные обнуляются. 2) Вдоль каждой строки происходит проход большим “скользящим окном”, длина которого превышает в 10 раз утроенную типичную ширину наклонной пачки сигналов. При этом каждый раз в “скользящем окне” остаются наиболее сильные сигналы, остальные обнуляются. 3) В каждом столбце матрицы сигналов остаются наиболее сильные сигналы таким образом, чтобы их суммарное количество не превышало утроенную максимальную поперечную высоту типичной наклонной пачки сигналов. Сигналы не удовлетворяющие этому условию обнуляются. 4) Фильтрация диагональной маской. По матрице сигналов осуществляется проход “скользящим окном” в виде специальной диагональной маски, которая оставляет только те сигналы, которые образуют связные наклонные линии. Остальные сигналы обнуляются (см. [7]).
Варианты реализации Прежде всего, можно заметить, что некоторые этапы алгоритма включается в себя множество одинаковых задач, которые можно выполнить независимо. Получается, что при большом объеме данных очень выгодно выполнять все задачи этапа параллельно, чтобы достигнуть наибольшей скорости. Рассмотрим варианты выполнения кода алгоритма, мы можем делать это на центральном процессоре компьютера (CPU), либо на графическом процессоре (GPU). Центральный процессор в среднем обладает от 2-8 ядрами, соответственно столько же параллельных задач можно запустить на нем. Тогда как графический процессор позволяет запустить миллионы параллельных задач одновременно. Таким образом, для ускорения выполнения задачи фильтрации подходит работа с выполнением программного кода на графическом процессоре. Рассмотрим это на примере технологии CUDA. Это достаточно современная технология для обработки больших объемов данных. Она активно применяется в целях машинного обучения, которое используется во всех поисковых системах для лучшего ранжирования данных. Технология представленная компанией Nvidia программно-аппаратная архитектура для расчётов на видеочипах CUDA неплохо подходит для решения многих задач, требующих большого количества параллельных расчётов. В плюсах у CUDA то, что она работает на всех современных решениях Nvidia, удобна в применении (по сравнению с другими моделями программирования GPGPU, но не CPU, тут пока есть место для улучшений) и уже довольно долгое время находится в разработке, что минимизирует количество недоработок и ошибок. CUDA использует и сама Nvidia, и одиночные исследователи в своих любительских проектах, и учебные заведения, и компании, производящие программно-аппаратные средства для множества сфер применения. Ведь CUDA — это технология, доступная любому разработчику ПО, неважно в какой сфере и компании он работает. Если применяемые в работе алгоритмы хорошо распараллеливаются в принципе, и GPU их ничем не ограничивает (точность вычислений, объём разделяемой памяти и т.п.), то нужно привыкнуть к иной парадигме программирования, присущей параллельным вычислениям, и тогда затраты времени на программирование на CUDA вернутся в виде хорошего прироста производительности.
Технология CUDA Введение CUDA, Compute Unified Device Architecture — программно-аппаратная архитектура, позволяющая производить вычисления с использованием графических процессоров NVIDIA, поддерживающих технологию GPGPU (произвольных вычислений на видеокартах). Впервые появились на рынке с выходом чипа NVIDIA восьмого поколения — G80 и присутствует во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce, Quadro и NVidia Tesla. CUDA SDK позволяет программистам реализовывать на специальном упрощённом диалекте языка программирования Си алгоритмы, выполнимые на графических процессорах NVIDIA, и включать специальные функции в текст программы на Cи. CUDA даёт разработчику возможность по своему усмотрению организовывать доступ к набору инструкций графического ускорителя и управлять его памятью, организовывать на нём сложные параллельные вычисления. Стоит отметить, что технология активно развивается в последнее время и находит применение в различных областях, где требуется работа с большими данными. На сегодняшний день продажи CUDA процессоров достигли миллионов, а разработчики программного обеспечения, ученые и исследователи широко используют CUDA в различных областях, включая обработку видео и изображений, вычислительную биологию и химию, моделирование динамики жидкостей, восстановление изображений, полученных путем компьютерной томографии, сейсмический анализ, трассировку лучей и многое другое.
Вычислительная модель GPU Рассмотрим вычислительную модель GPU более подробно: · является сопроцессором к CPU (называемому host); · обладает собственной памятью (DRAM); · обладает возможностью параллельного выполнения огромного количества отдельных нитей (threads).
При этом между нитями на CPU и нитями на GPU есть принципиальные различия - нити на GPU обладают крайне " небольшой стоимостью" - их создание и управление требует минимальных ресурсов (в отличии от CPU)для эффективной утилизации возможностей GPU нужно использовать многие тысячи отдельных нитей (для CPU обычно нужно не более 10-20 нитей). Сами программы пишутся на " расширенном" С, при этом их параллельная часть (ядра) выполняется на GPU, а обычная часть - на CPU. CUDA автоматически осуществляет разделением частей и управлением их запуском. CUDA использует большое число отдельных нитей для вычислений, часто каждому вычисляемому элементами соответствует одна нить. Все нити группируются в иерархию - grid/block/thread, как показано на (рис. 3).
Рис. 3 Иерархия потоков CUDA. Верхний уровень - grid - соответствует ядру и объединяет все нити выполняющие данное ядро. grid представляет собой одномерный или двухмерный массив блоков (block). Каждый блок (block) представляет из себя одно/двух/трехмерный массив нитей (threads). При этом каждый блок представляет собой полностью независимый набор взаимодействующих между собой нитей, нити из разных блоков не могут между собой взаимодействовать (см. [1]). Фактически блок соответствует независимо решаемой подзадаче, так например если нужно найти произведение двух матриц, то матрицу-результат можно разбить на отдельные подматрицы одинакового размера. Нахождение каждой такой подматрицы может происходить абсолютно независимо от нахождения остальных подматриц. Нахождение такой подматрицы - задача отдельного блока, внутри блока каждому элементу подматрицы соответствует отдельная нить. При этом нити внутри блока могут взаимодействовать между собой (т.е. совместно решать подзадачу) через общую память (shared memory)функцию синхронизации всех нитей блока (__synchronize. Подобная иерархия довольно естественна - с одной стороны хочется иметь возможность взаимодействия между отдельными нитями, а с другой - чем больше таких нитей, тем выше оказывается цена подобного взаимодействия. Поэтому исходная задача (применение ядра к входным данным) разбивается на ряд подзадач, каждая из которых решается абсолютно независимо (т.е. никакого взаимодействия между подзадачами нет) и в произвольном порядке. Сама же подзадача решается при помощи набора взаимодействующих между собой нитей. С аппаратной точки зрения все нити разбиваются на так называемые warp'ы - блоки подряд идущих нитей, которые одновременно (физически) выполняются и могут взаимодействовать друг с другом. Каждый блок нитей разбивается на несколько warp'ов, размер warp'а для всех существующих сейчас GPU равен 32 (см. [2]). Важным моментом является то, что нити фактически выполняют одну и ту же команды, но каждая со своими данными. Поэтому если внутри warp'а происходит ветвление (например в результате выполнения оператора if), то все нити warp'а выполняют все возникающие при этом ветви. Поэтому крайне желательно уменьшить ветвление в пределах каждого отдельного warp'а. Также используется понятие half-warp'а - это первая или вторая половина warp'а. Подобное разбиение warp'а на половины связано с тем, что обычно обращение к памяти делаются отдельно для каждого half-warp'а. Кроме иерархии нитей существует также несколько различных типов памяти. Быстродействие приложения очень сильно зависит от скорости работы с памятью. Именно поэтому в традиционных CPU большую часть кристалла занимают различные кэши, предназначенные для ускорения работы с памятью (в то время как для GPU основную часть кристалла занимают ALU).
2.4 Модель памяти CUDA Модель памяти в CUDA отличается возможностью побайтной адресации. Доступно довольно большое количество регистров на каждый потоковый процессор, до 1024 штук. Доступ к ним очень быстрый, хранить в них можно 32-битные целые или числа с плавающей точкой (см.[9]). Каждый поток имеет доступ к пяти типам памяти, все они изображены на (рис. 4).
Рис. 4 Типы памяти CUDA Глобальная память — самый большой объём памяти, доступный для всех мультипроцессоров на видеочипе, размер составляет от 256 мегабайт до 1.5 гигабайт на текущих решениях (и до 4 Гбайт на Tesla). Обладает высокой пропускной способностью, более 100 гигабайт/с для лучших решений Nvidia, но очень большими задержками в несколько сот тактов. Не кэшируется, поддерживает обобщённые инструкции load и store, и обычные указатели на память. Локальная память — это небольшой объём памяти, к которому имеет доступ только один потоковый процессор. Она относительно медленная — такая же, как и глобальная. Разделяемая память — это 16-килобайтный (в видеочипах нынешней архитектуры) блок памяти с общим доступом для всех потоковых процессоров в мультипроцессоре. Эта память весьма быстрая, такая же, как регистры. Она обеспечивает взаимодействие потоков, управляется разработчиком напрямую и имеет низкие задержки. Преимущества разделяемой памяти: использование в виде управляемого программистом кэша первого уровня, снижение задержек при доступе исполнительных блоков (ALU) к данным, сокращение количества обращений к глобальной памяти. Память констант — область памяти объемом 64 килобайта (то же — для нынешних GPU), доступная только для чтения всеми мультипроцессорами. Она кэшируется по 8 килобайт на каждый мультипроцессор. Довольно медленная — задержка в несколько сот тактов при отсутствии нужных данных в кэше. Текстурная память — блок памяти, доступный для чтения всеми мультипроцессорами. Выборка данных осуществляется при помощи текстурных блоков видеочипа, поэтому предоставляются возможности линейной интерполяции данных без дополнительных затрат. Кэшируется по 8 килобайт на каждый мультипроцессор. Медленная, как глобальная — сотни тактов задержки при отсутствии данных в кэше (см.[10]). Естественно, что глобальная, локальная, текстурная и память констант — это физически одна и та же память, известная как локальная видеопамять видеокарты. Их отличия в различных алгоритмах кэширования и моделях доступа. Центральный процессор может обновлять и запрашивать только внешнюю память: глобальную, константную и текстурную. Как изображено на (рис. 5).
Рис. 5 Обращение CPU к памяти GPU Из написанного выше понятно, что CUDA предполагает специальный подход к разработке, не совсем такой, как принят в программах для CPU. Нужно помнить о разных типах памяти, о том, что локальная и глобальная память не кэшируется и задержки при доступе к ней гораздо выше, чем у регистровой памяти, так как она физически находится в отдельных микросхемах. Типичный, но не обязательный шаблон решения задач: · задача разбивается на подзадачи; · входные данные делятся на блоки, которые вмещаются в разделяемую память; · каждый блок обрабатывается блоком потоков; · подблок подгружается в разделяемую память из глобальной; · над данными в разделяемой памяти проводятся соответствующие вычисления; · результаты копируются из разделяемой памяти обратно в глобальную.
Программирование CUDA При использовании традиционных API программист вне зависимости от сложности алгоритма всегда обязан конфигурировать все части графического конвейера. Этот факт существенно затрудняет использование GPU для решения задач общего назначения, так как даже простое сложение двух матриц требует выполнение ряда команд по подготовке и отрисовке изображений во внеэкранном буфере. В итоге на несколько строк шейдерной программы приходятся сотни строк дополнительного кода. При решении задач с небольшой размерностью эти дополнительные затраты способны свести на нет весь выигрыш от использования GPU (см.[8]). Модель программирования, используемая в CUDA отличается от традиционных API тем, что полностью скрывает графический конвейер от программиста, позволяя ему тем писать программы в более привычных для него «терминах» на расширенной вариации языка C. Кроме того, CUDA предоставляет программисту более удобную модель работы с памятью. Больше нет необходимости хранить данные в 128-битных текстурах, так как CUDA позволяет читать данные напрямую из памяти видеокарты. В состав NVIDIACUDA входят два API: высокого уровня (CUDA Runtime API) и низкого (CUDA Driver API) смотри (рис. 6).
Рис. 6 Схема CUDA API При необходимости задействовать низкоуровневые функции графического процессора программист всегда может отказаться от Runtime API в пользу Driver API (см.[12]). Первым шагом при переносе существующего алгоритма на CUDA непременно является его анализ, цель которого состоит в поиске «узкого места», нуждающегося в распараллеливании. Как правило, в алгоритме для CPU такие участки заключены в цикл или рекурсию. Полный перенос алгоритма на GPU не является возможным, так как графический процессор не имеет доступа ни к памяти компьютера, ни к устройствам ввода/вывода (за исключением буфера кадра, который может быть отображен в виде картинки на экране компьютера). При исполнении программы CPU по-прежнему отвечает за подготовку и постобработку данных, сама же трудоемкая работа ложится на GPU. Набор инструкций, исполняемый на графическом процессоре, называется ядром (kernel). Ядро, по сути, является развитием концепции шейдеров. За формирование и компиляцию ядер отвечает CPU. Видеочип просто принимает уже скомпилированное ядро и создает его копии для каждого элемента данных. Каждое из ядер исполняется в своем собственном потоке. Потоки в GPU могут исполняться лишь группами по 32 экземпляра (wrap). При этом общее число потоков необходимое для решения задачи может превосходить максимально допустимое для текущей видеокарты. Поэтому каждый такт аппаратное обеспечение выбирает, какой из wrap будет исполнен. Но если бы в CPU подобное переключение заняло бы сотни тактов, то GPU делает это почти мгновенно. В отличие от шейдеров, где все данные представлены в виде четырех компонентных векторов, данные в ядре скалярных. Такое представление более естественно для большинства неграфических задач. Модель программирования CUDA предполагает группировку потоков блоки–одно-, двух-или трехмерные матрицы. Взаимодействие между ними осуществляется при помощи разделяемой памяти. Также существуют точки синхронизации, позволяющие привести данные во всех потоках в актуальное состояние. Каждое из ядер исполняется над сеткой (grid) блоков. В каждый момент времени на GPU может исполняться лишь одна сетка. Подобная группировка позволяет достичь максимальной масштабируемости. Если у GPU недостаточно ресурсов для запуска всех блоков – они будут выполняться последовательно, друг за другом. Это позволяет разработчику не задумываться о мощности устройства, на котором будет запущено приложение. В каждом приложении, написанном на NVIDIACUDA вне зависимости от его назначения можно выделить ряд общих шагов: Подготовка памяти. ПосколькуGPU не имеет доступа к оперативной памяти программисту необходимо заранее позаботиться о том, что все ресурсы, необходимые для исполнения ядра приложения находятся в памяти видеокарты. Для этих целей используются три основных функции из CUDASDK: cudaMalloc, cudaMemcpy и cudaFree. Эти функции имеют то же назначение, что и стандартные malloc, memcpy и free, но, разумеется, все операции проводятся над видеопамятью. Так же стоит отметить, что функция cudaMemcpy имеет дополнительный параметр, обозначающий направление копирования информации (из CPU в GPU или наоборот). Конфигурация сетки ( grid ) и блоков ( blocks ). Процесс конфигурации крайне прост и заключается в задании размеров сетки и блоков. Основной задачей программиста на данном шаге является нахождение оптимального баланса между размером и количеством блоков. Увеличением количества потоков в блоке можно снизить количество обращений к глобальной памяти за счет увеличения интенсивности обмена данными между потоками через быструю разделяемую память. С другой стороны, количество регистров выделяемых на блок фиксировано и если количество потоков окажется сильно большим, то GPU начнет размещать данные в медленной локальной памяти, что существенно увеличит время исполнения ядра. NVIDIA рекомендует программистам использовать блоки по 128 или 256 потоков. В большинстве задач такое количество потоков в блоке позволяет достичь оптимальных задержек и количества регистров. Запуск ядра. Ядро вызывается как обычная функция в языке C. Единственное существенное отличие заключается в том, что при вызове ядра необходимо передать ранее определенные размерности сетки и блока. Получение результатов и освобождение памяти. После исполнения ядра необходимо скопировать результаты выполнения программы назад, в оперативную память при помощи функции cudaMemcpy с указанием обратного направления копирования (из GPU в CPU). Точно также, как и в любой C-программе для предотвращения утечек памяти необходимо освободить все выделенные ресурсы.
Спецификаторы доступа CUDA
На функции, выполняемые на GPU (__device__ и __global__) накладываются следующие ограничения: · нельзя брать их адрес (за исключением __global__ функций); · не поддерживается рекурсия; · не поддерживаются static-переменные внутри функции; · не поддерживается переменное число входных аргументов.
Для задания размещения в памяти GPU переменных используются следующие спецификаторы - __device__, __constant__ и __shared__. На их использование также накладывается ряд ограничений: · эти спецификаторы не могут быть применены к полям структуры (struct или union); · соответствующие переменные могут использоваться только в пределах одного файла, их нельзя объявлять как extern; · запись в переменные типа __constant__ может осуществляться только CPU при помощи специальных функций; · __shared__ переменные не могут инициализироваться при объявлении.
В язык добавлены следующие специальные переменные: · gridDim - размер grid'а (имеет тип dim3); · blockDim - размер блока (имеет тип dim3); · blockIdx - индекс текущего блока в grid'е (имеет тип uint3); · threadIdx - индекс текущей нити в блоке (имеет тип uint3); · warpSize - размер warp'а (имеет тип int).
В язык добавляются 1/2/3/4-мерные вектора из базовых типов - char1, char2, char3, char4, uchar1, uchar2, uchar3, uchar4, short1, short2, short3, short4, ushort1, ushort2, ushort3, ushort4, int1, int2, int3, int4, uint1, uint2, uint3, uint4, long1, long2, long3, long4, ulong1, ulong2, ulong3, ulong4, float1, float2, float3, float2, и double2. Обращение к компонентам вектора идет по именам - x, y, z и w. Для создания значений-векторов заданного типа служит конструкция вида make_< typeName>. Также для задания размерности служит тип dim3, основанный на типе uint3, но обладающий нормальным конструктором, инициализирующим все не заданные компоненты единицами. Для запуска ядра на GPU используется следующая конструкция: kernelName < < < Dg, Db, Ns, S> > > ( args ) Здесь kernelName это имя (адрес) соответствующей __global__ функции, Dg - переменная (или значение) типа dim3, задающая размерность и размер grid'a (в блоках), Db - переменная (или значение) типа dim3, задающая размерность и размер блока (в нитях), Ns - переменная (или значение) типа size_t, задающая дополнительный объем shared-памяти, которая должна быть динамически выделена (к уже статически выделенной shared-памяти), S - переменная (или значение) типа cudaStream_t задает поток (CUDA stream), в котором должен произойти вызов, по умолчанию используется поток 0. Через args обозначены аргументы вызова функции kernelName. Также в язык С добавлена функция __syncthreads, осуществляющая синхронизацию всех нитей блока. Управление из нее будет возвращено только тогда, когда все нити данного блока вызовут эту функцию. Т.е. когда весь код, идущий перед этим вызовом, уже выполнен (и, значит, на его результаты можно смело рассчитывать). Эта функция очень удобная для организации безконфликтной работы с shared-памятью. Также CUDA поддерживает все математические функции из стандартной библиотеки С, однако с точки зрения быстродействия лучше использовать их float-аналоги (а не double) - например sinf. Кроме этого CUDA предоставляет дополнительный набор математических функций (__sinf, __powf и т.д.) обеспечивающие более низкую точность, но заметно более высокое быстродействие чем sinf, powf и т.п. Таким образом, для использования технологии CUDA достаточно знать перечисленные особенности и можно использовать любой удобный для вас редактор кода, стоит отметить, что в Visual Studio встроена поддержка данной технологии, которая позволяет легко использовать все описанные возможности. На данный момент самой новой версией является CUDA Runtime API 8.0. нужно сказать, что ведется разработка в сторону поддержки все большего числа языков программирования, например для Java существует JCUDA.
Описание реализации Поставленная задача фильтрации была реализована двумя способами. Первый заключался в стандартном кодировании алгоритма на языке c++. Второй базируется на использовании технологии CUDA и кодировании на расширенном языке c++ специально для этой технологии. Особенность реализации на технологии CUDA заключается в том, что каждая подзадача алгоритма выполняется отдельным потоком на видеокарте. При этом все потоки асинхронны и не зависимы между собой. Каждая отдельная нить обращается к своему отдельному блоку памяти. Программа читает входные данные из файла, это описано в приложении Б, размер данных задается заранее. Также по выбору запускается фильтрация на процессоре, видеокарте, либо на том и другом с замером скорости работы. После данные сохраняются в файл, для каждого выбора отдельно. Программа с использованием технологии CUDA описана в приложении А, нужные методы для работы импортируются и запускаются из обычного.cpp файла. Поэтому переход на технологию CUDA практически не заметен. Рассмотрим последовательность действий алгоритма, все данные визуализированы в программе matlab. Для иллюстрации работы алгоритма рассмотрим фрагмент дефектограммы на (рис. 7).
Рис. 7 Фрагмент исходных данных На нем представлен фрагмент полезного сигнала. Как было описано ранее, темно синий цвет соответствует сигналу с слабой амплитудой (нулевой), более светлый оттенок синего это слабые шумы, а яркий желтый сигнал с сильной амплитудой от болтовых стыков, головки и подошвы рельса. На первом этапе прохода по строкам матрицы удаляются слабые сигналы в подстроке размером 20 значений. Матрица в gpu разбивается на 180 блоков по количеству строк. Затем внутри каждого блока вся его длина разбивается по 20 значений. Получается одновременно 90000 потоков, каждый из которых работает со своей выборкой данных. Выборка не пересекается. Результат показан на (рис. 8).
Рис. 8 Дефектограмма после 1го шага алгоритма На втором этапе прохода по строкам матрицы удаляются слабые сигналы в подстроке размером 200 значений. В этом случае разбиение схоже с первым этапом, только внутри каждого блока выборка уже больше, поэтому в целом получается 9000 потоков. Результат показан на (рис. 9).
Рис. 9 Дефектограмма после 2го шага алгоритма На третьем этапе проход осуществляется по столбцам матрицы, при этом удаляются слабые сигналы в массиве размером 50 значений. Матрица разбивается по числу столбцов на 10000 блоков. Затем внутри каждый блок делится на 4. Итого получается 40000 потоков, где каждый работает со своей частью матрицы. Выборка не пересекается. Результат показан на (рис. 10).
Рис. 10 Дефектограмма после 3го шага алгоритма На четвертом этапе проход осуществляется диагональной матрицей, при этом удаляются слабые сигналы, которые не образуют связной наклонной линии. Матрица внутри gpu разбивается на 500 блоков, размер каждого блока 180 строк и 20 столбцов. Результат показан на (рис. 11).
Рис. 11 Дефектограмма после 4го шага алгоритма На финальном этапе алгоритма остаются наиболее выраженные связные наборы сигналов. Результат фильтрации на CPU совпадает с представленными данными, разница заключается в скорости работы алгоритма. Результаты замеров
В таблице 2 приведены средние результаты, по итогу 100 тестов. Из результатов видно, что на малых значения технология CUDA дает преимущество в 50 раз, а при увеличении размера матрицы, это значение возрастает до 150 раз и больше. Это можно объяснить тем, что для выполнения вычислений на видеокарте требуется их подготовка, копирование в память графического процессора и обратно, поэтому при малых размерах исходных данных, преимущество параллельных вычислений может быть менее заметным. Таким образом, на довольно стандартном размере матрицы 180 х 10 000 в задачах фильтрации, работа с видеокартой позволяет значительно ускорить время работы алгоритма. Также стоит отметить, что исходя из возможностей видеокарты это не предельный размер для работы с технологией CUDA. Так как размер данных не большой, то можно заметить, что потери на подготовку данных для GPU малы в сравнении с общим временем выполнения. В итоге получается матрица, которая представляет собой готовый объект для изучения на предмет выявления дефектов. Большинство шумов исчезло, остались лишь связные наклонные пачки сигналов с наибольшей амплитудой, которые дают объективную картину состояния исследуемого объекта.
Заключение В ходе работы была показана актуальность обработки больших объемов данных с целью выявления дефектов на железнодорожных путях. Изучена современная технология CUDA для решения поставленной задачи фильтрации. Предложен алгоритм и методы его оптимизации. На основе полученных тестов было продемонстрировано, что работа с видеокартой может значительно ускорить выполнение алгоритма за счет большого количества параллельных потоков. В целом, выполнение задач с использованием видеокарты в современном информационном обществе находит все больше применения. Работа с большими данными требует больших вычислительных мощностей, которые может обеспечить видеокарта, либо несколько процессоров. На основе полученных тестов можно убедиться, что видеокарта дает заметное преимущество в обработке значительного числа информации. Также стоит заметить, что технология CUDA хорошо справляется с большим объемом арифметических операции, но менее применима к структурам, где требуется множество ветвлений.
Список литератур ы 1. Боресков А.В. Параллельные вычисления на GPU. Архитектура и программная модель CUDA: Учебное пособие. Издательство Московского университета, 2012. 86с. 2. Боресков А.В., Харламов А.А. Основы работы с технологией CUDA. 2010. 40 - 56с. 3. Сандерс Д., Кэндрот Эю. NVIDIA CUDA в примерах: введение в программирование графических процессоров – ДМК Пресс -2011. 4. Седжвик Р. Фундаментальные алгоритмы на C++. М., СПб., Киев: DiaSoft, 2001. 152с. 5. Перепёлкин Е.Е., Садовников Б.И., Иноземцева Н.Г. Вычисления на графических процессорах (GPU) в задачах математической и теоретической физики. Издательство: URSS, 2014. 15с. 6. Марков А.А., Кузнецова Е.А. Дефектоскопия рельсов. Формирование и анализ сигналов. Санкт-Петербург, Ультра Принт, 2014. 212с. 7. Тарабрин В.Ф., Зверев А. В., Горбунов О.Е., Кузьмин Е.В. О фильтрации данных..., В мире неразрушающего контроля 2 (64) июнь 2014, 5 - 9 с. 8. Официальная документация по технологии CUDA: http: //www.nvidia.ru. 1веб-страница. URL: http: //www.nvidia.ru/object/cuda_home_new_ru.html. 9. Описание технологии CUDA: http: //bourabai.kz. 1веб-страница. URL: http: //bourabai.kz/graphics/cuda.htm. 10. Неграфические вычисления CUDA: https: //www.ixbt.com. 1веб-страница. URL: https: //www.ixbt.com/video3/cuda-1.shtml. 11. Параллельные вычисления: https: //studwood.ru. 1веб-страница. URL: https: //studwood.ru/1622722/informatika/tehnologiya_nvidia_cuda. 12. Cuda - Основные примеры: http: //www.ksmlab.ru. 1веб-страница. URL: http: //www.ksmlab.ru/page/cuda-osnovy-primery. 13. Дорофеев А.Л., Кузаманов Ю.Г. Электромагнитная дефектоскопия. — 2-е изд., доп. и перераб. -М.: Машиностроение, 1980. —232 е.: ил. 14. Горделий В.И. Состояние и перспективы развития средств неразрушающего контроля рельсов // Нефть, газ, промышленность №6, 2005. 15. Правила технической эксплуатации железных дорог Российской Федерации. - министерство путей сообщения Российской Федерации. — М.: «Техинформ», - 2000г. 114 - 190 с. Приложение А. Листинг программы по технологии CUDA В данном файле представлен код, который выполняется на видеокарте, его можно отличить по приставке __device__ или __global__. Остальные же методы служат для копирования данных на видеокарту и обратно, выделение памяти, разметки процессов и непосредственно запуска ядра видеокарты.
Файл kernel. cu // Библиотеки для работы с CUDA Runtime API 8.0 #include " cuda_runtime.h" #include " device_launch_parameters.h"
#include < stdio.h> #include < iostream>
/* * Проверка на принадлежность элемента к массиву, метод выполняется на видеокарте */ __device__ bool contains(unsigned short int* indexes, int n, unsigned short index) { for (int i = 0; i < n; ++i) { if (indexes[i] == index) { return true; } } return false; }
/* * Фильтрация данных по столбцам, метод запускается на видеокарте */ __global__ void verticalFilteringWithSearching(unsigned short int* dev, int const n, int const height, int const sizeOfMask, int const keepCount) { int const startIndex = threadIdx.x + threadIdx.y * sizeOfMask + blockDim.x * blockIdx.x;
unsigned short int* maxIndexes = new unsigned short[keepCount];
for (int i = 0; i < keepCount; i++) maxIndexes[i] = -1;
for (int k = 0; k < keepCount; ++k) { short int maxValue = -1; int maxIndex = -1; for (int i = startIndex; i < startIndex + sizeOfMask; i += n) { unsigned short value = dev[i]; if (value > = maxValue & & ! contains(maxIndexes, keepCount, i)) { maxValue = value; maxIndex = i; } } maxIndexes[k] = maxIndex; }
for (int i = startIndex; i < startIndex + sizeOfMask; i += n) { if (! contains(maxIndexes, keepCount, i)) dev[i] = 0; } }
/* * Фильтрация данных по строкам, метод запускается на видеокарте */ __global__ void horizontalFilteringWithSearching(unsigned short int* dev, int const n, int const sizeOfMask, int const keepCount) { int const startIndex = n * blockIdx.y + sizeOfMask * threadIdx.x;
unsigned short int* maxIndexes = new unsigned short[keepCount];
for (int i = 0; i < keepCount; i++) maxIndexes[i] = -1;
for (int k = 0; k < keepCount; ++k) { short int maxValue = -1; int maxIndex = -1; for (int i = startIndex; i < sizeOfMask + startIndex; ++i) { unsigned short value = dev[i]; if (value > = maxValue & & ! contains(maxIndexes, keepCount, i)) { maxValue = value; maxIndex = i; } } maxIndexes[k] = maxIndex; }
for (int i = startIndex; i < startIndex + sizeOfMask; ++i) { if (! contains(maxIndexes, keepCount, i)) dev[i] = 0; } } /* * Метод фильтрации по строкам с использование технологии CUDA */ void cudaHorizontalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount) { cudaError_t cudaStatus; unsigned short int* dev = nullptr; int numBytes = n * height * sizeof(short int); if (n % sizeOfMask) { std:: cout < < " WARNING: size of mask is not the multiplicity for size line"; } std:: cout < < " Start cuda horizontal filter\n";
cudaStatus = cudaSetDevice(0); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaSetDevice failed! Do you have a CUDA-capable GPU installed? " ); } cudaStatus = cudaMalloc((void**)& dev, numBytes); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaMalloc failed! " ); }
dim3 threads = dim3(n / sizeOfMask, 1); dim3 blocks = dim3(1, height);
cudaMemcpy(dev, data, numBytes, cudaMemcpyHostToDevice);
horizontalFilteringWithSearching< < < blocks, threads > > > (dev, n, sizeOfMask, keepCount);
cudaMemcpy(data, dev, numBytes, cudaMemcpyDeviceToHost);
cudaFree(dev); }
/* * Метод фильтрации по столбцам с использование технологии CUDA */ void cudaVerticalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount) { unsigned short int* dev = nullptr; int numBytes = n * height * sizeof(short int); if (height % sizeOfMask) { std:: cout < < " WARNING: size of mask is not the multiplicity for size line"; } std:: cout < < " Start cuda vertical filter\n";
cudaMalloc((void**)& dev, numBytes);
dim3 threads = dim3(n / 100, height / sizeOfMask); dim3 blocks = dim3(100, 1);
cudaMemcpy(dev, data, numBytes, cudaMemcpyHostToDevice); verticalFilteringWithSearching < < < blocks, threads > > > (dev, n, height, sizeOfMask*n, keepCount);
cudaMemcpy(data, dev, numBytes, cudaMemcpyDeviceToHost);
cudaFree(dev); }
/* * Фильтрация данных по диагонали, метод запускается на видеокарте */ __global__ void diagonalFilteringWithSearching(unsigned short int* dev, int const n, int const sizeOfMask, int const keepCount) { int const startIndex = sizeOfMask * threadIdx.x;
unsigned short int* maxIndexes = new unsigned short[keepCount];
for (int i = 0; i < keepCount; i++) maxIndexes[i] = -1;
int maxIndex = -1; for (int k = 0; k < keepCount; ++k) { short int maxValue = -1; for (int i = startIndex; i < sizeOfMask + startIndex; ++i) { unsigned short value = dev[i]; if (value > = maxValue & & ! contains(maxIndexes, keepCount, i)) { maxValue = value; maxIndex = i; } } if (startIndex + sizeOfMask + 1 < n) { int maxIndexNeighbor = -1; short int maxValueNeighbor = -1; for (int i = startIndex; i < sizeOfMask + startIndex; ++i) { unsigned short value = dev[i]; if (value > = maxValueNeighbor & & ! contains(maxIndexes, keepCount, i)) { maxValueNeighbor = value; maxIndexNeighbor = i; } } if (abs(maxIndexNeighbor - maxIndex) < 10) { maxIndexes[k] = maxIndex; } } else { maxIndexes[k] = maxIndex; } }
for (int i = startIndex; i < startIndex + sizeOfMask; ++i) { if (! contains(maxIndexes, keepCount, i)) dev[i] = 0; } }
/* * Метод фильтрации по диагонали с использование технологии CUDA */ void cudaDiagonalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount) { cudaError_t cudaStatus; unsigned short int* dev = nullptr; int numBytes = n * height * sizeof(short int); if (n % sizeOfMask) { std:: cout < < " WARNING: size of mask is not the multiplicity for size line"; } std:: cout < < " Start cuda diagonal filter\n";
cudaStatus = cudaSetDevice(0); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaSetDevice failed! Do you have a CUDA-capable GPU installed? " ); }
cudaStatus = cudaMalloc((void**)& dev, numBytes); if (cudaStatus! = cudaSuccess) { fprintf(stderr, " cudaMalloc failed! " ); }
dim3 threads = dim3(n / sizeOfMask, 1); dim3 blocks = dim3(1, 1);
cudaMemcpy(dev, data, numBytes, cudaMemcpyHostToDevice);
diagonalFilteringWithSearching < < < blocks, threads > > > (dev, n, sizeOfMask, keepCount);
cudaMemcpy(data, dev, numBytes, cudaMemcpyDeviceToHost);
cudaFree(dev); }
/* * Получение информации о всех видеокартах компьютера */ void getInfo() { int deviceCount; cudaDeviceProp devProp;
cudaGetDeviceCount(& deviceCount);
printf(" Found %d devices\n", deviceCount);
for (int device = 0; device < deviceCount; device++) { cudaGetDeviceProperties(& devProp, device);
printf(" Device %d\n", device); printf(" Compute capability: %d.%d\n", devProp.major, devProp.minor); printf(" Name : %s\n", devProp.name); printf(" Total Global Memory: %d\n", devProp.totalGlobalMem); printf(" Shared memory per block: %d\n", devProp.sharedMemPerBlock); printf(" Registers per block: %d\n", devProp.regsPerBlock); printf(" Warp size : %d\n", devProp.warpSize); printf(" Max threads per block: %d\n", devProp.maxThreadsPerBlock); printf(" Total constant memory: %d\n", devProp.totalConstMem);
} }
Приложение Б. Листинг программы без технологии CUDA В данном файле представлен код, который считывает исходную матрицу значений, затем запускает фильтрацию на процессоре и видеокарте с замерами времени соответственно. Размер исходные данные задается заранее и не изменяется в процессе фильтрации. Методы для работы с технологией CUDA импортируются как внешние зависимости, поэтому компилируются отдельно от основного файла. После выполнения фильтрации данные сохраняются в файл, для каждого способа отдельно. Файл Main.cpp #include < iostream> #include < fstream> #include < sstream> #include < iterator> #include < vector> #include < ctime> #include < iomanip> #include < thread> #include < chrono>
// Подключение функций работающих с CUDA extern void cudaHorizontalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount); extern void cudaVerticalFilter(short int* data, int const n, int const height, int const sizeOfMask, int const keepCount); extern void diagonalFilteringWithSearching(unsigned short int* dev, int const n, int const sizeOfMask, int const keepCount); extern void getInfo();
// Имя файла для чтения данных std:: string const pathToData = " smp1.txt"; int const countLines = 180; int const countColumns = 10000;
// Переменные для выбора способа фильтрации std:: string onlyCPU = " cpu"; std:: string onlyGPU = " gpu"; std:: string any = " any";
int getRightSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount);
int getDownSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount);
std:: vector< short int> getKeepIndexesHorizontal(int keepCount, int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data);
std:: vector< short int> getKeepIndexesVertical(int keepCount, int i, int j, int downSide, std:: vector< std:: vector< short int> > & data); std:: vector< short int> getKeepIndexesDiagonal(int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data); /* * Получение данных из файла */ std:: vector< std:: vector< short int> > getDatafromFile() { std:: vector< std:: vector< short int> > data; std:: ifstream reader(pathToData, std:: ios_base:: in); std:: string str; if (! reader.is_open()) { std:: cerr < < " Error: file " < < pathToData < < " is not open\n"; return data; } int counter = 0; while (getline(reader, str) & & counter! = countLines) { counter++; std:: istringstream iss(str); std:: vector< std:: string> results((std:: istream_iterator< std:: string> (iss)), std:: istream_iterator< std:: string> ()); std:: vector< short int> dataLine; if (results.size() < countColumns) std:: cerr < < " Error: size of columns is less then need\n";
for (int i = 0; i < countColumns; ++i) { dataLine.push_back((short int)std:: stoi(results[i])); } data.push_back(dataLine); } reader.close(); return data; }
/* * Сохранение данных после фильтрации на процессоре */ void saveDataToFile(std:: string filename, std:: vector< std:: vector< short int> > data) { std:: ofstream file(filename, std:: ios:: out); if (! file.is_open()) { std:: cerr < < " Error: file " < < filename < < " is not open\n"; return; }
for (int i = 0; i < data.size(); ++i) { for (int j = 0; j < data[0].size(); ++j) { file < < std:: setw(4) < < data[i][j] < < " "; } file < < std:: endl; }
file.close(); }
/* * Сохранение данных после фильтрации на видеокарте */ void saveDataToFile(std:: string filename, short int* data, int n, int height) { std:: ofstream file(filename, std:: ios:: out); if (! file.is_open()) { std:: cerr < < " Error: file " < < filename < < " is not open\n"; return; }
for (int i = 0; i < height; ++i) { for (int j = 0; j < n; ++j) { file < < std:: setw(4) < < data[i * n + j] < < " "; } file < < " \n"; }
file.close(); }
/* * Проверка на принадлежность элемента к массиву */ bool contains(std:: vector< short int> a, int source) { for (int i = 0; i < a.size(); ++i) { if (a[i] == source) { return true; } } return false; }
/* * Фильтрация данных по столбцам на процессоре */ void verticalFiltering(std:: vector< std:: vector< short int> > & data, int sizeOfMask, int keepCount) { for (int i = 0; i < data[0].size(); ++i) { for (int j = 0; j < data.size(); j += sizeOfMask) { int downSide = getDownSide(sizeOfMask, i, j, data, keepCount); std:: vector< short int> keepsIndexes = getKeepIndexesVertical(keepCount, i, j, downSide, data);
for (int m = j; m < downSide; ++m) { if (! contains(keepsIndexes, m)) { data[m][i] = 0; } } } } }
/* * Фильтрация данных по строкам на процессоре */ void horizontalFiltering(std:: vector< std:: vector< short int> > & data, int sizeOfMask, int keepCount) { for (int i = 0; i < data.size(); ++i) { for (int j = 0; j < data[i].size(); j += sizeOfMask) { int rightSide = getRightSide(sizeOfMask, i, j, data, keepCount); std:: vector< short int> keepsIndexes = getKeepIndexesHorizontal(keepCount, i, j, rightSide, data);
for (int m = j; m < rightSide; ++m) { if (! contains(keepsIndexes, m)) { data[i][m] = 0; } } } } }
/* * Фильтрация данных по диагонали на процессоре */ void diagonalFiltering(std:: vector< std:: vector< short int> > & data, int sizeOfMask, int minSignal) { for (int i = 0; i < data.size(); ++i) { for (int j = 0; j < data[i].size(); j += sizeOfMask) { int notUsed = 0; int rightSide = getRightSide(sizeOfMask, i, j, data, notUsed); std:: vector< short int> keepsIndexes = getKeepIndexesDiagonal(i, j, rightSide, data);
for (int m = j; m < rightSide; ++m) { if (! contains(keepsIndexes, m)) { data[i][m] = 0; } } } } }
/* * Получение списка индексов сигналов в строке, которые нужно сохранить. Остальные обнуляются. */ std:: vector< short int> getKeepIndexesHorizontal(int keepCount, int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data) { std:: vector< short int> keepsIndexes; for (int l = 0; l < keepCount; ++l) { int max = -1; int indexOfMax = 0; for (int k = j; k < rightSide; ++k) { if (data[i][k] > = max & & ! contains(keepsIndexes, k)) { max = data[i][k]; indexOfMax = k; } } keepsIndexes.push_back(indexOfMax); } return keepsIndexes; }
/* * Получение списка индексов сигналов в столбце, которые нужно сохранить. Остальные обнуляются. */ std:: vector< short int> getKeepIndexesVertical(int keepCount, int i, int j, int downSide, std:: vector< std:: vector< short int> > & data) { std:: vector< short int> keepsIndexes; for (int l = 0; l < keepCount; ++l) { int max = -1; int indexOfMax = 0; for (int k = j; k < downSide; ++k) { if (data[k][i] > = max & & ! contains(keepsIndexes, k)) { max = data[k][i]; indexOfMax = k; } } keepsIndexes.push_back(indexOfMax); } return keepsIndexes; }
/* * Получение списка индексов сигналов в строке, которые нужно сохранить. Остальные обнуляются. */ std:: vector< short int> getKeepIndexesDiagonal(int i, int j, int rightSide, std:: vector< std:: vector< short int> > & data) { std:: vector< short int> keepsIndexes; int indexOfMax = 0; for (int l = 0; l < 1; ++l) { int max = -1; for (int k = j; k < rightSide; ++k) { if (data[i][k] > = max & & ! contains(keepsIndexes, k)) { max = data[i][k]; indexOfMax = k; } } } if (data[i].size() > j + 1) { int indexOfMaxNeighbor = 0; for (int k = j; k < rightSide; ++k) { int max = -1; if (data[i][k] > = max & & ! contains(keepsIndexes, k)) { max = data[i][k]; indexOfMaxNeighbor = k; } } if (abs(indexOfMaxNeighbor - indexOfMax) < 10) { keepsIndexes.push_back(indexOfMax); } } else { keepsIndexes.push_back(indexOfMax); } return keepsIndexes; }
/* * получение правой границы маски в фильтрации на процессоре */ int getRightSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount) { int rightSide; if (j + sizeOfMask < data[i].size()) { rightSide = j + sizeOfMask; if (rightSide + sizeOfMask > data[i].size()) { keepCount *= (int)(sizeOfMask + data[i].size() - rightSide) / sizeOfMask; rightSide = (int)data[i].size(); } } else { rightSide = (int)data[i].size(); } return rightSide; }
/* * Получение нижней границы маски в филтрации на процессоре */ int getDownSide(int sizeOfMask, int i, int j, std:: vector< std:: vector< short int> > & data, int& keepCount) { int downSide; if (j + sizeOfMask < data.size()) { downSide = j + sizeOfMask; if (downSide + sizeOfMask > data.size()) { keepCount *= (int)(sizeOfMask + data.size() - downSide) / sizeOfMask; downSide = (int)data.size(); } } else { downSide = (int)data.size(); } return downSide; }
/* * Фильтрация с использование технологии CUDA */ void cudaHandling(short int* data) { cudaHorizontalFilter(data, countColumns, countLines, 20, 5); cudaHorizontalFilter(data, countColumns, countLines, 200, 20); cudaVerticalFilter(data, countColumns, countLines, 20, 5); cudaDiagonalFilter(data, countColumns, countLines, 20, 5); }
/* * Фильтрация данных с помощью процессора */ void cpuHandling(std:: vector< std:: vector< short int> > & data) { horizontalFiltering(data, 20, 5); horizontalFiltering(data, 200, 20); verticalFiltering(data, 20, 5); diagonalFiltering(data, 20, 5); }
/* * Старт программы, чтение данных из файла, запуск методов фильтрации с сохранение данных в файлы */ int main() { time_t start, stop; double timeCPU = 0, timeGPU = 0; std:: string mode = any; std:: cout < < " ------------START READING DATA------------------------------\n";
std:: vector< std:: vector< short int> > vec = getDatafromFile(); if (vec.size()! = countLines || vec[0].size()! = countColumns) { std:: cout < < " ------------ERROR SIZE IS WRONG------------------------------\n"; return 1; } short int* data = new short int[countColumns * countLines]; for (int i = 0; i < countLines; ++i) { for (int j = 0; j < countColumns; ++j) { data[i * countColumns + j] = vec[i][j]; } }
std:: cout < < " ------------DATA SUCCESSFULLY OBTAINED------------------------------\n"; if (mode == onlyCPU || mode == any) { std:: cout < < " ------------CPU START------------------------------\n"; time(& start); timeCPU = clock(); cpuHandling(vec); time(& stop); timeCPU = clock() - timeCPU; saveDataToFile(" cpuresult.txt", vec); std:: cout < < " ------------CPU FINISH------------------------------\n"; }
if (mode == onlyGPU || mode == any) { std:: cout < < " ------------GPU START------------------------------\n"; time(& start); timeGPU = clock(); cudaHandling(data); time(& stop); timeGPU = clock() - timeGPU; saveDataToFile(" gpuresult.txt", data, countColumns, countLines); std:: cout < < " ------------GPU FINISH------------------------------\n"; }
std:: cout < < " CPU time: " < < timeCPU < < " ms\n" < < " GPU time: " < < timeGPU < < " ms\n";
std:: cout < < " ------------FINISH PROGRAM----------------------------------\n"; free(data); return 0; }
МИНОБРНАУКИ РОССИИ
Федеральное государственное бюджетное образовательное учреждение высшего образования «Ярославский государственный университет им. П.Г. Демидова»
Кафедра теоретической информатики
Сдано на кафедру «____» ________ 2018 г. Заведующий кафедрой, д. ф.- м. н., профессор ________ В.А. Соколов
Выпускная квалификационная работа
Фильтрация данных дефектограмм на основе технологии « CUDA »
по направлению подготовки 02.03.02 Фундаментальная информатика и информационные технологии
Научный руководитель д. ф.- м. н., профессор ________Е.В. Кузьмин «___» ________ 2018 г.
Студент группы ИТ-41БО ________А.О. Шуманов «___» _________ 2018 г.
Ярославль 2018 г. Реферат Объем 51 с., 3 гл., 11 рис., 2 табл., 15 источников, 2 прил. |
Последнее изменение этой страницы: 2019-03-20; Просмотров: 381; Нарушение авторского права страницы
