793 lines
56 KiB
TeX
793 lines
56 KiB
TeX
\documentclass[a4paper, final]{article}
|
||
%\usepackage{literat} % Нормальные шрифты
|
||
\usepackage[14pt]{extsizes} % для того чтобы задать нестандартный 14-ый размер шрифта
|
||
\usepackage{tabularx}
|
||
\usepackage[T2A]{fontenc}
|
||
\usepackage[utf8]{inputenc}
|
||
\usepackage[russian]{babel}
|
||
\usepackage{amsmath}
|
||
\usepackage[left=25mm, top=20mm, right=20mm, bottom=20mm, footskip=10mm]{geometry}
|
||
\usepackage{ragged2e} %для растягивания по ширине
|
||
\usepackage{setspace} %для межстрочно го интервала
|
||
\usepackage{moreverb} %для работы с листингами
|
||
\usepackage{indentfirst} % для абзацного отступа
|
||
\usepackage{moreverb} %для печати в листинге исходного кода программ
|
||
\usepackage{pdfpages} %для вставки других pdf файлов
|
||
\usepackage{tikz}
|
||
\usepackage{graphicx}
|
||
\usepackage{afterpage}
|
||
\usepackage{longtable}
|
||
\usepackage{float}
|
||
|
||
|
||
|
||
% \usepackage[paper=A4,DIV=12]{typearea}
|
||
\usepackage{pdflscape}
|
||
% \usepackage{lscape}
|
||
|
||
\usepackage{array}
|
||
\usepackage{multirow}
|
||
|
||
\renewcommand\verbatimtabsize{4\relax}
|
||
\renewcommand\listingoffset{0.2em} %отступ от номеров строк в листинге
|
||
\renewcommand{\arraystretch}{1.4} % изменяю высоту строки в таблице
|
||
\usepackage[font=small, singlelinecheck=false, justification=centering, format=plain, labelsep=period]{caption} %для настройки заголовка таблицы
|
||
\usepackage{listings} %листинги
|
||
\usepackage{xcolor} % цвета
|
||
\usepackage{hyperref}% для гиперссылок
|
||
\usepackage{enumitem} %для перечислений
|
||
|
||
\newcommand{\specialcell}[2][l]{\begin{tabular}[#1]{@{}l@{}}#2\end{tabular}}
|
||
|
||
|
||
\setlist[enumerate,itemize]{leftmargin=1.2cm} %отступ в перечислениях
|
||
|
||
\hypersetup{colorlinks,
|
||
allcolors=[RGB]{010 090 200}} %красивые гиперссылки (не красные)
|
||
|
||
% подгружаемые языки — подробнее в документации listings (это всё для листингов)
|
||
\lstloadlanguages{ SQL}
|
||
% включаем кириллицу и добавляем кое−какие опции
|
||
\lstset{tabsize=2,
|
||
breaklines,
|
||
basicstyle=\footnotesize,
|
||
columns=fullflexible,
|
||
flexiblecolumns,
|
||
numbers=left,
|
||
numberstyle={\footnotesize},
|
||
keywordstyle=\color{blue},
|
||
inputencoding=cp1251,
|
||
extendedchars=true
|
||
}
|
||
\lstdefinelanguage{MyC}{
|
||
language=SQL,
|
||
% ndkeywordstyle=\color{darkgray}\bfseries,
|
||
% identifierstyle=\color{black},
|
||
% morecomment=[n]{/**}{*/},
|
||
% commentstyle=\color{blue}\ttfamily,
|
||
% stringstyle=\color{red}\ttfamily,
|
||
% morestring=[b]",
|
||
% showstringspaces=false,
|
||
% morecomment=[l][\color{gray}]{//},
|
||
keepspaces=true,
|
||
escapechar=\%,
|
||
texcl=true
|
||
}
|
||
|
||
\textheight=24cm % высота текста
|
||
\textwidth=16cm % ширина текста
|
||
\oddsidemargin=0pt % отступ от левого края
|
||
\topmargin=-1.5cm % отступ от верхнего края
|
||
\parindent=24pt % абзацный отступ
|
||
\parskip=5pt % интервал между абзацами
|
||
\tolerance=2000 % терпимость к "жидким" строкам
|
||
\flushbottom % выравнивание высоты страниц
|
||
|
||
|
||
% Настройка листингов
|
||
\lstset{
|
||
language=C++,
|
||
extendedchars=\true,
|
||
inputencoding=utf8,
|
||
keepspaces=true,
|
||
% captionpos=b, % подписи листингов снизу
|
||
}
|
||
|
||
\begin{document} % начало документа
|
||
|
||
|
||
|
||
% НАЧАЛО ТИТУЛЬНОГО ЛИСТА
|
||
\begin{center}
|
||
\hfill \break
|
||
\hfill \break
|
||
\normalsize{МИНИСТЕРСТВО НАУКИ И ВЫСШЕГО ОБРАЗОВАНИЯ РОССИЙСКОЙ ФЕДЕРАЦИИ\\
|
||
федеральное государственное автономное образовательное учреждение высшего образования «Санкт-Петербургский политехнический университет Петра Великого»\\[10pt]}
|
||
\normalsize{Институт компьютерных наук и кибербезопасности}\\[10pt]
|
||
\normalsize{Высшая школа технологий искусственного интеллекта}\\[10pt]
|
||
\normalsize{Направление: 02.03.01 <<Математика и компьютерные науки>>}\\
|
||
|
||
\hfill \break
|
||
\hfill \break
|
||
\hfill \break
|
||
\large{Отчёт по по дисциплине}\\
|
||
\large{<<Курсовое проектирование по управлению ресурсами суперэвм>>}\\
|
||
\large{Решение задачи нахождения пути движения робота по полигону}\\
|
||
\hfill \break
|
||
|
||
% \hfill \break
|
||
\hfill \break
|
||
\end{center}
|
||
|
||
\small{
|
||
\begin{tabular}{lrrl}
|
||
\!\!\!Студент, & \hspace{2cm} & & \\
|
||
\!\!\!группы 5130201/20102 & \hspace{2cm} & \underline{\hspace{3cm}} &Тищенко А. А. \\\\
|
||
\!\!\!Преподаватель & \hspace{2cm} & \underline{\hspace{3cm}} & Курочкин М. А. \\\\
|
||
&&\hspace{4cm}
|
||
\end{tabular}
|
||
\begin{flushright}
|
||
<<\underline{\hspace{1cm}}>>\underline{\hspace{2.5cm}} 2025г.
|
||
\end{flushright}
|
||
}
|
||
|
||
\hfill \break
|
||
% \hfill \break
|
||
\begin{center} \small{Санкт-Петербург, 2025} \end{center}
|
||
\thispagestyle{empty} % выключаем отображение номера для этой страницы
|
||
|
||
% КОНЕЦ ТИТУЛЬНОГО ЛИСТА
|
||
\newpage
|
||
|
||
\tableofcontents
|
||
|
||
|
||
\newpage
|
||
|
||
\section*{Введение}
|
||
|
||
\addcontentsline{toc}{section}{Введение}
|
||
|
||
Многие задачи, возникающие на практике, требуют большого объема вычислений. Одним из вариантов решения сложных вычислительных задач является использование параллельного программирования. За последние несколько десятилетий стало очень распространено вычисление с помощью графических ускорителей — устройств с массивно-параллельной архитектурой. Производить вычисления общего
|
||
назначения можно на видеокартах архитектуры Nvidia CUDA. Сегодня спроектированы и испытаны многие компьютеры, которые используют в своей архитектуре
|
||
тот или иной вид параллельной обработки данных. Сложность работы программирования заключается в координации используемых ресурсов. Одним из примеров
|
||
массивных вычислительных систем является суперкомпьютерный центр «Политехнический». Часть узлов этого суперкомпьютера оборудована графическими ускорителями Nvidia Tesla K40X.
|
||
|
||
|
||
\newpage
|
||
|
||
\section {Постановка задачи}
|
||
В рамках данной работы необходимо изучить технологию параллельного программирования с использованием архитектуры Nvidia CUDA.
|
||
|
||
Также необходимо ознакомиться с принципом использования ресурсов суперкомпьютерного центра «Политехнический» для решения прикладной задачи.
|
||
|
||
Необходимо написать программу для решения поставленной практической задачи с использованием технологии Nvidia CUDA и провести исследование зависимости времени
|
||
выполнения программы от количества используемых ресурсов.
|
||
|
||
В рамках курсовой работы необходимо написать программу для построения пути движения робота по полигону.
|
||
|
||
\newpage
|
||
\section {Аппаратно-программная платформа Nvidia
|
||
CUDA}
|
||
CUDA (Compute Unified Device Architecture) — это архитектура параллельных
|
||
вычислений от Nvidia, позволяющая существенно увеличить вычислительную производительность благодаря использованию GPU (графических процессоров) фирмы Nvidia.
|
||
|
||
\subsection{Архитектура Nvidia CUDA}
|
||
|
||
Видеокарты Nvidia CUDA имеют иерархическую архитектуру:
|
||
\begin{itemize}
|
||
\item Процессор GPU представляет собой массив потоковых процессоров (Streaming
|
||
Processor Array);
|
||
\item Потоковый процессор состоит из кластеров текстурных процессоров (Texture
|
||
Processor Clusters);
|
||
\item Текстурные процессоры состоят из набора мультипроцессоров (Streaming
|
||
Multiprocessor);
|
||
\item Мультипроцессоры содержат несколько потоковых процессоров (Streaming
|
||
Processors) или ядер (CUDA cores).
|
||
\end{itemize}
|
||
|
||
Планирование выполнения команд происходит при помощи GigaThread Engine.
|
||
Он распределяет блоки потоков по мультипроцессорам. Общий вид GPU кардинально не меняется при переходе от одной микроархитектуры к другой, изменяется размер и скорость L2 кэша.
|
||
|
||
В архитектуре Nvidia CUDA применяется SIMT (Single Instruction Multiple
|
||
Thread) модель исполнения. Это модель является комбинацией из MIMD (Multiple
|
||
Instruction Multiple Data) и SIMD (Single Instruction Multiple Data). Вычисляемая
|
||
задача состоит из сетки блоков, а блок состоит из нитей (thread), при исполнении
|
||
нити разбиваются на варпы — группы по 32 нити. Все нити варпы выполняют в одно
|
||
время одну и ту же инструкцию.
|
||
|
||
\subsection{Вычислительные возможности Nvidia CUDA}
|
||
|
||
Видеокарты Nvidia CUDA разных микроархитектур обладают разным количеством ядер разного назначения. В этом разделе приведены краткие сведения о вычислительных возможностях различных поколений Nvidia CUDA.
|
||
|
||
\subsubsection*{Микроархитектура Fermi}
|
||
В каждом потоковом процессоре:
|
||
\begin{itemize}
|
||
\item 32 ядра CUDA для выполнения операций с целыми числами и с числами с
|
||
плавающей точкой;
|
||
\item 16 ядер загрузки/выгрузки данных;
|
||
\item 4 блока специального назначения (для вычисления сложных арифметических
|
||
функций).
|
||
\end{itemize}
|
||
|
||
Для распределения задач используется два планировщика варпов. Для хранения
|
||
используется:
|
||
\begin{itemize}
|
||
\item Регистровый файл, 128KB;
|
||
\item L1-cache, 16KB/48KB;
|
||
\item Разделяемая память, 48KB/16KB.
|
||
\end{itemize}
|
||
|
||
|
||
\subsubsection*{Микроархитектура Kepler}
|
||
В каждом потоковом процессоре:
|
||
\begin{itemize}
|
||
\item 192 ядра CUDA для выполнения операций с целыми числами и с числами с
|
||
плавающей точкой;
|
||
\item 64 блока для обработки чисел с двойной точностью;
|
||
\item 32 ядер загрузки/выгрузки данных;
|
||
\item 32 блока специального назначения (для вычисления сложных арифметических
|
||
функций).
|
||
\end{itemize}
|
||
|
||
Для распределения задач используется два планировщика варпов. Для хранения
|
||
используется:
|
||
\begin{itemize}
|
||
\item Регистровый файл, 256KB;
|
||
\item L1-cache, 16KB/48KB/32KB;
|
||
\item Разделяемая память, 48KB/16KB/32KB;
|
||
\item Константная память (Read-Only Data Cache), 48KB.
|
||
\end{itemize}
|
||
|
||
\subsubsection*{Микроархитектура Maxwell}
|
||
Каждый потоковый процессор в этой архитектуре состоит из четырех блоков. В
|
||
каждом блоке:
|
||
\begin{itemize}
|
||
\item Свой планировщик варпов;
|
||
\item Свой регистровый файл, 64KB;
|
||
\item 32 ядра CUDA;
|
||
\item 8 ядер загрузки/выгрузки данных;
|
||
\item 8 блока специального назначения (для вычисления сложных арифметических
|
||
функций).
|
||
\end{itemize}
|
||
Общими для четырех блоков являются L1/Texture кэш и разделенная память. В
|
||
отличие от предыдущих микроархитектур shared memory теперь является отдельным
|
||
блоком размером 96KB.
|
||
|
||
\subsubsection*{Микроархитектура Pascal}
|
||
В плане архитектуры Pascal отличается от Maxwell тем, что каждый потоковый
|
||
процессов содержит 2 блока.
|
||
|
||
\subsubsection*{Микроархитектура Tesla}
|
||
В Volta SM отказались от CUDA ядер и расщепили их на отдельные блоки (INT
|
||
+ FP32), что позволило использовать их одновременно и, соответственно, увеличило
|
||
общую производительность. Также в состав включили абсолютно новый компонент
|
||
под названием Tensor Core, нацеленный на увеличение производительности для глубокого обучения. По аналогии с Maxwell в SM есть 4 одинаковых блока, каждый из
|
||
которых содержит планировщик варпов, один модуль отправки команд, 8 ядер двойной точности, 32 целочисленных ядра, 32 ядра одинарной точности, 2 Tensor Core, 8
|
||
LD/ST и 1 SFU. Регистровый файл для каждого блока равен 64KB. Для Volta L1 кэш
|
||
и разделяемая память общие для всех блоков и снова объединены общим объемом в
|
||
128KB, где под разделяемую память можно отвести до 96 KB.
|
||
|
||
\subsubsection*{Вычислительные возможности Nvidia Tesla K40}
|
||
Программа в рамках данной работы исполнялась на видеокарте Nvidia Tesla
|
||
K40. Ниже приведены вычислительные возможности этой видеокарты:
|
||
|
||
\begin{itemize}
|
||
\item Размерность сетки блоков. Максимальное число измерений: 3; максимальное
|
||
число блоков по каждому измерению: $2^{31-1}$, 65535, 65535;
|
||
\item Размерность блока. Максимальное число измерений: 3; максимальное число
|
||
нитей по каждому измерению: 1024, 1024, 64; максимальное число нитей в
|
||
блоке: 1024;
|
||
\item Размер варпа: 64;
|
||
\item Размер регистрового файла. На блок: 64000 регистров; на нить: 255 регистров;
|
||
\item Размер разделяемой памяти на мультипроцессор: 48К;
|
||
\item Размер константной памяти на блок: 64KB.
|
||
\end{itemize}
|
||
|
||
|
||
\subsubsection{Потоковая модель}
|
||
Вычислительная архитектура CUDA основана на понятии мультипроцессора и
|
||
концепции SIMT (Single Instruction Multiple Threads). При выполнении многопоточной программы на видеокарте CUDA, все потоки разделяются на блоки, а внутри
|
||
блоков на варпы, где все потоки выполняют одну и ту же инструкцию. Группа блоков
|
||
выполняется на потоковом процессоре, распределением задач занимается планировщик. Программа, выполняющаяся на нескольких блоках одновременно называется
|
||
ядром (kernel).
|
||
|
||
Особенностью архитектуры CUDA является блочно-сеточная организация,
|
||
необычная для многопоточных приложений (Рис.~\ref{fig:streams}). При этом драйвер CUDA самостоятельно распределяет ресурсы устройства между потоками.
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.5\linewidth]{img/streams.png}
|
||
\caption{Организация потоков.}
|
||
\label{fig:streams}
|
||
\end{figure}
|
||
|
||
|
||
\subsection{Устройство памяти}
|
||
Видеокарта имеет собственную оперативную глобальную память, отдельную от
|
||
оперативной памяти CPU (на хосте). При выполнении кода на видеокарте (на устройстве) обращение может происходить только к памяти на видеокарте. Для перемещения данных с хоста на устройство и обратно используются служебные функции,
|
||
вызываемые с хоста. Помимо глобальной оперативной памяти на каждом мультипроцессоре есть свой кэш, своя текстурная, константная, разделяемая память и тд.
|
||
|
||
\subsection{Модели памяти}
|
||
В CUDA выделяют шесть видов памяти:
|
||
|
||
\begin{itemize}
|
||
\item регистры;
|
||
\item локальная память;
|
||
\item глобальная память;
|
||
\item разделяемая память;
|
||
\item константная память;
|
||
\item текстурная память.
|
||
\end{itemize}
|
||
|
||
На Рис.~\ref{fig:memory} представлено устройство памяти.
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.7\linewidth]{img/memory.png}
|
||
\caption{Устройство памяти.}
|
||
\label{fig:memory}
|
||
\end{figure}
|
||
|
||
\subsubsection*{Регистры}
|
||
На один мультипроцессор доступно 8192 32-разрядных регистров. Они распределяются между нитями в этом потоке. Обращение к этой памяти самое быстрое.
|
||
|
||
\subsubsection*{Глобальная память}
|
||
Глобальная память имеет большой объем. Она поддерживает произвольный доступ для всех мультипроцессоров, а также запись и чтение с хоста. Однако, эта
|
||
память очень медленная и не кэшируется, поэтому рекомендуется сократить количество обращений к этой памяти.
|
||
|
||
\subsubsection*{Локальная память}
|
||
Это небольшой объём памяти, к которому имеет доступ только один потоковый
|
||
процессор. Она относительно медленная — такая же, как и глобальная.
|
||
|
||
\subsubsection*{Разделяемая память}
|
||
Разделяемая память — это некэшируемая, но быстрая память. Ее и рекомендуется использовать как управляемый кэш. На один мультипроцессор доступно всего
|
||
16KB разделяемой памяти. Она обеспечивает взаимодействие потоков, управляется
|
||
разработчиком напрямую и имеет низкие задержки. Разделив это число на количество задач в блоке, получим максимальное количество разделяемой памяти, доступной на один поток.
|
||
|
||
\subsubsection*{Константная память}
|
||
Константная память — это тип памяти, который хранит неизменяемые данные, доступные на уровне сетки. Ее размер 64KB, физически она не отделена от глобальной памяти, но, используя системы кэшей и механизму широкого вещания
|
||
(broadcast), она может обеспечить прирост производительности за счет сокращения
|
||
трафика между процессором и памятью.
|
||
|
||
\subsubsection*{Текстурная память}
|
||
Текстурная память — это тип памяти, который похож на константную память,
|
||
поскольку через текстурный блок разрешены запросы только для чтения. Физически
|
||
текстурная память не отделена от глобальной. Как и константная память, позволяет
|
||
увеличить производительность за счет системы кэшей. Отличительной особенностью
|
||
является оптимизация текстурного кэша для двумерной пространственной локальности (данные расположены рядом в двумерном пространстве).
|
||
|
||
|
||
\subsection{Модель вычислений на GPU}
|
||
Программа, запускаемая на GPU с Nvidia CUDA, называется ядром (kernel).
|
||
Ядро запускается одновременно на сетке из блоков. Каждый из блоков, как было
|
||
сказано раньше, состоит из нескольких потоков. Количество блоков и потоков задается при вызове ядра. В каждом потоке выполняется один и тот же код. Добиться
|
||
разного поведения в потоках возможно при помощи информации о номере блока и
|
||
потока, которая доступна для каждого потока.
|
||
|
||
Все команды, пришедшие на GPU, будут исполняться в порядке общей для всех
|
||
потоков очереди. Последовательность исполнения может сгладить планировщик, который может запустить одновременно копирование с хоста, на хост и исполнение
|
||
ядра. А если ядро использует меньше 50\% мощности GPU, то запустить параллельно следующее ядро из
|
||
другого потока, если оно готово к запуску. Так, для оптимизации времени выполнения задач в разных потоках нужно учитывать, что на GPU команды будут
|
||
исполнятся в том порядке, в котором их вызвали в хост коде. Это значит, что не
|
||
всегда лучше заполнить задачами один поток, затем второй и т.д. Скорее более оптимальным подходом будет равномерный запуск задач по всем потокам. Для этого
|
||
можно сначала заполнить один поток командами, а затем другой. Тогда команды
|
||
второго потока будут ожидать окончания выполнения команд первого потока, или,
|
||
что более вероятно, начала копирования в память хоста из первого потока. Другой
|
||
вариант, когда команды будут распределяться по потокам поочередно, т.е. в первый
|
||
поток отправляется первая команда, затем во второй отправляется также первая, затем в первый вторая, и т.д. В результате такого равномерного распределения можно
|
||
добиться улучшения производительности за счет умения планировщика одновременно запускать операции копирования и ядра.
|
||
|
||
\subsection{Планировщик задач}
|
||
Поток (stream) в CUDA — логическая последовательность зависимых асинхронных операций, независимая от операций в других потоках. Потоки позволяют запускать CUDA команды на GPU в порядке, определенном в контексте одного потока.
|
||
С точки зрения GPU потоков не существует, и все команды, пришедшие на GPU,
|
||
будут исполняться в порядке общей для всех потоков очереди, знание этой особенности может помочь при оптимизации. Последовательность исполнения может сгладить планировщик, который может запустить одновременно копирование с хоста, на
|
||
хост и исполнение ядра. При вызове ядра можно указать поток, в который будет добавлено это ядро. По умолчанию все ядра добавляются в поток 0, который является
|
||
синхронным с хостом.
|
||
|
||
|
||
\subsection{Компиляция программы}
|
||
|
||
Программа для видеокарт Nvidia CUDA пишется на основе других языков, в
|
||
частности используется расширение языка C. Оно называется CUDA C. Для сборки
|
||
программы используется компилятор nvcc, который входит в пакет инструментов
|
||
разработчика. Этот пакет, а также библиотеку CUDA можно скачать с сайта Nvidia.
|
||
|
||
|
||
\newpage
|
||
\section{Суперкомпьютерный центр «Политехнический»}
|
||
|
||
\subsection{Состав}
|
||
|
||
Суперкомпьютерный центр «Политехнический» состоит из узлов трех типов:
|
||
|
||
\begin{itemize}
|
||
\item 668 узлов кластера «Политехник - РСК Торнадо»;
|
||
\item 288 узлов вычислителя с ультравысокой многопоточностью «Политехник -
|
||
РСК ПетаСтрим»;
|
||
\item 64 узла кластера «Политехник - NUMA».
|
||
\end{itemize}
|
||
|
||
\subsection{Характеристики}
|
||
|
||
\subsubsection*{Политехник - РСК Торнадо}
|
||
Кластер содержит узлы двух типов:
|
||
|
||
\begin{itemize}
|
||
\item 612 узлов с прямым жидкостным охлаждением серии «Торнадо»(производитель
|
||
РСК Технологии РФ), имеющие каждый два CPU Intel Xeon E5-2697 v3 (14
|
||
ядер, 2.6 ГГц) и 64 ГБ оперативной памяти DDR4;
|
||
\item 56 узлов с прямым жидкостным охлаждением серии Tornado содержащие каждый два CPU Intel Xeon E5-2697 v3 и два ускорителя вычислений NVIDIA Tesla
|
||
K40X, 64 ГБ оперативной памяти DDR4.
|
||
\end{itemize}
|
||
|
||
\subsubsection*{Политехник - РСК ПетаСтрим
|
||
}
|
||
Кластер содержит узлы двух типов:
|
||
|
||
\begin{itemize}
|
||
\item 288 однопроцессорных узлов с пиковой производительностью 1 ТФлопс каждый;
|
||
\item 17280 многопоточных ядер общего назначения (69120) потоков, поддерживающих векторную обработку данных посредством аппаратно реализованных
|
||
инструкций FMA (Fused Multiply-Accumulate);
|
||
\item оперативная память узла - 8 ГБ, GDDR5; суммарный объём оперативной памяти системы 2304 ГБ;
|
||
\item пропускная способность между двумя узлами модуля системы на тесте MPI
|
||
OSU или Intel MPI Benchmarks не менее 6 ГБ/с.
|
||
\end{itemize}
|
||
|
||
\subsubsection*{Политехник - NUMA}
|
||
Кластер содержит узлы двух типов:
|
||
|
||
\begin{itemize}
|
||
\item 64 вычислительных узла, каждый из которых включает:
|
||
\begin{itemize}
|
||
\item 3 CPU AMD Opteron 638;
|
||
\item Адаптер NumaConnect N313-48;
|
||
\item 192 ГБ оперативной памяти;
|
||
\end{itemize}
|
||
\item 192 процессора;
|
||
\item 3072 ядер х86.
|
||
\end{itemize}
|
||
|
||
\subsection{Технология подключения}
|
||
|
||
Для подключения зарегистрированного пользователя к СКЦ необходимо использовать SSH клиент. С помощью него получается доступ к удаленному терминалу
|
||
для работы с ресурсами СКЦ.
|
||
|
||
В рамках работы была использована следующая технология подключения:
|
||
\begin{enumerate}
|
||
\item Были получен приватный ключ от администрации СКЦ в виде файла.
|
||
|
||
\item При помощи команды ssh был произведен вход: ssh -v tm3u21@login1.hpc.spbstu.ru
|
||
-i /.ssh/id\_rsa, где tm3u21 - логин, login1.hpc.spbstu.ru - адрес, id\_rsa - приватный ключ.
|
||
|
||
\item Чтобы переслать файлы, использовалась команда \\
|
||
scp -r D:\textbackslash kernel.cu tm3u21@login1.hpc.spbstu.ru\:home/kernel.cu. \\ Где \texttt{"D:\textbackslash kernel.cu"} -- путь
|
||
до файла на локальном компьютере, tm3u21 - логин, login1.hpc.spbstu.ru - адрес,
|
||
home/kernel.cu - пусть сохранения файла на СКЦ.
|
||
\end{enumerate}
|
||
|
||
|
||
\newpage
|
||
\section{Постановка решаемой практической задачи}
|
||
|
||
Построение пути движения робота по полигону.
|
||
|
||
Дано:
|
||
\begin{itemize}
|
||
\item массив Р - полигон (nxn), -- целых чисел;
|
||
\item P1(x1,y1), P2(x2,y2) точки на полигоне;
|
||
\item V - Контуры, запрещенные для движения, граница контура имеет значение -1.
|
||
\end{itemize}
|
||
|
||
Надо:
|
||
\begin{itemize}
|
||
\item построить L - траекторию, соединяющую P1 и P2, либо вывести на экран сообщение о том, что такой траектории не существует.
|
||
\end{itemize}
|
||
|
||
Ограничения:
|
||
\begin{itemize}
|
||
\item L не включает точки полигона со значениями P(I,j)=-1
|
||
\item M -- Число контуров -- случайная величина;
|
||
\item контуры V могут пересекаться;
|
||
\end{itemize}
|
||
|
||
|
||
\newpage
|
||
\section{Алгоритм решения задачи}
|
||
Для решения задачи используется волновой алгоритм, также известный как алгоритм Ли. Алгоритм состоит из двух частей. В первой части от начальной точки к конечной распространяется волна. Во второй выполняется обратный ход, в процессе которого из ячеек волны формируется путь.
|
||
|
||
Волна, идущая от начальной к конечной точке, на каждом шаге первой части алгоритма пополняется свободными ячейками, которые, во-первых, еще не принадлежат волне, и, во-вторых, являются соседями ячеек, попавших в волну на предыдущем шаге. Непроходимые ячейки просто игнорируются. Новые ячейки волны заполняются минимальным значением соседской клетки, которая уже принадлежит волне, увеличенным на единицу. Начальная клетка заполняется нулём и является стартовой для распространения волны. Для определения соседних ячеек используется окрестность фон Неймана (Рис.~\ref{fig:fon}).
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.2\linewidth]{img/fon.png}
|
||
\caption{Окрестность фон Неймана.}
|
||
\label{fig:fon}
|
||
\end{figure}
|
||
|
||
На Рис.~\ref{fig:first-step}-\ref{fig:last-step} представлены иллюстрации работы первой части алгоритма.
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.45\linewidth]{img/first-step.png}
|
||
\caption{Первый шаг распространения волны.}
|
||
\label{fig:first-step}
|
||
\end{figure}
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.45\linewidth]{img/third-step.png}
|
||
\caption{Третий шаг распространения волны.}
|
||
\label{fig:third-step}
|
||
\end{figure}
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.45\linewidth]{img/last-step.png}
|
||
\caption{Последний шаг распространения волны.}
|
||
\label{fig:last-step}
|
||
\end{figure}
|
||
|
||
Восстановление пути происходит в обратном направлении: при выборе ячейки от финишной ячейки к стартовой на каждом шаге выбирается ячейка со значением на единицу меньше текущей ячейки (Рис.~\ref{fig:reverse}).
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=0.45\linewidth]{img/reverse.png}
|
||
\caption{Восстановление пути.}
|
||
\label{fig:reverse}
|
||
\end{figure}
|
||
|
||
Волновой алгоритм либо находит кратчайший путь от начальной к конечной точке, либо информирует о неудаче, если путь к конечной точке блокируется препятствиями.
|
||
|
||
\subsection{Метод распараллеливания алгоритма}
|
||
Рассмотрим два подхода к распараллеливанию этой задачи на CUDA: с использованием глобальной памяти и с использованием разделяемой (shared) памяти.
|
||
|
||
В обоих подходах восстановление пути происходит на CPU и не подвергается распараллеливанию.
|
||
|
||
\subsubsection{Глобальная память}
|
||
Используется одномерная сетка блоков и одномерная сетка потоков внутри блоков. Таким образом уникальный идентификатор потока вычисляется по формуле:
|
||
|
||
\texttt{tid = threadIdx.x + blockIdx.x * blockDim.x}
|
||
|
||
Каждый поток обрабатывает несколько ячеек исходной матрицы с шагом \texttt{blockDim.x * gridDim.x} (общее количество потоков во всех блоках). Каждая клетка исходной матрицы обрабатывается только одним потоком.
|
||
|
||
На каждом шаге алгоритма каждый поток пытается уменьшить значения своих клеток. Поток обращается к глобальной памяти, чтобы получить значения соседних клеток. Значение текущей клетки обновляется, если значение одной из соседних клеток, увеличенное на единицу, меньше текущего значения клетки. То есть по следующей формуле:
|
||
|
||
$$
|
||
cell = \min(cell, left + 1, right + 1, top + 1, bottom + 1),
|
||
$$
|
||
где $cell$ -- значение в текущей клетке, $left$ -- значение в соседней слева, $right$ -- значение в соседней справа, $top$ -- значение в соседней сверху, $bottom$ -- значение в соседней снизу. Клетки с препятствиями игнорируются и не участвуют в формуле.
|
||
|
||
Следующий шаг алгоритма запускается, если хотя бы одному потоку удалось обновить значение в какой-нибудь своей клетке. Алгоритм завершается, если больше не осталось незаполненных клеток и дальнейшее обновление расстояний невозможно. Для отслеживания изменений используется флаг, который также хранится в глобальной памяти.
|
||
|
||
\subsubsection{Разделяемая память}
|
||
Используется одномерная сетка блоков. Потоки внутри блоков организованы в двумерную сетку.
|
||
|
||
Каждый блок потоков обрабатывает несколько соразмерных ему подматриц исходной матрицы с шагом, равным количеству используемых блоков. Каждая подматрица обрабатывается только одним блоком. Посколько блок и поматрица имеют равные размеры, каждому потоку блока в соответствие ставится одна ячейка поматрицы. Таким образом каждая ячейка исходной матрицы обрабатывается только одним потоком.
|
||
|
||
На каждом шаге алгоритма каждый блок потоков копирует соответствующую подматрицу исходной матриц из глобальной памяти в разделяемую. Затем на подматрице запускается паралелльная версия волнового алгоритма, описанная в предыдущем разделе, с двумя отличиями:
|
||
\begin{itemize}
|
||
\item К глобальной памяти обращаются только те потоки, которые обрабатывают клетки на границах подматрицы. Потому что некоторые соседи граничных клеток находятся вне скопированной в разделяемую память подматрицы. Такой подход позволяет значительно сократить обращения к глобальной памяти при использовании больших размеров блоков.
|
||
\item Локальный флаг для отслеживания изменений хранится в разделяемой памяти каждого потока.
|
||
\end{itemize}
|
||
После выполнения волнового алгоритма на подматрице в разделяемой памяти, её значения копируются в соответствующую подматрицу исходной матрицы в глобальной памяти.
|
||
|
||
Таким образом, на каждом шаге алгоритма к волне прибавляются не просто отдельные ячейки, а целые подматрицы, соразмерные используемым блокам потоков.
|
||
|
||
Следующий шаг алгоритма запускается, если хотя бы одному блоку потоков удалось обновить значения в своей подматрице. Алгоритм завершается, если больше не осталось незаполненных подматриц и дальнейшее обновление расстояний невозможно. Для отслеживания изменений в подматрицах используется дополнительный флаг, хранящийся в глобальной памяти.
|
||
|
||
|
||
\newpage
|
||
\section{Описание эксперимента}
|
||
В этом разделе выполняется исследование времени решения задачи при изменении следующих параметров:
|
||
|
||
\begin{itemize}
|
||
\item Размеры матрицы: 500 × 500, 1000 × 1000, 1500 × 1500;
|
||
\item Количество блоков: 1, 10, 100, 1000, 10000;
|
||
\item Количество потоков: 1, 9, 100, 1024;
|
||
\item Используемая память: глобальная, глобальная и разделяемая.
|
||
\end{itemize}
|
||
|
||
Для каждой комбинации параметров было проведено 100 измерений с использованием событий CUDA, после чего вычислялось среднее значение.
|
||
|
||
\newpage
|
||
\section{Анализ результатов}
|
||
В таблицах 1, 2, 3 приведены результаты измерения времени в миллисекундах для глобальной памяти для размеров матрицы 100 × 100, 500 × 500, 1000 × 1000 соответственно. В таблицах 4, 5, 6 приведены результаты измерения времени в миллисекундах для разделяемой памяти для размеров матрицы 100 × 100, 500 × 500, 1000 × 1000 соответственно.
|
||
|
||
\begin{table}[h!]
|
||
\centering
|
||
\caption{Результаты измерения времени исполнения программы для матрицы 100 × 100 и глобальной памяти. Время указано в миллисекундах.}
|
||
\footnotesize
|
||
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||
\hline
|
||
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||
\hline
|
||
& 1 & 10 & 100 & 1000 & 10000 \\
|
||
\hline
|
||
1 & 171.85 & 120.98 & 19.41 & 12.84 & 23.68 \\
|
||
\hline
|
||
9 & 233.35 & 29.13 & 8.14 & 5.68 & 10.55 \\
|
||
\hline
|
||
100 & 27.57 & 7.79 & 3.90 & \textbf{3.79} & 8.97 \\
|
||
\hline
|
||
1024 & 7.96 & 4.04 & 4.16 & 6.47 & 32.24 \\
|
||
\hline
|
||
\end{tabularx}
|
||
\end{table}
|
||
|
||
\begin{table}[h!]
|
||
\centering
|
||
\caption{Результаты измерения времени исполнения программы для матрицы 500 × 500 и глобальной памяти. Время указано в миллисекундах.}
|
||
\footnotesize
|
||
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||
\hline
|
||
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||
\hline
|
||
& 1 & 10 & 100 & 1000 & 10000 \\
|
||
\hline
|
||
1 & 44648 & 15096 & 1694 & 1027 & 1140 \\
|
||
\hline
|
||
9 & 23257 & 4075 & 554 & 291 & 223 \\
|
||
\hline
|
||
100 & 3643 & 536 & 93 & 55 & 77 \\
|
||
\hline
|
||
1024 & 596 & 88 & \textbf{49} & 64 & 196 \\
|
||
\hline
|
||
\end{tabularx}
|
||
\end{table}
|
||
|
||
\begin{table}[h!]
|
||
\centering
|
||
\caption{Результаты измерения времени исполнения программы для матрицы 1000 × 1000 и глобальной памяти. Время указано в миллисекундах.}
|
||
\footnotesize
|
||
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||
\hline
|
||
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||
\hline
|
||
& 1 & 10 & 100 & 1000 & 10000 \\
|
||
\hline
|
||
1 & 189324 & 121341 & 12760 & 7010 & 6472 \\
|
||
\hline
|
||
9 & 164810 & 29369 & 4435 & 2208 & 1356 \\
|
||
\hline
|
||
100 & 28946 & 2985 & 628 & 329 & 360 \\
|
||
\hline
|
||
1024 & 3701 & 609 & \textbf{277} & 322 & 579 \\
|
||
\hline
|
||
\end{tabularx}
|
||
\end{table}
|
||
|
||
|
||
|
||
\begin{table}[h!]
|
||
\centering
|
||
\caption{Результаты измерения времени исполнения программы для матрицы 100 × 100 и разделяемой памяти. Время указано в миллисекундах.}
|
||
\footnotesize
|
||
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||
\hline
|
||
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||
\hline
|
||
& 1 & 10 & 100 & 1000 & 10000 \\
|
||
\hline
|
||
1 & 260.60 & 230.96 & 47.04 & 18.13 & 32.70 \\
|
||
\hline
|
||
9 & 28.65 & 19.92 & 7.27 & 3.10 & 4.38 \\
|
||
\hline
|
||
100 & 5.46 & 4.71 & \textbf{1.35} & 1.36 & 1.45 \\
|
||
\hline
|
||
1024 & 4.39 & 2.68 & 2.65 & 2.67 & 3.00 \\
|
||
\hline
|
||
\end{tabularx}
|
||
\end{table}
|
||
|
||
\begin{table}[h!]
|
||
\centering
|
||
\caption{Результаты измерения времени исполнения программы для матрицы 500 × 500 и разделяемой памяти. Время указано в миллисекундах.}
|
||
\footnotesize
|
||
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||
\hline
|
||
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||
\hline
|
||
& 1 & 10 & 100 & 1000 & 10000 \\
|
||
\hline
|
||
1 & 48837 & 31345 & 3566 & 1786 & 1724 \\
|
||
\hline
|
||
9 & 1120 & 2109 & 347 & 163 & 148 \\
|
||
\hline
|
||
100 & 166 & 166 & 65 & 37 & \textbf{17} \\
|
||
\hline
|
||
1024 & 99 & 64 & 34 & 112 & 114 \\
|
||
\hline
|
||
\end{tabularx}
|
||
\end{table}
|
||
|
||
\begin{table}[h!]
|
||
\centering
|
||
\caption{Результаты измерения времени исполнения программы для матрицы 1000 × 1000 и разделяемой памяти. Время указано в миллисекундах.}
|
||
\footnotesize
|
||
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||
\hline
|
||
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||
\hline
|
||
& 1 & 10 & 100 & 1000 & 10000 \\
|
||
\hline
|
||
1 & 127512 & 48471 & 24677 & 14065 & 12445 \\
|
||
\hline
|
||
9 & 5775 & 16595 & 2062 & 1138 & 727 \\
|
||
\hline
|
||
100 & 656 & 961 & 491 & 124 & \textbf{103} \\
|
||
\hline
|
||
1024 & 398 & 190 & 222 & 384 & 618 \\
|
||
\hline
|
||
\end{tabularx}
|
||
\end{table}
|
||
|
||
\newpage
|
||
В целом, увеличение количества потоков приводит к уменьшению времени выполнения вплоть до некоторого предела, после которого время начинает возрастать. Это связано с тем, что слишком большое число потоков приводит к неэффективному использованию ресурсов GPU, вызывая задержки из-за управления бездействующими потоками. В среднем лучше всего себя показывают конфигурации, где количестов потоков близко к количеству элементов в матрице.
|
||
|
||
Были выделены лучшие конфигурации для матриц с разными размерами:
|
||
\begin{itemize}
|
||
\item Для матрицы 100 × 100:
|
||
\begin{itemize}
|
||
\item Глобальная память: 100 блоков, 100 потоков, 3.79 мс.
|
||
\item Разделяемая память: 100 блоков, 100 потоков, 1.35 мс.
|
||
\end{itemize}
|
||
\item Для матрицы 500 × 500:
|
||
\begin{itemize}
|
||
\item Глобальная память: 100 блоков, 1024 потока, 49 мс.
|
||
\item Разделяемая память: 10 000 блоков, 100 потоков, 17 мс.
|
||
\end{itemize}
|
||
\item Для матрицы 1000 × 1000:
|
||
\begin{itemize}
|
||
\item Глобальная память: 100 блоков, 1024 потока, 277 мс.
|
||
\item Разделяемая память: 10 000 блоков, 100 потоков, 103 мс.
|
||
\end{itemize}
|
||
\end{itemize}
|
||
|
||
На Рис.~\ref{fig:plot} отображена зависимость времения от разного числа потоков в определенной выборке блоков.
|
||
|
||
\begin{figure}[h!]
|
||
\centering
|
||
\includegraphics[width=1\linewidth]{img/plot.png}
|
||
\caption{Зависимость времени выполнения программы от разного числа потоков в блоке для матрицы 1000 x 1000 и глобальной памяти. Используется логарифмический масштаб по обеим осям.}
|
||
\label{fig:plot}
|
||
\end{figure}
|
||
|
||
Использование разделяемой памяти в большинстве случаев позволяет значительно ускорить выполнение алгоритма по сравнению с использованием только глобальной памяти (в среднем на 60\%). Особенно заметно это на больших размерах матриц, где доступ к глобальной памяти становится узким местом.
|
||
|
||
Алгоритму с разделяемой памятью важно, чтобы в каждом блоке было как можно больше потоков. Это связано с тем, что чем больше потоков в блоке, тем больше ячеек матрицы можно обработать одновременно, используя быструю разделяемую память вместо глобальной. При небольшом количестве потоков в блоке преимуществ разделяемой памяти становится меньше, так как потоки вынуждены чаще обращаться к глобальной памяти, что снижает общую производительность.
|
||
|
||
|
||
\newpage
|
||
\phantom{text}
|
||
\newpage
|
||
\section*{Заключение}
|
||
\addcontentsline{toc}{section}{Заключение}
|
||
|
||
В рамках курсовой работы было изучена технология параллельного программирования на основе архитектуры Nvidia CUDA.
|
||
|
||
Для задачи построения пути движения робота по полигону был разработан параллельный асинхронный алгоритм, алгоритм был реализован на языке CUDA C. Программа была
|
||
запущена на ресурсах суперкомпьютерного центра «Политехнический». Для запуска использовался узел типа «Торнадо» с видеокартой NVIDIA Tesla K40X. Запуск
|
||
программы проводился на одном узле с использованием одной видеокарты.
|
||
|
||
Было измерено время работы программы при различной степени распараллеливания, разных размерах матриц и разной памяти. Использование оптимальной конфигурации позволило
|
||
уменьшить время выполнения в 100 раз относительно наихудшей конфигурации для
|
||
матрциы 100 × 100, в 1000 раз для 500 × 500 и в 1500 раз для 1000 × 1000.
|
||
|
||
Реализация алгоритма с использованием разделяемой памяти показала значительно более высокую эффективность. Время выполнения алгоритма снизилось в среднем на 60\%.
|
||
|
||
В рамках курсовой работы была написана программа размером 280 строк. Работа на СКЦ «Политехнический» шла две недели, за это время было сделано примерно 100 запусков задач на исполнение.
|
||
Для сборки использовался компилятор NVCC версии 11.6u2.
|
||
|
||
|
||
\newpage
|
||
\section*{Список литературы}
|
||
\addcontentsline{toc}{section}{Список литературы}
|
||
|
||
\vspace{-1.5cm}
|
||
\begin{thebibliography}{0}
|
||
\bibitem{mayers}
|
||
Сандерс, Д. Технология CUDA в примерах: введение в программирование графических процессоров -- Москва: изд. ДМК Пресс, 2013 г -- 232 с.
|
||
\end{thebibliography}
|
||
|
||
\end{document} |