\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}