Отчет по OpenCL. Отчет по лабораторной работе "Сравнение производительности различных программноаппаратных средств для задачи определения доступности узлов полигона с использованием программных платформ Opencl и Nvidia cuda"
 Скачать 1.76 Mb.
|
|
САНКТ-ПЕТЕРБУРГСКИЙ ПОЛИТЕХНИЧЕСКИЙ УНИВЕРСИТЕТ им. ПЕТРА ВЕЛИКОГО Институт прикладной математики и механики Высшая школа прикладной математики и вычислительной физики ОТЧЕТ ПО ЛАБОРАТОРНОЙ РАБОТЕ ”Сравнение производительности различных программно-аппаратных средств для задачи определения доступности узлов полигона с использованием программных платформ OpenCL и Nvidia CUDA” 4 курс, группа 3630201/60101 Выполнил И.И.Любимов Преподаватель М.А. Курочкин « » 2020. Санкт-Петербург 2020 Содержание Введение 3 1 Постановка задачи 4 2 Суперкомпьютерный центр «Политехнический» 5 3 Nvidia CUDA 6 3.1 Потоковая модель . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 6 3.2 Устройство памяти 8 3.2.1 Локальная память . . . . . . . . . . . . . . . . . . . . . . . . . . . 9 3.2.2 Константная память . . . . . . . . . . . . . . . . . . . . . . . . . . 9 3.2.3 Текстурная память 10 3.2.4 Регистровая память . . . . . . . . . . . . . . . . . . . . . . . . . . 10 3.2.5 Разделяемая память . . . . . . . . . . . . . . . . . . . . . . . . . . 10 3.3 Ядро . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 11 3.4 Основные понятия: 11 4 OpenCL . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 13 4.1 Потоковая модель . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 13 4.2 Устройство памяти 14 4.3 Ядро . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 14 5 Сравнение OpenCL и CUDA . . . . . . . . . . . . . . . . . . . . . . 16 6 Компиляторы gcc и icc для процессора Intel . . . . . . . . . . 17 6.1 Компилятор gcc . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 18 6.2 Компилятор ICC . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 19 7 Технология cuBLAS для архитектуры Nvidia . . . . . . . . . 20 8 Постановка задачи . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22 9 Алгоритм решения задачи . . . . . . . . . . . . . . . . . . . . . . . . 23 10 Описание эксперимента . . . . . . . . . . . . . . . . . . . . . . . . . . 24 10.1 Эксперимент 1 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 24 10.2 Эксперимент 2 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 24 10.3 Эксперимент 3 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 25 11 Анализ результатов эксперимента . . . . . . . . . . . . . . . . . . 26 11.1 Эксперимент 1 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 26 11.2 Эксперимент 2 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 27 11.3 Эксперимент 3 . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 28 Заключение . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 29 2 Введение В последнее время все большее распространение получают технологии, позволя- ющие использовать графические процессоры видеокарт (GPU) для различного рода вычислений, проводимых на ЭВМ. Данная технология получила название GPGPU - неспециализированные универсальные вычисления на графических процессорах. Вычисление на GPU не предполагает полного переноса нагрузки с центрально- го процессора на графический. Центральный процессор при этом остается частич- но загруженным, а использование графического процессора вместе с центральным позволяет повысить производительность, то есть сократить время выполнения по- ставленной задачи. Выгода от использования графического процессора кроется в архитектуре самих CPU и GPU. Ядра CPU созданы для выполнения одного потока последовательных инструк- ций с максимальной производительностью, а ядра GPU - для быстрого исполне- ния очень большого числа параллельно выполняемых потоков инструкций. В этом и заключается принципиальное отличие графических процессоров от центральных: CPU представляет собой универсальный процессор общего назначения, оптимизи- рованный для достижения высокой производительности единственного потока ко- манд, а GPU устроен иначе. Он проектировался для выполнения огромного количе- ства параллельных потоков команд, причем эти потоки распараллелены изначально, нет накладных расходов на распараллеливание инструкций в графическом процессо- ре. Такое распараллеливание обеспечивается большим количеством исполнительных блоков, заложенным в GPU аппаратно, которые и позволяют одновременно обраба- тывать несколько тысяч потоков команд. На сегодняшний день существует несколько реализаций технологии GPGPU. Ком- пания NVIDIA представила программно-аппаратную архитектуру параллельных вы- числений CUDА (Compute Unified Device Architecture), позволяющую использовать графические процессоры своих видеокарт для проведения вычислений. В основе CUDA лежит упрощенный диалект языка С. Архитектура CUDA SDK обеспечивает программистам реализацию алгоритмов, выполнимых на графических процессорах NVIDIA, и включение специальных функций в текст программы на языке С. Для успешной трансляции кода на этом языке в состав CUDA SDK так же входит соб- ственный С-компиляторкомпании Nvidia. Так же, в последние драйверы графических процессоров NVIDIA включена под- держка OpenCL (Open Computing Language) - это низкоуровневый API для гетеро- генных вычислений, который работает на графических процессорах на базе CUDA. Используя OpenCL API, разработчики могут запускать вычислительные ядра, напи- санные с использованием ограниченного подмножества языка программирования C на графическом процессоре. Выполнение расчётов на GPU показывает отличные результаты в алгоритмах, использующих параллельную обработку данных. То есть, когда одну и ту же после- довательность математических операций применяют к большому объёму данных. При этом лучшие результаты достигаются, если отношение числа арифметических инструкций к числу обращений к памяти достаточно велико. Это предъявляет мень- шие требования к управлению исполнением (flow control), а высокая плотность ма- тематики и большой объёмданных отменяет необходимость в больших кэшах, как на CPU. Целью данной работы является изучение аппаратно-программных платформ Nvidia CUDA и OpenCL, их сравнение, а также исследование скорости выполнения практической задачи в зависимости от параметров распараллеливания 3 1 Постановка задачи • Изучить аппаратно-программную платформу Nvidia CUDA для распараллели- вания вычислении; • Изучить аппаратно-программную платформу OpenCL для распараллеливания вычислении; • Сравнить аппаратно-программные платформы OpenCL и CUDA; • Решить практическую задачу с помощью Nvidia CUDA и с помощью OpenCL; • Исследовать скорость работы задачи в зависимости от параметров распарал- леливания; • Проанализировать результаты исследования 4 2 Суперкомпьютерный центр «Политехнический» В качестве вычислительного ресурса для решения поставленной задачи использует- ся суперкомпьютерный комплекс Санкт-Петербургского Политехнического Универ- ситета. Цель создания Суперкомпьютерного центра (СКЦ) "Политехнический ускорение выполнения и повышение эффективности фундаментальных и прикладных научных исследований, обеспечивающих прорывные инженерные решения. СКЦ "Политехнический"располагает высокопроизводительными вычислительны- ми системами разной архитектуры с общей пиковой производительностью более 1.2 ПФлопс, что выводит его на вторую позицию среди аналогичных центров универси- тетов и исследовательских организаций России. Эти ресурсы доступны отечествен- ным и зарубежным ученым, инженерам в рамках совместных исследовательских про- ектов и разработок. В центре представлены три суперкомпьютера: • "Политехник - РСК Торнадоластер с пиковой производительностью 943 Тфлопс; содержит 668 двухпроцессорных узлов (Intel Xeon E5 2697 v3), из них - 56 узлов имеют два ускорителя вычислений NVIDIA K40; • "Политехник - РСК ПетаСтримассивно-параллельный компьютер с ультравы- сокой многопоточностью; единственная в России система, способная поддер- жать более 70 тыс. потоков; пиковая производительность 291 Тфлопс; содержит 288 узлов на сопроцессорах Intel Xeon Phi; • "Политехник - NUMAассивно-параллельная система с кеш-когерентной гло- бально адресуемой памятью объемом более 12 ТБ; содержит 64 узла, 192 про- цессора; пиковая производительность 30 ТФлопс. Все вычислительные системы работают с общей системой хранения данных (Seagate ClusterStore, 1.1ПБ), имеют единую систему управления и мониторинга. В качестве вычислительного ресурса выбран кластер "Политехник - РСК Торна- до". Для проведения тестов использовался один узел "Политехник - РСК Торнадо"с NVIDIA Tesla K40. Tesla K40 обладает вдвое большим объемом памяти и на 40% быстрее предше- ственника Tesla K20X, а также в 10 раз быстрее самого мощного на сегодня CPU, являясь первым в мире и самым мощным ускорителем, оптимизированным для ана- лиза больших объемов данных и масштабных научных исследований. Ключевые особенности ускорителя Tesla K40: • 12ГБ скоростной памяти GDDR5 для обработки вдвое больших объемов данных и их быстрого анализа; • 2880 параллельных ядер CUDA обеспечивают ускорение работы приложений до 10 раз по сравнению с CPU; • динамический параллелизм позволяет GPU-потокам динамически создавать новые потоки для быстрой и легкой работы с адаптивными и динамическими структурами данных; • поддержка интерфейса PCIe 3-го поколения ускоряет передачу данных вдвое по сравнению с PCIe 2. 5 3 Nvidia CUDA CUDA (Compute Unified Device Architecture) – программно-аппаратная архитек- тура параллельных вычислений, которая позволяет существенно увеличить вычис- лительную производительность благодаря использованию графических процессоров фирмы Nvidia. Процессор GPU представляет собой массив потоковых процессоров (StreamingProcessor Array), что состоит из кластеров текстурных процессоров (TCP -Texture Processor Clusters). TPC в свою очередь состоит из набора мультипроцессоров (SM - Streaming Multiprocessor), в каждом из которых несколько потоковых процессоров (SP - Streaming Processors) или ядер (CUDA cores). GigaThread Engine - это планировщик блоков потоков. Он распределяет блоки потоков по SM. Общий вид GPU кардинально не меняется при переходе от одной микроархитектуры к другой, изменяется размер и скорость L2 кэша. Несмотря на то, что набор компонентов, входящих в SM несильно изменяется, внутренняя компоновка и число этих компонентов может сильно раз- ниться в разных микроархитектурах. 3.1 Потоковая модель В архитектуре Nvidia CUDA применяется SIMT (Single Instruction MultipleThread) модель исполнения. Это модель является комбинацией из MIMD (Multiple Instruction Multiple Data) и SIMD (Single Instruction Multiple Data). Вычисляемая задача состоит из сетки блоков, а блок состоит из нитей (thread). Планировщик блоков запускает блоки на выполнение на определенном SM. На одном SM может выполнятся несколько блоков, но один блок может исполнятся только на одном SM. Но не все нити блока сразу запускаются на исполнение, они разбиваются на варпы - группы по 32 потока. Нити в варпе исполняемые на SM исполняют одну и ту же инструкцию одновременно, отсюда и происходит название модели SIMT. Из определения SIMT модели можно выделить одну проблему, а именно испол- нение потоков одного варпа, содержащих ветвления. При исполнениитаких варпов сначала исполняются инструкции одной ветви оператора, а затем другой. В самой актуальной на данный момент микроархитектуры Volta, этот механизм изменился, но все еще согласуется с SIMT моделью. CUDA SDK позволяет программистам реализовывать на специальном упрощён- ном диалекте языка программирования Си алгоритмы, выполнимые на графических процессорах Nvidia, и включать специальные функции в текст программы на Си. Архитектура CUDA (рис.1) даёт разработчику возможность по своему усмотрению организовывать доступ к набору инструкций графического ускорителя и управлять его памятью. 6 Рис. 1: Логическое представление GPU (архитектура CUDA). Платформа CUDA впервые появились на рынке с выходом чипа NVIDIA восьмого поколения G80 и стала присутствовать во всех последующих сериях графических чипов, которые используются в семействах ускорителей GeForce, Quadro и NVidia Tesla. Первая серия оборудования, поддерживающая CUDA SDK, G8x, имела 32-битный векторный процессор одинарной точности, использующий CUDA SDK как API (CUDA поддерживает тип double языка Си, однако сейчас его точность понижена до 32- битного с плавающей запятой). Более поздние процессоры GT200 имеют поддержку 64-битной точности (только для SFU), но производительность значительно хуже, чем для 32-битной точности (из-за того, что SFU всего два на каждый потоковый мульти- процессор, а скалярных процессоров — восемь). Графический процессор организует аппаратную многопоточность, что позволяет задействовать все ресурсы графиче- ского процессора. Таким образом, открывается перспектива переложить функции физического ускорителя на графический ускоритель (пример реализации — PhysX). Также открываются широкие возможности использования графического оборудова- ния компьютера для выполнения сложных неграфических вычислений: например, в вычислительной биологии и в иных отраслях науки. CUDA. Общие положения. • GPU – сопроцессор для CPU (хоста); • У GPU есть собственная память (device memory); • GPU способен одновременно обрабатывать множество процессов (threads) дан- ных одним и тем же алгоритмом; • Для осуществления расчётов при помощи GPU хост должен осуществить за- пуск вычислительного ядра (kernel), который определяет конфигурацию GPU в вычислениях и способ получения результатов (алгоритм); • Процессы GPU (в отличие от CPU) очень просты и многочисленны (∼ 1000 для полной загрузки GPU). 7 Преимущества: По сравнению с традиционным подходом к организации вычислений общего на- значения посредством возможностей графических API, у архитектуры CUDA отме- чают следующие преимущества в этой области: • Интерфейс программирования приложений CUDA (CUDA API) основан на стандартном языке программирования Си с некоторыми ограничениями. По мнению разработчиков, это должно упростить и сгладить процесс изучения архитектуры CUDA; • Разделяемая между потоками память (shared memory) размером в 16 Кб может быть использована под организованный пользователем кэш с более широкой полосой пропускания, чем при выборке из обычных текстур; • Более эффективные транзакции между памятью центрального процессора и видеопамятью; • Полная аппаратная поддержка целочисленных и побитовых операций; • Поддержка компиляции GPU кода средствами открытого LLVM. Ограничения Все функции, выполнимые на устройстве, не поддерживают рекурсии (в версии CUDA Toolkit 3.1 поддерживает указатели и рекурсию), ни одна из размерностей массива блоков не должна превышать 65535 и некоторые другие ограничения. 3.2 Устройство памяти Устройства CUDA имеют несколько различных областей памяти: глобальную, ло- кальную, текстурную, константную, разделяемую и регистровую память. Каждый тип памяти на устройстве имеет свои преимущества и недостатки. Неправильное использование доступной памяти в приложении может лишить желаемой произво- дительности. Особенности памяти • Данные, хранящиеся впамяти регистров, видны только потоку, который их на- писал, и сохраняются только в течение времени жизни этого потока; • Локальная память имеет те же правила области, что и регистровая память, но работает медленнее; • Данные, хранящиеся в разделяемой памяти, видны всем потокам в этом бло- ке и сохраняются в течение всего блока. Этот тип памяти позволяет потокам общаться и обмениваться данными друг с другом; • Константная память используется для данных, которые не изменятсяв ходе вы- полнения ядра, и доступна только для чтения. Использование константной, а не глобальной памяти может уменьшить требуемую пропускную способность па- мяти, однако этот прирост производительностиможет быть реализован только в том случае, если варп потоков читает одно и то же место; • Память текстур - это еще одна разновидность постоянной памяти на устрой- стве, когда все чтения в варпе физически смежны, использованиетекстурной памяти может уменьшить трафик и повысить производительность по сравне- нию с глобальной памятью. 8 • Данные, хранящиеся в глобальной памяти, видны всем потокам в приложении (включая хост) и сохраняются в течение всего времени выделения хоста. На рисунке 2 представлена сравнительная таблица архитектур памяти. Регистро- вая память располагается на кристалле, кэширование не поддерживается, доступна для чтения и записи для одной нити, время жизни - пока жива нить. Локальная память расположена вне кристалла, не кэшируется, доступ для чтения и записи для одной нити, время жизни - пока жива нить. Разделяемая память располагается на красталле, кэширование не поддерживается, доступна для чтения и записи для всех нитей в блоке, время жизни - пока жив блок. Глобальная память располагается вне кристалла, не кэшируется, доступна на чтение и запись для всех нитей и хоста, время жизни - пока работает приложение. Константная память расположена вне кристалла, кэшируется, доступна на чтение для всех нитей и хоста, время жизни - пока работает приложение. Текстурная память расположена вне кристалла, кэшируется, доступна на чтение для всех нитей и хоста, время жизни - пока работает приложение. Рис. 2: Архитектуры памяти. 3.2.1 Локальная память Локальная память - это не физический тип памяти, а абстракция глобальной па- мяти. Его область является локальной для потока, и он находитсявне микросхемы, что делает его доступным как глобальная память. Локальная память используется только для хранения автоматических переменных. Компилятор использует локальную память, когда определяет, что в регистре недостаточно места для хранения переменной. Автоматические переменные, кото- рые являются большими структурами или массивами, также обычно помещаются в локальную память. 3.2.2 Константная память Константная память - это тип памяти, который хранит неизменяемые данные, до- ступные на уровне сетки. Ее размер 64KB, физически она не отделенаот глобальной памяти, но используя системы кэшей и механизму широкоговещания (broadcast) она может обеспечить прирост производительности засчет сокращения трафика между процессором и памятью. Механизм широкого вещания позволяет за одну операцию чтения обеспечить дан- ными полуварп (группу из 16 потоков одного варпа), если эти потоки запрашивают одни и те же данные. Написание кода с правильным применением этого механизма способно уменьшить количество трафика в 1610 раз. Обратной стороной такого под- ходя является то, что одновременно можно отправить только один запрос на чтение. 9 Это может негативно отразиться на производительности, если потоки в полуварпе будут читать данные из разных адресов. 3.2.3 Текстурная память Текстурная память - это тип памяти который похож на константную память, по- скольку через текстурный блок разрешены запросы только для чтения. Физически текстурная память не отделена от глобальной. Как и константная память, позволяет увеличить производительность за счет системы кэшей. Отличительной особенностью является оптимизация текстурного кэша для двумерной пространственной локаль- ности (данные расположены рядом в двумерном пространстве). Текстурная память может привязываться к 1D, 2D и 3D массивам. Для обращения к ней используются не обычные скобки, а специальные функции, которые говорят компилятору, что эти операции чтения нужно проводить через текстурный блок. 3.2.4 Регистровая память В большинстве случаев для доступа к регистру требуется ноль тактов на инструк- цию. Однако задержки могут возникать из-за зависимостей чтения после записи и конфликтов банков. Задержка чтения после записи составляетпримерно 24 такта. Для более новых устройств CUDA, имеющих 32 ядрана мультипроцессор, может по- требоваться до 768 потоков, чтобы полностью скрыть задержку. В дополнение к задержке чтения после записи, давление в регистре может силь- но ухудшить производительность приложения. Давление в регистре возникает, когда для данной задачи недостаточно регистров. Когда это происходит, данные "перете- кают"с использованием локальной памяти. 3.2.5 Разделяемая память Распространенная проблема возникает, когда память используется совместно: при наличии всей памяти, доступной для всех потоков, будет много потоков, одновре- менно получающих доступ к данным. Чтобы устранить это потенциальное узкое ме- сто, разделяемая память делится на 32 логических банка. Последовательные участки памяти назначаются последовательным банкам (см. рисунок 3). Общий размер разделяемой памяти может быть установлен на 16 КБ, 32КБ или 48 КБ (при этом оставшийся объем автоматически используется для кэша L1), как показано на рисунке 3. По умолчанию общая память имеет значение 48 КБ (при этом для кэша L1 остается 16 КБ). Благодаря архитектуре Kepler каждый банк имеет полосу пропускания 64 бита за такт. Более старая архитектура Fermi работала по-другому, но фактически обес- печивала половину этой пропускной способности. В варпе 32 потока и ровно 32 банка общей памяти. Поскольку каждый банк обслу- живает только один запрос за цикл, несколько одновременных обращений к одному и тому же банку приведут к так называемому конфликту банков. 10 Рис. 3: Общая память и кэш L1. 3.3 Ядро Ядро CUDA компилируется вместе с остальными исходниками программы. В ядрах CUDA есть свои инструменты вычисления индекса потока, они представлены в виде констант. Поскольку размер вычислительной сетки задается в блоках, а размер блока в потоках, то и константы соответствуют этой логике. В потоке доступны: индекс блока, индекс потока в блоке, размер сетки в блоках, размер блока в потоках. В CUDA также присутствует барьер синхронизации. Он ожидает пока все пото- ки в блоке выполнятся до него. Он может быть использован, чтобы гарантировать запись/чтение всеми потоками данных из разделяемой памяти. В CUDA нет барьера для глобальной памяти, поскольку считается, что проще запустить два ядра подряд, чем одно с глобальным барьером. 3.4 Основные понятия: • Устройство – GPU. Выполняет роль «подчиненного» - делает то, что ему гово- рит CPU. • Хост – GPU. Выполняет управляющую роль – запускает задачи на устройстве, выделяет память на устройстве, перемещает данные с/на устройство. CUDA поддерживает любой язык программирования. При компиляции програм- мы компилятор CUDA генерирует код отдельно для хоста и отдельно для устройства. Основные этапы CUDA-программы: 1. Хост выделяет нужное количество памяти на устройстве; 11 2. Хост копирует данные из своей памяти в память устройства; 3. Хост стартует выполнение определенных ядер на устройстве; 4. Устройство выполняет задачу; 5. Хост копирует данные с результатом из памяти устройства в свою память. Следует отметить, что ядра представляют из себя обычные последовательное про- граммы – то есть создание и запуск потоков не происходит в коде самих ядер, этим занимается хост. Технология CUDA использовалась в данной работе для того, чтобы распарал- лелить алгоритм. Также CUDA позволяет значительно увеличить вычислительную производительность благодаря использованию GPU (графических процессоров). Для написания кода GPU необходимо использовать: 1. Спецификаторы функций. Они предназначены для указания, как и откуда должна выполняться функция. К таковым относятся • _ _ host _ _ - выполняется на CPU (центральный процессор), вызывается с CPU • _ _ global _ _ - выполняется на GPU, вызывается с CPU • _ _ device _ _ - выполняется на GPU, вызывается с GPU 2. Спецификаторы переменных, которые служат для указания типа используемой памяти GPU. 3. Спецификаторы запуска ядра GPU, которые служат для описания колическтва сеток, блоков и памяти, которые необходимо выделить при расчете на GPU. Синтаксис запуска ядра: funcName« param2), где: • gridSize – количество сеток с блоками, выделенных для расчетов • blockSize – количество блоков в сетке, выделенных для расчетов. 4. Встроенные переменные: • gridDim – размер сетки, которая имеет тип dim3. Позволяет узнать размер сетки выделенного при текущем вызове ядра. • blockDim – размер блока, которая имеет тип dim3. Позволяет узнать размер блока выделенного при текущем вызове ядра. • blockIdx – индекс текущего блока в вычислении на GPU, имеет тип unit3. • threadIdx – индекс текущего треда в вычислении на GPU, имеет тип unit3. 12 4 OpenCL OpenCL (Open Computing Language) является платформой для написания программ, которые выполняются на гетерогенных платформах, состоящих из центральных про- цессоров (CPU), графических процессоров (GPU), цифровых сигнальных процессо- ров (DSP), ПЛИС (FPGA) и других процессоры или аппаратные ускорители. OpenCL описывает язык программирования (на основе C99 и C++ 11) для программирования этих устройств и интерфейсы прикладного программирования (API) для управления платформой и выполнения программ на вычислительных устройствах. Основные концепции OpenCL основывается на трех основных концепциях: • host (хост) - управляющее по отношению к GPU или любому другому девайсу устройство, чаще CPU. Соответственно, такие понятия как хост-код и память хоста означают, что такой код будет исполняться на CPU, а память будет хра- ниться в RAM. • device (девайс) - устройство, на котором запускаются вычислительные ядра. Аналогично хост-коду и памяти существует понятие о девайс-коде и памяти. Девайс-код и память должны исполняться и храниться на девайсе. Важно, что CPU также может выступать в качестве девайса, но с ограниченным функци- оналом. • kernel (ядро) - эта функция, которая будет запущена на девайсе, вызывается только с хоста, помечается квалификатором __kernel. 4.1 Потоковая модель Вычислительное ядро (kernel) содержит код, исполняемый каждым элементом (work item) вычислительной сетки NDRange (N-dimensional range) (см. рисунок 4). Сетки могут быть 1, 2, 3-мерные. Размеры сетки в каждом измерении задают количество элементов в этом измерении. Общее число элементов, исполняющих ядро, может быть посчитано как произведение размеров сетки в 3-х измерениях. Рис. 4: Группировка элементов в OpenCL. Элементы на устройстве запускаются не штучно, а группами (work groups). Груп- пы могут быть 1, 2, 3-мерные. Размер групп обязательно должен быть кратен разме- ру сетки. Группы являются важной частью OpenCL, т.к. только элементы из одной группы имеют механизмы взаимодействия между собой: локальная память и ба- рьер синхронизации. Максимальное размеры сетки, количество доступных измере- ний, максимальный размер группы, максимальное количество групп являются пара- 13 метрами, зависимыми от конкретного OpenCL устройства. Их можно получить непо- средственно во время исполнения программы. Это важно, поскольку это позволяет запускать ядра с наиболее подходящей группировкой элементов, что положительно сказывается на производительности. 4.2 Устройство памяти OpenCL определяет четырехуровневую иерархию памяти для вычислительного устрой- ства (см. рисунок 5): • Глобальная память (global memory): область памяти предоставляет доступ на чтение/запись всем элементам во всех группах, работающие на любом устрой- стве в контексте. Элементы могут читать или писать в любой элемент объекта памяти. Чтение и запись в глобальную память могут быть кэшированы в зави- симости от возможностей устройства. • Постоянная память (constant memory): область глобальной памяти, которая остается постоянной во время выполнения экземпляра ядра. Хост выделяет и инициализирует объекты памяти, помещенные в константную память. • Локальная память (local memory): область памяти, локальная для рабочей группы. Эта область памяти может быть использована для выделения пере- менных, которые являются общими для всех элементов в этой группе. • Приватная память (private memory): область памяти, частная для элемента. Переменные, определенные в приватной памяти элемента не видна другому элементу. Рис. 5: Иерархия памяти OpenCL в устройствах. 4.3 Ядро Функция, которая исполняется всеми элементами в сетке. Ядра и все вспомогатель- ные функции, исполняемые на устройстве, описываются либо в отдельном файле и подгружаются в программу, либо в виде строки в программе. Во время работы про- грамма компилирует исходный код, и используетего для запуска ядер на устройствах. Также существует возможность заранее скомпилировать код ядер и в программу подгружать уже бинарный файл. 14 В OpenCL ядрах можно использовать встроенные функции для определения гло- бального (относительно всей сетки), локального (относительно группы) индекс эле- мента, индекс группы по каждому из измерений. Также возможно получить число используемых измерений, размеры сетки и группы, и число групп. Эти функции позволяют потоку определить участки обрабатываемой памяти. Поскольку последовательность исполнения элементов не определена, но они мо- гут обращаться к локальной или глобальной памяти, существует необходимость га- рантировать, что все элементы произвели чтение или запись. Для этого используется барьер синхронизации с флагами для глобальной и локальной синхронизации. 15 5 Сравнение OpenCL и CUDA Сводная таблица характеристик и свойств OpenCL и CUDA приведена на рис. 6 Рис. 6: Таблица сравнения OpenCL и CUDA. 16 6 Компиляторы gcc и icc для процессора Intel Для разработки программ на C или C++ необходим компилятор – он транслиру- ет исходный код в исполняемый файл, который затем можно запускать. В настоящий момент существуют множество различных компиляторов, которые могут отличаться по разным аспектам, в частности, по реализации стандартов. Стадии компиляции программы на C/C++: 1. Препроцессинг. Препроцессор – это макро процессор, который преобразует про- грамму для дальнейшего компилирования. На данной стадии происходит рабо- та с препроцессорными директивами, например удаление комментариев и встав- ка кода из заголовочных файлов. 2. Компиляция. На данном этапе компилятор выполняет свою главную работу – компилирует, то есть преобразует полученный на прошлом шаге код без дерик- тив, в ассемблерный код. Это промежуточный шаг между языками высокого уровня и машинным языком. 3. Ассемблирование. Так как х86 процессоры исполняют команды бинарного кода, необходимо перевести ассемблерный код в машинный с помощью ассемблера. Ассемблер преобразовывает ассемблерный код в магинный, сохраняя его в объ- ектном файле. Объектный файл – созданный ассемблером промежуточный файл, хранящий кусок машинного кода. Этот кусок машинного кода, который еще не был связан с другими кусками в конечную выполянемую программу, называется объект- ным кодом. 4. Компоновка. Компоновщик (линковщик, линкер) связывает все объектные фай- лы и статические библиотеки в единый исполняемый файл, который может быть запущен. 5. Загрузка. Последний этап процесса компиляции программы – вызов загруз- чика для загрузки программы в память. На данной стадии также возможна подгрузка динамических библиотек. 17 6.1 Компилятор gcc GCC — своббодно доступный оптимизирующий компилятор для языков C /C++. Программа gcc, запускаемая из командной строки, представляет собой надстрой- ку над группой компиляторов. В зависимости от расширений имен файлов, переда- ваемых в качестве параметров, и дополнительных опций, gcc запускает необходимые препроцессоры, компиляторы и линковщики. Файл с расширением .cc, .C или .cpp рассматривается как файл на языкек C++, файлы с расширениями .c и .cu рассматриваются как программы на языке C, а файлы с расширением .o и .obj считаются объектными. Чтобы начать процесс компиляции достаточно вызвать команду gcc следующим образом: gcc – абсолютный или относи- тельный путь к файлу, а Рис. 7: Опции компилятора GCC. 18 6.2 Компилятор ICC Intel C++ compiler — оптимизирующий компилятор, разрабатываемый фирмой Intel для процессоров семейств х86, х86-64 и IA-64. Главным достоинством компиля- тора являются выполняемые им высокоуровневые и целевые оптимизации под про- цессоры Intel. Компилятор работает под ОС Linux, Windows, MacOS X. Основные возможности: • Высокоуровневая оптимизация • Межпроцедурная оптимизация • Автоматическое распараллеливание кода • Веторизация для SSE, SSE2, SSE3, SSE4 • Оптимизация с учетом профильной информации Компилятор поддерживает стандарт OpenMP 3.0 для написания параллельных программ, а также содержит модификацию OpenMP под названием Cluster OpenMP, при помощи которой можно запускать приложения, написанные в соответствии с OpenMP на кластерах, использующих MPI. Intel C++ compiler использует фронтэнд – часть компилятора, занимающуюся синтаксическим анализом компилируемой программы – от Edison Design Group. Этот же фронтэнд используется и в других достаточно популярных компиляторах. К недостаткам Linux–версии компилятора можно отнести несовместимость с GNU– расширениями языка Си (поддерживаемые в GCC), что может вызвать проблемы при компиляции некоторых программ. Поддерживаются языки C, C++, Fortran. На рисунке 8 приведены наиболее часто употребляемые опции компилятора ICC. Рис. 8: Опции компилятора ICC. 19 7 Технология cuBLAS для архитектуры Nvidia Библиотека cuBLAS представляет собой реализацию BLAS (базовых подпро- грамм линейной алгебры) поверх среды выполнения NVIDIA CUDA. Это позволяет пользователю получить доступ к вычислительным ресурсам графического процессо- ра NVIDIA (GPU). Библиотека cuBLAS предоставляет три набора API: • API-интерфейс cuBLAS, который в этом документе называется просто API- интерфейсом cuBLAS (начиная с CUDA 6.0), • API CUBLASXT (начиная с CUDA 6.0) и • API cuBLASLt (начиная с CUDA 10.1) Чтобы использовать API cuBLAS, приложение должно выделить необходимые матрицы и векторы в пространстве памяти GPU, заполнить их данными, вызвать последовательность требуемых функций cuBLAS, а затем загрузить результаты из пространства памяти GPU обратно на хост. API cuBLAS также предоставляет вспо- могательные функции для записи и извлечения данных из графического процессора. Чтобы использовать API CUBLASXT, приложение должно хранить данные на хосте, а библиотека позаботится о распределении операции на один или несколь- ко графических процессоров, присутствующих в системе, в зависимости от запроса пользователя. CuBLASLt - это облегченная библиотека, предназначенная для операций GEneral Matrix-to-matrix Multiply (GEMM) с новым гибким API. Эта библиотека добавля- ет гибкость в расположении матричных данных, типах ввода, типах вычислений, а также в выборе алгоритмических реализаций и эвристики посредством програм- мируемости параметров. После того, как пользователь определил набор опций для предполагаемой операции GEMM, эти опции могут многократно использоваться для разных входов. Это аналогично тому, как cuFFT и FFTW сначала создают план и повторно используют для БПФ одинакового размера и типа с разными входными данными. Ниже приведен пример использование cuBLAS при решении задачи "1-based indexing" #include #include #include #include #include "cublas_v2.h" #define M 6 #define N 5 #define IDX2F(i,j,ld) ((((j)-1)*(ld))+((i)-1)) static __inline__ void modify (cublasHandle_t handle, float *m, int ldm, int n, int p, int q, float alpha, float beta){ cublasSscal (handle, n-q+1, &alpha, &m[IDX2F(p,q,ldm)], ldm); cublasSscal (handle, ldm-p+1, &beta, &m[IDX2F(p,q,ldm)], 1); } int main (void){ cudaError_t cudaStat; cublasStatus_t stat; cublasHandle_t handle; int i, j; 20 float* devPtrA; float* a = 0; a = (float *)malloc (M * N * sizeof (*a)); if (!a) { printf ("host memory allocation failed"); return EXIT_FAILURE; } for (j = 1; j <= N; j++) { for (i = 1; i <= M; i++) { a[IDX2F(i,j,M)] = (float)((i-1) * M + j); } } cudaStat = cudaMalloc ((void**)&devPtrA, M*N*sizeof(*a)); if (cudaStat != cudaSuccess) { printf ("device memory allocation failed"); return EXIT_FAILURE; } stat = cublasCreate(&handle); if (stat != CUBLAS_STATUS_SUCCESS) { printf ("CUBLAS initialization failed\n"); return EXIT_FAILURE; } stat = cublasSetMatrix (M, N, sizeof(*a), a, M, devPtrA, M); if (stat != CUBLAS_STATUS_SUCCESS) { printf ("data download failed"); cudaFree (devPtrA); cublasDestroy(handle); return EXIT_FAILURE; } modify (handle, devPtrA, M, N, 2, 3, 16.0f, 12.0f); stat = cublasGetMatrix (M, N, sizeof(*a), devPtrA, M, a, M); if (stat != CUBLAS_STATUS_SUCCESS) { printf ("data upload failed"); cudaFree (devPtrA); cublasDestroy(handle); return EXIT_FAILURE; } cudaFree (devPtrA); cublasDestroy(handle); for (j = 1; j <= N; j++) { for (i = 1; i <= M; i++) { printf ("%7.0f", a[IDX2F(i,j,M)]); } printf ("\n"); } free(a); return EXIT_SUCCESS; } 21 8 Постановка задачи Название определение маршрута из одной точки в другую при случайном распределении доступных узлов полигона до движения робота. Дано • карта высот M ap[i][j], представляющая собой двумерный массив, каждый из узлов которой показывает высоту данной точки; • максимальное значение высоты maxDelta, которое способен преодолеть робот при переходе из одной точки к другой. Необходимо: пометить каждую клетку на возможность доступа из нее в одну из четырех бли- жайших ячеек (верхнюю, левую, нижнюю и правую) и построить маршрут из за- данной точки А в заданную точку В. Соседняя клетка считается доступной, если разница высот между выбранной и соседней клетками меньше некоторого заданного значения. Ограничения карта квадратная, максимальный размер карты равен 25000 точек, узлы содержат целые числа в диапазоне (-50, 49), узлы A и B не совпадают 22 9 Алгоритм решения задачи Заведём массив d[], в котором для каждой вершины v будем хранить текущую длину d[v] кратчайшего пути из s в v. Изначально d[s] = 0, а для всех остальных вершин эта длина равна бесконечности: d[v] = ∞, v 6= s Кроме того, для каждой вершины v будем хранить, помечена она ещё или нет, т.е. заведём булевский массив u[]. Изначально все вершины не помечены, т.е. u[v] = false Сам алгоритм Дейкстры состоит из n итераций. На очередной итерации выбира- ется вершина v с наименьшей величиной d[v] среди ещё не помеченных, т.е.: d[v] = min p: u[p]=false d[p] Выбранная таким образом вершина v отмечается помеченной. Далее, на теку- щей итерации, из вершины v производятся релаксации: просматриваются все рёбра (v, to), исходящие из вершины v, и для каждой такой вершины to алгоритм пытается улучшить значение d[to]. Пусть длина текущего ребра равна len, тогда в виде кода релаксация выглядит как: d[to] = min(d[to], d[v] + len) На этом текущая итерация заканчивается, алгоритм переходит к следующей ите- рации (снова выбирается вершина с наименьшей величиной d, из неё производятся релаксации, и т.д.). При этом в конце концов, после n итераций, все вершины графа станут помеченными, и алгоритм свою работу завершает. Утверждается, что найден- ные значения d[v] и есть искомые длины кратчайших путей из s в v. Стоит заметить, что, если не все вершины графа достижимы из вершины s, то значения d[v] для них так и останутся бесконечными. Понятно, что несколько послед- них итераций алгоритма будут как раз выбирать эти вершины, но никакой полезной работы производить эти итерации не будут (поскольку бесконечное расстояние не сможет прорелаксировать другие, даже тоже бесконечные расстояния). Поэтому ал- горитм можно сразу останавливать, как только в качестве выбранной вершины бе- рётся вершина с бесконечным расстоянием. В данном случае длина вершины будет вычисляться динамически, то есть если следующая вершина доступна, т.к. разница высот не больше maxDelta, то len = 1, в противном случае len = inf Последовательная часть выделение и освобождение памяти, загрузка ивыгрузка данных с устройства, ини- циализация переменных размерностей, генерация карты высот, вывод результатов работы алгоритма. Параллельная часть выбор следующей вершины и релаксация ребер. 23 10 Описание эксперимента Целью данного исследования является выявление зависимости времени выпол- нения параллельного участка кода от типа используемой технологии (OpenCL или CUDA), от типа используемой памяти (глобальная или текстурная), размера поли- гона, а так же от числа потоков, на которых происходит выполнение алгоритма. 10.1 Эксперимент 1 Установление зависимости времени выполнения алгоритмов от числа потоков при фиксированном размере стороны полигона равной 25000. Для сбора данных было проведено 10 испытаний для каждого исследуемого числа потоков, после чего было посчитано среднее значение времени выполнения и были построены графики. Число потоков регулировалось размером формируемого для вычислений блоком (константа BLOCKSIZE). Результаты представлены на рис. 9 Рис. 9: Результаты эксперимента 1. 10.2 Эксперимент 2 Установление зависимости времени выполнения алгоритмов от размера стороны полигона при фиксированном размере блока равном 32. Для сбора данных было проведено 10 испытаний для каждого исследуемого размера полигона, после чего было посчитано среднее значение времени выполнения и были построены графики. Размер полигона регулировался константой FIELD_SIDE. Результаты представлены на рис. 10 Рис. 10: Результаты эксперимента 2. 24 10.3 Эксперимент 3 Установление зависимости времени выполнения алгоритмов от используемого компилятора при фиксированном размере блока равном 32 и фиксированном разме- ре поля равном 25000x25000. Для сбора данных было проведено 10 испытаний для каждого исследуемого размера полигона, после чего было посчитано среднее значе- ние времени выполнения и были построены графики. Результаты представлены на рис. 11 Рис. 11: Результаты эксперимента 3. 25 11 Анализ результатов эксперимента 11.1 Эксперимент 1 На рисунке 12 представлен график результатов первого эксперимента. По оси абсцисс указана размерность блоков, задающих количество потоков, по оси орди- нат - время выполнения расчетов в секундах. Размер полигона фиксирован и равен 25000*25000. Рис. 12: Результаты эксперимента 1. Из рисунка видно, что скорость выполнения вычислений растет по мере увеличе- ния числа потоков, однако для алгоритма, исполняемого с использованием техноло- гии OpenCL прирост производительности уменьшается быстрее, чем для алгоритмов, использующих технологию Nvidia CUDA. Пиковая производительность алгоритмов, использующих технологию Nvidia CUDA достигается при размере блока равном 16, а далее с увеличением размера блока по- казатель ухудшается в связи с затратами на выделение дополнительных ресурсов. Из-за этого при размере блока равном 32 результаты работы алгоритмов практиче- ски идентичны. Наблюдается прирост производительности при использовании текстурной памяти по сравнению с использованием глобальной памяти, однако данные показатели будут присутствовать не на всех типах графических ускорителей (эксперимент проведен с использованием Nvidia Tesla K40x), а так же зависит от типа и размерности решаемой задачи. 26 11.2 Эксперимент 2 На рисунке 13 представлен график результатов второго эксперимента. По оси абсцисс указано время выполнения расчетов в секундах, по оси ординат – размер сто- роны полигона. Размер блока фиксирован и равен 32. В виде столбцов представлены результаты выполнения задачи с использованием OpenCL для большей наглядно- сти при сравнении с выполнением задачи с различной моделью памяти при помощи Nvidia CUDA. Рис. 13: Результаты эксперимента 2. Из рисунка видно, что OpenCL дает преимущество при расчете задачи по срав- нению с использованием технологии CUDA с глобальной памятью на 1,7%, одна- ко использование технологии CUDA с текстурной памятью для данной задачи дает прирост производительности порядка 10%. Для всех трех алгоритмов наблюдается квадратичная зависимость скорости вычисления от размера полигона, что является ожидаемым для алгоритма поиска пути. 27 11.3 Эксперимент 3 На рисунке 13 представлен график результатов третьего эксперимента. По оси абсцисс указано время выполнения расчетов в секундах, по оси ординат – размер стороны полигона. Размер блока фиксирован и равен 32. Результаты представлены на рис. 14 Рис. 14: Результаты эксперимента 3. Из рисунка видно, что ICC дает преимущество при расчете задачи по сравнению с использованием компилятора GCC с глобальной памятью на 10,3%, это объясня- ется тем, что компилятор ICC оптимизирует код, исполняемый на процессоре, так как в кластере "Политехник – РСК Торнадо"используются процессоры Intel Xeon, а ICC поддерживает стандарт OpenMP, что в итоге позволяет дополнительно рас- параллелить однопоточную часть задачи. Для всех трех алгоритмов наблюдается квадратичная зависимость скорости вычисления от размера полигона, что является ожидаемым для алгоритма поиска пути. 28 Заключение В результате выполнения данной работы были изучены аппаратная ипрограмм- ная части CUDA и OpenCL, технология написания параллельных программ с ис- пользованием GPU, суперкомпьютера и технологии CUDA и OpenCL. На основе изученных материалов был реализован и распараллелен с помощью CUDA и OpenCL алгоритм определения кратчайшего пути на карте высот до движе- ния робота. Для распараллеливания использовалась глобальная и текстурная память (для Nvidia CUDA). Удалось исследовать зависимость времени работы алгоритма в зависимости от количества потоков, размерности полигона, а так же от типа используемой памяти и используемого компилятора (на примере использования технологии CUDA). Ско- рость выполнения вычислений растет по мере увеличения числа потоков, однако для алгоритма, исполняемого с использованием технологии OpenCL прирост производи- тельности уменьшается быстрее, чем для алгоритмов, использующих технологию Nvidia CUDA. Пиковая производительность алгоритмов, использующих технологию Nvidia CUDA достигается при размере блока равном 16, а далее с увеличением размера блока по- казатель ухудшается в связи с затратами на выделение дополнительных ресурсов. Из-за этого при размере блока равном 32 результаты работы алгоритмов практиче- ски идентичны. Также в силу того, что на используемом кластере суперкомпьютера использует- ся процессоры Intel серии Xeon, использование компилятора ICC позволяет также дополнительно распараллелить часть исполняемого на процессоре кода, что влечет за собой уменьшение времени исполнения программы. Такая функциональность с компилятором GCC возможно только при ручной оптимизации кода и нативного использования технологии OpenMP. Для всех трех алгоритмов наблюдается квадратичная зависимость скорости вы- числения от размера полигона. OpenCL дает преимущество при расчете задачи по сравнению с использованием технологии CUDA с глобальной памятью на 1,7%, одна- ко использование технологии CUDA с текстурной памятью для данной задачи дает прирост производительности порядка 10%. 29 |