\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*{Заключение}
    \addcontentsline{toc}{section}{Заключение}


\newpage
\section*{Список литературы}
\addcontentsline{toc}{section}{Список литературы}

\vspace{-1.5cm}
\begin{thebibliography}{0}
    \bibitem{mayers}
    Майерс, Г. Искусство тестирования программ. -- Санкт-Петербург: Диалектика, 2012 г.
\end{thebibliography}

\end{document}